From 692c0e055ab53918616694d8c7877394bdd68993 Mon Sep 17 00:00:00 2001 From: sriraj Date: Sat, 22 Jun 2019 19:50:53 -0400 Subject: Fix SM6_TITANX config to use pascal file --- configs/tested-cfgs/SM6_TITANX/gpgpusim.config | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config index 2fe898a..a94b66a 100644 --- a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config +++ b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config @@ -135,7 +135,7 @@ # interconnection -network_mode 1 --inter_config_file config_fermi_islip.icnt +-inter_config_file config_pascal_islip.icnt # memory partition latency config -rop_latency 120 -- cgit v1.3 From 779a636d32914e30ca039a043d0f3abc88d2f1ea Mon Sep 17 00:00:00 2001 From: chris Date: Thu, 4 Jul 2019 00:18:58 +0000 Subject: added temp. solution to bypass function pointers in order for pytorch to run --- src/cuda-sim/ptx.y | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index 45392fb..260564f 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -287,6 +287,7 @@ ptr_align_spec: ALIGN_DIRECTIVE INT_OPERAND statement_block: LEFT_BRACE statement_list RIGHT_BRACE statement_list: directive_statement { add_directive(); } + | statement_list prototype_block {printf("Prototype statement detected. WARNING: this is not supported yet on GPGPU-SIM\n"); } | instruction_statement { add_instruction(); } | statement_list directive_statement { add_directive(); } | statement_list instruction_statement { add_instruction(); } @@ -403,6 +404,23 @@ initializer_list: LEFT_BRACE literal_list RIGHT_BRACE { add_array_initializer(); literal_list: literal_operand | literal_list COMMA literal_operand; +// TODO: This is currently hardcoded to handle and ignore one specific case +// that all prototype statements follow in the PTX from Pytorch. As a +// workaround, this parses and ignores both the prototype declaration +// and calling of the prototype (which conveniently comes right after the +// declaration for all cases.) This should be changed to handle both +// declaring the prototype, and actually calling it. +prototype_block: prototype_decl prototype_call + +prototype_decl: IDENTIFIER COLON CALLPROTOTYPE_DIRECTIVE LEFT_PAREN prototype_param RIGHT_PAREN IDENTIFIER LEFT_PAREN prototype_param RIGHT_PAREN SEMI_COLON + +prototype_call: OPCODE LEFT_PAREN IDENTIFIER RIGHT_PAREN COMMA operand COMMA LEFT_PAREN IDENTIFIER RIGHT_PAREN COMMA IDENTIFIER SEMI_COLON + | OPCODE IDENTIFIER COMMA LEFT_PAREN IDENTIFIER RIGHT_PAREN COMMA IDENTIFIER SEMI_COLON + +prototype_param: /* empty */ + | PARAM_DIRECTIVE B64_TYPE IDENTIFIER + | PARAM_DIRECTIVE B32_TYPE IDENTIFIER + instruction_statement: instruction SEMI_COLON | IDENTIFIER COLON { add_label($1); } | pred_spec instruction SEMI_COLON; -- cgit v1.3 From d212d7e5fdcc9f8e10779d5cfb398a451f8ad033 Mon Sep 17 00:00:00 2001 From: boyealKim Date: Sun, 21 Jul 2019 02:28:59 +0900 Subject: take account of shfl latency --- configs/tested-cfgs/SM2_GTX480/gpgpusim.config | 6 +++--- configs/tested-cfgs/SM6_TITANX/gpgpusim.config | 6 +++--- configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 6 +++--- src/cuda-sim/cuda-sim.cc | 29 +++++++++++++------------- src/gpgpu-sim/shader.cc | 9 ++++---- 5 files changed, 29 insertions(+), 27 deletions(-) diff --git a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config index cf3627b..7ff836a 100644 --- a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config +++ b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config @@ -46,9 +46,9 @@ -gpgpu_num_dp_units 0 # Instruction latencies and initiation intervals -# "ADD,MAX,MUL,MAD,DIV" --ptx_opcode_latency_int 4,13,4,5,145 --ptx_opcode_initiation_int 1,2,2,1,8 +# "ADD,MAX,MUL,MAD,DIV,SHFL" +-ptx_opcode_latency_int 4,13,4,5,145,32 +-ptx_opcode_initiation_int 1,2,2,1,8,4 -ptx_opcode_latency_fp 4,13,4,5,39 -ptx_opcode_initiation_fp 1,2,1,1,4 -ptx_opcode_latency_dp 8,19,8,8,330 diff --git a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config index 2fe898a..da1af48 100644 --- a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config +++ b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config @@ -56,11 +56,11 @@ # Instruction latencies and initiation intervals -# "ADD,MAX,MUL,MAD,DIV" +# "ADD,MAX,MUL,MAD,DIV,SHFL" # All Div operations are executed on SFU unit # Throughput (initiation latency) are adopted from CUDA SDK document V8, section 5.4.1, Table 2 --ptx_opcode_latency_int 4,13,4,5,145 --ptx_opcode_initiation_int 1,1,1,1,4 +-ptx_opcode_latency_int 4,13,4,5,145,32 +-ptx_opcode_initiation_int 1,1,1,1,4,4 -ptx_opcode_latency_fp 4,13,4,5,39 -ptx_opcode_initiation_fp 1,2,1,1,4 -ptx_opcode_latency_dp 8,19,8,8,330 diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config index ebd442f..f2301e8 100644 --- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config +++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config @@ -66,10 +66,10 @@ # Instruction latencies and initiation intervals # "ADD,MAX,MUL,MAD,DIV" # All Div operations are executed on SFU unit -# Throughput (initiation latency) are adopted from +# Throughput (initiation latency except shfl) are adopted from # http://on-demand.gputechconf.com/gtc/2018/presentation/s8122-dissecting-the-volta-gpu-architecture-through-microbenchmarking.pdf --ptx_opcode_latency_int 4,13,4,5,145 --ptx_opcode_initiation_int 2,2,2,2,8 +-ptx_opcode_latency_int 4,13,4,5,145,32 +-ptx_opcode_initiation_int 2,2,2,2,8,4 -ptx_opcode_latency_fp 4,13,4,5,39 -ptx_opcode_initiation_fp 2,2,2,2,4 -ptx_opcode_latency_dp 8,19,8,8,330 diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index f7bb9cc..7376cc6 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -71,9 +71,9 @@ 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, - "Opcode latencies for integers " - "Default 1,1,19,25,145", - "1,1,19,25,145"); + "Opcode latencies for integers " + "Default 1,1,19,25,145,32", + "1,1,19,25,145,32"); option_parser_register(opp, "-ptx_opcode_latency_fp", OPT_CSTR, &opcode_latency_fp, "Opcode latencies for single precision floating points " "Default 1,1,1,1,30", @@ -91,9 +91,9 @@ void ptx_opcocde_latency_options (option_parser_t opp) { "Default 64", "64"); option_parser_register(opp, "-ptx_opcode_initiation_int", OPT_CSTR, &opcode_initiation_int, - "Opcode initiation intervals for integers " - "Default 1,1,4,4,32", - "1,1,4,4,32"); + "Opcode initiation intervals for integers " + "Default 1,1,4,4,32,4", + "1,1,4,4,32,4"); option_parser_register(opp, "-ptx_opcode_initiation_fp", OPT_CSTR, &opcode_initiation_fp, "Opcode initiation intervals for single precision floating points " "Default 1,1,1,1,5", @@ -648,12 +648,12 @@ void ptx_instruction::set_bar_type() void ptx_instruction::set_opcode_and_latency() { - unsigned int_latency[5]; + unsigned int_latency[6]; unsigned fp_latency[5]; unsigned dp_latency[5]; unsigned sfu_latency; unsigned tensor_latency; - unsigned int_init[5]; + unsigned int_init[6]; unsigned fp_init[5]; unsigned dp_init[5]; unsigned sfu_init; @@ -664,10 +664,11 @@ void ptx_instruction::set_opcode_and_latency() * [2] MUL * [3] MAD * [4] DIV + * [5] SHFL */ - sscanf(opcode_latency_int, "%u,%u,%u,%u,%u", + sscanf(opcode_latency_int, "%u,%u,%u,%u,%u,%u", &int_latency[0],&int_latency[1],&int_latency[2], - &int_latency[3],&int_latency[4]); + &int_latency[3],&int_latency[4],&int_latency[5]); sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u", &fp_latency[0],&fp_latency[1],&fp_latency[2], &fp_latency[3],&fp_latency[4]); @@ -678,9 +679,9 @@ void ptx_instruction::set_opcode_and_latency() &sfu_latency); sscanf(opcode_latency_tensor, "%u", &tensor_latency); - sscanf(opcode_initiation_int, "%u,%u,%u,%u,%u", + sscanf(opcode_initiation_int, "%u,%u,%u,%u,%u,%u", &int_init[0],&int_init[1],&int_init[2], - &int_init[3],&int_init[4]); + &int_init[3],&int_init[4],&int_init[5]); sscanf(opcode_initiation_fp, "%u,%u,%u,%u,%u", &fp_init[0],&fp_init[1],&fp_init[2], &fp_init[3],&fp_init[4]); @@ -873,8 +874,8 @@ void ptx_instruction::set_opcode_and_latency() op=TENSOR_CORE_OP; break; case SHFL_OP: - latency = 32; - initiation_interval = 4; + latency = int_latency[5]; + initiation_interval = int_init[5]; break; default: break; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 007ad42..96ba385 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -3002,7 +3002,7 @@ void shader_core_config::set_pipeline_latency() { //calculate the max latency based on the input - unsigned int_latency[5]; + unsigned int_latency[6]; unsigned fp_latency[5]; unsigned dp_latency[5]; unsigned sfu_latency; @@ -3014,10 +3014,11 @@ void shader_core_config::set_pipeline_latency() { * [2] MUL * [3] MAD * [4] DIV + * [5] SHFL */ - sscanf(opcode_latency_int, "%u,%u,%u,%u,%u", + sscanf(opcode_latency_int, "%u,%u,%u,%u,%u,%u", &int_latency[0],&int_latency[1],&int_latency[2], - &int_latency[3],&int_latency[4]); + &int_latency[3],&int_latency[4],&int_latency[5]); sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u", &fp_latency[0],&fp_latency[1],&fp_latency[2], &fp_latency[3],&fp_latency[4]); @@ -3034,7 +3035,7 @@ void shader_core_config::set_pipeline_latency() { max_sfu_latency = std::max(dp_latency[4],sfu_latency); //assume that the max operation has the max latency max_sp_latency = fp_latency[1]; - max_int_latency = int_latency[1]; + max_int_latency = std::max(int_latency[1],int_latency[5]); max_dp_latency = dp_latency[1]; max_tensor_core_latency = tensor_latency; -- cgit v1.3 From f829a12add380c78e50ecc86357d4a1b513f21f0 Mon Sep 17 00:00:00 2001 From: boyealkim Date: Thu, 25 Jul 2019 20:42:13 +0900 Subject: fix stream_merge front algo --- src/stream_manager.cc | 50 +++++++++++++++++++++++++++++++------------------- src/stream_manager.h | 12 ++++++++++-- 2 files changed, 41 insertions(+), 21 deletions(-) diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 6cd62a2..6e02dcd 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -194,14 +194,17 @@ bool stream_operation::do_operation( gpgpu_sim *gpu ) m_stream->record_next_done(); } break; - case stream_wait_event: { + case stream_wait_event: //only allows next op to go if event is done //otherwise stays in the stream queue printf("stream wait event processing...\n"); - if(m_event->done()) + if(m_event->done()){ printf("stream wait event done\n"); m_stream->record_next_done(); } + else{ + return false; + } break; default: abort(); @@ -232,6 +235,7 @@ stream_manager::stream_manager( gpgpu_sim *gpu, bool cuda_launch_blocking ) m_service_stream_zero = false; m_cuda_launch_blocking = cuda_launch_blocking; pthread_mutex_init(&m_lock,NULL); + m_last_stream = m_streams.begin(); } bool stream_manager::operation( bool * sim) @@ -330,22 +334,31 @@ stream_operation stream_manager::front() m_service_stream_zero = false; } } - if(!m_service_stream_zero) { - std::list::iterator s; - for( s=m_streams.begin(); s != m_streams.end(); s++) { - CUstream_st *stream = *s; - if( !stream->busy() && !stream->empty() ) { - result = stream->next(); - if( result.is_kernel() ) { - unsigned grid_id = result.get_kernel()->get_uid(); - m_grid_id_to_stream[grid_id] = stream; - } - break; - } - } - } + std::list::iterator s = m_last_stream; + if(m_last_stream == m_streams.end()){ + s = m_streams.begin(); + } + else{ + s++; + } + for(size_t ii = 0 ; ii < m_streams.size(); ii++, s++) { + if(s == m_streams.end()){ + s = m_streams.begin(); + } + m_last_stream = s; + CUstream_st *stream = *s; + if( !stream->busy() && !stream->empty() ) { + result = stream->next(); + if( result.is_kernel() ) { + unsigned grid_id = result.get_kernel()->get_uid(); + m_grid_id_to_stream[grid_id] = stream; + } + break; + } + } + } return result; } @@ -370,7 +383,8 @@ void stream_manager::destroy_stream( CUstream_st *stream ) break; } } - delete stream; + delete stream; + m_last_stream = m_streams.begin(); pthread_mutex_unlock(&m_lock); } @@ -384,7 +398,6 @@ bool stream_manager::concurrent_streams_empty() for( s=m_streams.begin(); s!=m_streams.end();++s ) { struct CUstream_st *stream = *s; if( !stream->empty() ) { - //stream->print(stdout); result = false; break; } @@ -437,7 +450,6 @@ void stream_manager::print_impl( FILE *fp) void stream_manager::push( stream_operation op ) { struct CUstream_st *stream = op.get_stream(); - // block if stream 0 (or concurrency disabled) and pending concurrent operations exist bool block= !stream || m_cuda_launch_blocking; while(block) { diff --git a/src/stream_manager.h b/src/stream_manager.h index 91d1b36..ae1af9f 100644 --- a/src/stream_manager.h +++ b/src/stream_manager.h @@ -163,7 +163,6 @@ public: void print( FILE *fp ) const; struct CUstream_st *get_stream() { return m_stream; } void set_stream( CUstream_st *stream ) { m_stream = stream; } - private: struct CUstream_st *m_stream; @@ -179,9 +178,9 @@ private: const char *m_symbol; size_t m_offset; + struct CUevent_st *m_event; bool m_sim_mode; kernel_info_t *m_kernel; - struct CUevent_st *m_event; }; struct CUevent_st { @@ -193,6 +192,7 @@ public: m_updates = 0; m_wallclock = 0; m_gpu_tot_sim_cycle = 0; + m_issued = 0; m_done = false; } void update( double cycle, time_t clk ) @@ -207,11 +207,18 @@ public: unsigned num_updates() const { return m_updates; } bool done() const { return m_done; } time_t clock() const { return m_wallclock; } + void issue(){ + m_issued++; + } + unsigned int num_issued() const{ + return m_issued; + } private: int m_uid; bool m_blocking; bool m_done; int m_updates; + unsigned int m_issued; time_t m_wallclock; double m_gpu_tot_sim_cycle; @@ -268,6 +275,7 @@ private: CUstream_st m_stream_zero; bool m_service_stream_zero; pthread_mutex_t m_lock; + std::list::iterator m_last_stream; }; #endif -- cgit v1.3 From 8e9f0375a6a8bed217b0634a08f7f1f0cd149744 Mon Sep 17 00:00:00 2001 From: boyealkim Date: Thu, 25 Jul 2019 20:31:21 +0900 Subject: fix stream fix bug and fix cudaStreamWaitEvent impl --- libcuda/cuda_runtime_api.cc | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 44f0f4e..7005eef 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1696,6 +1696,7 @@ __host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t s if( !e ) return g_last_cudaError = cudaErrorUnknown; struct CUstream_st *s = (struct CUstream_st *)stream; stream_operation op(e,s); + e->issue(); g_stream_manager->push(op); return g_last_cudaError = cudaSuccess; } @@ -1708,9 +1709,13 @@ __host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEven //reference: https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDART__STREAM_gfe68d207dc965685d92d3f03d77b0876.html CUevent_st *e = get_event(event); if( !e ){ - printf("GPGPU-Sim API: Warning: cudaEventRecord has not been called on event before calling cudaStreamWaitEvent.\nNothing to be done.\n"); - return g_last_cudaError = cudaSuccess; + printf("GPGPU-Sim API: Error at cudaStreamWaitEvent. Event is not created .\n"); + return g_last_cudaError = cudaErrorInvalidResourceHandle; } + else if(e->num_issued() == 0){ + printf("GPGPU-Sim API: Warning: cudaEventRecord has not been called on event before calling cudaStreamWaitEvent.\nNothing to be done.\n"); + return g_last_cudaError = cudaSuccess; + } if (!stream){ g_stream_manager->pushCudaStreamWaitEventToAllStreams(e, flags); } else { -- cgit v1.3 From b13557baa4ea6d1084bdcc9b28b9eb660c394427 Mon Sep 17 00:00:00 2001 From: boyealkim Date: Fri, 26 Jul 2019 11:55:45 +0900 Subject: fix indentation --- libcuda/cuda_runtime_api.cc | 10 ++++---- src/stream_manager.cc | 58 ++++++++++++++++++++++----------------------- src/stream_manager.h | 17 ++++++------- 3 files changed, 40 insertions(+), 45 deletions(-) diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 7005eef..6a7826e 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1709,13 +1709,13 @@ __host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEven //reference: https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDART__STREAM_gfe68d207dc965685d92d3f03d77b0876.html CUevent_st *e = get_event(event); if( !e ){ - printf("GPGPU-Sim API: Error at cudaStreamWaitEvent. Event is not created .\n"); - return g_last_cudaError = cudaErrorInvalidResourceHandle; + printf("GPGPU-Sim API: Error at cudaStreamWaitEvent. Event is not created .\n"); + return g_last_cudaError = cudaErrorInvalidResourceHandle; } - else if(e->num_issued() == 0){ - printf("GPGPU-Sim API: Warning: cudaEventRecord has not been called on event before calling cudaStreamWaitEvent.\nNothing to be done.\n"); + else if(e->num_issued() == 0){ + printf("GPGPU-Sim API: Warning: cudaEventRecord has not been called on event before calling cudaStreamWaitEvent.\nNothing to be done.\n"); return g_last_cudaError = cudaSuccess; - } + } if (!stream){ g_stream_manager->pushCudaStreamWaitEventToAllStreams(e, flags); } else { diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 6e02dcd..d0eecdd 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -198,13 +198,13 @@ bool stream_operation::do_operation( gpgpu_sim *gpu ) //only allows next op to go if event is done //otherwise stays in the stream queue printf("stream wait event processing...\n"); - if(m_event->done()){ + if(m_event->done()){ printf("stream wait event done\n"); m_stream->record_next_done(); } - else{ - return false; - } + else{ + return false; + } break; default: abort(); @@ -235,7 +235,7 @@ stream_manager::stream_manager( gpgpu_sim *gpu, bool cuda_launch_blocking ) m_service_stream_zero = false; m_cuda_launch_blocking = cuda_launch_blocking; pthread_mutex_init(&m_lock,NULL); - m_last_stream = m_streams.begin(); + m_last_stream = m_streams.begin(); } bool stream_manager::operation( bool * sim) @@ -336,29 +336,25 @@ stream_operation stream_manager::front() } if(!m_service_stream_zero) { - std::list::iterator s = m_last_stream; - if(m_last_stream == m_streams.end()){ - s = m_streams.begin(); - } - else{ - s++; - } - for(size_t ii = 0 ; ii < m_streams.size(); ii++, s++) { - if(s == m_streams.end()){ - s = m_streams.begin(); - } - m_last_stream = s; - CUstream_st *stream = *s; - if( !stream->busy() && !stream->empty() ) { - result = stream->next(); - if( result.is_kernel() ) { - unsigned grid_id = result.get_kernel()->get_uid(); - m_grid_id_to_stream[grid_id] = stream; - } - break; - } - } - } + std::list::iterator s = m_last_stream; + if(m_last_stream == m_streams.end()){ s = m_streams.begin(); } + else{ s++; } + for(size_t ii = 0 ; ii < m_streams.size(); ii++, s++) { + if(s == m_streams.end()){ + s = m_streams.begin(); + } + m_last_stream = s; + CUstream_st *stream = *s; + if( !stream->busy() && !stream->empty() ) { + result = stream->next(); + if( result.is_kernel() ) { + unsigned grid_id = result.get_kernel()->get_uid(); + m_grid_id_to_stream[grid_id] = stream; + } + break; + } + } + } return result; } @@ -383,8 +379,8 @@ void stream_manager::destroy_stream( CUstream_st *stream ) break; } } - delete stream; - m_last_stream = m_streams.begin(); + delete stream; + m_last_stream = m_streams.begin(); pthread_mutex_unlock(&m_lock); } @@ -398,6 +394,7 @@ bool stream_manager::concurrent_streams_empty() for( s=m_streams.begin(); s!=m_streams.end();++s ) { struct CUstream_st *stream = *s; if( !stream->empty() ) { + //stream->print(stdout); result = false; break; } @@ -450,6 +447,7 @@ void stream_manager::print_impl( FILE *fp) void stream_manager::push( stream_operation op ) { struct CUstream_st *stream = op.get_stream(); + // block if stream 0 (or concurrency disabled) and pending concurrent operations exist bool block= !stream || m_cuda_launch_blocking; while(block) { diff --git a/src/stream_manager.h b/src/stream_manager.h index ae1af9f..eb4c6ff 100644 --- a/src/stream_manager.h +++ b/src/stream_manager.h @@ -163,6 +163,7 @@ public: void print( FILE *fp ) const; struct CUstream_st *get_stream() { return m_stream; } void set_stream( CUstream_st *stream ) { m_stream = stream; } + private: struct CUstream_st *m_stream; @@ -178,9 +179,9 @@ private: const char *m_symbol; size_t m_offset; - struct CUevent_st *m_event; bool m_sim_mode; kernel_info_t *m_kernel; + struct CUevent_st *m_event; }; struct CUevent_st { @@ -192,7 +193,7 @@ public: m_updates = 0; m_wallclock = 0; m_gpu_tot_sim_cycle = 0; - m_issued = 0; + m_issued = 0; m_done = false; } void update( double cycle, time_t clk ) @@ -207,18 +208,14 @@ public: unsigned num_updates() const { return m_updates; } bool done() const { return m_done; } time_t clock() const { return m_wallclock; } - void issue(){ - m_issued++; - } - unsigned int num_issued() const{ - return m_issued; - } + void issue(){ m_issued++; } + unsigned int num_issued() const{ return m_issued; } private: int m_uid; bool m_blocking; bool m_done; int m_updates; - unsigned int m_issued; + unsigned int m_issued; time_t m_wallclock; double m_gpu_tot_sim_cycle; @@ -275,7 +272,7 @@ private: CUstream_st m_stream_zero; bool m_service_stream_zero; pthread_mutex_t m_lock; - std::list::iterator m_last_stream; + std::list::iterator m_last_stream; }; #endif -- cgit v1.3 From f3ec23390a0798eab1426adf962487680ea89e93 Mon Sep 17 00:00:00 2001 From: boyealkim Date: Sat, 27 Jul 2019 03:54:44 +0900 Subject: fix event done impl --- src/stream_manager.cc | 2 +- src/stream_manager.h | 80 +++++++++++++++++++++++++-------------------------- 2 files changed, 41 insertions(+), 41 deletions(-) diff --git a/src/stream_manager.cc b/src/stream_manager.cc index d0eecdd..0d67e10 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -198,7 +198,7 @@ bool stream_operation::do_operation( gpgpu_sim *gpu ) //only allows next op to go if event is done //otherwise stays in the stream queue printf("stream wait event processing...\n"); - if(m_event->done()){ + if(m_event->num_updates()>=m_cnt){ printf("stream wait event done\n"); m_stream->record_next_done(); } diff --git a/src/stream_manager.h b/src/stream_manager.h index eb4c6ff..e73d71a 100644 --- a/src/stream_manager.h +++ b/src/stream_manager.h @@ -43,6 +43,45 @@ // unsigned m_pending_streams; //}; +struct CUevent_st { +public: + CUevent_st( bool blocking ) + { + m_uid = ++m_next_event_uid; + m_blocking = blocking; + m_updates = 0; + m_wallclock = 0; + m_gpu_tot_sim_cycle = 0; + m_issued = 0; + m_done = false; + } + void update( double cycle, time_t clk ) + { + m_updates++; + m_wallclock=clk; + m_gpu_tot_sim_cycle=cycle; + m_done = true; + } + //void set_done() { assert(!m_done); m_done=true; } + int get_uid() const { return m_uid; } + unsigned num_updates() const { return m_updates; } + bool done() const { return m_updates==m_issued; } + time_t clock() const { return m_wallclock; } + void issue(){ m_issued++; } + unsigned int num_issued() const{ return m_issued; } +private: + int m_uid; + bool m_blocking; + bool m_done; + int m_updates; + unsigned int m_issued; + time_t m_wallclock; + double m_gpu_tot_sim_cycle; + + static int m_next_event_uid; +}; + + enum stream_operation_type { stream_no_op, stream_memcpy_host_to_device, @@ -107,6 +146,7 @@ public: m_kernel=NULL; m_type=stream_wait_event; m_event=e; + m_cnt = m_event->num_issued(); m_stream=stream; m_done=false; } @@ -163,7 +203,6 @@ public: void print( FILE *fp ) const; struct CUstream_st *get_stream() { return m_stream; } void set_stream( CUstream_st *stream ) { m_stream = stream; } - private: struct CUstream_st *m_stream; @@ -183,45 +222,6 @@ private: kernel_info_t *m_kernel; struct CUevent_st *m_event; }; - -struct CUevent_st { -public: - CUevent_st( bool blocking ) - { - m_uid = ++m_next_event_uid; - m_blocking = blocking; - m_updates = 0; - m_wallclock = 0; - m_gpu_tot_sim_cycle = 0; - m_issued = 0; - m_done = false; - } - void update( double cycle, time_t clk ) - { - m_updates++; - m_wallclock=clk; - m_gpu_tot_sim_cycle=cycle; - m_done = true; - } - //void set_done() { assert(!m_done); m_done=true; } - int get_uid() const { return m_uid; } - unsigned num_updates() const { return m_updates; } - bool done() const { return m_done; } - time_t clock() const { return m_wallclock; } - void issue(){ m_issued++; } - unsigned int num_issued() const{ return m_issued; } -private: - int m_uid; - bool m_blocking; - bool m_done; - int m_updates; - unsigned int m_issued; - time_t m_wallclock; - double m_gpu_tot_sim_cycle; - - static int m_next_event_uid; -}; - struct CUstream_st { public: CUstream_st(); -- cgit v1.3 From 690774e8f70624ef759f0f76f88d804f89db8f46 Mon Sep 17 00:00:00 2001 From: Roy Spliet Date: Thu, 25 Jul 2019 16:43:06 +0100 Subject: libopencl: Stub clGetProgramBuildInfo() Assumes everything is ok, even if it isn't. Allows for some Rodinia benchmarks to fail slightly later. v2: Use available case macros. Signed-off-by: Roy Spliet --- libopencl/opencl_runtime_api.cc | 29 +++++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 97a54d8..a21e11d 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -1255,6 +1255,35 @@ clGetProgramInfo(cl_program program, return CL_SUCCESS; } +extern CL_API_ENTRY cl_int CL_API_CALL +clGetProgramBuildInfo (cl_program program, + cl_device_id device, + cl_program_build_info param_name, + size_t param_value_size, + void * param_value, + size_t * param_value_size_ret) CL_API_SUFFIX__VERSION_1_0 +{ + char *buf = (char*)param_value; + + switch( param_name ) { + case CL_PROGRAM_BUILD_STATUS: + CL_CASE( cl_build_status, CL_BUILD_SUCCESS ); + break; + case CL_PROGRAM_BUILD_OPTIONS: + case CL_PROGRAM_BUILD_LOG: + CL_STRING_CASE( "" ); + break; + case CL_PROGRAM_BINARY_TYPE: + CL_CASE( cl_program_binary_type, CL_PROGRAM_BINARY_TYPE_EXECUTABLE ); + break; + default: + return CL_INVALID_VALUE; + break; + } + + return CL_SUCCESS; +} + extern CL_API_ENTRY cl_int CL_API_CALL clEnqueueCopyBuffer(cl_command_queue command_queue, cl_mem src_buffer, -- cgit v1.3 From 6cdb7b18f70f8a1da200020f1b72f56fc33864d8 Mon Sep 17 00:00:00 2001 From: Roy Spliet Date: Thu, 25 Jul 2019 18:20:56 +0100 Subject: libopencl: Perform PDOM analysis Signed-off-by: Roy Spliet --- libopencl/opencl_runtime_api.cc | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index a21e11d..23e52f6 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -950,6 +950,17 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); } kernel_info_t *grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); + + //do dynamic PDOM analysis for performance simulation scenario + std::string kname = grid->name(); + function_info *kernel_func_info = grid->entry(); + if (kernel_func_info->is_pdom_set()) { + printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kname.c_str() ); + } else { + printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", kname.c_str() ); + kernel_func_info->do_pdom(); + kernel_func_info->set_pdom(); + } if ( g_ptx_sim_mode ) gpgpu_opencl_ptx_sim_main_func( grid ); else -- cgit v1.3 From 55a05c11b0af3159d81bce823c1fda4fb3cd04e3 Mon Sep 17 00:00:00 2001 From: Roy Spliet Date: Thu, 25 Jul 2019 18:34:44 +0100 Subject: libopencl: Calculate a valid offset in bind_args() v2: Adhere to alignment requirements. v3: Small style fix. Signed-off-by: Roy Spliet --- libopencl/opencl_runtime_api.cc | 14 +++++++++++++- 1 file changed, 13 insertions(+), 1 deletion(-) diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 23e52f6..752bfdf 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -263,15 +263,27 @@ void _cl_kernel::SetKernelArg( cl_int _cl_kernel::bind_args( gpgpu_ptx_sim_arg_list_t &arg_list ) { + size_t offset = 0; + assert( arg_list.empty() ); unsigned k=0; std::map::iterator i; for( i = m_args.begin(); i!=m_args.end(); i++ ) { if( i->first != k ) return CL_INVALID_KERNEL_ARGS; + arg_info arg = i->second; - gpgpu_ptx_sim_arg param( arg.m_arg_value, arg.m_arg_size, 0); + const symbol *sym = m_kernel_impl->get_arg(i->first); + const type_info_key &t = sym->type()->get_key(); + + int align = (t.get_alignment_spec() == -1) ? arg.m_arg_size : t.get_alignment_spec(); + if( offset % align ) + offset += (align - (offset % align)); + + gpgpu_ptx_sim_arg param( arg.m_arg_value, arg.m_arg_size, offset ); arg_list.push_front( param ); + + offset += arg.m_arg_size; k++; } return CL_SUCCESS; -- cgit v1.3 From 555cce395bb8457a209b26f6f7ff62cc09839fdf Mon Sep 17 00:00:00 2001 From: bftf Date: Mon, 11 Nov 2019 12:29:15 -0800 Subject: Enabled CUDA 10. Added implementations for __cudaPushCallConfiguration and __cudaPopCallConfiguration, neither of which are documented in the CUDA API but they preceed any cudaLaunchKernel() call and push the call configuration on a stack which is then popped when cudaLaunchKernel() is called. Currently the max stack size is 1, an assert will be triggered when this is violated. --- CHANGES | 1 + Makefile | 1 + libcuda/cuda_api.h | 3 ++- libcuda/cuda_runtime_api.cc | 54 +++++++++++++++++++++++++++++++++++++++++++-- linux-so-version.txt | 2 ++ setup_environment | 5 +++-- 6 files changed, 61 insertions(+), 5 deletions(-) diff --git a/CHANGES b/CHANGES index 2a4222b..31ae672 100644 --- a/CHANGES +++ b/CHANGES @@ -37,6 +37,7 @@ Version 4.0.0 (development branch) versus 3.2.3 3- Addig new system stats: gpu occupancy, L2BW, etc -Library: 1 Enabled CUTLASS Library on GPGPU-Sim +2 Enabled CUDA 10 -Regression: 1- Added TensorCore Regression Kernel -Configs: diff --git a/Makefile b/Makefile index 3db8ce8..8df3b57 100644 --- a/Makefile +++ b/Makefile @@ -164,6 +164,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.0; fi if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.1; fi if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.2 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.2; fi + if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.0; fi $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\ diff --git a/libcuda/cuda_api.h b/libcuda/cuda_api.h index 3808e8a..b304a59 100644 --- a/libcuda/cuda_api.h +++ b/libcuda/cuda_api.h @@ -234,10 +234,11 @@ typedef struct CUgraphicsResource_st *CUgraphicsResource; /**< CUDA graphics int typedef unsigned long long CUtexObject; /**< An opaque value that represents a CUDA texture object */ typedef unsigned long long CUsurfObject; /**< An opaque value that represents a CUDA surface object */ +#if CUDA_VERSION < 1000 typedef struct CUuuid_st { /**< CUDA definition of UUID */ char bytes[16]; } CUuuid; - +#endif #if __CUDA_API_VERSION >= 4010 diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 44f0f4e..fd092e5 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1448,6 +1448,50 @@ __host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, si return g_last_cudaError = cudaSuccess; } + +#if CUDA_VERSION >= 1000 +/* +* CUDA 10 requires a new CUDA kernel launch sequence +* A call to __cudaPushCallConfiguration() preceeds any call to cudaLaunchKernel() +* __cudaPushCallConfiguration is undocumented in the API but it simply sets up a buffer with the arguments which is accessed in cudaLaunchKernel() +* __cudaPopCallConfiguration is undocumented in the API but it simply pops the configuration set in cudaLaunchKernel() +* +* pushing more than 1 configuration without popping is currently not implemented in GPGPU-Sim and will result in an assert error +*/ +namespace g_cudaPushArgsBuffer +{ + bool g_is_initialized = false; + dim3 g_gridDim; + dim3 g_blockDim; + size_t g_sharedMem; + cudaStream_t g_stream; +} + +__host__ cudaError_t CUDARTAPI __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream) +{ + assert(g_cudaPushArgsBuffer::g_is_initialized == false); + printf("Pushing cuda call configuration \n"); + g_cudaPushArgsBuffer::g_is_initialized = true; + g_cudaPushArgsBuffer::g_gridDim = gridDim; + g_cudaPushArgsBuffer::g_blockDim = blockDim; + g_cudaPushArgsBuffer::g_sharedMem = sharedMem; + g_cudaPushArgsBuffer::g_stream = stream; + + return cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI __cudaPopCallConfiguration() +{ + printf("Inside __cudaPopCallConfiguration\n"); + assert(g_cudaPushArgsBuffer::g_is_initialized == true); + printf("Poping cuda call configuration \n"); + g_cudaPushArgsBuffer::g_is_initialized = false; + return cudaSuccess; +} + +#endif // #if CUDA_VERSION >= 1000 + + __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, size_t offset) { if(g_debug_execution >= 3){ @@ -1551,8 +1595,14 @@ __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 grid } CUctx_st *context = GPGPUSim_Context(); function_info *entry = context->get_kernel(hostFun); - - cudaConfigureCall(gridDim, blockDim, sharedMem, stream); + +#if CUDA_VERSION >= 1000 + assert(g_cudaPushArgsBuffer::g_is_initialized == false); + cudaConfigureCall(g_cudaPushArgsBuffer::g_gridDim, g_cudaPushArgsBuffer::g_blockDim, g_cudaPushArgsBuffer::g_sharedMem, g_cudaPushArgsBuffer::g_stream); +#else + cudaConfigureCall(gridDim, blockDim, sharedMem, stream); +#endif // #if CUDA_VERSION >= 1000 + for(unsigned i = 0; i < entry->num_args(); i++){ std::pair p = entry->get_param_config(i); cudaSetupArgument(args[i], p.first, p.second); diff --git a/linux-so-version.txt b/linux-so-version.txt index a7c2d3c..357c6cd 100644 --- a/linux-so-version.txt +++ b/linux-so-version.txt @@ -6,3 +6,5 @@ libcudart.so.9.2{ }; libcuda.so.1{ }; +libcudart.so.10.0{ +}; \ No newline at end of file diff --git a/setup_environment b/setup_environment index b420584..154dea0 100644 --- a/setup_environment +++ b/setup_environment @@ -51,9 +51,10 @@ CC_VERSION=`gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($i,/^[0 CUDA_VERSION_STRING=`$CUDA_INSTALL_PATH/bin/nvcc --version | awk '/release/ {print $5;}' | sed 's/,//'`; export CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'` -if [ $CUDA_VERSION_NUMBER -gt 9100 -o $CUDA_VERSION_NUMBER -lt 2030 ]; then +if [ $CUDA_VERSION_NUMBER -gt 10000 -o $CUDA_VERSION_NUMBER -lt 2030 ]; then echo "ERROR ** GPGPU-Sim version $GPGPUSIM_VERSION_STRING not tested with CUDA version $CUDA_VERSION_STRING (please see README)"; - return + echo $CUDA_VERSION_NUMBER + return fi if [ $CUDA_VERSION_NUMBER -ge 6000 ]; then -- cgit v1.3 From 57d2df04da439b1055590f10cff5f5f50791b9d9 Mon Sep 17 00:00:00 2001 From: bftf Date: Tue, 12 Nov 2019 19:14:16 -0800 Subject: Updated the ifdef flag which enables/disables the relevant code blocks for cuda 10; while assuring backwards compatibility for older cuda versions --- libcuda/cuda_api.h | 4 ++-- libcuda/cuda_runtime_api.cc | 10 +++++----- 2 files changed, 7 insertions(+), 7 deletions(-) diff --git a/libcuda/cuda_api.h b/libcuda/cuda_api.h index b304a59..0ded242 100644 --- a/libcuda/cuda_api.h +++ b/libcuda/cuda_api.h @@ -234,11 +234,11 @@ typedef struct CUgraphicsResource_st *CUgraphicsResource; /**< CUDA graphics int typedef unsigned long long CUtexObject; /**< An opaque value that represents a CUDA texture object */ typedef unsigned long long CUsurfObject; /**< An opaque value that represents a CUDA surface object */ -#if CUDA_VERSION < 1000 +#if CUDART_VERSION < 10000 typedef struct CUuuid_st { /**< CUDA definition of UUID */ char bytes[16]; } CUuuid; -#endif +#endif // #if CUDART_VERSION < 10000 #if __CUDA_API_VERSION >= 4010 diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index fd092e5..71b8722 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1449,7 +1449,7 @@ __host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, si } -#if CUDA_VERSION >= 1000 +#if CUDART_VERSION >= 10000 /* * CUDA 10 requires a new CUDA kernel launch sequence * A call to __cudaPushCallConfiguration() preceeds any call to cudaLaunchKernel() @@ -1489,7 +1489,7 @@ __host__ cudaError_t CUDARTAPI __cudaPopCallConfiguration() return cudaSuccess; } -#endif // #if CUDA_VERSION >= 1000 +#endif // #if CUDART_VERSION >= 10000 __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, size_t offset) @@ -1596,12 +1596,12 @@ __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 grid CUctx_st *context = GPGPUSim_Context(); function_info *entry = context->get_kernel(hostFun); -#if CUDA_VERSION >= 1000 +#if CUDART_VERSION >= 10000 assert(g_cudaPushArgsBuffer::g_is_initialized == false); cudaConfigureCall(g_cudaPushArgsBuffer::g_gridDim, g_cudaPushArgsBuffer::g_blockDim, g_cudaPushArgsBuffer::g_sharedMem, g_cudaPushArgsBuffer::g_stream); #else - cudaConfigureCall(gridDim, blockDim, sharedMem, stream); -#endif // #if CUDA_VERSION >= 1000 + cudaConfigureCall(gridDim, blockDim, sharedMem, stream); +#endif // #if CUDART_VERSION >= 10000 for(unsigned i = 0; i < entry->num_args(); i++){ std::pair p = entry->get_param_config(i); -- cgit v1.3