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') 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') 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') 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') 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') 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') 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') 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') 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') 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') 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') 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 38f811ab4d094e56b2065e1c8cb39327ad9f157b Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Thu, 30 Jun 2016 08:48:40 -0400 Subject: ADD: delete streams created by cta when deleting kernel --- src/abstract_hardware_model.cc | 12 ++++++++++++ src/abstract_hardware_model.h | 1 + 2 files changed, 13 insertions(+) (limited to 'src') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index f3c5b21..4db5f2f 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -574,6 +574,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * kernel_info_t::~kernel_info_t() { assert( m_active_threads.empty() ); + destroy_cta_streams(); delete m_param_mem; } @@ -668,6 +669,17 @@ void kernel_info_t::print_parent_info() { } } +void kernel_info_t::destroy_cta_streams() { + printf("Destroy streams for kernel %d: ", get_uid()); size_t stream_size = 0; + for(auto s = m_cta_streams.begin(); s != m_cta_streams.end(); s++) { + stream_size += s->second.size(); + for(auto ss = s->second.begin(); ss != s->second.end(); ss++) + g_stream_manager->destroy_stream(*ss); + s->second.clear(); + } + printf("size %lu\n", stream_size); + m_cta_streams.clear(); +} simt_stack::simt_stack( unsigned wid, unsigned warpSize) { diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index a0bf80d..3a268ad 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -282,6 +282,7 @@ public: 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 destroy_cta_streams(); void print_parent_info(); private: -- 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') 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') 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') 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 3321626d5e858df8e2154bf4e7a1bacda76658e7 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Thu, 13 Nov 2014 17:16:44 -0500 Subject: ADD: support concurrent kernels on one shader --- src/gpgpu-sim/gpu-sim.cc | 113 ++++++++++++++++++++++++++++++++++++++++++++--- src/gpgpu-sim/shader.cc | 79 ++++++++++++++++++++++++--------- src/gpgpu-sim/shader.h | 23 +++++++++- 3 files changed, 186 insertions(+), 29 deletions(-) (limited to 'src') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 126e007..8a5d581 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -367,6 +367,10 @@ void shader_core_config::reg_options(class OptionParser * opp) "For complete list of prioritization values see shader.h enum scheduler_prioritization_type" "Default: gto", "gto"); + + option_parser_register(opp, "-gpgpu_concurrent_kernel_sm", OPT_BOOL, &gpgpu_concurrent_kernel_sm, + "Support concurrent kernels on a SM (default = enabled)", + "1"); } void gpgpu_sim_config::reg_options(option_parser_t opp) @@ -1075,7 +1079,98 @@ void shader_core_ctx::mem_instruction_stats(const warp_inst_t &inst) abort(); } } +//Jin: concurrent kernels on one SM +bool shader_core_ctx::can_issue_1block(kernel_info_t & kernel) { + + if(m_config->max_cta(kernel) < 1) + return false; + + return occupy_shader_resource_1block(kernel, false); +} + +int shader_core_ctx::find_available_hwtid(unsigned int cta_size) { + + unsigned int step; + for(step = 0; step < m_config->n_thread_per_shader; + step += cta_size) { + + unsigned int hw_tid; + for(hw_tid = step; hw_tid < step + cta_size; + hw_tid++) { + if(m_active_threads.test(hw_tid)) + break; + } + if(hw_tid == step + cta_size) //consecutive non-active + break; + } + if(step >= m_config->n_thread_per_shader) //didn't find + return -1; + else + return step; +} + +bool shader_core_ctx::occupy_shader_resource_1block(kernel_info_t & k, bool occupy) { + unsigned threads_per_cta = k.threads_per_cta(); + const class function_info *kernel = k.entry(); + unsigned int padded_cta_size = threads_per_cta; + unsigned int warp_size = m_config->warp_size; + if (padded_cta_size%warp_size) + padded_cta_size = ((padded_cta_size/warp_size)+1)*(warp_size); + + if(m_occupied_n_threads + padded_cta_size > m_config->n_thread_per_shader) + return false; + + if(find_available_hwtid(padded_cta_size) == -1) + return false; + + const struct gpgpu_ptx_sim_kernel_info *kernel_info = ptx_sim_kernel_info(kernel); + + if(m_occupied_shmem + kernel_info->smem > m_config->gpgpu_shmem_size) + return false; + + unsigned int used_regs = padded_cta_size * ((kernel_info->regs+3)&~3); + if(m_occupied_regs + used_regs > m_config->gpgpu_shader_registers) + return false; + + if(m_occupied_ctas +1 > m_config->max_cta_per_core) + return false; + + if(occupy) { + m_occupied_n_threads += padded_cta_size; + m_occupied_shmem += kernel_info->smem; + m_occupied_regs += (padded_cta_size * ((kernel_info->regs+3)&~3)); + m_occupied_ctas++; + + printf("GPGPU-Sim uArch: Shader %d occupied %d threads, %d shared mem, %d registers, %d ctas\n", + m_sid, m_occupied_n_threads, m_occupied_shmem, m_occupied_regs, m_occupied_ctas); + } + + return true; +} +void shader_core_ctx::release_shader_resource_1block(kernel_info_t & k) { + unsigned threads_per_cta = k.threads_per_cta(); + const class function_info *kernel = k.entry(); + unsigned int padded_cta_size = threads_per_cta; + unsigned int warp_size = m_config->warp_size; + if (padded_cta_size%warp_size) + padded_cta_size = ((padded_cta_size/warp_size)+1)*(warp_size); + + assert(m_occupied_n_threads >= padded_cta_size); + m_occupied_n_threads -= padded_cta_size; + + const struct gpgpu_ptx_sim_kernel_info *kernel_info = ptx_sim_kernel_info(kernel); + + assert(m_occupied_shmem >= (unsigned int)kernel_info->smem); + m_occupied_shmem -= kernel_info->smem; + + unsigned int used_regs = padded_cta_size * ((kernel_info->regs+3)&~3); + assert(m_occupied_regs >= used_regs); + m_occupied_regs -= used_regs; + + assert(m_occupied_ctas >= 1); + m_occupied_ctas--; +} //////////////////////////////////////////////////////////////////////////////////////////////// @@ -1088,11 +1183,14 @@ void shader_core_ctx::mem_instruction_stats(const warp_inst_t &inst) void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) { - set_max_cta(kernel); +// set_max_cta(kernel); + kernel.inc_running(); + assert(occupy_shader_resource_1block(kernel, true)); // find a free CTA context unsigned free_cta_hw_id=(unsigned)-1; - for (unsigned i=0;imax_cta_per_core;i++ ) { if( m_cta_status[i]==0 ) { free_cta_hw_id=i; break; @@ -1109,8 +1207,11 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) int padded_cta_size = cta_size; if (cta_size%m_config->warp_size) padded_cta_size = ((cta_size/m_config->warp_size)+1)*(m_config->warp_size); - unsigned start_thread = free_cta_hw_id * padded_cta_size; - unsigned end_thread = start_thread + cta_size; + unsigned int start_thread = find_available_hwtid(padded_cta_size); + assert((int)start_thread != -1); + unsigned int end_thread = start_thread + cta_size; +// unsigned start_thread = free_cta_hw_id * padded_cta_size; +// unsigned end_thread = start_thread + cta_size; // reset the microarchitecture state of the selected hardware thread and warp contexts reinit(start_thread, end_thread,false); @@ -1138,7 +1239,9 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) m_n_active_cta++; shader_CTA_count_log(m_sid, 1); - printf("GPGPU-Sim uArch: core:%3d, cta:%2u initialized @(%lld,%lld)\n", m_sid, free_cta_hw_id, gpu_sim_cycle, gpu_tot_sim_cycle ); + printf("GPGPU-Sim uArch: core:%3d, cta:%2u, start_tid:%4u, end_tid:%4u, initialized @(%lld,%lld)\n", + m_sid, free_cta_hw_id, start_thread, end_thread, gpu_sim_cycle, gpu_tot_sim_cycle ); + } /////////////////////////////////////////////////////////////////////////////////////////// diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index de5bcf6..8cba31e 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -296,6 +296,12 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, m_last_inst_gpu_sim_cycle = 0; m_last_inst_gpu_tot_sim_cycle = 0; + + //Jin: for concurrent kernels on a SM + m_occupied_n_threads = 0; + m_occupied_shmem = 0; + m_occupied_regs = 0; + m_occupied_ctas = 0; } void shader_core_ctx::reinit(unsigned start_thread, unsigned end_thread, bool reset_not_completed ) @@ -303,6 +309,13 @@ void shader_core_ctx::reinit(unsigned start_thread, unsigned end_thread, bool re if( reset_not_completed ) { m_not_completed = 0; m_active_threads.reset(); + + //Jin: for concurrent kernels on a SM + m_occupied_n_threads = 0; + m_occupied_shmem = 0; + m_occupied_regs = 0; + m_occupied_ctas = 0; + } for (unsigned i = start_thread; iget_kernel())); m_not_completed -= 1; m_active_threads.reset(tid); assert( m_thread[tid]!= NULL ); @@ -1917,7 +1930,7 @@ void ldst_unit::cycle() } } -void shader_core_ctx::register_cta_thread_exit( unsigned cta_num ) +void shader_core_ctx::register_cta_thread_exit( unsigned cta_num, kernel_info_t * kernel) { assert( m_cta_status[cta_num] > 0 ); m_cta_status[cta_num]--; @@ -1925,23 +1938,33 @@ void shader_core_ctx::register_cta_thread_exit( unsigned cta_num ) m_n_active_cta--; m_barriers.deallocate_barrier(cta_num); shader_CTA_count_unlog(m_sid, 1); + printf("GPGPU-Sim uArch: Shader %d finished CTA #%d (%lld,%lld), %u CTAs running\n", m_sid, cta_num, gpu_sim_cycle, gpu_tot_sim_cycle, m_n_active_cta ); + if( m_n_active_cta == 0 ) { - assert( m_kernel != NULL ); - m_kernel->dec_running(); - printf("GPGPU-Sim uArch: Shader %u empty (release kernel %u \'%s\').\n", m_sid, m_kernel->get_uid(), - 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 %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 ); - } - } - m_kernel=NULL; + printf("GPGPU-Sim uArch: Shader %u empty (last released kernel %u \'%s\').\n", m_sid, kernel->get_uid(), + kernel->name().c_str() ); fflush(stdout); + + //Shader can only be empty when no more cta are dispatched + assert(m_kernel == NULL || !m_gpu->kernel_more_cta_left(m_kernel)); + m_kernel = NULL; } + + //Jin: for concurrent kernels on sm + release_shader_resource_1block(*kernel); + kernel->dec_running(); + if( !m_gpu->kernel_more_cta_left(kernel) ) { + if( !kernel->running() ) { + printf("GPGPU-Sim uArch: GPU detected kernel %u \'%s\' finished on shader %u.\n", kernel->get_uid(), + kernel->name().c_str(), m_sid ); + if(m_kernel == kernel) + m_kernel = NULL; + m_gpu->set_kernel_done( kernel ); + } + } + } } @@ -3239,15 +3262,27 @@ unsigned simt_core_cluster::issue_block2core() unsigned num_blocks_issued=0; for( unsigned i=0; i < m_config->n_simt_cores_per_cluster; i++ ) { unsigned core = (i+m_cta_issue_next_core+1)%m_config->n_simt_cores_per_cluster; - if( m_core[core]->get_not_completed() == 0 ) { - if( m_core[core]->get_kernel() == NULL ) { - kernel_info_t *k = m_gpu->select_kernel(); - if( k ) - m_core[core]->set_kernel(k); - } - } + kernel_info_t *kernel = m_core[core]->get_kernel(); - if( m_gpu->kernel_more_cta_left(kernel) && (m_core[core]->get_n_active_cta() < m_config->max_cta(*kernel)) ) { + + //Jin: check if to fetch the next kernel + if( !m_gpu->kernel_more_cta_left(kernel) ) { + if(m_config->gpgpu_concurrent_kernel_sm || //concurrent kernel on sm + + //otherwise wait till current kernel finishes + (!m_config->gpgpu_concurrent_kernel_sm && + m_core[core]->get_not_completed() == 0) ) + { + kernel_info_t *k = m_gpu->select_kernel(); + if( k ) + m_core[core]->set_kernel(k); + } + } + + kernel = m_core[core]->get_kernel(); + if( m_gpu->kernel_more_cta_left(kernel) && +// (m_core[core]->get_n_active_cta() < m_config->max_cta(*kernel)) ) { + m_core[core]->can_issue_1block(*kernel)) { m_core[core]->issue_block2core(*kernel); num_blocks_issued++; m_cta_issue_next_core=core; diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 38d09e9..fcbc8aa 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1327,6 +1327,9 @@ struct shader_core_config : public core_config int simt_core_sim_order; unsigned mem2device(unsigned memid) const { return memid + n_simt_clusters; } + + //Jin: concurrent kernel on sm + bool gpgpu_concurrent_kernel_sm; }; struct shader_core_stats_pod { @@ -1574,6 +1577,7 @@ public: void cycle(); void reinit(unsigned start_thread, unsigned end_thread, bool reset_not_completed ); void issue_block2core( class kernel_info_t &kernel ); + void cache_flush(); void accept_fetch_response( mem_fetch *mf ); void accept_ldst_unit_response( class mem_fetch * mf ); @@ -1582,7 +1586,7 @@ public: { assert(k); m_kernel=k; - k->inc_running(); +// k->inc_running(); printf("GPGPU-Sim uArch: Shader %d bind to kernel %u \'%s\'\n", m_sid, m_kernel->get_uid(), m_kernel->name().c_str() ); } @@ -1749,7 +1753,7 @@ public: virtual void checkExecutionStatusAndUpdate(warp_inst_t &inst, unsigned t, unsigned tid); address_type next_pc( int tid ) const; void fetch(); - void register_cta_thread_exit( unsigned cta_num ); + void register_cta_thread_exit( unsigned cta_num, kernel_info_t * kernel ); void decode(); @@ -1831,6 +1835,20 @@ public: // is that the dynamic_warp_id is a running number unique to every warp // run on this shader, where the warp_id is the static warp slot. unsigned m_dynamic_warp_id; + + //Jin: concurrent kernels on a sm +public: + bool can_issue_1block(kernel_info_t & kernel); + bool occupy_shader_resource_1block(kernel_info_t & kernel, bool occupy); + void release_shader_resource_1block(kernel_info_t & kernel); + int find_available_hwtid(unsigned int cta_size); +private: + unsigned int m_occupied_n_threads; + unsigned int m_occupied_shmem; + unsigned int m_occupied_regs; + unsigned int m_occupied_ctas; + + }; class simt_core_cluster { @@ -1851,6 +1869,7 @@ public: bool icnt_injection_buffer_full(unsigned size, bool write); void icnt_inject_request_packet(class mem_fetch *mf); + // for perfect memory interface bool response_queue_full() { return ( m_response_fifo.size() >= m_config->n_simt_ejection_buffer_size ); -- cgit v1.3 From 69bb1082de9df29d1d7b40486301049767e607b0 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Thu, 13 Nov 2014 18:47:03 -0500 Subject: BUG: for concurrent kernels on same shader, should select kernel from the distributor directly --- src/gpgpu-sim/gpu-sim.cc | 4 ++++ src/gpgpu-sim/shader.cc | 36 +++++++++++++++++++++--------------- 2 files changed, 25 insertions(+), 15 deletions(-) (limited to 'src') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 8a5d581..10c4ccf 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -523,6 +523,10 @@ 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()) + 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]) ){ diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 8cba31e..cd38cb7 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -3263,23 +3263,29 @@ unsigned simt_core_cluster::issue_block2core() for( unsigned i=0; i < m_config->n_simt_cores_per_cluster; i++ ) { unsigned core = (i+m_cta_issue_next_core+1)%m_config->n_simt_cores_per_cluster; - kernel_info_t *kernel = m_core[core]->get_kernel(); - - //Jin: check if to fetch the next kernel - if( !m_gpu->kernel_more_cta_left(kernel) ) { - if(m_config->gpgpu_concurrent_kernel_sm || //concurrent kernel on sm - - //otherwise wait till current kernel finishes - (!m_config->gpgpu_concurrent_kernel_sm && - m_core[core]->get_not_completed() == 0) ) - { - kernel_info_t *k = m_gpu->select_kernel(); - if( k ) - m_core[core]->set_kernel(k); - } + kernel_info_t * kernel; + //Jin: fetch kernel according to concurrent kernel setting + if(m_config->gpgpu_concurrent_kernel_sm) {//concurrent kernel on sm + //always select latest issued kernel + kernel_info_t *k = m_gpu->select_kernel(); + kernel = k; + } + else { + //first select core kernel, if no more cta, get a new kernel + //only when core completes + kernel = m_core[core]->get_kernel(); + if( !m_gpu->kernel_more_cta_left(kernel) ) { + //wait till current kernel finishes + if(m_core[core]->get_not_completed() == 0) + { + kernel_info_t *k = m_gpu->select_kernel(); + if( k ) + m_core[core]->set_kernel(k); + kernel = k; + } + } } - kernel = m_core[core]->get_kernel(); if( m_gpu->kernel_more_cta_left(kernel) && // (m_core[core]->get_n_active_cta() < m_config->max_cta(*kernel)) ) { m_core[core]->can_issue_1block(*kernel)) { -- 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') 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') 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 9c0384c1626205b4a9ed97f998b3d4d8e7758198 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 14 Nov 2014 22:11:52 -0500 Subject: BUG: concurrent kernels on same SM may occupy warps from running CTAs --- src/gpgpu-sim/gpu-sim.cc | 27 +++++++++++++++++++++------ src/gpgpu-sim/shader.cc | 7 +++++-- src/gpgpu-sim/shader.h | 8 +++++--- 3 files changed, 31 insertions(+), 11 deletions(-) (limited to 'src') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 10c4ccf..81c9c9a 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1092,7 +1092,7 @@ bool shader_core_ctx::can_issue_1block(kernel_info_t & kernel) { return occupy_shader_resource_1block(kernel, false); } -int shader_core_ctx::find_available_hwtid(unsigned int cta_size) { +int shader_core_ctx::find_available_hwtid(unsigned int cta_size, bool occupy) { unsigned int step; for(step = 0; step < m_config->n_thread_per_shader; @@ -1101,7 +1101,7 @@ int shader_core_ctx::find_available_hwtid(unsigned int cta_size) { unsigned int hw_tid; for(hw_tid = step; hw_tid < step + cta_size; hw_tid++) { - if(m_active_threads.test(hw_tid)) + if(m_occupied_hwtid.test(hw_tid)) break; } if(hw_tid == step + cta_size) //consecutive non-active @@ -1109,8 +1109,14 @@ int shader_core_ctx::find_available_hwtid(unsigned int cta_size) { } if(step >= m_config->n_thread_per_shader) //didn't find return -1; - else + else { + if(occupy) { + for(unsigned hw_tid = step; hw_tid < step + cta_size; + hw_tid++) + m_occupied_hwtid.set(hw_tid); + } return step; + } } bool shader_core_ctx::occupy_shader_resource_1block(kernel_info_t & k, bool occupy) { @@ -1124,7 +1130,7 @@ bool shader_core_ctx::occupy_shader_resource_1block(kernel_info_t & k, bool occu if(m_occupied_n_threads + padded_cta_size > m_config->n_thread_per_shader) return false; - if(find_available_hwtid(padded_cta_size) == -1) + if(find_available_hwtid(padded_cta_size, false) == -1) return false; const struct gpgpu_ptx_sim_kernel_info *kernel_info = ptx_sim_kernel_info(kernel); @@ -1152,7 +1158,7 @@ bool shader_core_ctx::occupy_shader_resource_1block(kernel_info_t & k, bool occu return true; } -void shader_core_ctx::release_shader_resource_1block(kernel_info_t & k) { +void shader_core_ctx::release_shader_resource_1block(unsigned hw_ctaid, kernel_info_t & k) { unsigned threads_per_cta = k.threads_per_cta(); const class function_info *kernel = k.entry(); unsigned int padded_cta_size = threads_per_cta; @@ -1163,6 +1169,13 @@ void shader_core_ctx::release_shader_resource_1block(kernel_info_t & k) { assert(m_occupied_n_threads >= padded_cta_size); m_occupied_n_threads -= padded_cta_size; + int start_thread = m_occupied_cta_to_hwtid[hw_ctaid]; + + for(unsigned hwtid = start_thread; hwtid < start_thread + padded_cta_size; + hwtid++) + m_occupied_hwtid.reset(hwtid); + m_occupied_cta_to_hwtid.erase(hw_ctaid); + const struct gpgpu_ptx_sim_kernel_info *kernel_info = ptx_sim_kernel_info(kernel); assert(m_occupied_shmem >= (unsigned int)kernel_info->smem); @@ -1211,9 +1224,11 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) int padded_cta_size = cta_size; if (cta_size%m_config->warp_size) padded_cta_size = ((cta_size/m_config->warp_size)+1)*(m_config->warp_size); - unsigned int start_thread = find_available_hwtid(padded_cta_size); + unsigned int start_thread = find_available_hwtid(padded_cta_size, true); assert((int)start_thread != -1); unsigned int end_thread = start_thread + cta_size; + assert(m_occupied_cta_to_hwtid.find(free_cta_hw_id) == m_occupied_cta_to_hwtid.end()); + m_occupied_cta_to_hwtid[free_cta_hw_id]= start_thread; // unsigned start_thread = free_cta_hw_id * padded_cta_size; // unsigned end_thread = start_thread + cta_size; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 8ce2146..b9caf18 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -302,6 +302,8 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, m_occupied_shmem = 0; m_occupied_regs = 0; m_occupied_ctas = 0; + m_occupied_hwtid.reset(); + m_occupied_cta_to_hwtid.clear(); } void shader_core_ctx::reinit(unsigned start_thread, unsigned end_thread, bool reset_not_completed ) @@ -315,6 +317,8 @@ void shader_core_ctx::reinit(unsigned start_thread, unsigned end_thread, bool re m_occupied_shmem = 0; m_occupied_regs = 0; m_occupied_ctas = 0; + m_occupied_hwtid.reset(); + m_occupied_cta_to_hwtid.clear(); } for (unsigned i = start_thread; im_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; } @@ -1980,7 +1983,7 @@ void shader_core_ctx::register_cta_thread_exit( unsigned cta_num, kernel_info_t } //Jin: for concurrent kernels on sm - release_shader_resource_1block(*kernel); + release_shader_resource_1block(cta_num, *kernel); kernel->dec_running(); if( !m_gpu->kernel_more_cta_left(kernel) ) { if( !kernel->running() ) { diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 882868e..bdd8dbe 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1766,7 +1766,7 @@ public: virtual void checkExecutionStatusAndUpdate(warp_inst_t &inst, unsigned t, unsigned tid); address_type next_pc( int tid ) const; void fetch(); - void register_cta_thread_exit( unsigned cta_num, kernel_info_t * kernel ); + void register_cta_thread_exit(unsigned cta_num, kernel_info_t * kernel ); void decode(); @@ -1853,13 +1853,15 @@ public: public: bool can_issue_1block(kernel_info_t & kernel); bool occupy_shader_resource_1block(kernel_info_t & kernel, bool occupy); - void release_shader_resource_1block(kernel_info_t & kernel); - int find_available_hwtid(unsigned int cta_size); + void release_shader_resource_1block(unsigned hw_ctaid, kernel_info_t & kernel); + int find_available_hwtid(unsigned int cta_size, bool occupy); private: unsigned int m_occupied_n_threads; unsigned int m_occupied_shmem; unsigned int m_occupied_regs; unsigned int m_occupied_ctas; + std::bitset m_occupied_hwtid; + std::map m_occupied_cta_to_hwtid; }; -- 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') 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') 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 60fa05d4de0f3c926f9ab4f687b5d0748ec19285 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sun, 16 Nov 2014 01:27:47 -0500 Subject: ADD: add kernel launching latency from stream to distributor --- src/abstract_hardware_model.cc | 4 ++++ src/abstract_hardware_model.h | 1 + src/gpgpu-sim/gpu-sim.cc | 6 ++++++ src/gpgpusim_entrypoint.cc | 1 + src/stream_manager.cc | 39 +++++++++++++++++++++++++-------------- 5 files changed, 37 insertions(+), 14 deletions(-) (limited to 'src') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index be5c5b9..819ad35 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -550,6 +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 +unsigned g_kernel_launch_latency; unsigned kernel_info_t::m_next_uid = 1; @@ -569,6 +571,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * //Jin: parent and child kernel management for CDP m_parent_kernel = NULL; + //Jin: launch latency management + m_launch_latency = g_kernel_launch_latency; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index e6ef521..ab77d63 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -298,6 +298,7 @@ public: unsigned long long launch_cycle; unsigned long long start_cycle; unsigned long long end_cycle; + unsigned m_launch_latency; }; struct core_config { diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index a9da1c9..7fb9ab3 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -443,6 +443,12 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) &Trace::sampling_memory_partition, "The memory partition which is printed using MEMPART_DPRINTF. Default -1 (i.e. all)", "-1"); ptx_file_line_stats_options(opp); + + //Jin: kernel launch latency + extern unsigned g_kernel_launch_latency; + option_parser_register(opp, "-gpgpu_kernel_launch_latency", OPT_INT32, + &g_kernel_launch_latency, "Kernel launch latency in cycles. Default: 0", + "0"); } ///////////////////////////////////////////////////////////////////////////// diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index fb17eed..04845e7 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -153,6 +153,7 @@ void *gpgpu_sim_thread_concurrent(void*) printf("GPGPU-Sim: ** STOP simulation thread (no work) **\n"); fflush(stdout); } + g_the_gpu->print_stats(); if(sim_cycles) { g_the_gpu->update_stats(); print_simulation_time(); diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 3459c6d..3b6cbd5 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -160,20 +160,31 @@ bool stream_operation::do_operation( gpgpu_sim *gpu ) m_stream->record_next_done(); break; case stream_kernel_launch: - if( gpu->can_start_kernel() ) { - gpu->set_cache_config(m_kernel->name()); - 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 ); + if( m_sim_mode ) { //Functional Sim + 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(); + } + gpu->set_cache_config(m_kernel->name()); + gpu->functional_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; + else { //Performance Sim + if( gpu->can_start_kernel() && m_kernel->m_launch_latency == 0) { + 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(); + } + gpu->set_cache_config(m_kernel->name()); + gpu->launch( m_kernel ); + } + else { + if(m_kernel->m_launch_latency) + m_kernel->m_launch_latency--; + if(g_debug_execution >= 3) + printf("kernel %d: \'%s\', latency %u not ready to transfer to GPU hardware scheduler\n", + m_kernel->get_uid(), m_kernel->name().c_str(), m_kernel->m_launch_latency); + return false; + } } break; case stream_event: { @@ -255,7 +266,7 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) //Jin: should check children kernels for CDP if(kernel->is_finished()) { // std::ofstream kernel_stat("kernel_stat.txt", std::ofstream::out | std::ofstream::app); -// kernel_stat<< " kernel " << grid_uid; +// kernel_stat<< " kernel " << grid_uid << ": " << kernel->name(); // if(kernel->get_parent()) // kernel_stat << ", parent " << kernel->get_parent()->get_uid() << // ", launch " << kernel->launch_cycle; -- 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') 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') 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 9f958e424f2fc952970794efc7647ceae1674d97 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Wed, 6 Jul 2016 02:40:10 -0400 Subject: MOD: modify to new structure name gpgpu_ptx_sim_info --- src/gpgpu-sim/gpu-sim.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'src') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 7fb9ab3..02c9c09 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1148,7 +1148,7 @@ bool shader_core_ctx::occupy_shader_resource_1block(kernel_info_t & k, bool occu if(find_available_hwtid(padded_cta_size, false) == -1) return false; - const struct gpgpu_ptx_sim_kernel_info *kernel_info = ptx_sim_kernel_info(kernel); + const struct gpgpu_ptx_sim_info *kernel_info = ptx_sim_kernel_info(kernel); if(m_occupied_shmem + kernel_info->smem > m_config->gpgpu_shmem_size) return false; @@ -1191,7 +1191,7 @@ void shader_core_ctx::release_shader_resource_1block(unsigned hw_ctaid, kernel_i m_occupied_hwtid.reset(hwtid); m_occupied_cta_to_hwtid.erase(hw_ctaid); - const struct gpgpu_ptx_sim_kernel_info *kernel_info = ptx_sim_kernel_info(kernel); + const struct gpgpu_ptx_sim_info *kernel_info = ptx_sim_kernel_info(kernel); assert(m_occupied_shmem >= (unsigned int)kernel_info->smem); m_occupied_shmem -= kernel_info->smem; -- 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') 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') 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 From ad4b448e846b1cd3ac244f02fc1e21ab37eaec55 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Tue, 30 Aug 2016 20:30:17 -0700 Subject: BUG: concurrent kernel on the same SMX does not work with non-legacy local memory mapping, turn off by default --- src/gpgpu-sim/gpu-sim.cc | 114 ++++++++++++++++++++++++++++------------------- src/gpgpu-sim/shader.cc | 4 +- 2 files changed, 72 insertions(+), 46 deletions(-) (limited to 'src') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 363fe5a..58a5d16 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -369,8 +369,8 @@ void shader_core_config::reg_options(class OptionParser * opp) "gto"); option_parser_register(opp, "-gpgpu_concurrent_kernel_sm", OPT_BOOL, &gpgpu_concurrent_kernel_sm, - "Support concurrent kernels on a SM (default = enabled)", - "1"); + "Support concurrent kernels on a SM (default = disabled)", + "0"); } void gpgpu_sim_config::reg_options(option_parser_t opp) @@ -1102,13 +1102,18 @@ void shader_core_ctx::mem_instruction_stats(const warp_inst_t &inst) abort(); } } -//Jin: concurrent kernels on one SM bool shader_core_ctx::can_issue_1block(kernel_info_t & kernel) { - - if(m_config->max_cta(kernel) < 1) - return false; - return occupy_shader_resource_1block(kernel, false); + //Jin: concurrent kernels on one SM + if(m_config->gpgpu_concurrent_kernel_sm) { + if(m_config->max_cta(kernel) < 1) + return false; + + return occupy_shader_resource_1block(kernel, false); + } + else { + return (get_n_active_cta() < m_config->max_cta(kernel)); + } } int shader_core_ctx::find_available_hwtid(unsigned int cta_size, bool occupy) { @@ -1178,34 +1183,37 @@ bool shader_core_ctx::occupy_shader_resource_1block(kernel_info_t & k, bool occu } void shader_core_ctx::release_shader_resource_1block(unsigned hw_ctaid, kernel_info_t & k) { - unsigned threads_per_cta = k.threads_per_cta(); - const class function_info *kernel = k.entry(); - unsigned int padded_cta_size = threads_per_cta; - unsigned int warp_size = m_config->warp_size; - if (padded_cta_size%warp_size) - padded_cta_size = ((padded_cta_size/warp_size)+1)*(warp_size); - - assert(m_occupied_n_threads >= padded_cta_size); - m_occupied_n_threads -= padded_cta_size; - - int start_thread = m_occupied_cta_to_hwtid[hw_ctaid]; - - for(unsigned hwtid = start_thread; hwtid < start_thread + padded_cta_size; - hwtid++) - m_occupied_hwtid.reset(hwtid); - m_occupied_cta_to_hwtid.erase(hw_ctaid); - - const struct gpgpu_ptx_sim_info *kernel_info = ptx_sim_kernel_info(kernel); - - assert(m_occupied_shmem >= (unsigned int)kernel_info->smem); - m_occupied_shmem -= kernel_info->smem; - unsigned int used_regs = padded_cta_size * ((kernel_info->regs+3)&~3); - assert(m_occupied_regs >= used_regs); - m_occupied_regs -= used_regs; - - assert(m_occupied_ctas >= 1); - m_occupied_ctas--; + if(m_config->gpgpu_concurrent_kernel_sm) { + unsigned threads_per_cta = k.threads_per_cta(); + const class function_info *kernel = k.entry(); + unsigned int padded_cta_size = threads_per_cta; + unsigned int warp_size = m_config->warp_size; + if (padded_cta_size%warp_size) + padded_cta_size = ((padded_cta_size/warp_size)+1)*(warp_size); + + assert(m_occupied_n_threads >= padded_cta_size); + m_occupied_n_threads -= padded_cta_size; + + int start_thread = m_occupied_cta_to_hwtid[hw_ctaid]; + + for(unsigned hwtid = start_thread; hwtid < start_thread + padded_cta_size; + hwtid++) + m_occupied_hwtid.reset(hwtid); + m_occupied_cta_to_hwtid.erase(hw_ctaid); + + const struct gpgpu_ptx_sim_info *kernel_info = ptx_sim_kernel_info(kernel); + + assert(m_occupied_shmem >= (unsigned int)kernel_info->smem); + m_occupied_shmem -= kernel_info->smem; + + unsigned int used_regs = padded_cta_size * ((kernel_info->regs+3)&~3); + assert(m_occupied_regs >= used_regs); + m_occupied_regs -= used_regs; + + assert(m_occupied_ctas >= 1); + m_occupied_ctas--; + } } //////////////////////////////////////////////////////////////////////////////////////////////// @@ -1219,14 +1227,23 @@ void shader_core_ctx::release_shader_resource_1block(unsigned hw_ctaid, kernel_i void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) { -// set_max_cta(kernel); + + if(!m_config->gpgpu_concurrent_kernel_sm) + set_max_cta(kernel); + else + assert(occupy_shader_resource_1block(kernel, true)); + kernel.inc_running(); - assert(occupy_shader_resource_1block(kernel, true)); // find a free CTA context unsigned free_cta_hw_id=(unsigned)-1; -// for (unsigned i=0;imax_cta_per_core;i++ ) { + + unsigned max_cta_per_core; + if(!m_config->gpgpu_concurrent_kernel_sm) + max_cta_per_core = kernel_max_cta_per_shader; + else + max_cta_per_core = m_config->max_cta_per_core; + for (unsigned i=0;iwarp_size) padded_cta_size = ((cta_size/m_config->warp_size)+1)*(m_config->warp_size); - unsigned int start_thread = find_available_hwtid(padded_cta_size, true); - assert((int)start_thread != -1); - unsigned int end_thread = start_thread + cta_size; - assert(m_occupied_cta_to_hwtid.find(free_cta_hw_id) == m_occupied_cta_to_hwtid.end()); - m_occupied_cta_to_hwtid[free_cta_hw_id]= start_thread; -// unsigned start_thread = free_cta_hw_id * padded_cta_size; -// unsigned end_thread = start_thread + cta_size; + + unsigned int start_thread, end_thread; + + if(!m_config->gpgpu_concurrent_kernel_sm) { + start_thread = free_cta_hw_id * padded_cta_size; + end_thread = start_thread + cta_size; + } + else { + start_thread = find_available_hwtid(padded_cta_size, true); + assert((int)start_thread != -1); + end_thread = start_thread + cta_size; + assert(m_occupied_cta_to_hwtid.find(free_cta_hw_id) == m_occupied_cta_to_hwtid.end()); + m_occupied_cta_to_hwtid[free_cta_hw_id]= start_thread; + } // reset the microarchitecture state of the selected hardware thread and warp contexts reinit(start_thread, end_thread,false); diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 59a2d8b..d17e51d 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1978,7 +1978,9 @@ void shader_core_ctx::register_cta_thread_exit( unsigned cta_num, kernel_info_t fflush(stdout); //Shader can only be empty when no more cta are dispatched - assert(m_kernel == NULL || !m_gpu->kernel_more_cta_left(m_kernel)); + if(kernel != m_kernel) { + assert(m_kernel == NULL || !m_gpu->kernel_more_cta_left(m_kernel)); + } m_kernel = NULL; } -- cgit v1.3 From 0421e45d3f8796bc9bf1ff3b62c05173f7862463 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Mon, 5 Sep 2016 18:09:12 -0400 Subject: MOD: modify Makefile to make CUDART_VERSION available to gpu-sim.cc --- src/gpgpu-sim/Makefile | 1 + 1 file changed, 1 insertion(+) (limited to 'src') diff --git a/src/gpgpu-sim/Makefile b/src/gpgpu-sim/Makefile index bead38a..f10a8a4 100644 --- a/src/gpgpu-sim/Makefile +++ b/src/gpgpu-sim/Makefile @@ -59,6 +59,7 @@ ifneq ($(GPGPUSIM_POWER_MODEL),) endif OPTFLAGS += -g3 -fPIC +OPTFLAGS += -DCUDART_VERSION=$(CUDART_VERSION) CPP = g++ $(SNOW) OEXT = o -- cgit v1.3 From a0de8a0a75f3640e10a27cb91822a11089f22b99 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 May 2017 16:52:24 -0400 Subject: Fix next block addr to link predicate ret block to consecutive block The block containing predicate ret instruction should add the consecutive block to its successor_ids set. next_addr should be assigned with current instruction address add instruction size instead of 1. Signed-off-by: Mengchi Zhang --- src/cuda-sim/ptx_ir.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src') diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 1805ce9..8ebdcf8 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -498,7 +498,7 @@ void function_info::connect_basic_blocks( ) //iterate across m_basic_blocks of f if( pI->has_pred() ) { printf("GPGPU-Sim PTX: Warning detected predicated return/exit.\n"); // if predicated, add link to next block - unsigned next_addr = pI->get_m_instr_mem_index() + 1; + unsigned next_addr = pI->get_m_instr_mem_index() + pI->inst_size(); if( next_addr < m_instr_mem_size && m_instr_mem[next_addr] ) { basic_block_t *next_bb = m_instr_mem[next_addr]->get_bb(); (*bb_itr)->successor_ids.insert(next_bb->bb_id); -- cgit v1.3 From 196703487e33ec383dab2a0cededb2289e342083 Mon Sep 17 00:00:00 2001 From: tgrogers Date: Tue, 16 May 2017 13:37:29 -0400 Subject: Changing the version detection to be much more detailed. Now the git commit # and branch will be embedded in the built executable and print out when gpgpu-sim runs --- src/cuda-sim/Makefile | 2 +- src/cuda-sim/cuda-sim.cc | 5 ++--- version | 1 - version_detection.mk | 9 +++++++-- 4 files changed, 10 insertions(+), 7 deletions(-) (limited to 'src') diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile index f479294..5132bcc 100644 --- a/src/cuda-sim/Makefile +++ b/src/cuda-sim/Makefile @@ -46,7 +46,7 @@ OPT := -O3 -g3 -Wall -Wno-unused-function -Wno-sign-compare ifeq ($(DEBUG),1) OPT := -g3 -Wall -Wno-unused-function -Wno-sign-compare endif -OPT += -I$(CUDA_INSTALL_PATH)/include -I$(OUTPUT_DIR)/ -I. +OPT += -I$(CUDA_INSTALL_PATH)/include -I$(OUTPUT_DIR)/ -I. -I$(SIM_OBJ_FILES_DIR) OPT += -fPIC ifeq ($(TRACE),1) diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 7ec3ce9..d4ace76 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1623,14 +1623,13 @@ kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, } #include "../../version" +#include "detailed_version" void print_splash() { static int splash_printed=0; if ( !splash_printed ) { - unsigned build=0; - sscanf(g_gpgpusim_build_string, "$Change"": %u $", &build); - fprintf(stdout, "\n\n *** %s [build %u] ***\n\n\n", g_gpgpusim_version_string, build ); + fprintf(stdout, "\n\n *** %s [build %s] ***\n\n\n", g_gpgpusim_version_string, g_gpgpusim_build_string ); splash_printed=1; } } diff --git a/version b/version index e565a98..c70c6ac 100644 --- a/version +++ b/version @@ -1,2 +1 @@ const char *g_gpgpusim_version_string = "GPGPU-Sim Simulator Version 3.2.2 "; -const char *g_gpgpusim_build_string = "$Change$"; diff --git a/version_detection.mk b/version_detection.mk index 8796d5c..00a8b1f 100644 --- a/version_detection.mk +++ b/version_detection.mk @@ -30,7 +30,13 @@ ifeq ($(GPGPUSIM_ROOT),) else GPGPUSIM_VERSION=$(shell cat $(GPGPUSIM_ROOT)/version | awk '/Version/ {print $$8}' ) -GPGPUSIM_BUILD=$(shell cat $(GPGPUSIM_ROOT)/version | awk '/Change/ {print $$6}' ) + +#Detect Git branch and commit # +GIT_BRANCH := $(shell git branch | grep "\*" | sed -re 's/\*\s+(.*)/\1/') +GIT_COMMIT := $(shell git log -n 1 | head -1 | sed -re 's/commit (.*)/\1/') +GIT_FILES_CHANGED := $(shell git diff --numstat --cached && git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1/') +GPGPUSIM_BUILD := "gpgpu-sim_$(GIT_BRANCH)_$(GIT_COMMIT)_modified_$(GIT_FILES_CHANGED)" +$(shell mkdir -p $(SIM_OBJ_FILES_DIR)/libcuda && echo "const char *g_gpgpusim_build_string=\"$(GPGPUSIM_BUILD)\";" > $(SIM_OBJ_FILES_DIR)/detailed_version) endif # Detect CUDA Runtime Version @@ -42,4 +48,3 @@ CC_VERSION := $(shell gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(mat # Detect Support for C++11 (C++0x) from GCC Version GNUC_CPP0X := $(shell gcc --version | perl -ne 'if (/gcc\s+\(.*\)\s+([0-9.]+)/){ if($$1 >= 4.3) {$$n=1} else {$$n=0;} } END { print $$n; }') - -- cgit v1.3 From f17b0d38c7654bf3a83e0bd93fcf87c6ab233b6b Mon Sep 17 00:00:00 2001 From: tgrogers Date: Thu, 6 Jul 2017 00:40:01 -0400 Subject: Adding the correct dependency for the detailed_version file. In order to updatet the built number output when we run gpgpu-sim we need to recompile cuda-sim everytime the detailed_version has changed --- src/cuda-sim/Makefile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src') diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile index 5132bcc..999dad7 100644 --- a/src/cuda-sim/Makefile +++ b/src/cuda-sim/Makefile @@ -142,7 +142,7 @@ $(OUTPUT_DIR)/ptx_parser.o: $(OUTPUT_DIR)/ptx.tab.c $(OUTPUT_DIR)/ptx_parser_dec $(OUTPUT_DIR)/ptxinfo.tab.o: $(OUTPUT_DIR)/ptx.tab.c $(OUTPUT_DIR)/ptx-stats.o: $(OUTPUT_DIR)/ptx.tab.c $(OUTPUT_DIR)/ptx_sim.o: $(OUTPUT_DIR)/ptx.tab.c -$(OUTPUT_DIR)/cuda-sim.o: $(OUTPUT_DIR)/ptx.tab.c +$(OUTPUT_DIR)/cuda-sim.o: $(OUTPUT_DIR)/ptx.tab.c $(SIM_OBJ_FILES_DIR)/detailed_version $(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 -- cgit v1.3 From 28a0a22a65d9c1ebe1558c838ff92bf17a22073b Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Wed, 12 Jul 2017 15:29:16 -0400 Subject: Fixing BankGroup Indexing Bug --- src/gpgpu-sim/gpu-sim.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src') diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index a2d1b9b..7d92c66 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -195,7 +195,7 @@ struct memory_config { for (i=0; nbkt>0; i++) { nbkt = nbkt>>1; } - bk_tag_length = i; + bk_tag_length = i-1; assert(nbkgrp>0 && "Number of bank groups cannot be zero"); tRCDWR = tRCD-(WL+1); tRTW = (CL+(BL/data_command_freq_ratio)+2-WL); -- cgit v1.3 From f5abe6f9573bfc1ccbdf52b4dc960599c974884f Mon Sep 17 00:00:00 2001 From: tgrogers Date: Wed, 19 Jul 2017 13:46:35 -0400 Subject: Updaing the interconnect simulator to properly check for dependencies --- src/intersim2/Makefile | 10 ++++++++++ 1 file changed, 10 insertions(+) (limited to 'src') diff --git a/src/intersim2/Makefile b/src/intersim2/Makefile index ef948d6..bd42000 100644 --- a/src/intersim2/Makefile +++ b/src/intersim2/Makefile @@ -125,6 +125,14 @@ endif # rules to compile simulator +$(OBJDIR)/Makefile.makedepend: depend + +ALL_SRCS = $(CPP_SRCS) +ALL_SRCS += $(shell ls **/*.cpp) + +depend: + touch $(OBJDIR)/Makefile.makedepend + makedepend -f$(OBJDIR)/Makefile.makedepend -I$(INCPATH) -p$(OBJDIR)/ $(ALL_SRCS) 2> /dev/null ${LEX_OBJS}: $(OBJDIR)/lex.yy.c $(OBJDIR)/y.tab.h $(CC) $(CPPFLAGS) -c $< -o $@ @@ -173,3 +181,5 @@ $(OBJDIR)/y.tab.c $(OBJDIR)/y.tab.h: config.y $(OBJDIR)/lex.yy.c: config.l $(LEX) -o$@ $< + +include $(OBJDIR)/Makefile.makedepend -- cgit v1.3