From e6d1487acd58c7db0fc260447ac256a4f71f6916 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 26 Sep 2014 12:41:58 -0400 Subject: ADD: support ptxinfo for sm_35 and cuda 6.5 --- src/cuda-sim/ptx.l | 1 + src/cuda-sim/ptx.y | 2 ++ src/cuda-sim/ptxinfo.y | 4 ++++ 3 files changed, 7 insertions(+) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index dfed936..58bdf3d 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -180,6 +180,7 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; \.union TC; return UNION_DIRECTIVE; /* not in PTX 2.1 */ \.version TC; return VERSION_DIRECTIVE; \.visible TC; return VISIBLE_DIRECTIVE; +\.weak TC; return WEAK_DIRECTIVE; \.address_size TC; return ADDRESS_SIZE_DIRECTIVE; \.constptr TC; return CONSTPTR_DIRECTIVE; /* Ptx plus directive for pointer to constant memory */ diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index 79faddf..2f85213 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -71,6 +71,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %token VERSION_DIRECTIVE %token ADDRESS_SIZE_DIRECTIVE %token VISIBLE_DIRECTIVE +%token WEAK_DIRECTIVE %token IDENTIFIER %token INT_OPERAND %token FLOAT_OPERAND @@ -243,6 +244,7 @@ function_decl_header: ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".ent | FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } | VISIBLE_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } | EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; g_func_decl=1; func_header(".func"); } + | WEAK_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } ; param_list: /*empty*/ diff --git a/src/cuda-sim/ptxinfo.y b/src/cuda-sim/ptxinfo.y index 294412d..1233da4 100644 --- a/src/cuda-sim/ptxinfo.y +++ b/src/cuda-sim/ptxinfo.y @@ -85,6 +85,7 @@ line: HEADER INFO COLON line_info line_info: function_name | function_info { ptxinfo_addinfo(); } + | gmem_info ; function_name: FUNC QUOTE IDENTIFIER QUOTE { ptxinfo_function($3); } @@ -95,6 +96,9 @@ function_info: info | function_info COMMA info ; +gmem_info: INT_OPERAND BYTES GMEM + ; + info: USED INT_OPERAND REGS { ptxinfo_regs($2); } | tuple LMEM { ptxinfo_lmem(g_declared,g_system); } | tuple SMEM { ptxinfo_smem(g_declared,g_system); } -- cgit v1.3 From 5cc6392a98bf736a05ffa7c7c557c4556753c8c2 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Wed, 1 Oct 2014 23:23:51 -0400 Subject: ADD: handle child kernel name in mov instruction. ADD: detect call cudaGetParameterBufferV2 and call cudaLaunchDeviceV2 --- libcuda/cuda_runtime_api.cc | 1 - src/cuda-sim/instructions.cc | 13 ++++++++++++- 2 files changed, 12 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 5dc9cc1..5310a52 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -266,7 +266,6 @@ struct CUctx_st { m_kernel_lookup[hostFun] = f; } else { - m_code[fat_cubin_handle]->dump(); printf("Warning: cannot find deviceFun %s\n", deviceFun); m_kernel_lookup[hostFun] = NULL; } diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index cf7f04a..254427b 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -148,6 +148,9 @@ ptx_reg_t ptx_thread_info::get_operand_value( const operand_info &op, operand_in result.u64 = op.get_symbol()->get_address(); } else if ( op.is_local() ) { result.u64 = op.get_symbol()->get_address(); + } else if ( op.is_function_address() ) { + result.u64 = op.get_symbol()->get_pc()->get_start_PC(); + printf("Get pc for kernel function %u\n", op.get_symbol()->get_pc()->get_start_PC()); } else { const char *name = op.name().c_str(); printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown operand type for %s\n", name ); @@ -1407,7 +1410,15 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) if( fname == "vprintf" ) { gpgpusim_cuda_vprintf(pI, thread, target_func); return; - } + } + else if(fname == "cudaGetParameterBufferV2") { + printf("calling cudaGetParameterBufferV2\n"); + return; + } + else if(fname == "cudaLaunchDeviceV2") { + printf("calling cudaLaunchDeviceV2\n"); + return; + } // read source arguements into register specified in declaration of function arg_buffer_list_t arg_values; -- cgit v1.3 From db5fe9600dd57ce59864b0f22c9d5407231ab5b1 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 3 Oct 2014 19:08:39 -0400 Subject: ADD: initial support for instruction group used by CDP --- cuobjdump_to_ptxplus/ptx_parser.h | 5 +++++ src/cuda-sim/ptx.y | 4 ++-- src/cuda-sim/ptx_ir.cc | 40 +++++++++++++++++++++++++++++++++++++++ src/cuda-sim/ptx_ir.h | 9 +++++++++ src/cuda-sim/ptx_parser.cc | 11 +++++++++++ src/cuda-sim/ptx_parser.h | 4 ++++ 6 files changed, 71 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim') diff --git a/cuobjdump_to_ptxplus/ptx_parser.h b/cuobjdump_to_ptxplus/ptx_parser.h index 1c96b46..418a733 100644 --- a/cuobjdump_to_ptxplus/ptx_parser.h +++ b/cuobjdump_to_ptxplus/ptx_parser.h @@ -110,6 +110,11 @@ void add_alignment_spec( int ) {PTX_PARSE_DPRINTF(" ");} void add_pragma( const char *a ) {PTX_PARSE_DPRINTF(" ");} void add_constptr(const char* identifier1, const char* identifier2, int offset) {PTX_PARSE_DPRINTF(" ");} +//Jin: handle instructino group for cdp +void start_inst_group(){PTX_PARSE_DPRINTF(" ");}; +void end_inst_group(){PTX_PARSE_DPRINTF(" ");}; + + /*non-dummy stuff below this point*/ extern cuobjdumpInstList *g_headerList; diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index 2f85213..c8208ea 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -270,8 +270,8 @@ statement_list: directive_statement { add_directive(); } | instruction_statement { add_instruction(); } | statement_list directive_statement { add_directive(); } | statement_list instruction_statement { add_instruction(); } - | statement_list statement_block - | statement_block + | statement_list {start_inst_group();} statement_block {end_inst_group();} + | {start_inst_group();} statement_block {end_inst_group();} ; directive_statement: variable_declaration SEMI_COLON diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 751b3f4..915c623 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -90,6 +90,11 @@ symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol m_const_next = 0; m_global_next = 0x100; m_local_next = 0; + m_tex_next = 0; + + //Jin: handle instruction group for cdp + m_inst_group_id = 0; + m_parent = parent; if ( m_parent ) { m_shared_next = m_parent->m_shared_next; @@ -170,6 +175,41 @@ void symbol_table::add_function( function_info *func, const char *filename, unsi m_symbols[ func->get_name() ] = s; } +//Jin: handle instruction group for cdp +symbol_table* symbol_table::start_inst_group() { + char inst_group_name[1024]; + snprintf(inst_group_name, 1024, "%s_inst_group_%u", m_scope_name.c_str(), m_inst_group_id); + + //previous added + assert(m_inst_group_symtab.find(std::string(inst_group_name)) == m_inst_group_symtab.end()); + symbol_table *sym_table = new symbol_table(inst_group_name, 3/*inst group*/, this ); + + sym_table->m_global_next = m_global_next; + sym_table->m_shared_next = m_shared_next; + sym_table->m_local_next = m_local_next; + sym_table->m_reg_allocator = m_reg_allocator; + sym_table->m_tex_next = m_tex_next; + sym_table->m_const_next = m_const_next; + + m_inst_group_symtab[std::string(inst_group_name)] = sym_table; + + return sym_table; +} + +symbol_table * symbol_table::end_inst_group() { + symbol_table * sym_table = m_parent; + + sym_table->m_global_next = m_global_next; + sym_table->m_shared_next = m_shared_next; + sym_table->m_local_next = m_local_next; + sym_table->m_reg_allocator = m_reg_allocator; + sym_table->m_tex_next = m_tex_next; + sym_table->m_const_next = m_const_next; + sym_table->m_inst_group_id++; + + return sym_table; +} + void register_ptx_function( const char *name, function_info *impl ); // either libcuda or libopencl bool symbol_table::add_function_decl( const char *name, int entry_point, function_info **func_info, symbol_table **sym_table ) diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 601a13d..7325e5f 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -330,6 +330,11 @@ public: iterator const_iterator_end() { return m_consts.end();} void dump(); + + //Jin: handle instruction group for cdp + symbol_table* start_inst_group(); + symbol_table* end_inst_group(); + private: unsigned m_reg_allocator; unsigned m_shared_next; @@ -347,6 +352,10 @@ private: std::list m_consts; std::map m_function_info_lookup; std::map m_function_symtab_lookup; + + //Jin: handle instruction group for cdp + unsigned m_inst_group_id; + std::map m_inst_group_symtab; }; class operand_info { diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 824714a..39257da 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -187,6 +187,17 @@ void add_function_name( const char *name ) g_global_symbol_table->add_function( g_func_info, g_filename, ptx_lineno ); } +//Jin: handle instruction group for cdp +void start_inst_group() { + PTX_PARSE_DPRINTF("start_instruction_group"); + g_current_symbol_table = g_current_symbol_table->start_inst_group(); +} + +void end_inst_group() { + PTX_PARSE_DPRINTF("end_instruction_group"); + g_current_symbol_table = g_current_symbol_table->end_inst_group(); +} + void add_directive() { PTX_PARSE_DPRINTF("add_directive"); diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index fef7635..32f3903 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -94,6 +94,10 @@ void change_operand_neg( ); void set_immediate_operand_type( ); void version_header(double a); +//Jin: handle instructino group for cdp +void start_inst_group(); +void end_inst_group(); + #define NON_ARRAY_IDENTIFIER 1 #define ARRAY_IDENTIFIER_NO_DIM 2 #define ARRAY_IDENTIFIER 3 -- cgit v1.3 From 70e02ee5283cb96f0edcb46a15edf0ab6e1d0697 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Mon, 13 Oct 2014 18:58:37 -0400 Subject: ADD: add cudaGetParameterBufferV2 and add cudaLaunchDeviceV2 implementation. Kernel launch to stream not yet implemented --- src/cuda-sim/Makefile | 3 +- src/cuda-sim/cuda-sim.cc | 10 +++ src/cuda-sim/cuda_device_runtime.cc | 175 ++++++++++++++++++++++++++++++++++++ src/cuda-sim/cuda_device_runtime.h | 7 ++ src/cuda-sim/instructions.cc | 12 ++- src/cuda-sim/ptx_ir.h | 4 +- src/cuda-sim/ptx_parser.cc | 14 +++ 7 files changed, 219 insertions(+), 6 deletions(-) create mode 100644 src/cuda-sim/cuda_device_runtime.cc create mode 100644 src/cuda-sim/cuda_device_runtime.h (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile index 166e256..f479294 100644 --- a/src/cuda-sim/Makefile +++ b/src/cuda-sim/Makefile @@ -62,7 +62,7 @@ ifeq ($(GNUC_CPP0X),1) endif endif -OBJS := $(OUTPUT_DIR)/ptx_parser.o $(OUTPUT_DIR)/ptx_loader.o $(OUTPUT_DIR)/cuda_device_printf.o $(OUTPUT_DIR)/instructions.o $(OUTPUT_DIR)/cuda-sim.o $(OUTPUT_DIR)/ptx_ir.o $(OUTPUT_DIR)/ptx_sim.o $(OUTPUT_DIR)/memory.o $(OUTPUT_DIR)/ptx-stats.o $(OUTPUT_DIR)/decuda_pred_table/decuda_pred_table.o $(OUTPUT_DIR)/ptx.tab.o $(OUTPUT_DIR)/lex.ptx_.o $(OUTPUT_DIR)/ptxinfo.tab.o $(OUTPUT_DIR)/lex.ptxinfo_.o +OBJS := $(OUTPUT_DIR)/ptx_parser.o $(OUTPUT_DIR)/ptx_loader.o $(OUTPUT_DIR)/cuda_device_printf.o $(OUTPUT_DIR)/instructions.o $(OUTPUT_DIR)/cuda-sim.o $(OUTPUT_DIR)/ptx_ir.o $(OUTPUT_DIR)/ptx_sim.o $(OUTPUT_DIR)/memory.o $(OUTPUT_DIR)/ptx-stats.o $(OUTPUT_DIR)/decuda_pred_table/decuda_pred_table.o $(OUTPUT_DIR)/ptx.tab.o $(OUTPUT_DIR)/lex.ptx_.o $(OUTPUT_DIR)/ptxinfo.tab.o $(OUTPUT_DIR)/lex.ptxinfo_.o $(OUTPUT_DIR)/cuda_device_runtime.o OPT += -DCUDART_VERSION=$(CUDART_VERSION) @@ -145,5 +145,6 @@ $(OUTPUT_DIR)/ptx_sim.o: $(OUTPUT_DIR)/ptx.tab.c $(OUTPUT_DIR)/cuda-sim.o: $(OUTPUT_DIR)/ptx.tab.c $(OUTPUT_DIR)/lex.ptxinfo_.o: $(OUTPUT_DIR)/ptx.tab.c $(OUTPUT_DIR)/lex.ptx_.o: $(OUTPUT_DIR)/ptx.tab.c +$(OUTPUT_DIR)/cuda_device_runtime.o: $(OUTPUT_DIR)/ptx.tab.c include $(OUTPUT_DIR)/Makefile.makedepend diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 715be98..4933029 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1065,6 +1065,16 @@ void function_info::add_param_data( unsigned argn, struct gpgpu_ptx_sim_arg *arg } } +unsigned function_info::get_args_aligned_size() { + unsigned int align_size = 4; // a word + unsigned int total_size = 0; + for(unsigned int i = 0; i < num_args(); i++) { + total_size += ((m_args[i]->get_size_in_bytes() + align_size - 1) / align_size) * align_size; + } + return total_size; +} + + void function_info::finalize( memory_space *param_mem ) { unsigned param_address = 0; diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc new file mode 100644 index 0000000..937eec8 --- /dev/null +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -0,0 +1,175 @@ +//Jin: cuda_device_runtime.cc +//Defines CUDA device runtime APIs for CDP support + +#include +#include + +#define __CUDA_RUNTIME_API_H__ + +#include +#include +#include "../gpgpu-sim/gpu-sim.h" +#include "cuda-sim.h" +#include "ptx_ir.h" +#include "../stream_manager.h" +#include "cuda_device_runtime.h" + +#define DEV_RUNTIME_REPORT(a) \ + if( g_debug_execution ) { \ + std::cout << __FILE__ << ", " << __LINE__ << ": " << a << "\n"; \ + std::cout.flush(); \ + } + +std::map g_cuda_device_launch_map; +struct CUstream_st * g_device_default_stream = NULL; +extern stream_manager *g_stream_manager; + +//Handling device runtime api: +//void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) +void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func) +{ + DEV_RUNTIME_REPORT("Calling cudaGetParameterBufferV2"); + + unsigned n_return = target_func->has_return(); + assert(n_return); + unsigned n_args = target_func->num_args(); + assert( n_args == 4 ); + + function_info * child_kernel_entry; + struct dim3 gridDim, blockDim; + unsigned int sharedMem; + + for( unsigned arg=0; arg < n_args; arg ++ ) { + const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# + const symbol *formal_param = target_func->get_arg(arg); //cudaGetParameterBufferV2_param_# + unsigned size=formal_param->get_size_in_bytes(); + assert( formal_param->is_param_local() ); + assert( actual_param_op.is_param_local() ); + addr_t from_addr = actual_param_op.get_symbol()->get_address(); + + if(arg == 0) {//function_info* for the child kernel + unsigned long long buf; + assert(size == sizeof(function_info *)); + thread->m_local_mem->read(from_addr, size, &buf); + child_kernel_entry = (function_info *)buf; + assert(child_kernel_entry); + DEV_RUNTIME_REPORT("child kernel name " << child_kernel_entry->get_name()); + } + else if(arg == 1) { //dim3 gridDim for the child kernel + assert(size == sizeof(struct dim3)); + thread->m_local_mem->read(from_addr, size, & gridDim); + DEV_RUNTIME_REPORT("grid (" << gridDim.x << ", " << gridDim.y << ", " << gridDim.z << ")"); + } + else if(arg == 2) { //dim3 blockDim for the child kernel + assert(size == sizeof(struct dim3)); + thread->m_local_mem->read(from_addr, size, & blockDim); + DEV_RUNTIME_REPORT("block (" << blockDim.x << ", " << blockDim.y << ", " << blockDim.z << ")"); + } + else if(arg == 3) { //unsigned int sharedMem + assert(size == sizeof(unsigned int)); + thread->m_local_mem->read(from_addr, size, & sharedMem); + DEV_RUNTIME_REPORT("shared memory " << sharedMem); + } + } + + //get total child kernel argument size and malloc buffer in global memory + unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); + void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size); + DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer); + + //create child kernel_info_t and index it with parameter_buffer address + kernel_info_t * child_grid = new kernel_info_t(gridDim, blockDim, child_kernel_entry); + assert(g_cuda_device_launch_map.find(param_buffer) == g_cuda_device_launch_map.end()); + g_cuda_device_launch_map[param_buffer] = child_grid; + + //copy the buffer address to retval0 + const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 + const symbol *formal_return = target_func->get_return_var(); //void * + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaGetParameterBufferV2 return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(void *)); + addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); + thread->m_local_mem->write(ret_param_addr, return_size, ¶m_buffer, NULL, NULL); + +} + +//Handling device runtime api: +//cudaError_t cudaLaunchDeviceV2(void *parameterBuffer, cudaStream_t stream) +void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func) { + DEV_RUNTIME_REPORT("Calling cudaLaunchDeviceV2"); + + unsigned n_return = target_func->has_return(); + assert(n_return); + unsigned n_args = target_func->num_args(); + assert( n_args == 2 ); + + kernel_info_t * child_grid = NULL; + function_info * child_kernel_entry = NULL; + void * parameter_buffer; + struct CUstream_st * child_stream; + for( unsigned arg=0; arg < n_args; arg ++ ) { + const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# + const symbol *formal_param = target_func->get_arg(arg); //cudaLaunchDeviceV2_param_# + unsigned size=formal_param->get_size_in_bytes(); + assert( formal_param->is_param_local() ); + assert( actual_param_op.is_param_local() ); + addr_t from_addr = actual_param_op.get_symbol()->get_address(); + + if(arg == 0) {//paramter buffer for child kernel (in global memory) + //get parameter_buffer from the cudaDeviceLaunchV2_param0 + assert(size == sizeof(void *)); + thread->m_local_mem->read(from_addr, size, ¶meter_buffer); + assert((size_t)parameter_buffer >= GLOBAL_HEAP_START); + DEV_RUNTIME_REPORT("Parameter buffer locating at global memory " << parameter_buffer); + + //get child grid info through parameter_buffer address + assert(g_cuda_device_launch_map.find(parameter_buffer) != g_cuda_device_launch_map.end()); + child_grid = g_cuda_device_launch_map[parameter_buffer]; + child_kernel_entry = child_grid->entry(); + DEV_RUNTIME_REPORT("find child kernel " << child_kernel_entry->get_name()); + + //copy data in parameter_buffer to child kernel param memory + unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); + DEV_RUNTIME_REPORT("child_kernel_arg_size " << child_kernel_arg_size); + memory_space *child_kernel_param_mem = child_grid->get_param_memory(); + size_t param_start_address = 0; + for(unsigned n = 0; n < child_kernel_arg_size; n++) { + unsigned char one_byte; + thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 1, &one_byte); + child_kernel_param_mem->write(param_start_address + n, 1, &one_byte, NULL, NULL); + } + } + else if(arg == 1) { //cudaStream for the child kernel + assert(size == sizeof(cudaStream_t)); + thread->m_local_mem->read(from_addr, size, &child_stream); + if(child_stream == 0) { //default stream on device + if(!g_device_default_stream) { + //g_device_default_stream = new struct CUstream_st(); + //g_stream_manager->add_stream(g_device_default_stream); + } + child_stream = g_device_default_stream; + } +// DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid()); + } + + } + + //launch child kernel + stream_operation op(child_grid, g_ptx_sim_mode, child_stream); +// g_stream_manager->push(op); +// g_cuda_device_launch_map.erase(parameter_buffer); + + //set retval0 + const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 + const symbol *formal_return = target_func->get_return_var(); //cudaError_t + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaLaunchDeviceV2 return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size + && return_size == sizeof(cudaError_t)); + cudaError_t error = cudaSuccess; + addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); + thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); + +} diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h new file mode 100644 index 0000000..1b10407 --- /dev/null +++ b/src/cuda-sim/cuda_device_runtime.h @@ -0,0 +1,7 @@ +//Jin: cuda_device_runtime.h +//Defines CUDA device runtime APIs for CDP support + +#pragma once + +void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); +void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 254427b..1c47ad3 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -41,6 +41,9 @@ #include "../gpgpu-sim/gpu-sim.h" #include "../gpgpu-sim/shader.h" +//Jin: include device runtime for CDP +#include "cuda_device_runtime.h" + #include unsigned ptx_instruction::g_num_ptx_inst_uid=0; @@ -149,8 +152,7 @@ ptx_reg_t ptx_thread_info::get_operand_value( const operand_info &op, operand_in } else if ( op.is_local() ) { result.u64 = op.get_symbol()->get_address(); } else if ( op.is_function_address() ) { - result.u64 = op.get_symbol()->get_pc()->get_start_PC(); - printf("Get pc for kernel function %u\n", op.get_symbol()->get_pc()->get_start_PC()); + result.u64 = (size_t)op.get_symbol()->get_pc(); } else { const char *name = op.name().c_str(); printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown operand type for %s\n", name ); @@ -1411,12 +1413,14 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) gpgpusim_cuda_vprintf(pI, thread, target_func); return; } + + //Jin: handle device runtime apis for CDP else if(fname == "cudaGetParameterBufferV2") { - printf("calling cudaGetParameterBufferV2\n"); + gpgpusim_cuda_getParameterBufferV2(pI, thread, target_func); return; } else if(fname == "cudaLaunchDeviceV2") { - printf("calling cudaLaunchDeviceV2\n"); + gpgpusim_cuda_launchDeviceV2(pI, thread, target_func); return; } diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 7325e5f..a7ca27e 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -699,7 +699,7 @@ public: } bool is_immediate_address() const { - return m_immediate_address; + return m_immediate_address; } bool is_literal() const { return m_type == int_t || @@ -1209,6 +1209,8 @@ public: { return m_args.size(); } + unsigned get_args_aligned_size(); + const symbol* get_arg( unsigned n ) const { assert( n < m_args.size() ); diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 39257da..baa3bcd 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -120,6 +120,20 @@ symbol_table *init_parser( const char *ptx_filename ) #define DEF(X,Y) g_ptx_token_decode[X] = Y; #include "ptx_parser_decode.def" #undef DEF + g_ptx_token_decode[undefined_space] = "undefined_space"; + g_ptx_token_decode[undefined_space] = "undefined_space=0"; + g_ptx_token_decode[reg_space] = "reg_space"; + g_ptx_token_decode[local_space] = "local_space"; + g_ptx_token_decode[shared_space] = "shared_space"; + g_ptx_token_decode[param_space_unclassified] = "param_space_unclassified"; + g_ptx_token_decode[param_space_kernel] = "param_space_kernel"; + g_ptx_token_decode[param_space_local] = "param_space_local"; + g_ptx_token_decode[const_space] = "const_space"; + g_ptx_token_decode[tex_space] = "tex_space"; + g_ptx_token_decode[surf_space] = "surf_space"; + g_ptx_token_decode[global_space] = "global_space"; + g_ptx_token_decode[generic_space] = "generic_space"; + g_ptx_token_decode[instruction_space] = "instruction_space"; return g_global_symbol_table; } -- cgit v1.3 From 582106ec595bb7745db092f2f4aa2b0fb9521b16 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sat, 18 Oct 2014 23:07:53 -0400 Subject: MOD: add child kernel stream and scheduling support --- src/abstract_hardware_model.cc | 77 +++++++++++++++++++++++++++++++++++++ src/abstract_hardware_model.h | 39 +++++++++++++++++++ src/cuda-sim/cuda-sim.cc | 9 ++++- src/cuda-sim/cuda_device_runtime.cc | 31 +++++++++------ src/cuda-sim/ptx_sim.h | 3 ++ src/gpgpu-sim/gpu-sim.cc | 4 ++ src/gpgpu-sim/gpu-sim.h | 20 ++++++++++ src/gpgpusim_entrypoint.cc | 9 +++++ src/stream_manager.cc | 39 ++++++++++--------- 9 files changed, 199 insertions(+), 32 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 84d165c..b19b9d1 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -565,6 +565,10 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_num_cores_running=0; m_uid = m_next_uid++; m_param_mem = new memory_space_impl<8192>("param",64*1024); + + //Jin: parent and child kernel management for CDP + m_parent_kernel = NULL; + } kernel_info_t::~kernel_info_t() @@ -578,6 +582,79 @@ std::string kernel_info_t::name() const return m_kernel_entry->get_name(); } +//Jin: parent and child kernel management for CDP +void kernel_info_t::set_parent(kernel_info_t * parent, + dim3 parent_ctaid, dim3 parent_tid) { + m_parent_kernel = parent; + m_parent_ctaid = parent_ctaid; + m_parent_tid = parent_tid; + parent->set_child(this); +} + +void kernel_info_t::set_child(kernel_info_t * child) { + m_child_kernels.push_back(child); +} + +bool kernel_info_t::is_finished() { + if(done() && children_all_finished()) + return true; + else + return false; +} + +bool kernel_info_t::children_all_finished() { + for(auto child = m_child_kernels.begin(); child != m_child_kernels.end(); child++) + { + if(!(*child)->is_finished()) + return false; + } + return true; +} + +void kernel_info_t::notify_parent_finished() { + if(m_parent_kernel) { + g_stream_manager->register_finished_kernel(m_parent_kernel->get_uid()); + } +} + +CUstream_st * kernel_info_t::create_stream_cta(dim3 ctaid) { + assert(get_default_stream_cta(ctaid)); + CUstream_st * stream = new CUstream_st(); + g_stream_manager->add_stream(stream); + assert(m_cta_streams.find(ctaid) != m_cta_streams.end()); + assert(m_cta_streams[ctaid].size() >= 1); //must have default stream + m_cta_streams[ctaid].push_back(stream); + + return stream; +} + +CUstream_st * kernel_info_t::get_default_stream_cta(dim3 ctaid) { + if(m_cta_streams.find(ctaid) != m_cta_streams.end()) { + assert(m_cta_streams[ctaid].size() >= 1); //already created, must have default stream + return *(m_cta_streams[ctaid].begin()); + } + else { + m_cta_streams[ctaid] = std::list(); + CUstream_st * stream = new CUstream_st(); + g_stream_manager->add_stream(stream); + m_cta_streams[ctaid].push_back(stream); + return stream; + } +} + +bool kernel_info_t::cta_has_stream(dim3 ctaid, CUstream_st* stream) { + if(m_cta_streams.find(ctaid) == m_cta_streams.end()) + return false; + + std::list &stream_list = m_cta_streams[ctaid]; + if(std::find(stream_list.begin(), stream_list.end(), stream) + == stream_list.end()) + return false; + else + return true; +} + + simt_stack::simt_stack( unsigned wid, unsigned warpSize) { m_warp_id=wid; diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index ba4ea29..16f4b31 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -154,15 +154,35 @@ enum _memory_op_t { #include #include #include +#include #if !defined(__VECTOR_TYPES_H__) struct dim3 { unsigned int x, y, z; }; #endif +struct dim3comp { + bool operator() (const dim3 & a, const dim3 & b) const + { + if(a.z < b.z) + return true; + else if(a.y < b.y) + return true; + else if (a.x < b.x) + return true; + else + return false; + } +}; void increment_x_then_y_then_z( dim3 &i, const dim3 &bound); +//Jin: child kernel information for CDP +#include "stream_manager.h" +class stream_manager; +struct CUstream_st; +extern stream_manager * g_stream_manager; + class kernel_info_t { public: // kernel_info_t() @@ -250,6 +270,25 @@ private: std::list m_active_threads; class memory_space *m_param_mem; + +public: + //Jin: parent and child kernel management for CDP + void set_parent(kernel_info_t * parent, dim3 parent_ctaid, dim3 parent_tid); + void set_child(kernel_info_t * child); + bool is_finished(); + bool children_all_finished(); + void notify_parent_finished(); + CUstream_st * create_stream_cta(dim3 ctaid); + CUstream_st * get_default_stream_cta(dim3 ctaid); + bool cta_has_stream(dim3 ctaid, CUstream_st* stream); + +private: + kernel_info_t * m_parent_kernel; + dim3 m_parent_ctaid; + dim3 m_parent_tid; + std::list m_child_kernels; //child kernel launched + std::map< dim3, std::list, dim3comp > m_cta_streams; //streams created in each CTA + }; struct core_config { diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 4933029..482e3f6 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1762,8 +1762,13 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) extern stream_manager *g_stream_manager; //openCL kernel simulation calls don't register the kernel so we don't register its exit - if(!openCL) - g_stream_manager->register_finished_kernel(kernel.get_uid()); + if(!openCL) { + + //Jin: for CDP, children should be finished first + if(kernel.is_finished()) { + g_stream_manager->register_finished_kernel(kernel.get_uid()); + } + } //******PRINTING******* printf( "GPGPU-Sim: Done functional simulation (%u instructions simulated).\n", g_ptx_sim_num_insn ); diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 937eec8..4d8d1e1 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -21,8 +21,7 @@ } std::map g_cuda_device_launch_map; -struct CUstream_st * g_device_default_stream = NULL; -extern stream_manager *g_stream_manager; +extern stream_manager * g_stream_manager; //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -78,7 +77,13 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer); //create child kernel_info_t and index it with parameter_buffer address - kernel_info_t * child_grid = new kernel_info_t(gridDim, blockDim, child_kernel_entry); + kernel_info_t * child_grid = new kernel_info_t(gridDim, blockDim, child_kernel_entry); + kernel_info_t & parent_grid = thread->get_kernel(); + DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << + thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z << + "), thread (" << thread->get_tid().x << ", " << thread->get_tid().y << ", " << thread->get_tid().z << + ")"); + child_grid->set_parent(&parent_grid, thread->get_ctaid(), thread->get_tid()); assert(g_cuda_device_launch_map.find(param_buffer) == g_cuda_device_launch_map.end()); g_cuda_device_launch_map[param_buffer] = child_grid; @@ -137,28 +142,30 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * for(unsigned n = 0; n < child_kernel_arg_size; n++) { unsigned char one_byte; thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 1, &one_byte); + std::cout << "one byte " << std::hex << one_byte << "\n"; child_kernel_param_mem->write(param_start_address + n, 1, &one_byte, NULL, NULL); } } else if(arg == 1) { //cudaStream for the child kernel assert(size == sizeof(cudaStream_t)); thread->m_local_mem->read(from_addr, size, &child_stream); - if(child_stream == 0) { //default stream on device - if(!g_device_default_stream) { - //g_device_default_stream = new struct CUstream_st(); - //g_stream_manager->add_stream(g_device_default_stream); - } - child_stream = g_device_default_stream; + + kernel_info_t & parent_kernel = thread->get_kernel(); + if(child_stream == 0) { //default stream on device for current CTA + child_stream = parent_kernel.get_default_stream_cta(thread->get_ctaid()); + } + else { + assert(parent_kernel.cta_has_stream(thread->get_ctaid(), child_stream)); } -// DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid()); + DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid() << " " << child_stream); } } //launch child kernel stream_operation op(child_grid, g_ptx_sim_mode, child_stream); -// g_stream_manager->push(op); -// g_cuda_device_launch_map.erase(parameter_buffer); + g_stream_manager->push(op); + g_cuda_device_launch_map.erase(parameter_buffer); //set retval0 const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h index f926e6d..ea171c5 100644 --- a/src/cuda-sim/ptx_sim.h +++ b/src/cuda-sim/ptx_sim.h @@ -418,6 +418,9 @@ public: void or_reduction(unsigned ctaid, unsigned barid, bool value) {m_core->or_reduction(ctaid,barid,value);} void popc_reduction(unsigned ctaid, unsigned barid, bool value) {m_core->popc_reduction(ctaid,barid,value);} + //Jin: get corresponding kernel grid for CDP purpose + kernel_info_t & get_kernel() { return m_kernel; } + public: addr_t m_last_effective_address; bool m_branch_taken; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index eafb909..47dbc89 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -622,6 +622,10 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config ) *active_sms=0; last_liveness_message_time = 0; + + //Jin: functional simulation for CDP + m_functional_sim = false; + m_functional_sim_kernel = NULL; } int gpgpu_sim::shared_mem_size() const diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 33fffd3..a2d1b9b 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -492,6 +492,7 @@ private: std::string executed_kernel_info_string(); //< format the kernel information into a string for stat printout void clear_executed_kernel_info(); //< clear the kernel information after stat printout + public: unsigned long long gpu_sim_insn; unsigned long long gpu_tot_sim_insn; @@ -504,6 +505,25 @@ public: void change_cache_config(FuncCache cache_config); void set_cache_config(std::string kernel_name); + //Jin: functional simulation for CDP +private: + //set by stream operation every time a functoinal simulation is done + bool m_functional_sim; + kernel_info_t * m_functional_sim_kernel; + +public: + bool is_functional_sim() { return m_functional_sim; } + kernel_info_t * get_functional_kernel() { return m_functional_sim_kernel; } + void functional_launch(kernel_info_t * k) { + m_functional_sim = true; + m_functional_sim_kernel = k; + } + void finish_functional_sim(kernel_info_t * k) { + assert(m_functional_sim); + assert(m_functional_sim_kernel == k); + m_functional_sim = false; + m_functional_sim_kernel = NULL; + } }; #endif diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 6ba38eb..31f3a41 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -127,6 +127,15 @@ void *gpgpu_sim_thread_concurrent(void*) if(g_stream_manager->operation(&sim_cycles) && !g_the_gpu->active()) break; + //functional simulation + if( g_the_gpu->is_functional_sim()) { + kernel_info_t * kernel = g_the_gpu->get_functional_kernel(); + assert(kernel); + gpgpu_cuda_ptx_sim_main_func(*kernel); + g_the_gpu->finish_functional_sim(kernel); + } + + //performance simulation if( g_the_gpu->active() ) { g_the_gpu->cycle(); sim_cycles = true; diff --git a/src/stream_manager.cc b/src/stream_manager.cc index dd42f0a..1f93c03 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -155,7 +155,7 @@ void stream_operation::do_operation( gpgpu_sim *gpu ) gpu->set_cache_config(m_kernel->name()); printf("kernel \'%s\' transfer to GPU hardware scheduler\n", m_kernel->name().c_str() ); if( m_sim_mode ) - gpgpu_cuda_ptx_sim_main_func( *m_kernel ); + gpu->functional_launch( m_kernel ); else gpu->launch( m_kernel ); } @@ -212,11 +212,9 @@ bool stream_manager::operation( bool * sim) bool stream_manager::check_finished_kernel() { - - unsigned grid_uid = m_gpu->finished_kernel(); - bool check=register_finished_kernel(grid_uid); - return check; - + unsigned grid_uid = m_gpu->finished_kernel(); + bool check=register_finished_kernel(grid_uid); + return check; } bool stream_manager::register_finished_kernel(unsigned grid_uid) @@ -226,13 +224,17 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) CUstream_st *stream = m_grid_id_to_stream[grid_uid]; kernel_info_t *kernel = stream->front().get_kernel(); assert( grid_uid == kernel->get_uid() ); - stream->record_next_done(); - m_grid_id_to_stream.erase(grid_uid); - delete kernel; - return true; - }else{ - return false; + + //Jin: should check children kernels for CDP + if(kernel->is_finished()) { + stream->record_next_done(); + m_grid_id_to_stream.erase(grid_uid); + kernel->notify_parent_finished(); + delete kernel; + return true; + } } + return false; } @@ -259,21 +261,22 @@ stream_operation stream_manager::front() { // called by gpu simulation thread stream_operation result; - if( concurrent_streams_empty() ) - m_service_stream_zero = true; +// if( concurrent_streams_empty() ) + m_service_stream_zero = true; if( m_service_stream_zero ) { - if( !m_stream_zero.empty() ) { - if( !m_stream_zero.busy() ) { + if( !m_stream_zero.empty() && !m_stream_zero.busy() ) { result = m_stream_zero.next(); if( result.is_kernel() ) { unsigned grid_id = result.get_kernel()->get_uid(); m_grid_id_to_stream[grid_id] = &m_stream_zero; } - } } else { m_service_stream_zero = false; } - } else { + } + + if(!m_service_stream_zero) + { std::list::iterator s; for( s=m_streams.begin(); s != m_streams.end(); s++) { CUstream_st *stream = *s; -- cgit v1.3 From d31c2da2b34d3d7a63d9980335c85d5e1a19ad02 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sun, 19 Oct 2014 01:02:20 -0400 Subject: BUG: parameter alignment --- src/cuda-sim/cuda-sim.cc | 26 +++++++++++++++++++++----- src/cuda-sim/cuda_device_runtime.cc | 1 - 2 files changed, 21 insertions(+), 6 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 482e3f6..4281913 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1066,12 +1066,24 @@ void function_info::add_param_data( unsigned argn, struct gpgpu_ptx_sim_arg *arg } unsigned function_info::get_args_aligned_size() { - unsigned int align_size = 4; // a word + + unsigned param_address = 0; unsigned int total_size = 0; - for(unsigned int i = 0; i < num_args(); i++) { - total_size += ((m_args[i]->get_size_in_bytes() + align_size - 1) / align_size) * align_size; + for( std::map::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { + param_info &p = i->second; + std::string name = p.get_name(); + symbol *param = m_symtab->lookup(name.c_str()); + + size_t arg_size = p.get_size(); // size of param in bytes + total_size = (total_size + arg_size - 1) / arg_size * arg_size; //aligned + p.add_offset(total_size); + param_address += total_size; + param->set_address(param_address); + total_size += arg_size; } - return total_size; + + return (total_size + 3) / 4; //final size aligned to word + } @@ -1097,13 +1109,17 @@ void function_info::finalize( memory_space *param_mem ) size = (size<(p.get_size()/8))?size:(p.get_size()/8); } // copy the parameter over word-by-word so that parameter that crosses a memory page can be copied over + //Jin: copy parameter using aligned rules const size_t word_size = 4; + param_address = (param_address + size - 1) / size * size; //aligned with size for (size_t idx = 0; idx < size; idx += word_size) { const char *pdata = reinterpret_cast(param_value.pdata) + idx; // cast to char * for ptr arithmetic param_mem->write(param_address + idx, word_size, pdata,NULL,NULL); } + unsigned offset = p.get_offset(); + assert(offset == param_address); param->set_address(param_address); - param_address += size; + param_address += size; } } diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 4d8d1e1..148471b 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -142,7 +142,6 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * for(unsigned n = 0; n < child_kernel_arg_size; n++) { unsigned char one_byte; thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 1, &one_byte); - std::cout << "one byte " << std::hex << one_byte << "\n"; child_kernel_param_mem->write(param_start_address + n, 1, &one_byte, NULL, NULL); } } -- cgit v1.3 From 5fc7bf5872aad126a09cad4b385054c4b3a094aa Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Mon, 20 Oct 2014 17:19:32 -0400 Subject: BUG: do not handle cudaGetParameterBufferV2 and cudaLaunchDeviceV2 as call.uni in reconvergence --- src/abstract_hardware_model.h | 2 ++ src/cuda-sim/cuda-sim.cc | 6 ++++-- src/cuda-sim/ptx_ir.cc | 3 +++ 3 files changed, 9 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 16f4b31..11fee10 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -864,6 +864,7 @@ public: m_mem_accesses_created=false; m_cache_hit=false; m_is_printf=false; + m_is_cdp = false; } virtual ~warp_inst_t(){ } @@ -1016,6 +1017,7 @@ protected: unsigned cycles; // used for implementing initiation interval delay bool m_isatomic; bool m_is_printf; + bool m_is_cdp; unsigned m_warp_id; unsigned m_dynamic_warp_id; const core_config *m_config; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 4281913..58bd4e0 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -638,7 +638,7 @@ void ptx_instruction::set_opcode_and_latency() case MEMBAR_OP: op = MEMORY_BARRIER_OP; break; case CALL_OP: { - if(m_is_printf) + if(m_is_printf || m_is_cdp) op = ALU_OP; else op = CALL_OPS; @@ -646,7 +646,7 @@ void ptx_instruction::set_opcode_and_latency() } case CALLP_OP: { - if(m_is_printf) + if(m_is_printf || m_is_cdp) op = ALU_OP; else op = CALL_OPS; @@ -1219,6 +1219,8 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) bool skip = false; int op_classification = 0; addr_t pc = next_instr(); + if(pc == 440) + pc = 440; assert( pc == inst.pc ); // make sure timing model and functional model are in sync const ptx_instruction *pI = m_func_info->get_instruction(pc); set_npc( pc + pI->inst_size() ); diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 915c623..861f0dc 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1241,6 +1241,9 @@ ptx_instruction::ptx_instruction( int opcode, if (fname =="vprintf"){ m_is_printf = true; } + if (fname == "cudaGetParameterBufferV2" + || fname == "cudaLaunchDeviceV2") + m_is_cdp = true; } } -- cgit v1.3 From 6286547cfdc5d14c84568505db267f5e8dd9841f Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Mon, 20 Oct 2014 23:23:00 -0400 Subject: BUG: multiple child kernels finish --- src/abstract_hardware_model.cc | 22 ++++++++++++++++++---- src/abstract_hardware_model.h | 2 ++ src/cuda-sim/cuda-sim.cc | 6 +----- src/stream_manager.cc | 3 ++- 4 files changed, 23 insertions(+), 10 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index b19b9d1..f3c5b21 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -595,6 +595,12 @@ void kernel_info_t::set_child(kernel_info_t * child) { m_child_kernels.push_back(child); } +void kernel_info_t::remove_child(kernel_info_t * child) { + assert(std::find(m_child_kernels.begin(), m_child_kernels.end(), child) + != m_child_kernels.end()); + m_child_kernels.remove(child); +} + bool kernel_info_t::is_finished() { if(done() && children_all_finished()) return true; @@ -603,16 +609,15 @@ bool kernel_info_t::is_finished() { } bool kernel_info_t::children_all_finished() { - for(auto child = m_child_kernels.begin(); child != m_child_kernels.end(); child++) - { - if(!(*child)->is_finished()) + if(!m_child_kernels.empty()) return false; - } + return true; } void kernel_info_t::notify_parent_finished() { if(m_parent_kernel) { + m_parent_kernel->remove_child(this); g_stream_manager->register_finished_kernel(m_parent_kernel->get_uid()); } } @@ -654,6 +659,15 @@ bool kernel_info_t::cta_has_stream(dim3 ctaid, CUstream_st* stream) { return true; } +void kernel_info_t::print_parent_info() { + if(m_parent_kernel) { + printf("Parent %d: \'%s\', Block (%d, %d, %d), Thread (%d, %d, %d)\n", + m_parent_kernel->get_uid(), m_parent_kernel->name().c_str(), + m_parent_ctaid.x, m_parent_ctaid.y, m_parent_ctaid.z, + m_parent_tid.x, m_parent_tid.y, m_parent_tid.z); + } +} + simt_stack::simt_stack( unsigned wid, unsigned warpSize) { diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 11fee10..a0bf80d 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -275,12 +275,14 @@ public: //Jin: parent and child kernel management for CDP void set_parent(kernel_info_t * parent, dim3 parent_ctaid, dim3 parent_tid); void set_child(kernel_info_t * child); + void remove_child(kernel_info_t * child); bool is_finished(); bool children_all_finished(); void notify_parent_finished(); CUstream_st * create_stream_cta(dim3 ctaid); CUstream_st * get_default_stream_cta(dim3 ctaid); bool cta_has_stream(dim3 ctaid, CUstream_st* stream); + void print_parent_info(); private: kernel_info_t * m_parent_kernel; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 58bd4e0..8dd1078 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1781,11 +1781,7 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) //openCL kernel simulation calls don't register the kernel so we don't register its exit if(!openCL) { - - //Jin: for CDP, children should be finished first - if(kernel.is_finished()) { - g_stream_manager->register_finished_kernel(kernel.get_uid()); - } + g_stream_manager->register_finished_kernel(kernel.get_uid()); } //******PRINTING******* diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 1f93c03..687d544 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -153,7 +153,8 @@ void stream_operation::do_operation( gpgpu_sim *gpu ) case stream_kernel_launch: if( gpu->can_start_kernel() ) { gpu->set_cache_config(m_kernel->name()); - printf("kernel \'%s\' transfer to GPU hardware scheduler\n", m_kernel->name().c_str() ); + printf("kernel %d: \'%s\' transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); + m_kernel->print_parent_info(); if( m_sim_mode ) gpu->functional_launch( m_kernel ); else -- cgit v1.3 From 3bb32d87175d873e7089ad50a0069acc195edb34 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Tue, 21 Oct 2014 01:24:26 -0400 Subject: ADD: add support for cudaStreamCreateWithFlags --- src/cuda-sim/cuda_device_runtime.cc | 61 ++++++++++++++++++++++++++++++++++++- src/cuda-sim/cuda_device_runtime.h | 1 + src/cuda-sim/instructions.cc | 4 +++ src/cuda-sim/ptx_ir.cc | 3 +- 4 files changed, 67 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 148471b..be87f6a 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -152,11 +152,12 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * kernel_info_t & parent_kernel = thread->get_kernel(); if(child_stream == 0) { //default stream on device for current CTA child_stream = parent_kernel.get_default_stream_cta(thread->get_ctaid()); + DEV_RUNTIME_REPORT("launching child kernel to default stream of the cta " << child_stream->get_uid() << ": " << child_stream); } else { assert(parent_kernel.cta_has_stream(thread->get_ctaid(), child_stream)); + DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid() << ": " << child_stream); } - DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid() << " " << child_stream); } } @@ -179,3 +180,61 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); } + + +//Handling device runtime api: +//cudaError_t cudaStreamCreateWithFlags ( cudaStream_t* pStream, unsigned int flags) +//flags can only be cudaStreamNonBlocking +void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func) { + DEV_RUNTIME_REPORT("Calling cudaStreamCreateWithFlags"); + + unsigned n_return = target_func->has_return(); + assert(n_return); + unsigned n_args = target_func->num_args(); + assert( n_args == 2 ); + + size_t generic_pStream_addr; + addr_t pStream_addr; + unsigned int flags; + for( unsigned arg=0; arg < n_args; arg ++ ) { + const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# + const symbol *formal_param = target_func->get_arg(arg); //cudaStreamCreateWithFlags_param_# + unsigned size=formal_param->get_size_in_bytes(); + assert( formal_param->is_param_local() ); + assert( actual_param_op.is_param_local() ); + addr_t from_addr = actual_param_op.get_symbol()->get_address(); + + if(arg == 0) {//cudaStream_t * pStream, address of cudaStream_t + assert(size == sizeof(cudaStream_t *)); + thread->m_local_mem->read(from_addr, size, &generic_pStream_addr); + + //pStream should be non-zero address in local memory + pStream_addr = generic_to_local(thread->get_hw_sid(), thread->get_hw_tid(), generic_pStream_addr); + + DEV_RUNTIME_REPORT("pStream locating at local memory " << pStream_addr); + } + else if(arg == 1) { //unsigned int flags, should be cudaStreamNonBlocking + assert(size == sizeof(unsigned int)); + thread->m_local_mem->read(from_addr, size, &flags); + assert(flags == cudaStreamNonBlocking); + } + } + + //create stream and write back to param0 + CUstream_st * stream = thread->get_kernel().create_stream_cta(thread->get_ctaid()); + DEV_RUNTIME_REPORT("Create stream " << stream->get_uid() << ": " << stream); + thread->m_local_mem->write(pStream_addr, sizeof(cudaStream_t), &stream, NULL, NULL); + + //set retval0 + const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 + const symbol *formal_return = target_func->get_return_var(); //cudaError_t + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaStreamCreateWithFlags return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size + && return_size == sizeof(cudaError_t)); + cudaError_t error = cudaSuccess; + addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); + thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); + +} diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index 1b10407..7447d5a 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -5,3 +5,4 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); +void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 1c47ad3..5d909d3 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1423,6 +1423,10 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) gpgpusim_cuda_launchDeviceV2(pI, thread, target_func); return; } + else if(fname == "cudaStreamCreateWithFlags") { + gpgpusim_cuda_streamCreateWithFlags(pI, thread, target_func); + return; + } // read source arguements into register specified in declaration of function arg_buffer_list_t arg_values; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 861f0dc..bdc8381 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1242,7 +1242,8 @@ ptx_instruction::ptx_instruction( int opcode, m_is_printf = true; } if (fname == "cudaGetParameterBufferV2" - || fname == "cudaLaunchDeviceV2") + || fname == "cudaLaunchDeviceV2" + || fname == "cudaStreamCreateWithFlags") m_is_cdp = true; } -- cgit v1.3 From f372a4c641b9e6d38470ead6ae25743d26c5fed1 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Thu, 23 Oct 2014 15:33:59 -0400 Subject: BUG: kernels should return to stream if not pushed to concurrent kernel pool --- src/cuda-sim/cuda_device_runtime.cc | 8 +++++--- src/gpgpu-sim/shader.cc | 3 ++- src/stream_manager.cc | 37 ++++++++++++++++++++++++++++++++----- src/stream_manager.h | 3 ++- 4 files changed, 41 insertions(+), 10 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index be87f6a..c53ea04 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -122,7 +122,7 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * addr_t from_addr = actual_param_op.get_symbol()->get_address(); if(arg == 0) {//paramter buffer for child kernel (in global memory) - //get parameter_buffer from the cudaDeviceLaunchV2_param0 + //get parameter_buffer from the cudaLaunchDeviceV2_param0 assert(size == sizeof(void *)); thread->m_local_mem->read(from_addr, size, ¶meter_buffer); assert((size_t)parameter_buffer >= GLOBAL_HEAP_START); @@ -152,11 +152,13 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * kernel_info_t & parent_kernel = thread->get_kernel(); if(child_stream == 0) { //default stream on device for current CTA child_stream = parent_kernel.get_default_stream_cta(thread->get_ctaid()); - DEV_RUNTIME_REPORT("launching child kernel to default stream of the cta " << child_stream->get_uid() << ": " << child_stream); + DEV_RUNTIME_REPORT("launching child kernel " << child_grid->get_uid() << + " to default stream of the cta " << child_stream->get_uid() << ": " << child_stream); } else { assert(parent_kernel.cta_has_stream(thread->get_ctaid(), child_stream)); - DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid() << ": " << child_stream); + DEV_RUNTIME_REPORT("launching child kernel " << child_grid->get_uid() << + " to stream " << child_stream->get_uid() << ": " << child_stream); } } diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index ff2fac7..de5bcf6 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1934,7 +1934,8 @@ void shader_core_ctx::register_cta_thread_exit( unsigned cta_num ) m_kernel->name().c_str() ); if( !m_gpu->kernel_more_cta_left(m_kernel) ) { if( !m_kernel->running() ) { - printf("GPGPU-Sim uArch: GPU detected kernel \'%s\' finished on shader %u.\n", m_kernel->name().c_str(), m_sid ); + printf("GPGPU-Sim uArch: GPU detected kernel %u \'%s\' finished on shader %u.\n", m_kernel->get_uid(), + m_kernel->name().c_str(), m_sid ); m_gpu->set_kernel_done( m_kernel ); } } diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 687d544..546ab3b 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -95,6 +95,15 @@ stream_operation CUstream_st::next() return result; } +void CUstream_st::cancel_front() +{ + pthread_mutex_lock(&m_lock); + assert(m_pending); + m_pending = false; + pthread_mutex_unlock(&m_lock); + +} + void CUstream_st::print(FILE *fp) { pthread_mutex_lock(&m_lock); @@ -111,10 +120,10 @@ void CUstream_st::print(FILE *fp) } -void stream_operation::do_operation( gpgpu_sim *gpu ) +bool stream_operation::do_operation( gpgpu_sim *gpu ) { if( is_noop() ) - return; + return true; assert(!m_done && m_stream); if(g_debug_execution >= 3) @@ -153,16 +162,23 @@ void stream_operation::do_operation( gpgpu_sim *gpu ) case stream_kernel_launch: if( gpu->can_start_kernel() ) { gpu->set_cache_config(m_kernel->name()); - printf("kernel %d: \'%s\' transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); + if(g_debug_execution >= 3) + printf("kernel %d: \'%s\' transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); m_kernel->print_parent_info(); if( m_sim_mode ) gpu->functional_launch( m_kernel ); else gpu->launch( m_kernel ); } + else { + if(g_debug_execution >= 3) + printf("kernel %d: \'%s\' not ready to transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); + return false; + } break; case stream_event: { - printf("event update\n"); + if(g_debug_execution >= 3) + printf("event update\n"); time_t wallclock = time((time_t *)NULL); m_event->update( gpu_tot_sim_cycle, wallclock ); m_stream->record_next_done(); @@ -173,6 +189,7 @@ void stream_operation::do_operation( gpgpu_sim *gpu ) } m_done=true; fflush(stdout); + return true; } void stream_operation::print( FILE *fp ) const @@ -204,7 +221,16 @@ bool stream_manager::operation( bool * sim) bool check=check_finished_kernel(); if(check)m_gpu->print_stats(); stream_operation op =front(); - op.do_operation( m_gpu ); + if(!op.do_operation( m_gpu )) //not ready to execute + { + //cancel operation + if( op.is_kernel() ) { + unsigned grid_uid = op.get_kernel()->get_uid(); + m_grid_id_to_stream.erase(grid_uid); + } + op.get_stream()->cancel_front(); + + } pthread_mutex_unlock(&m_lock); //pthread_mutex_lock(&m_lock); // simulate a clock cycle on the GPU @@ -228,6 +254,7 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) //Jin: should check children kernels for CDP if(kernel->is_finished()) { + printf("kernel %d finishes, retires from stream %d\n", grid_uid, stream->get_uid()); stream->record_next_done(); m_grid_id_to_stream.erase(grid_uid); kernel->notify_parent_finished(); diff --git a/src/stream_manager.h b/src/stream_manager.h index 701b33c..222a1b2 100644 --- a/src/stream_manager.h +++ b/src/stream_manager.h @@ -150,7 +150,7 @@ public: bool is_noop() const { return m_type == stream_no_op; } bool is_done() const { return m_done; } kernel_info_t *get_kernel() { return m_kernel; } - void do_operation( gpgpu_sim *gpu ); + bool do_operation( gpgpu_sim *gpu ); void print( FILE *fp ) const; struct CUstream_st *get_stream() { return m_stream; } void set_stream( CUstream_st *stream ) { m_stream = stream; } @@ -218,6 +218,7 @@ public: void push( const stream_operation &op ); void record_next_done(); stream_operation next(); + void cancel_front(); //front operation fails, cancle the pending status stream_operation &front() { return m_operations.front(); } void print( FILE *fp ); unsigned get_uid() const { return m_uid; } -- cgit v1.3 From cf507ddb207337bc7a67ded7c4438f0cb0bed26f Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sun, 2 Nov 2014 19:24:37 -0500 Subject: BUG: PTX section id. ADD: cudaDeviceSetLimit. BUG: parameter addresses for child kernels in CDP. BUG: .weak .entry and .weak .global directives in ptx file. BUG: empty_protected() for stream manager causes deadlock, change to empty() --- libcuda/cuda_runtime_api.cc | 5 +++++ src/cuda-sim/cuda-sim.cc | 9 +++------ src/cuda-sim/cuda_device_runtime.cc | 9 +++++---- src/cuda-sim/ptx.y | 2 ++ src/gpgpusim_entrypoint.cc | 2 +- 5 files changed, 16 insertions(+), 11 deletions(-) (limited to 'src/cuda-sim') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 5310a52..30bf823 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2018,6 +2018,11 @@ __host__ cudaError_t CUDARTAPI cudaFuncSetCacheConfig(const char *func, enum cud context->get_device()->get_gpgpu()->set_cache_config(context->get_kernel(func)->get_name(), (FuncCache)cacheConfig); return g_last_cudaError = cudaSuccess; } + +//Jin: hack for cdp +__host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t value) { + return g_last_cudaError = cudaSuccess; +} #endif #endif diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 8dd1078..980afc8 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1074,15 +1074,14 @@ unsigned function_info::get_args_aligned_size() { std::string name = p.get_name(); symbol *param = m_symtab->lookup(name.c_str()); - size_t arg_size = p.get_size(); // size of param in bytes + size_t arg_size = p.get_size() / 8; // size of param in bytes total_size = (total_size + arg_size - 1) / arg_size * arg_size; //aligned p.add_offset(total_size); - param_address += total_size; - param->set_address(param_address); + param->set_address(param_address + total_size); total_size += arg_size; } - return (total_size + 3) / 4; //final size aligned to word + return (total_size + 3) / 4 * 4; //final size aligned to word } @@ -1219,8 +1218,6 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) bool skip = false; int op_classification = 0; addr_t pc = next_instr(); - if(pc == 440) - pc = 440; assert( pc == inst.pc ); // make sure timing model and functional model are in sync const ptx_instruction *pI = m_func_info->get_instruction(pc); set_npc( pc + pI->inst_size() ); diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index c53ea04..2a90cba 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -139,10 +139,11 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * DEV_RUNTIME_REPORT("child_kernel_arg_size " << child_kernel_arg_size); memory_space *child_kernel_param_mem = child_grid->get_param_memory(); size_t param_start_address = 0; - for(unsigned n = 0; n < child_kernel_arg_size; n++) { - unsigned char one_byte; - thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 1, &one_byte); - child_kernel_param_mem->write(param_start_address + n, 1, &one_byte, NULL, NULL); + //copy in word + for(unsigned n = 0; n < child_kernel_arg_size; n += 4) { + unsigned int oneword; + thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 4, &oneword); + child_kernel_param_mem->write(param_start_address + n, 4, &oneword, NULL, NULL); } } else if(arg == 1) { //cudaStream for the child kernel diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index c8208ea..e29b973 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -241,6 +241,7 @@ function_ident_param: IDENTIFIER { add_function_name($1); } LEFT_PAREN {func_hea function_decl_header: ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".entry"); } | VISIBLE_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".entry"); } + | WEAK_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".entry"); } | FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } | VISIBLE_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } | EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; g_func_decl=1; func_header(".func"); } @@ -323,6 +324,7 @@ var_spec: space_spec | type_spec | align_spec | EXTERN_DIRECTIVE { add_extern_spec(); } + | WEAK_DIRECTIVE ; align_spec: ALIGN_DIRECTIVE INT_OPERAND { add_alignment_spec($2); } diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 31f3a41..fb17eed 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -100,7 +100,7 @@ void *gpgpu_sim_thread_concurrent(void*) printf("GPGPU-Sim: *** simulation thread starting and spinning waiting for work ***\n"); fflush(stdout); } - while( g_stream_manager->empty_protected() && !g_sim_done ) + while( g_stream_manager->empty() && !g_sim_done ) ; if(g_debug_execution >= 3) { printf("GPGPU-Sim: ** START simulation thread (detected work) **\n"); -- cgit v1.3 From 3656376ccd83d5c514389c4c3818f1969bc30e0c Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 1 Jul 2016 05:29:58 -0400 Subject: MOD: schedule one child kernel each cycle --- src/cuda-sim/cuda_device_runtime.cc | 214 ++++++++++++++++++++++++------------ src/cuda-sim/cuda_device_runtime.h | 2 + src/gpgpu-sim/gpu-sim.cc | 4 + src/stream_manager.cc | 2 +- 4 files changed, 148 insertions(+), 74 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 2a90cba..df2653e 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -20,8 +20,45 @@ std::cout.flush(); \ } -std::map g_cuda_device_launch_map; -extern stream_manager * g_stream_manager; +class device_launch_config_t { + +public: + device_launch_config_t() {} + + device_launch_config_t(dim3 _grid_dim, + dim3 _block_dim, + unsigned int _shared_mem, + function_info * _entry): + grid_dim(_grid_dim), + block_dim(_block_dim), + shared_mem(_shared_mem), + entry(_entry) {} + + dim3 grid_dim; + dim3 block_dim; + unsigned int shared_mem; + function_info * entry; + +}; + +class device_launch_operation_t { + +public: + device_launch_operation_t() {} + device_launch_operation_t(kernel_info_t *_grid, + CUstream_st * _stream) : + grid(_grid), stream(_stream) {} + + kernel_info_t * grid; //a new child grid + + CUstream_st * stream; + +}; + + +std::map g_cuda_device_launch_param_map; +std::list g_cuda_device_launch_op; +extern stream_manager *g_stream_manager; //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -35,8 +72,8 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i assert( n_args == 4 ); function_info * child_kernel_entry; - struct dim3 gridDim, blockDim; - unsigned int sharedMem; + struct dim3 grid_dim, block_dim; + unsigned int shared_mem; for( unsigned arg=0; arg < n_args; arg ++ ) { const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# @@ -48,54 +85,48 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i if(arg == 0) {//function_info* for the child kernel unsigned long long buf; - assert(size == sizeof(function_info *)); + assert(size == sizeof(function_info *)); thread->m_local_mem->read(from_addr, size, &buf); child_kernel_entry = (function_info *)buf; assert(child_kernel_entry); DEV_RUNTIME_REPORT("child kernel name " << child_kernel_entry->get_name()); } - else if(arg == 1) { //dim3 gridDim for the child kernel - assert(size == sizeof(struct dim3)); - thread->m_local_mem->read(from_addr, size, & gridDim); - DEV_RUNTIME_REPORT("grid (" << gridDim.x << ", " << gridDim.y << ", " << gridDim.z << ")"); + else if(arg == 1) { //dim3 grid_dim for the child kernel + assert(size == sizeof(struct dim3)); + thread->m_local_mem->read(from_addr, size, & grid_dim); + DEV_RUNTIME_REPORT("grid (" << grid_dim.x << ", " << grid_dim.y << ", " << grid_dim.z << ")"); } - else if(arg == 2) { //dim3 blockDim for the child kernel - assert(size == sizeof(struct dim3)); - thread->m_local_mem->read(from_addr, size, & blockDim); - DEV_RUNTIME_REPORT("block (" << blockDim.x << ", " << blockDim.y << ", " << blockDim.z << ")"); + else if(arg == 2) { //dim3 block_dim for the child kernel + assert(size == sizeof(struct dim3)); + thread->m_local_mem->read(from_addr, size, & block_dim); + DEV_RUNTIME_REPORT("block (" << block_dim.x << ", " << block_dim.y << ", " << block_dim.z << ")"); } - else if(arg == 3) { //unsigned int sharedMem - assert(size == sizeof(unsigned int)); - thread->m_local_mem->read(from_addr, size, & sharedMem); - DEV_RUNTIME_REPORT("shared memory " << sharedMem); + else if(arg == 3) { //unsigned int shared_mem + assert(size == sizeof(unsigned int)); + thread->m_local_mem->read(from_addr, size, & shared_mem); + DEV_RUNTIME_REPORT("shared memory " << shared_mem); } } - //get total child kernel argument size and malloc buffer in global memory - unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); - void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size); - DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer); - - //create child kernel_info_t and index it with parameter_buffer address - kernel_info_t * child_grid = new kernel_info_t(gridDim, blockDim, child_kernel_entry); - kernel_info_t & parent_grid = thread->get_kernel(); - DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << - thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z << - "), thread (" << thread->get_tid().x << ", " << thread->get_tid().y << ", " << thread->get_tid().z << - ")"); - child_grid->set_parent(&parent_grid, thread->get_ctaid(), thread->get_tid()); - assert(g_cuda_device_launch_map.find(param_buffer) == g_cuda_device_launch_map.end()); - g_cuda_device_launch_map[param_buffer] = child_grid; - - //copy the buffer address to retval0 + //get total child kernel argument size and malloc buffer in global memory + unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); + void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size); + DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer); + + //store param buffer address and launch config + device_launch_config_t device_launch_config(grid_dim, block_dim, shared_mem, child_kernel_entry); + assert(g_cuda_device_launch_param_map.find(param_buffer) == g_cuda_device_launch_param_map.end()); + g_cuda_device_launch_param_map[param_buffer] = device_launch_config; + + //copy the buffer address to retval0 const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 const symbol *formal_return = target_func->get_return_var(); //void * - unsigned int return_size = formal_return->get_size_in_bytes(); - DEV_RUNTIME_REPORT("cudaGetParameterBufferV2 return value has size of " << return_size); - assert(actual_return_op.is_param_local()); - assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(void *)); + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaGetParameterBufferV2 return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(void *)); addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); - thread->m_local_mem->write(ret_param_addr, return_size, ¶m_buffer, NULL, NULL); + thread->m_local_mem->write(ret_param_addr, return_size, ¶m_buffer, NULL, NULL); } @@ -109,10 +140,13 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * unsigned n_args = target_func->num_args(); assert( n_args == 2 ); - kernel_info_t * child_grid = NULL; - function_info * child_kernel_entry = NULL; + kernel_info_t * device_grid = NULL; + function_info * device_kernel_entry = NULL; void * parameter_buffer; struct CUstream_st * child_stream; + device_launch_config_t config; + device_launch_operation_t device_launch_op; + for( unsigned arg=0; arg < n_args; arg ++ ) { const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# const symbol *formal_param = target_func->get_arg(arg); //cudaLaunchDeviceV2_param_# @@ -123,64 +157,80 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * if(arg == 0) {//paramter buffer for child kernel (in global memory) //get parameter_buffer from the cudaLaunchDeviceV2_param0 - assert(size == sizeof(void *)); + assert(size == sizeof(void *)); thread->m_local_mem->read(from_addr, size, ¶meter_buffer); assert((size_t)parameter_buffer >= GLOBAL_HEAP_START); DEV_RUNTIME_REPORT("Parameter buffer locating at global memory " << parameter_buffer); //get child grid info through parameter_buffer address - assert(g_cuda_device_launch_map.find(parameter_buffer) != g_cuda_device_launch_map.end()); - child_grid = g_cuda_device_launch_map[parameter_buffer]; - child_kernel_entry = child_grid->entry(); - DEV_RUNTIME_REPORT("find child kernel " << child_kernel_entry->get_name()); - - //copy data in parameter_buffer to child kernel param memory - unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); - DEV_RUNTIME_REPORT("child_kernel_arg_size " << child_kernel_arg_size); - memory_space *child_kernel_param_mem = child_grid->get_param_memory(); + assert(g_cuda_device_launch_param_map.find(parameter_buffer) != g_cuda_device_launch_param_map.end()); + config = g_cuda_device_launch_param_map[parameter_buffer]; + //device_grid = op.grid; + device_kernel_entry = config.entry; + DEV_RUNTIME_REPORT("find device kernel " << device_kernel_entry->get_name()); + + //copy data in parameter_buffer to device kernel param memory + unsigned device_kernel_arg_size = device_kernel_entry->get_args_aligned_size(); + DEV_RUNTIME_REPORT("device_kernel_arg_size " << device_kernel_arg_size); + memory_space *device_kernel_param_mem; + + //create child kernel_info_t and index it with parameter_buffer address + device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry); + kernel_info_t & parent_grid = thread->get_kernel(); + DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << + thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z << + "), thread (" << thread->get_tid().x << ", " << thread->get_tid().y << ", " << thread->get_tid().z << + ")"); + device_grid->set_parent(&parent_grid, thread->get_ctaid(), thread->get_tid()); + device_launch_op = device_launch_operation_t(device_grid, NULL); + device_kernel_param_mem = device_grid->get_param_memory(); //kernel param + size_t param_start_address = 0; //copy in word - for(unsigned n = 0; n < child_kernel_arg_size; n += 4) { + for(unsigned n = 0; n < device_kernel_arg_size; n += 4) { unsigned int oneword; thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 4, &oneword); - child_kernel_param_mem->write(param_start_address + n, 4, &oneword, NULL, NULL); + device_kernel_param_mem->write(param_start_address + n, 4, &oneword, NULL, NULL); } } else if(arg == 1) { //cudaStream for the child kernel - assert(size == sizeof(cudaStream_t)); - thread->m_local_mem->read(from_addr, size, &child_stream); + assert(size == sizeof(cudaStream_t)); + thread->m_local_mem->read(from_addr, size, &child_stream); + kernel_info_t & parent_kernel = thread->get_kernel(); if(child_stream == 0) { //default stream on device for current CTA child_stream = parent_kernel.get_default_stream_cta(thread->get_ctaid()); - DEV_RUNTIME_REPORT("launching child kernel " << child_grid->get_uid() << + DEV_RUNTIME_REPORT("launching child kernel " << device_grid->get_uid() << " to default stream of the cta " << child_stream->get_uid() << ": " << child_stream); } else { assert(parent_kernel.cta_has_stream(thread->get_ctaid(), child_stream)); - DEV_RUNTIME_REPORT("launching child kernel " << child_grid->get_uid() << + DEV_RUNTIME_REPORT("launching child kernel " << device_grid->get_uid() << " to stream " << child_stream->get_uid() << ": " << child_stream); } + + device_launch_op.stream = child_stream; + } } - } + //launch child kernel - stream_operation op(child_grid, g_ptx_sim_mode, child_stream); - g_stream_manager->push(op); - g_cuda_device_launch_map.erase(parameter_buffer); + g_cuda_device_launch_op.push_back(device_launch_op); + g_cuda_device_launch_param_map.erase(parameter_buffer); //set retval0 const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 const symbol *formal_return = target_func->get_return_var(); //cudaError_t - unsigned int return_size = formal_return->get_size_in_bytes(); - DEV_RUNTIME_REPORT("cudaLaunchDeviceV2 return value has size of " << return_size); - assert(actual_return_op.is_param_local()); - assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaLaunchDeviceV2 return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(cudaError_t)); cudaError_t error = cudaSuccess; addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); - thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); + thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); } @@ -208,7 +258,7 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_ addr_t from_addr = actual_param_op.get_symbol()->get_address(); if(arg == 0) {//cudaStream_t * pStream, address of cudaStream_t - assert(size == sizeof(cudaStream_t *)); + assert(size == sizeof(cudaStream_t *)); thread->m_local_mem->read(from_addr, size, &generic_pStream_addr); //pStream should be non-zero address in local memory @@ -217,7 +267,7 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_ DEV_RUNTIME_REPORT("pStream locating at local memory " << pStream_addr); } else if(arg == 1) { //unsigned int flags, should be cudaStreamNonBlocking - assert(size == sizeof(unsigned int)); + assert(size == sizeof(unsigned int)); thread->m_local_mem->read(from_addr, size, &flags); assert(flags == cudaStreamNonBlocking); } @@ -231,13 +281,31 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_ //set retval0 const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 const symbol *formal_return = target_func->get_return_var(); //cudaError_t - unsigned int return_size = formal_return->get_size_in_bytes(); - DEV_RUNTIME_REPORT("cudaStreamCreateWithFlags return value has size of " << return_size); - assert(actual_return_op.is_param_local()); - assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaStreamCreateWithFlags return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(cudaError_t)); cudaError_t error = cudaSuccess; addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); - thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); + thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); } + + +void launch_one_device_kernel() { + if(!g_cuda_device_launch_op.empty()) { + device_launch_operation_t &op = g_cuda_device_launch_op.front(); + + stream_operation stream_op = stream_operation(op.grid, g_ptx_sim_mode, op.stream); + g_stream_manager->push(stream_op); + g_cuda_device_launch_op.pop_front(); + } +} + +void launch_all_device_kernels() { + while(!g_cuda_device_launch_op.empty()) { + launch_one_device_kernel(); + } +} + diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index 7447d5a..385d605 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -6,3 +6,5 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); +void launch_all_device_kernels(); +void launch_one_device_kernel(); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 47dbc89..126e007 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -61,6 +61,7 @@ #include "power_stat.h" #include "visualizer.h" #include "stats.h" +#include "../cuda-sim/cuda_device_runtime.h" #ifdef GPGPUSIM_POWER_MODEL #include "power_interface.h" @@ -1370,6 +1371,9 @@ void gpgpu_sim::cycle() } try_snap_shot(gpu_sim_cycle); spill_log_to_file (stdout, 0, gpu_sim_cycle); + + //launch device kernel + launch_one_device_kernel(); } } diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 546ab3b..f90d9be 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -217,8 +217,8 @@ stream_manager::stream_manager( gpgpu_sim *gpu, bool cuda_launch_blocking ) bool stream_manager::operation( bool * sim) { - pthread_mutex_lock(&m_lock); bool check=check_finished_kernel(); + pthread_mutex_lock(&m_lock); if(check)m_gpu->print_stats(); stream_operation op =front(); if(!op.do_operation( m_gpu )) //not ready to execute -- cgit v1.3 From dafeb411265dbc0228889fe97d85b00f71363f10 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 1 Jul 2016 06:54:46 -0400 Subject: MOD: compute child parameter size --- src/cuda-sim/cuda-sim.cc | 9 +++++++-- src/cuda-sim/ptx_ir.cc | 1 + src/cuda-sim/ptx_ir.h | 3 +++ 3 files changed, 11 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 980afc8..f5d8a88 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1066,7 +1066,10 @@ void function_info::add_param_data( unsigned argn, struct gpgpu_ptx_sim_arg *arg } unsigned function_info::get_args_aligned_size() { - + + if(m_args_aligned_size >= 0) + return m_args_aligned_size; + unsigned param_address = 0; unsigned int total_size = 0; for( std::map::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { @@ -1081,7 +1084,9 @@ unsigned function_info::get_args_aligned_size() { total_size += arg_size; } - return (total_size + 3) / 4 * 4; //final size aligned to word + m_args_aligned_size = (total_size + 3) / 4 * 4; //final size aligned to word + + return m_args_aligned_size; } diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index bdc8381..176eb14 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1292,6 +1292,7 @@ function_info::function_info(int entry_point ) m_kernel_info.regs = 0; m_kernel_info.smem = 0; m_local_mem_framesize = 0; + m_args_aligned_size = -1; } unsigned function_info::print_insn( unsigned pc, FILE * fp ) const diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index a7ca27e..a54ae41 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -291,6 +291,9 @@ private: std::list m_initializer; static unsigned sm_next_uid; + + //parameter size for device kernels + int m_args_aligned_size; }; class symbol_table { -- cgit v1.3 From 0b37da2434beb66713754869a1de775b82a72283 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 1 Jul 2016 07:00:20 -0400 Subject: ADD: launch all device kernels at once in functional simulator --- src/cuda-sim/cuda-sim.cc | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index f5d8a88..276cb9d 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -48,6 +48,7 @@ #include "../gpgpusim_entrypoint.h" #include "decuda_pred_table/decuda_pred_table.h" #include "../stream_manager.h" +#include "cuda_device_runtime.h" int gpgpu_ptx_instruction_classification; void ** g_inst_classification_stat = NULL; @@ -1776,13 +1777,15 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) g_the_gpu->getShaderCoreConfig()->warp_size ); cta.execute(); + + launch_all_device_kernels(); } //registering this kernel as done - extern stream_manager *g_stream_manager; //openCL kernel simulation calls don't register the kernel so we don't register its exit if(!openCL) { + extern stream_manager *g_stream_manager; g_stream_manager->register_finished_kernel(kernel.get_uid()); } -- cgit v1.3 From 8ef2e4eb13093c59190439800fdd0cc552a3779e Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 14 Nov 2014 18:45:46 -0500 Subject: ADD: add cdp latency --- src/abstract_hardware_model.h | 8 ++++++-- src/cuda-sim/cuda-sim.cc | 21 ++++++++++++++++----- src/cuda-sim/ptx_ir.cc | 10 ++++++---- src/gpgpu-sim/shader.cc | 26 ++++++++++++++++++++++++++ src/gpgpu-sim/shader.h | 13 +++++++++++++ 5 files changed, 67 insertions(+), 11 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 3a268ad..45334b6 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -867,7 +867,7 @@ public: m_mem_accesses_created=false; m_cache_hit=false; m_is_printf=false; - m_is_cdp = false; + m_is_cdp = 0; } virtual ~warp_inst_t(){ } @@ -1020,7 +1020,6 @@ protected: unsigned cycles; // used for implementing initiation interval delay bool m_isatomic; bool m_is_printf; - bool m_is_cdp; unsigned m_warp_id; unsigned m_dynamic_warp_id; const core_config *m_config; @@ -1041,6 +1040,11 @@ protected: std::list m_accessq; static unsigned sm_next_uid; + + //Jin: cdp support +public: + int m_is_cdp; + }; void move_warp( warp_inst_t *&dst, warp_inst_t *&src ); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 276cb9d..9ecd92b 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -64,6 +64,8 @@ unsigned gpgpu_param_num_shaders = 0; char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp; char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp; +char *cdp_latency_str; +unsigned cdp_latency[4]; void ptx_opcocde_latency_options (option_parser_t opp) { option_parser_register(opp, "-ptx_opcode_latency_int", OPT_CSTR, &opcode_latency_int, @@ -90,6 +92,11 @@ void ptx_opcocde_latency_options (option_parser_t opp) { "Opcode initiation intervals for double precision floating points " "Default 8,8,8,8,130", "8,8,8,8,130"); + option_parser_register(opp, "-cdp_latency", OPT_CSTR, &cdp_latency_str, + "CDP API latency " + "Default 1,7200,19320,1680", + "1,7200,19320,1680"); } static address_type get_converge_point(address_type pc); @@ -609,6 +616,8 @@ void ptx_instruction::set_opcode_and_latency() sscanf(opcode_initiation_dp, "%u,%u,%u,%u,%u", &dp_init[0],&dp_init[1],&dp_init[2], &dp_init[3],&dp_init[4]); + sscanf(cdp_latency_str, "%u,%u,%u,%u", + &cdp_latency[0],&cdp_latency[1],&cdp_latency[2], &cdp_latency[3]); if(!m_operands.empty()){ std::vector::iterator it; @@ -639,19 +648,21 @@ void ptx_instruction::set_opcode_and_latency() case MEMBAR_OP: op = MEMORY_BARRIER_OP; break; case CALL_OP: { - if(m_is_printf || m_is_cdp) + if(m_is_printf || m_is_cdp) { op = ALU_OP; + } else op = CALL_OPS; break; } case CALLP_OP: { - if(m_is_printf || m_is_cdp) + if(m_is_printf || m_is_cdp) { op = ALU_OP; - else - op = CALL_OPS; - break; + } + else + op = CALL_OPS; + break; } case RET_OP: case RETP_OP: op = RET_OPS;break; case ADD_OP: case ADDP_OP: case ADDC_OP: case SUB_OP: case SUBC_OP: diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 176eb14..4931213 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1241,10 +1241,12 @@ ptx_instruction::ptx_instruction( int opcode, if (fname =="vprintf"){ m_is_printf = true; } - if (fname == "cudaGetParameterBufferV2" - || fname == "cudaLaunchDeviceV2" - || fname == "cudaStreamCreateWithFlags") - m_is_cdp = true; + if(fname == "cudaGetParameterBufferV2") + m_is_cdp = 1; + if(fname == "cudaStreamCreateWithFlags") + m_is_cdp = 2; + if(fname == "cudaLaunchDeviceV2") + m_is_cdp = 3; } } diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index cd38cb7..e85c4a8 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -841,6 +841,13 @@ void scheduler_unit::cycle() unsigned max_issue = m_shader->m_config->gpgpu_max_insn_issue_per_warp; while( !warp(warp_id).waiting() && !warp(warp_id).ibuffer_empty() && (checked < max_issue) && (checked <= issued) && (issued < max_issue) ) { const warp_inst_t *pI = warp(warp_id).ibuffer_next_inst(); + //Jin: handle cdp latency; + if(pI->m_is_cdp && warp(warp_id).m_cdp_latency > 0) { + assert(warp(warp_id).m_cdp_dummy); + warp(warp_id).m_cdp_latency--; + break; + } + bool valid = warp(warp_id).ibuffer_next_valid(); bool warp_inst_issued = false; unsigned pc,rpc; @@ -875,6 +882,25 @@ void scheduler_unit::cycle() bool sp_pipe_avail = m_sp_out->has_free(); bool sfu_pipe_avail = m_sfu_out->has_free(); if( sp_pipe_avail && (pI->op != SFU_OP) ) { + + //Jin: special for CDP api + if(pI->m_is_cdp && !warp(warp_id).m_cdp_dummy) { + assert(warp(warp_id).m_cdp_latency == 0); + + extern unsigned cdp_latency[3]; + if(pI->m_is_cdp != 3) + warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1]; + else //cudaLaunchDeviceV2 + warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1] + + cdp_latency[pI->m_is_cdp] * active_mask.count(); + warp(warp_id).m_cdp_dummy = true; + break; + } + else if(pI->m_is_cdp && warp(warp_id).m_cdp_dummy) { + assert(warp(warp_id).m_cdp_latency == 0); + warp(warp_id).m_cdp_dummy = false; + } + // always prefer SP pipe for operations that can use both SP and SFU pipelines m_shader->issue_warp(*m_sp_out,pI,active_mask,warp_id); issued++; diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index fcbc8aa..882868e 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -108,6 +108,10 @@ public: m_last_fetch=0; m_next=0; m_inst_at_barrier=NULL; + + //Jin: cdp support + m_cdp_latency = 0; + m_cdp_dummy = false; } void init( address_type start_pc, unsigned cta_id, @@ -124,6 +128,10 @@ public: n_completed -= active.count(); // active threads are not yet completed m_active_threads = active; m_done_exit=false; + + //Jin: cdp support + m_cdp_latency = 0; + m_cdp_dummy = false; } bool functional_done() const; @@ -260,6 +268,11 @@ private: unsigned m_stores_outstanding; // number of store requests sent but not yet acknowledged unsigned m_inst_in_pipeline; + + //Jin: cdp support +public: + unsigned int m_cdp_latency; + bool m_cdp_dummy; }; -- cgit v1.3 From bbcb492c0f6d887c4034bd15adf57420dd735c5e Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 14 Nov 2014 20:01:10 -0500 Subject: ADD: add separate cdp latency --- src/cuda-sim/cuda-sim.cc | 16 +++++++++------- src/cuda-sim/ptx_ir.cc | 6 +++--- src/gpgpu-sim/shader.cc | 7 ++++--- 3 files changed, 16 insertions(+), 13 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 9ecd92b..c87e3e4 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -65,7 +65,7 @@ unsigned gpgpu_param_num_shaders = 0; char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp; char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp; char *cdp_latency_str; -unsigned cdp_latency[4]; +unsigned cdp_latency[5]; void ptx_opcocde_latency_options (option_parser_t opp) { option_parser_register(opp, "-ptx_opcode_latency_int", OPT_CSTR, &opcode_latency_int, @@ -93,10 +93,11 @@ void ptx_opcocde_latency_options (option_parser_t opp) { "Default 8,8,8,8,130", "8,8,8,8,130"); option_parser_register(opp, "-cdp_latency", OPT_CSTR, &cdp_latency_str, - "CDP API latency " - "Default 1,7200,19320,1680", - "1,7200,19320,1680"); + "CDP API latency " + "Default 7200,8000,100,12000,1600", + "7200,8000,100,12000,1600"); } static address_type get_converge_point(address_type pc); @@ -616,8 +617,9 @@ void ptx_instruction::set_opcode_and_latency() sscanf(opcode_initiation_dp, "%u,%u,%u,%u,%u", &dp_init[0],&dp_init[1],&dp_init[2], &dp_init[3],&dp_init[4]); - sscanf(cdp_latency_str, "%u,%u,%u,%u", - &cdp_latency[0],&cdp_latency[1],&cdp_latency[2], &cdp_latency[3]); + sscanf(cdp_latency_str, "%u,%u,%u,%u,%u", + &cdp_latency[0],&cdp_latency[1],&cdp_latency[2], + &cdp_latency[3],&cdp_latency[4]); if(!m_operands.empty()){ std::vector::iterator it; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 4931213..783c885 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1241,12 +1241,12 @@ ptx_instruction::ptx_instruction( int opcode, if (fname =="vprintf"){ m_is_printf = true; } - if(fname == "cudaGetParameterBufferV2") - m_is_cdp = 1; if(fname == "cudaStreamCreateWithFlags") + m_is_cdp = 1; + if(fname == "cudaGetParameterBufferV2") m_is_cdp = 2; if(fname == "cudaLaunchDeviceV2") - m_is_cdp = 3; + m_is_cdp = 4; } } diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index e85c4a8..8ce2146 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -887,12 +887,13 @@ void scheduler_unit::cycle() if(pI->m_is_cdp && !warp(warp_id).m_cdp_dummy) { assert(warp(warp_id).m_cdp_latency == 0); - extern unsigned cdp_latency[3]; - if(pI->m_is_cdp != 3) + extern unsigned cdp_latency[5]; + if(pI->m_is_cdp == 1) warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1]; - else //cudaLaunchDeviceV2 + else //cudaLaunchDeviceV2 and cudaGetParameterBufferV2 warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1] + cdp_latency[pI->m_is_cdp] * active_mask.count(); + printf("set latency %d\n", warp(warp_id).m_cdp_latency); warp(warp_id).m_cdp_dummy = true; break; } -- cgit v1.3 From a31392370bc7b23e115160b170e95b117597ebc4 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sat, 15 Nov 2014 02:09:24 -0500 Subject: ADD: add stats for kernel launching and complete cycle --- src/abstract_hardware_model.h | 6 ++++++ src/cuda-sim/cuda_device_runtime.cc | 2 +- src/gpgpu-sim/gpu-sim.cc | 2 ++ src/stream_manager.cc | 9 ++++++++- 4 files changed, 17 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 45334b6..e6ef521 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -284,6 +284,7 @@ public: bool cta_has_stream(dim3 ctaid, CUstream_st* stream); void destroy_cta_streams(); void print_parent_info(); + kernel_info_t * get_parent() { return m_parent_kernel; } private: kernel_info_t * m_parent_kernel; @@ -292,6 +293,11 @@ private: std::list m_child_kernels; //child kernel launched std::map< dim3, std::list, dim3comp > m_cta_streams; //streams created in each CTA +//Jin: kernel timing +public: + unsigned long long launch_cycle; + unsigned long long start_cycle; + unsigned long long end_cycle; }; struct core_config { diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index df2653e..1b8c8d9 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -176,6 +176,7 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * //create child kernel_info_t and index it with parameter_buffer address device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry); + device_grid->launch_cycle = gpu_sim_cycle; kernel_info_t & parent_grid = thread->get_kernel(); DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z << @@ -184,7 +185,6 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * device_grid->set_parent(&parent_grid, thread->get_ctaid(), thread->get_tid()); device_launch_op = device_launch_operation_t(device_grid, NULL); device_kernel_param_mem = device_grid->get_param_memory(); //kernel param - size_t param_start_address = 0; //copy in word for(unsigned n = 0; n < device_kernel_arg_size; n += 4) { diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 81c9c9a..27362cf 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -531,6 +531,7 @@ kernel_info_t *gpgpu_sim::select_kernel() unsigned idx = (n+m_last_issued_kernel+1)%m_config.max_concurrent_kernel; if( kernel_more_cta_left(m_running_kernels[idx]) ){ m_last_issued_kernel=idx; + m_running_kernels[idx]->start_cycle = gpu_sim_cycle; // record this kernel for stat print if it is the first time this kernel is selected for execution unsigned launch_uid = m_running_kernels[idx]->get_uid(); if (std::find(m_executed_kernel_uids.begin(), m_executed_kernel_uids.end(), launch_uid) == m_executed_kernel_uids.end()) { @@ -560,6 +561,7 @@ void gpgpu_sim::set_kernel_done( kernel_info_t *kernel ) std::vector::iterator k; for( k=m_running_kernels.begin(); k!=m_running_kernels.end(); k++ ) { if( *k == kernel ) { + kernel->end_cycle = gpu_sim_cycle; *k = NULL; break; } diff --git a/src/stream_manager.cc b/src/stream_manager.cc index f90d9be..5bd7737 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -254,7 +254,14 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) //Jin: should check children kernels for CDP if(kernel->is_finished()) { - printf("kernel %d finishes, retires from stream %d\n", grid_uid, stream->get_uid()); +// std::ofstream kernel_stat("kernel_stat.txt", std::ofstream::out | std::ofstream::app); +// kernel_stat<< " kernel " << grid_uid; +// if(kernel->get_parent()) +// kernel_stat << ", parent " << kernel->get_parent()->get_uid() << +// ", launch " << kernel->launch_cycle; +// kernel_stat<< ", start " << kernel->start_cycle << +// ", end " << kernel->end_cycle << ", retire " << gpu_sim_cycle << "\n"; +// printf("kernel %d finishes, retires from stream %d\n", grid_uid, stream->get_uid()); stream->record_next_done(); m_grid_id_to_stream.erase(grid_uid); kernel->notify_parent_finished(); -- cgit v1.3 From b70a5a69fa14cc01f707b910f8f021e36067922f Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sat, 15 Nov 2014 05:44:03 -0500 Subject: ADD: print kernel parameter size footprint. BUG: concurrent kernels on same shader, should use hw_cta_id to store shared mem info --- src/abstract_hardware_model.cc | 2 ++ src/cuda-sim/cuda-sim.cc | 3 ++- src/cuda-sim/cuda_device_runtime.cc | 7 ++++++- src/gpgpu-sim/gpu-sim.cc | 23 +++++++++++++++-------- src/stream_manager.cc | 6 ++++-- 5 files changed, 29 insertions(+), 12 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 4db5f2f..be5c5b9 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -618,6 +618,8 @@ bool kernel_info_t::children_all_finished() { void kernel_info_t::notify_parent_finished() { if(m_parent_kernel) { + extern unsigned long long g_total_param_size; + g_total_param_size -= ((m_kernel_entry->get_args_aligned_size() + 255)/256*256); m_parent_kernel->remove_child(this); g_stream_manager->register_finished_kernel(m_parent_kernel->get_uid()); } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index c87e3e4..3f5af7e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1499,7 +1499,8 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, unsigned max_cta_per_sm = num_threads/cta_size; // e.g., 256 / 48 = 5 assert( max_cta_per_sm > 0 ); - unsigned sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid; + //unsigned sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid; + unsigned sm_idx = hw_cta_id*gpgpu_param_num_shaders + sid; if ( shared_memory_lookup.find(sm_idx) == shared_memory_lookup.end() ) { if ( g_debug_execution >= 1 ) { diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 1b8c8d9..66cd063 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -59,6 +59,8 @@ public: std::map g_cuda_device_launch_param_map; std::list g_cuda_device_launch_op; extern stream_manager *g_stream_manager; +unsigned long long g_total_param_size = 0; +unsigned long long g_max_total_param_size = 0; //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -111,7 +113,10 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i //get total child kernel argument size and malloc buffer in global memory unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size); + g_total_param_size += ((child_kernel_arg_size + 255) / 256 * 256); DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer); + if(g_total_param_size > g_max_total_param_size) + g_max_total_param_size = g_total_param_size; //store param buffer address and launch config device_launch_config_t device_launch_config(grid_dim, block_dim, shared_mem, child_kernel_entry); @@ -176,7 +181,7 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * //create child kernel_info_t and index it with parameter_buffer address device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry); - device_grid->launch_cycle = gpu_sim_cycle; + device_grid->launch_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; kernel_info_t & parent_grid = thread->get_kernel(); DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z << diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 27362cf..a9da1c9 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -524,20 +524,26 @@ bool gpgpu_sim::get_more_cta_left() const kernel_info_t *gpgpu_sim::select_kernel() { if(m_running_kernels[m_last_issued_kernel] && - !m_running_kernels[m_last_issued_kernel]->no_more_ctas_to_run()) + !m_running_kernels[m_last_issued_kernel]->no_more_ctas_to_run()) { + unsigned launch_uid = m_running_kernels[m_last_issued_kernel]->get_uid(); + if(std::find(m_executed_kernel_uids.begin(), m_executed_kernel_uids.end(), launch_uid) == m_executed_kernel_uids.end()) { + m_running_kernels[m_last_issued_kernel]->start_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; + m_executed_kernel_uids.push_back(launch_uid); + m_executed_kernel_names.push_back(m_running_kernels[m_last_issued_kernel]->name()); + } return m_running_kernels[m_last_issued_kernel]; + } for(unsigned n=0; n < m_running_kernels.size(); n++ ) { unsigned idx = (n+m_last_issued_kernel+1)%m_config.max_concurrent_kernel; if( kernel_more_cta_left(m_running_kernels[idx]) ){ m_last_issued_kernel=idx; - m_running_kernels[idx]->start_cycle = gpu_sim_cycle; + m_running_kernels[idx]->start_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; // record this kernel for stat print if it is the first time this kernel is selected for execution unsigned launch_uid = m_running_kernels[idx]->get_uid(); - if (std::find(m_executed_kernel_uids.begin(), m_executed_kernel_uids.end(), launch_uid) == m_executed_kernel_uids.end()) { - m_executed_kernel_uids.push_back(launch_uid); - m_executed_kernel_names.push_back(m_running_kernels[idx]->name()); - } + assert(std::find(m_executed_kernel_uids.begin(), m_executed_kernel_uids.end(), launch_uid) == m_executed_kernel_uids.end()); + m_executed_kernel_uids.push_back(launch_uid); + m_executed_kernel_names.push_back(m_running_kernels[idx]->name()); return m_running_kernels[idx]; } @@ -561,7 +567,7 @@ void gpgpu_sim::set_kernel_done( kernel_info_t *kernel ) std::vector::iterator k; for( k=m_running_kernels.begin(); k!=m_running_kernels.end(); k++ ) { if( *k == kernel ) { - kernel->end_cycle = gpu_sim_cycle; + kernel->end_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; *k = NULL; break; } @@ -942,7 +948,8 @@ void gpgpu_sim::gpu_print_stat() printf("gpu_tot_ipc = %12.4f\n", (float)(gpu_tot_sim_insn+gpu_sim_insn) / (gpu_tot_sim_cycle+gpu_sim_cycle)); printf("gpu_tot_issued_cta = %lld\n", gpu_tot_issued_cta + m_total_cta_launched); - + extern unsigned long long g_max_total_param_size; + fprintf(statfout, "max_total_param_size = %llu\n", g_max_total_param_size); // performance counter for stalls due to congestion. printf("gpu_stall_dramfull = %d\n", gpu_stall_dramfull); diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 5bd7737..3459c6d 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -219,7 +219,7 @@ bool stream_manager::operation( bool * sim) { bool check=check_finished_kernel(); pthread_mutex_lock(&m_lock); - if(check)m_gpu->print_stats(); +// if(check)m_gpu->print_stats(); stream_operation op =front(); if(!op.do_operation( m_gpu )) //not ready to execute { @@ -260,8 +260,10 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) // kernel_stat << ", parent " << kernel->get_parent()->get_uid() << // ", launch " << kernel->launch_cycle; // kernel_stat<< ", start " << kernel->start_cycle << -// ", end " << kernel->end_cycle << ", retire " << gpu_sim_cycle << "\n"; +// ", end " << kernel->end_cycle << ", retire " << gpu_sim_cycle + gpu_tot_sim_cycle << "\n"; // printf("kernel %d finishes, retires from stream %d\n", grid_uid, stream->get_uid()); +// kernel_stat.flush(); +// kernel_stat.close(); stream->record_next_done(); m_grid_id_to_stream.erase(grid_uid); kernel->notify_parent_finished(); -- cgit v1.3 From 06c5a3ddc979a78dd38bf0f74170e4919cd32fab Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sun, 3 Jul 2016 06:46:15 -0400 Subject: BUG: wrong declaration for m_args_aligned_size --- src/cuda-sim/ptx_ir.h | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index a54ae41..55b01fd 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -292,8 +292,6 @@ private: std::list m_initializer; static unsigned sm_next_uid; - //parameter size for device kernels - int m_args_aligned_size; }; class symbol_table { @@ -1302,6 +1300,9 @@ private: static std::vector s_g_pc_to_insn; // a direct mapping from PC to instruction static unsigned sm_next_uid; + + //parameter size for device kernels + int m_args_aligned_size; }; class arg_buffer_t { -- cgit v1.3 From b64b0f41e217db350637aff6de0ddb18c8e7ad49 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sun, 3 Jul 2016 07:10:38 -0400 Subject: BUG: extra bracket --- src/cuda-sim/cuda_device_runtime.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 66cd063..12c83d2 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -215,9 +215,9 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * " to stream " << child_stream->get_uid() << ": " << child_stream); } - device_launch_op.stream = child_stream; - } + device_launch_op.stream = child_stream; } + } -- cgit v1.3 From 2af85e353bccad8c8536c1d5f039361884b96872 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Wed, 6 Jul 2016 05:18:44 -0400 Subject: ADD: add knob to enable CDP in gpgpusim config --- libcuda/cuda_runtime_api.cc | 14 +++++++++++++- src/abstract_hardware_model.cc | 3 ++- src/cuda-sim/ptx_loader.cc | 6 +++++- src/gpgpu-sim/gpu-sim.cc | 4 ++++ 4 files changed, 24 insertions(+), 3 deletions(-) (limited to 'src/cuda-sim') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 30bf823..aa9a3eb 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1336,7 +1336,11 @@ void extract_code_using_cuobjdump(){ system(command); // Running cuobjdump using dynamic link to current process // Needs the option '-all' to extract PTX from CDP-enabled binary - snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all %s > %s", app_binary.c_str(), fname); + extern bool g_cdp_enabled; + if(!g_cdp_enabled) + snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass %s > %s", app_binary.c_str(), fname); + else + snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all %s > %s", app_binary.c_str(), fname); bool parse_output = true; int result = system(command); if(result) { @@ -1545,6 +1549,14 @@ cuobjdumpPTXSection* findPTXSectionInList(std::list sectionli if((ptxsection=dynamic_cast(*iter)) != NULL){ if(ptxsection->getIdentifier() == identifier) return ptxsection; + else { + extern bool g_cdp_enabled; + if(g_cdp_enabled) { + printf("Warning: __cudaRegisterFatBinary needs %s, but find PTX section with %s\n", + identifier.c_str(), ptxsection->getIdentifier().c_str()); + return ptxsection; + } + } } } return NULL; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 819ad35..fe6f8ab 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -550,7 +550,8 @@ void warp_inst_t::completed( unsigned long long cycle ) const ptx_file_line_stats_add_latency(pc, latency * active_count()); } -//Jin: kernel launch latency +//Jin: CDP support +bool g_cdp_enabled; unsigned g_kernel_launch_latency; unsigned kernel_info_t::m_next_uid = 1; diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index f7bf70e..a646408 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -217,7 +217,11 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num #if CUDART_VERSION >= 3000 if (sm_version == 0) sm_version = 20; - snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version); + extern bool g_cdp_enabled; + if(!g_cdp_enabled) + snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version); + else + snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version); #endif snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file /dev/null 2> %s", diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 02c9c09..0b4b2f6 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -449,6 +449,10 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-gpgpu_kernel_launch_latency", OPT_INT32, &g_kernel_launch_latency, "Kernel launch latency in cycles. Default: 0", "0"); + extern bool g_cdp_enabled; + option_parser_register(opp, "-gpgpu_cdp_enabled", OPT_BOOL, + &g_cdp_enabled, "Turn on CDP", + "0"); } ///////////////////////////////////////////////////////////////////////////// -- cgit v1.3 From 623a88e5d5c6c3edb94404ef6e5ea100caec9deb Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Mon, 29 Aug 2016 18:10:00 -0700 Subject: MOD: Add macros to turn off cuda_device_runtime for CUDA < 5.0 --- src/cuda-sim/cuda-sim.cc | 2 ++ src/cuda-sim/cuda_device_runtime.cc | 10 +++++++--- src/cuda-sim/cuda_device_runtime.h | 3 ++- src/cuda-sim/instructions.cc | 2 ++ src/gpgpu-sim/gpu-sim.cc | 2 ++ src/gpgpu-sim/shader.cc | 2 +- 6 files changed, 16 insertions(+), 5 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 3f5af7e..ec51779 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1792,7 +1792,9 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) ); cta.execute(); +#if (CUDART_VERSION >= 5000) launch_all_device_kernels(); +#endif } //registering this kernel as done diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 12c83d2..4a8ffe5 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -1,9 +1,15 @@ //Jin: cuda_device_runtime.cc //Defines CUDA device runtime APIs for CDP support + #include #include +unsigned long long g_total_param_size = 0; +unsigned long long g_max_total_param_size = 0; + + +#if (CUDART_VERSION >= 5000) #define __CUDA_RUNTIME_API_H__ #include @@ -59,8 +65,6 @@ public: std::map g_cuda_device_launch_param_map; std::list g_cuda_device_launch_op; extern stream_manager *g_stream_manager; -unsigned long long g_total_param_size = 0; -unsigned long long g_max_total_param_size = 0; //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -313,4 +317,4 @@ void launch_all_device_kernels() { launch_one_device_kernel(); } } - +#endif diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index 385d605..6dbcd71 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -1,6 +1,6 @@ //Jin: cuda_device_runtime.h //Defines CUDA device runtime APIs for CDP support - +#if (CUDART_VERSION >= 5000) #pragma once void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); @@ -8,3 +8,4 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void launch_all_device_kernels(); void launch_one_device_kernel(); +#endif diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 5d909d3..e68f9fd 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1414,6 +1414,7 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) return; } +#if (CUDART_VERSION >= 5000) //Jin: handle device runtime apis for CDP else if(fname == "cudaGetParameterBufferV2") { gpgpusim_cuda_getParameterBufferV2(pI, thread, target_func); @@ -1427,6 +1428,7 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) gpgpusim_cuda_streamCreateWithFlags(pI, thread, target_func); return; } +#endif // read source arguements into register specified in declaration of function arg_buffer_list_t arg_values; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 0b4b2f6..363fe5a 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1513,8 +1513,10 @@ void gpgpu_sim::cycle() try_snap_shot(gpu_sim_cycle); spill_log_to_file (stdout, 0, gpu_sim_cycle); +#if (CUDART_VERSION >= 5000) //launch device kernel launch_one_device_kernel(); +#endif } } diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index b9caf18..59a2d8b 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -846,7 +846,7 @@ void scheduler_unit::cycle() while( !warp(warp_id).waiting() && !warp(warp_id).ibuffer_empty() && (checked < max_issue) && (checked <= issued) && (issued < max_issue) ) { const warp_inst_t *pI = warp(warp_id).ibuffer_next_inst(); //Jin: handle cdp latency; - if(pI->m_is_cdp && warp(warp_id).m_cdp_latency > 0) { + if(pI && pI->m_is_cdp && warp(warp_id).m_cdp_latency > 0) { assert(warp(warp_id).m_cdp_dummy); warp(warp_id).m_cdp_latency--; break; -- cgit v1.3