From 2a6788b59055b5ce694882a282af0cc6311854d4 Mon Sep 17 00:00:00 2001 From: Nick Date: Mon, 26 Aug 2019 13:42:10 -0400 Subject: Fix a bunch of outstanding warnings and undefined behavior --- src/gpgpu-sim/addrdec.cc | 2 +- src/gpgpu-sim/addrdec.h | 2 +- src/gpgpu-sim/dram.cc | 9 ++++----- src/gpgpu-sim/gpu-cache.cc | 6 ++---- src/gpgpu-sim/gpu-sim.cc | 10 +++++----- src/gpgpu-sim/l2cache.cc | 6 +++--- src/gpgpu-sim/local_interconnect.cc | 2 +- src/gpgpu-sim/scoreboard.cc | 4 ++-- src/gpgpu-sim/shader.cc | 14 +++++++------- src/gpgpu-sim/shader.h | 6 +++--- 10 files changed, 29 insertions(+), 32 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/src/gpgpu-sim/addrdec.cc b/src/gpgpu-sim/addrdec.cc index ca88ec9..09bbc3c 100644 --- a/src/gpgpu-sim/addrdec.cc +++ b/src/gpgpu-sim/addrdec.cc @@ -182,7 +182,7 @@ void linear_to_raw_address_translation::addrdec_tlx(new_addr_type addr, addrdec_ } assert(tlx->chip < m_n_channel); - assert(tlx->sub_partition < m_n_channel*m_n_sub_partition_in_channel); + assert(tlx->sub_partition < m_n_channel * m_n_sub_partition_in_channel); return; break; } diff --git a/src/gpgpu-sim/addrdec.h b/src/gpgpu-sim/addrdec.h index a5333fb..c9a1420 100644 --- a/src/gpgpu-sim/addrdec.h +++ b/src/gpgpu-sim/addrdec.h @@ -92,7 +92,7 @@ private: new_addr_type sub_partition_id_mask; unsigned int gap; - int m_n_channel; + unsigned m_n_channel; int m_n_sub_partition_in_channel; }; diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc index d443d79..9c33822 100644 --- a/src/gpgpu-sim/dram.cc +++ b/src/gpgpu-sim/dram.cc @@ -482,7 +482,6 @@ void dram_t::cycle() bool memory_pending_rw_found=false; for (unsigned j=0;jnbk;j++) { - unsigned grp = get_bankgrp_number(j); if (bk[j]->mrq && (((bk[j]->curr_row == bk[j]->mrq->row) && (bk[j]->mrq->rw == READ) && (bk[j]->state == BANK_ACTIVE)) @@ -817,10 +816,10 @@ void dram_t::visualize() const void dram_t::print_stat( FILE* simFile ) { - fprintf(simFile,"DRAM (%llu): n_cmd=%llu n_nop=%llu n_act=%llu n_pre=%llu n_ref=%llu n_req=%llu n_rd=%llu n_write=%llu bw_util=%.4g ", + fprintf(simFile,"DRAM (%u): n_cmd=%llu n_nop=%llu n_act=%llu n_pre=%llu n_ref=%llu n_req=%llu n_rd=%llu n_write=%llu bw_util=%.4g ", id, n_cmd, n_nop, n_act, n_pre, n_ref, n_req, n_rd, n_wr, (float)bwutil/n_cmd); - fprintf(simFile, "mrqq: %d %.4g mrqsmax=%d ", max_mrqs, (float)ave_mrqs/n_cmd, max_mrqs_temp); + fprintf(simFile, "mrqq: %d %.4g mrqsmax=%llu ", max_mrqs, (float)ave_mrqs/n_cmd, max_mrqs_temp); fprintf(simFile, "\n"); fprintf(simFile, "dram_util_bins:"); for (unsigned i=0;i<10;i++) fprintf(simFile, " %d", dram_util_bins[i]); @@ -899,10 +898,10 @@ void dram_t::set_dram_power_stats( unsigned &cmd, unsigned dram_t::get_bankgrp_number(unsigned i) { if(m_config->dram_bnkgrp_indexing_policy == HIGHER_BITS) { //higher bits - return i>>m_config->bk_tag_length; + return i >> m_config->bk_tag_length; } else if (m_config->dram_bnkgrp_indexing_policy == LOWER_BITS) { //lower bits - return i&((m_config->nbkgrp-1)); + return i & ((m_config->nbkgrp - 1)); } else { assert(1); diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc index 1705821..dec61db 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -777,7 +777,7 @@ void cache_stats::print_stats(FILE *fout, const char *cache_name) const{ } for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { if(total_access[type] > 0) - fprintf(fout, "\t%s[%s][%s] = %llu\n", + fprintf(fout, "\t%s[%s][%s] = %u\n", m_cache_name.c_str(), mem_access_type_str((enum mem_access_type)type), "TOTAL_ACCESS", @@ -790,7 +790,7 @@ void cache_stats::print_fail_stats(FILE *fout, const char *cache_name) const{ for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { for (unsigned fail = 0; fail < NUM_CACHE_RESERVATION_FAIL_STATUS; ++fail) { if(m_fail_stats[type][fail] > 0){ - fprintf(fout, "\t%s[%s][%s] = %u\n", + fprintf(fout, "\t%s[%s][%s] = %llu\n", m_cache_name.c_str(), mem_access_type_str((enum mem_access_type)type), cache_fail_status_str((enum cache_reservation_fail_reason)fail), @@ -1417,8 +1417,6 @@ data_cache::wr_miss_wa_lazy_fetch_on_read( new_addr_type addr, { new_addr_type block_addr = m_config.block_addr(addr); - new_addr_type mshr_addr = m_config.mshr_addr(mf->get_addr()); - //if the request writes to the whole cache line/sector, then, write and set cache line Modified. //and no need to send read request to memory or reserve mshr diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index e4ae04f..1f9a422 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1101,8 +1101,8 @@ void gpgpu_sim::gpu_print_stat() printf("gpu_tot_sim_insn = %lld\n", gpu_tot_sim_insn+gpu_sim_insn); 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); - printf("gpu_occupancy = %.4f\% \n", gpu_occupancy.get_occ_fraction() * 100); - printf("gpu_tot_occupancy = %.4f\% \n", (gpu_occupancy + gpu_tot_occupancy).get_occ_fraction() * 100); + printf("gpu_occupancy = %.4f%% \n", gpu_occupancy.get_occ_fraction() * 100); + printf("gpu_tot_occupancy = %.4f%% \n", (gpu_occupancy + gpu_tot_occupancy).get_occ_fraction() * 100); fprintf(statfout, "max_total_param_size = %llu\n", gpgpu_ctx->device_runtime->g_max_total_param_size); @@ -1343,7 +1343,7 @@ bool shader_core_ctx::occupy_shader_resource_1block(kernel_info_t & k, bool occu m_occupied_regs += (padded_cta_size * ((kernel_info->regs+3)&~3)); m_occupied_ctas++; - SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Occupied %d threads, %d shared mem, %d registers, %d ctas\n", + SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Occupied %u threads, %u shared mem, %u registers, %u ctas\n", m_occupied_n_threads, m_occupied_shmem, m_occupied_regs, m_occupied_ctas); } @@ -1460,7 +1460,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) nthreads_in_block += ptx_sim_init_thread(kernel,&m_thread[i],m_sid,i,cta_size-(i-start_thread),m_config->n_thread_per_shader,this,free_cta_hw_id,warp_id,m_cluster->get_gpu()); m_threadState[i].m_active = true; // load thread local memory and register file - if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaidcheckpoint_CTA_t ) + if(m_gpu->resume_option == 1 && kernel.get_uid() == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t ) { char fname[2048]; snprintf(fname,2048,"checkpoint_files/thread_%d_%d_reg.txt",i%cta_size,ctaid ); @@ -1475,7 +1475,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) assert( nthreads_in_block > 0 && nthreads_in_block <= m_config->n_thread_per_shader); // should be at least one, but less than max m_cta_status[free_cta_hw_id]=nthreads_in_block; - if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaidcheckpoint_CTA_t ) + if(m_gpu->resume_option == 1 && kernel.get_uid() == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t ) { char f1name[2048]; snprintf(f1name,2048,"checkpoint_files/shared_mem_%d.txt", ctaid); diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc index 6540b52..862461f 100644 --- a/src/gpgpu-sim/l2cache.cc +++ b/src/gpgpu-sim/l2cache.cc @@ -82,7 +82,7 @@ void memory_partition_unit::handle_memcpy_to_gpu( size_t addr, unsigned global_s unsigned p = global_sub_partition_id_to_local_id(global_subpart_id); std::string mystring = mask.to_string(); - MEMPART_DPRINTF("Copy Engine Request Received For Address=%llx, local_subpart=%u, global_subpart=%u, sector_mask=%s \n", addr, p, global_subpart_id, mystring.c_str()); + MEMPART_DPRINTF("Copy Engine Request Received For Address=%zx, local_subpart=%u, global_subpart=%u, sector_mask=%s \n", addr, p, global_subpart_id, mystring.c_str()); m_sub_partition[p]->force_l2_tag_update(addr,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle, mask); } @@ -622,7 +622,7 @@ std::vector memory_sub_partition::breakdown_request_to_sector_reques } } else { - printf("Invalid sector received, address = 0x%06x, sector mask = %s, data size = %d", + printf("Invalid sector received, address = 0x%06llx, sector mask = %s, data size = %d", mf->get_addr(), mf->get_access_sector_mask(), mf->get_data_size()); assert(0 && "Undefined sector mask is received"); } @@ -657,7 +657,7 @@ std::vector memory_sub_partition::breakdown_request_to_sector_reques byte_sector_mask <<= SECTOR_SIZE; } } else { - printf("Invalid sector received, address = 0x%06x, sector mask = %d, byte mask = , data size = %d", + printf("Invalid sector received, address = 0x%06llx, sector mask = %d, byte mask = , data size = %u", mf->get_addr(), mf->get_access_sector_mask().count(), mf->get_data_size()); assert(0 && "Undefined data size is received"); } diff --git a/src/gpgpu-sim/local_interconnect.cc b/src/gpgpu-sim/local_interconnect.cc index 1416b2c..bb09d44 100644 --- a/src/gpgpu-sim/local_interconnect.cc +++ b/src/gpgpu-sim/local_interconnect.cc @@ -231,7 +231,7 @@ LocalInterconnect::LocalInterconnect(const struct inct_config& m_localinct_confi } LocalInterconnect::~LocalInterconnect(){ - for (int i=0; i inst_regs; - for(int iii=0;iiioutcount;iii++) + for(unsigned iii=0; iii < inst->outcount; iii++) inst_regs.insert(inst->out[iii]); - for(int jjj=0;jjjincount;jjj++) + for(unsigned jjj=0;jjjincount;jjj++) inst_regs.insert(inst->in[jjj]); if(inst->pred > 0) inst_regs.insert(inst->pred); diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index c697450..c365ebb 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -87,7 +87,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, shader_core_stats *stats ) : core_t( gpu, NULL, config->warp_size, config->n_thread_per_shader ), m_barriers( this, config->max_warps_per_shader, config->max_cta_per_core, config->max_barriers_per_cta, config->warp_size ), - m_dynamic_warp_id(0), m_active_warps(0) + m_active_warps(0), m_dynamic_warp_id(0) { m_cluster = cluster; m_config = config; @@ -164,7 +164,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, NUM_CONCRETE_SCHEDULERS; assert ( scheduler != NUM_CONCRETE_SCHEDULERS ); - for (int i = 0; i < m_config->gpgpu_num_sched_per_core; i++) { + for (unsigned i = 0; i < m_config->gpgpu_num_sched_per_core; i++) { switch( scheduler ) { case CONCRETE_SCHEDULER_LRR: @@ -263,7 +263,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, //distribute i's evenly though schedulers; schedulers[i%m_config->gpgpu_num_sched_per_core]->add_supervised_warp_id(i); } - for ( int i = 0; i < m_config->gpgpu_num_sched_per_core; ++i ) { + for ( unsigned i = 0; i < m_config->gpgpu_num_sched_per_core; ++i ) { schedulers[i]->done_adding_supervised_warps(); } @@ -474,7 +474,7 @@ void shader_core_ctx::init_warps( unsigned cta_id, unsigned start_thread, unsign } m_simt_stack[i]->launch(start_pc,active_threads); - if(m_gpu->resume_option==1 && kernel_id==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaidcheckpoint_CTA_t ) + if(m_gpu->resume_option == 1 && kernel_id == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t ) { char fname[2048]; snprintf(fname,2048,"checkpoint_files/warp_%d_%d_simt.txt",i%warp_per_cta,ctaid ); @@ -868,7 +868,7 @@ void shader_core_ctx::func_exec_inst( warp_inst_t &inst ) void shader_core_ctx::issue_warp( register_set& pipe_reg_set, const warp_inst_t* next_inst, const active_mask_t &active_mask, unsigned warp_id, unsigned sch_id ) { - warp_inst_t** pipe_reg = pipe_reg = pipe_reg_set.get_free(m_config->sub_core_model, sch_id); + warp_inst_t** pipe_reg = pipe_reg_set.get_free(m_config->sub_core_model, sch_id); assert(pipe_reg); m_warp[warp_id].ibuffer_free(); @@ -2134,7 +2134,7 @@ ldst_unit::ldst_unit( mem_fetch_interface *icnt, if(m_config->m_L1D_config.l1_latency > 0) { - for(int i=0; im_L1D_config.l1_latency; i++ ) + for(unsigned i = 0; i < m_config->m_L1D_config.l1_latency; i++ ) l1_latency_queue.push_back((mem_fetch*)NULL); } } @@ -2446,7 +2446,7 @@ void shader_core_ctx::register_cta_thread_exit( unsigned cta_num, kernel_info_t m_barriers.deallocate_barrier(cta_num); shader_CTA_count_unlog(m_sid, 1); - SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Finished CTA #%d (%lld,%lld), %u CTAs running\n", + SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Finished CTA #%u (%lld,%lld), %u CTAs running\n", cta_num, m_gpu->gpu_sim_cycle, m_gpu->gpu_tot_sim_cycle, m_n_active_cta); if( m_n_active_cta == 0 ) { diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index b0d7f7f..dbe2285 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1392,9 +1392,9 @@ class shader_core_config : public core_config If we won't remove it, old regression will be broken. So to support the legacy config files it's best to handle in this way. */ - int num_config_to_read=N_PIPELINE_STAGES-2*(!gpgpu_tensor_core_avail); + int num_config_to_read= N_PIPELINE_STAGES - 2 * (!gpgpu_tensor_core_avail); - for (unsigned i = 0; i Date: Fri, 30 Aug 2019 20:17:01 -0400 Subject: Making the max cycles/inst counters 64 bits --- src/gpgpu-sim/gpu-sim.cc | 4 ++-- src/gpgpu-sim/gpu-sim.h | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 1720836..d236c74 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -475,10 +475,10 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) m_shader_config.reg_options(opp); m_memory_config.reg_options(opp); power_config::reg_options(opp); - option_parser_register(opp, "-gpgpu_max_cycle", OPT_INT32, &gpu_max_cycle_opt, + option_parser_register(opp, "-gpgpu_max_cycle", OPT_INT64, &gpu_max_cycle_opt, "terminates gpu simulation early (0 = no limit)", "0"); - option_parser_register(opp, "-gpgpu_max_insn", OPT_INT32, &gpu_max_insn_opt, + option_parser_register(opp, "-gpgpu_max_insn", OPT_INT64, &gpu_max_insn_opt, "terminates gpu simulation early (0 = no limit)", "0"); option_parser_register(opp, "-gpgpu_max_cta", OPT_INT32, &gpu_max_cta_opt, diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 76c7a06..fba770d 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -368,8 +368,8 @@ private: double l2_period; // GPGPU-Sim timing model options - unsigned gpu_max_cycle_opt; - unsigned gpu_max_insn_opt; + unsigned long long gpu_max_cycle_opt; + unsigned long long gpu_max_insn_opt; unsigned gpu_max_cta_opt; char *gpgpu_runtime_stat; bool gpgpu_flush_l1_cache; -- cgit v1.3 From 48887cfe0261abde6de23fc5d1d76694426b7e8e Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 12 Sep 2019 02:01:01 -0400 Subject: Remove GPGPUsim_ctx_ptr() Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 116 ++++++++++++++++++++--------- libcuda/gpgpu_context.h | 5 ++ libopencl/opencl_runtime_api.cc | 8 +- src/cuda-sim/cuda-sim.cc | 2 +- src/gpgpu-sim/gpu-sim.cc | 4 +- src/gpgpusim_entrypoint.cc | 161 +++++++++++++++++++--------------------- src/gpgpusim_entrypoint.h | 6 -- 7 files changed, 167 insertions(+), 135 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 9459b64..5cefe60 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -150,9 +150,6 @@ #endif -extern void synchronize(); -extern void exit_simulation(); - /*DEVICE_BUILTIN*/ struct cudaArray { @@ -242,8 +239,8 @@ struct _cuda_device_id *gpgpu_context::GPGPUSim_Init() prop->maxThreadsPerMultiProcessor = the_gpu->threads_per_core(); #endif the_gpu->set_prop(prop); - GPGPUsim_ctx_ptr()->the_cude_device = new _cuda_device_id(the_gpu); - the_device = GPGPUsim_ctx_ptr()->the_cude_device; + the_gpgpusim->the_cude_device = new _cuda_device_id(the_gpu); + the_device = the_gpgpusim->the_cude_device; } start_sim_thread(1); return the_device; @@ -1876,6 +1873,28 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroyInternal(cudaStream_t stream, gp return g_last_cudaError = cudaSuccess; } +__host__ cudaError_t CUDARTAPI cudaStreamSynchronizeInternal(cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } +#if (CUDART_VERSION >= 3000) + if( stream == NULL ) + ctx->synchronize(); + return g_last_cudaError = cudaSuccess; + stream->synchronize(); +#else + printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); +#endif + return g_last_cudaError = cudaSuccess; +} + void __cudaRegisterTextureInternal( void **fatCubinHandle, const struct textureReference *hostVar, @@ -2065,6 +2084,53 @@ __host__ cudaError_t CUDARTAPI cudaStreamWaitEventInternal(cudaStream_t stream, return g_last_cudaError = cudaSuccess; } +__host__ cudaError_t CUDARTAPI cudaThreadExitInternal(gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + ctx->exit_simulation(); + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaThreadSynchronizeInternal(gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + //Called on host side + ctx->synchronize(); + return g_last_cudaError = cudaSuccess; +} + +cudaError_t CUDARTAPI cudaDeviceSynchronizeInternal(gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + //Blocks until the device has completed all preceding requested tasks + ctx->synchronize(); + return g_last_cudaError = cudaSuccess; +} + /******************************************************************************* * * * * @@ -2592,18 +2658,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) __host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } -#if (CUDART_VERSION >= 3000) - if( stream == NULL ) - synchronize(); - return g_last_cudaError = cudaSuccess; - stream->synchronize(); -#else - printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); -#endif - return g_last_cudaError = cudaSuccess; + return cudaStreamSynchronizeInternal(stream); } __host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream) @@ -2722,22 +2777,13 @@ __host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float *ms, cudaEvent_t start __host__ cudaError_t CUDARTAPI cudaThreadExit(void) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - exit_simulation(); - return g_last_cudaError = cudaSuccess; + return cudaThreadExitInternal(); } __host__ cudaError_t CUDARTAPI cudaThreadSynchronize(void) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - //Called on host side - synchronize(); - return g_last_cudaError = cudaSuccess; -}; + return cudaThreadSynchronizeInternal(); +} int CUDARTAPI __cudaSynchronizeThreads(void**, void*) { @@ -3441,15 +3487,11 @@ cudaError_t cudaDeviceReset ( void ) { } return g_last_cudaError = cudaSuccess; } -cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - //Blocks until the device has completed all preceding requested tasks - synchronize(); - return g_last_cudaError = cudaSuccess; -} +cudaError_t CUDARTAPI cudaDeviceSynchronize(void) +{ + return cudaDeviceSynchronizeInternal(); +} void __cudaRegisterShared( void **fatCubinHandle, diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 45c5cdd..61d7507 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -52,6 +52,10 @@ class gpgpu_context { cuda_device_runtime* device_runtime; ptx_stats* stats; // member function list + void synchronize(); + void exit_simulation(); + void print_simulation_time(); + int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ); void cuobjdumpParseBinary(unsigned int handle); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ); @@ -60,6 +64,7 @@ class gpgpu_context { void print_ptx_file( const char *p, unsigned source_num, const char *filename ); class symbol_table* init_parser(const char*); class gpgpu_sim *gpgpu_ptx_sim_init_perf(); + void start_sim_thread(int api); struct _cuda_device_id *GPGPUSim_Init(); void ptx_reg_options(option_parser_t opp); const ptx_instruction* pc_to_instruction(unsigned pc); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 7f029c7..b032c05 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -647,13 +647,13 @@ unsigned _cl_kernel::sm_context_uid = 0; class _cl_device_id *GPGPUSim_Init() { static _cl_device_id *the_device = NULL; + gpgpu_context *ctx; + ctx = GPGPU_Context(); if( !the_device ) { - gpgpu_context *ctx; - ctx = GPGPU_Context(); gpgpu_sim *the_gpu = ctx->gpgpu_ptx_sim_init_perf(); the_device = new _cl_device_id(the_gpu); } - start_sim_thread(2); + ctx->start_sim_thread(2); return the_device; } @@ -960,7 +960,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, if ( ctx->func_sim->g_ptx_sim_mode ) ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid ); else - gpgpu_opencl_ptx_sim_main_perf( grid ); + ctx->gpgpu_opencl_ptx_sim_main_perf( grid ); return CL_SUCCESS; } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index fc1ea8e..7a130ea 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2203,7 +2203,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL //g_simulation_starttime is initilized by gpgpu_ptx_sim_init_perf() in gpgpusim_entrypoint.cc upon starting gpgpu-sim time_t end_time, elapsed_time, days, hrs, minutes, sec; end_time = time((time_t *)NULL); - elapsed_time = MAX(end_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1); + elapsed_time = MAX(end_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1); //calculating and printing simulation time in terms of days, hours, minutes and seconds diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index d236c74..56ea8c4 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1144,7 +1144,7 @@ void gpgpu_sim::gpu_print_stat() time_t curr_time; time(&curr_time); - unsigned long long elapsed_time = MAX( curr_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1 ); + unsigned long long elapsed_time = MAX( curr_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1 ); printf( "gpu_total_sim_rate=%u\n", (unsigned)( ( gpu_tot_sim_insn + gpu_sim_insn ) / elapsed_time ) ); //shader_print_l1_miss_stat( stdout ); @@ -1717,7 +1717,7 @@ void gpgpu_sim::cycle() time_t days, hrs, minutes, sec; time_t curr_time; time(&curr_time); - unsigned long long elapsed_time = MAX(curr_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1); + unsigned long long elapsed_time = MAX(curr_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1); if ( (elapsed_time - last_liveness_message_time) >= m_config.liveness_message_freq && DTRACE(LIVENESS) ) { days = elapsed_time/(3600*24); hrs = elapsed_time/3600 - 24*days; diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 5ce40c6..846773d 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -43,38 +43,28 @@ static int sg_argc = 3; static const char *sg_argv[] = {"", "-config","gpgpusim.config"}; -GPGPUsim_ctx* the_gpgpusim = NULL; - -GPGPUsim_ctx* GPGPUsim_ctx_ptr(){ - if(the_gpgpusim == NULL) - the_gpgpusim = GPGPU_Context()->the_gpgpusim; - - return the_gpgpusim; -} - -static void print_simulation_time(); - -void *gpgpu_sim_thread_sequential(void*) +void * gpgpu_sim_thread_sequential(void * ctx_ptr) { + gpgpu_context * ctx = (gpgpu_context *)ctx_ptr; // at most one kernel running at a time bool done; do { - sem_wait(&(GPGPUsim_ctx_ptr()->g_sim_signal_start)); + sem_wait(&(ctx->the_gpgpusim->g_sim_signal_start)); done = true; - if( GPGPUsim_ctx_ptr()->g_the_gpu->get_more_cta_left() ) { + if( ctx->the_gpgpusim->g_the_gpu->get_more_cta_left() ) { done = false; - GPGPUsim_ctx_ptr()->g_the_gpu->init(); - while( GPGPUsim_ctx_ptr()->g_the_gpu->active() ) { - GPGPUsim_ctx_ptr()->g_the_gpu->cycle(); - GPGPUsim_ctx_ptr()->g_the_gpu->deadlock_check(); + ctx->the_gpgpusim->g_the_gpu->init(); + while( ctx->the_gpgpusim->g_the_gpu->active() ) { + ctx->the_gpgpusim->g_the_gpu->cycle(); + ctx->the_gpgpusim->g_the_gpu->deadlock_check(); } - GPGPUsim_ctx_ptr()->g_the_gpu->print_stats(); - GPGPUsim_ctx_ptr()->g_the_gpu->update_stats(); - print_simulation_time(); + ctx->the_gpgpusim->g_the_gpu->print_stats(); + ctx->the_gpgpusim->g_the_gpu->update_stats(); + ctx->print_simulation_time(); } - sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_finish)); + sem_post(&(ctx->the_gpgpusim->g_sim_signal_finish)); } while(!done); - sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit)); + sem_post(&(ctx->the_gpgpusim->g_sim_signal_exit)); return NULL; } @@ -86,8 +76,9 @@ static void termination_callback() fflush(stdout); } -void *gpgpu_sim_thread_concurrent(void*) +void *gpgpu_sim_thread_concurrent(void * ctx_ptr) { + gpgpu_context * ctx = (gpgpu_context *)ctx_ptr; atexit(termination_callback); // concurrent kernel execution simulation thread do { @@ -95,19 +86,19 @@ void *gpgpu_sim_thread_concurrent(void*) printf("GPGPU-Sim: *** simulation thread starting and spinning waiting for work ***\n"); fflush(stdout); } - while( GPGPUsim_ctx_ptr()->g_stream_manager->empty_protected() && !GPGPUsim_ctx_ptr()->g_sim_done ) + while( ctx->the_gpgpusim->g_stream_manager->empty_protected() && !ctx->the_gpgpusim->g_sim_done ) ; if(g_debug_execution >= 3) { printf("GPGPU-Sim: ** START simulation thread (detected work) **\n"); - GPGPUsim_ctx_ptr()->g_stream_manager->print(stdout); + ctx->the_gpgpusim->g_stream_manager->print(stdout); fflush(stdout); } - pthread_mutex_lock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); - GPGPUsim_ctx_ptr()->g_sim_active = true; - pthread_mutex_unlock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); + pthread_mutex_lock(&(ctx->the_gpgpusim->g_sim_lock)); + ctx->the_gpgpusim->g_sim_active = true; + pthread_mutex_unlock(&(ctx->the_gpgpusim->g_sim_lock)); bool active = false; bool sim_cycles = false; - GPGPUsim_ctx_ptr()->g_the_gpu->init(); + ctx->the_gpgpusim->g_the_gpu->init(); do { // check if a kernel has completed // launch operation on device if one is pending and can be run @@ -119,82 +110,82 @@ void *gpgpu_sim_thread_concurrent(void*) // another kernel, the gpu is not re-initialized and the inter-kernel // behaviour may be incorrect. Check that a kernel has finished and // no other kernel is currently running. - if(GPGPUsim_ctx_ptr()->g_stream_manager->operation(&sim_cycles) && !GPGPUsim_ctx_ptr()->g_the_gpu->active()) + if(ctx->the_gpgpusim->g_stream_manager->operation(&sim_cycles) && !ctx->the_gpgpusim->g_the_gpu->active()) break; //functional simulation - if( GPGPUsim_ctx_ptr()->g_the_gpu->is_functional_sim()) { - kernel_info_t * kernel = GPGPUsim_ctx_ptr()->g_the_gpu->get_functional_kernel(); + if( ctx->the_gpgpusim->g_the_gpu->is_functional_sim()) { + kernel_info_t * kernel = ctx->the_gpgpusim->g_the_gpu->get_functional_kernel(); assert(kernel); - GPGPUsim_ctx_ptr()->gpgpu_ctx->func_sim->gpgpu_cuda_ptx_sim_main_func(*kernel); - GPGPUsim_ctx_ptr()->g_the_gpu->finish_functional_sim(kernel); + ctx->the_gpgpusim->gpgpu_ctx->func_sim->gpgpu_cuda_ptx_sim_main_func(*kernel); + ctx->the_gpgpusim->g_the_gpu->finish_functional_sim(kernel); } //performance simulation - if( GPGPUsim_ctx_ptr()->g_the_gpu->active() ) { - GPGPUsim_ctx_ptr()->g_the_gpu->cycle(); + if( ctx->the_gpgpusim->g_the_gpu->active() ) { + ctx->the_gpgpusim->g_the_gpu->cycle(); sim_cycles = true; - GPGPUsim_ctx_ptr()->g_the_gpu->deadlock_check(); + ctx->the_gpgpusim->g_the_gpu->deadlock_check(); }else { - if(GPGPUsim_ctx_ptr()->g_the_gpu->cycle_insn_cta_max_hit()){ - GPGPUsim_ctx_ptr()->g_stream_manager->stop_all_running_kernels(); - GPGPUsim_ctx_ptr()->g_sim_done = true; - GPGPUsim_ctx_ptr()->break_limit = true; + if(ctx->the_gpgpusim->g_the_gpu->cycle_insn_cta_max_hit()){ + ctx->the_gpgpusim->g_stream_manager->stop_all_running_kernels(); + ctx->the_gpgpusim->g_sim_done = true; + ctx->the_gpgpusim->break_limit = true; } } - active=GPGPUsim_ctx_ptr()->g_the_gpu->active() || !(GPGPUsim_ctx_ptr()->g_stream_manager->empty_protected()); + active=ctx->the_gpgpusim->g_the_gpu->active() || !(ctx->the_gpgpusim->g_stream_manager->empty_protected()); - } while( active && !GPGPUsim_ctx_ptr()->g_sim_done); + } while( active && !ctx->the_gpgpusim->g_sim_done); if(g_debug_execution >= 3) { printf("GPGPU-Sim: ** STOP simulation thread (no work) **\n"); fflush(stdout); } if(sim_cycles) { - GPGPUsim_ctx_ptr()->g_the_gpu->print_stats(); - GPGPUsim_ctx_ptr()->g_the_gpu->update_stats(); - print_simulation_time(); + ctx->the_gpgpusim->g_the_gpu->print_stats(); + ctx->the_gpgpusim->g_the_gpu->update_stats(); + ctx->print_simulation_time(); } - pthread_mutex_lock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); - GPGPUsim_ctx_ptr()->g_sim_active = false; - pthread_mutex_unlock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); - } while( !GPGPUsim_ctx_ptr()->g_sim_done ); + pthread_mutex_lock(&(ctx->the_gpgpusim->g_sim_lock)); + ctx->the_gpgpusim->g_sim_active = false; + pthread_mutex_unlock(&(ctx->the_gpgpusim->g_sim_lock)); + } while( !ctx->the_gpgpusim->g_sim_done ); printf("GPGPU-Sim: *** simulation thread exiting ***\n"); fflush(stdout); - if(GPGPUsim_ctx_ptr()->break_limit) { + if(ctx->the_gpgpusim->break_limit) { printf("GPGPU-Sim: ** break due to reaching the maximum cycles (or instructions) **\n"); exit(1); } - sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit)); + sem_post(&(ctx->the_gpgpusim->g_sim_signal_exit)); return NULL; } -void synchronize() +void gpgpu_context::synchronize() { printf("GPGPU-Sim: synchronize waiting for inactive GPU simulation\n"); - GPGPUsim_ctx_ptr()->g_stream_manager->print(stdout); + the_gpgpusim->g_stream_manager->print(stdout); fflush(stdout); // sem_wait(&g_sim_signal_finish); bool done = false; do { - pthread_mutex_lock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); - done = ( GPGPUsim_ctx_ptr()->g_stream_manager->empty() && !GPGPUsim_ctx_ptr()->g_sim_active ) || GPGPUsim_ctx_ptr()->g_sim_done; - pthread_mutex_unlock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); + pthread_mutex_lock(&(the_gpgpusim->g_sim_lock)); + done = ( the_gpgpusim->g_stream_manager->empty() && !the_gpgpusim->g_sim_active ) || the_gpgpusim->g_sim_done; + pthread_mutex_unlock(&(the_gpgpusim->g_sim_lock)); } while (!done); printf("GPGPU-Sim: detected inactive GPU simulation thread\n"); fflush(stdout); // sem_post(&g_sim_signal_start); } -void exit_simulation() +void gpgpu_context::exit_simulation() { - GPGPUsim_ctx_ptr()->g_sim_done=true; + the_gpgpusim->g_sim_done=true; printf("GPGPU-Sim: exit_simulation called\n"); fflush(stdout); - sem_wait(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit)); + sem_wait(&(the_gpgpusim->g_sim_signal_exit)); printf("GPGPU-Sim: simulation thread signaled exit\n"); fflush(stdout); } @@ -211,8 +202,8 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() func_sim->ptx_opcocde_latency_options(opp); icnt_reg_options(opp); - GPGPUsim_ctx_ptr()->g_the_gpu_config = new gpgpu_sim_config(this); - GPGPUsim_ctx_ptr()->g_the_gpu_config->reg_options(opp); // register GPU microrachitecture options + the_gpgpusim->g_the_gpu_config = new gpgpu_sim_config(this); + the_gpgpusim->g_the_gpu_config->reg_options(opp); // register GPU microrachitecture options option_parser_cmdline(opp, sg_argc, sg_argv); // parse configuration options fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n"); @@ -220,37 +211,37 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() // Set the Numeric locale to a standard locale where a decimal point is a "dot" not a "comma" // so it does the parsing correctly independent of the system environment variables assert(setlocale(LC_NUMERIC,"C")); - GPGPUsim_ctx_ptr()->g_the_gpu_config->init(); + the_gpgpusim->g_the_gpu_config->init(); - GPGPUsim_ctx_ptr()->g_the_gpu = new gpgpu_sim(*(GPGPUsim_ctx_ptr()->g_the_gpu_config), this); - GPGPUsim_ctx_ptr()->g_stream_manager = new stream_manager((GPGPUsim_ctx_ptr()->g_the_gpu), func_sim->g_cuda_launch_blocking); + the_gpgpusim->g_the_gpu = new gpgpu_sim(*(the_gpgpusim->g_the_gpu_config), this); + the_gpgpusim->g_stream_manager = new stream_manager((the_gpgpusim->g_the_gpu), func_sim->g_cuda_launch_blocking); - GPGPUsim_ctx_ptr()->g_simulation_starttime = time((time_t *)NULL); + the_gpgpusim->g_simulation_starttime = time((time_t *)NULL); - sem_init(&(GPGPUsim_ctx_ptr()->g_sim_signal_start),0,0); - sem_init(&(GPGPUsim_ctx_ptr()->g_sim_signal_finish),0,0); - sem_init(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit),0,0); + sem_init(&(the_gpgpusim->g_sim_signal_start),0,0); + sem_init(&(the_gpgpusim->g_sim_signal_finish),0,0); + sem_init(&(the_gpgpusim->g_sim_signal_exit),0,0); - return GPGPUsim_ctx_ptr()->g_the_gpu; + return the_gpgpusim->g_the_gpu; } -void start_sim_thread(int api) +void gpgpu_context::start_sim_thread(int api) { - if( GPGPUsim_ctx_ptr()->g_sim_done ) { - GPGPUsim_ctx_ptr()->g_sim_done = false; + if( the_gpgpusim->g_sim_done ) { + the_gpgpusim->g_sim_done = false; if( api == 1 ) { - pthread_create(&(GPGPUsim_ctx_ptr()->g_simulation_thread),NULL,gpgpu_sim_thread_concurrent,NULL); + pthread_create(&(the_gpgpusim->g_simulation_thread),NULL,gpgpu_sim_thread_concurrent,(void *)this); } else { - pthread_create(&(GPGPUsim_ctx_ptr()->g_simulation_thread),NULL,gpgpu_sim_thread_sequential,NULL); + pthread_create(&(the_gpgpusim->g_simulation_thread),NULL,gpgpu_sim_thread_sequential,(void *)this); } } } -void print_simulation_time() +void gpgpu_context::print_simulation_time() { time_t current_time, difference, d, h, m, s; current_time = time((time_t *)NULL); - difference = MAX(current_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1); + difference = MAX(current_time - the_gpgpusim->g_simulation_starttime, 1); d = difference/(3600*24); h = difference/3600 - 24*d; @@ -260,18 +251,18 @@ void print_simulation_time() fflush(stderr); printf("\n\ngpgpu_simulation_time = %u days, %u hrs, %u min, %u sec (%u sec)\n", (unsigned)d, (unsigned)h, (unsigned)m, (unsigned)s, (unsigned)difference ); - printf("gpgpu_simulation_rate = %u (inst/sec)\n", (unsigned)(GPGPUsim_ctx_ptr()->g_the_gpu->gpu_tot_sim_insn / difference) ); - const unsigned cycles_per_sec = (unsigned)(GPGPUsim_ctx_ptr()->g_the_gpu->gpu_tot_sim_cycle / difference); + printf("gpgpu_simulation_rate = %u (inst/sec)\n", (unsigned)(the_gpgpusim->g_the_gpu->gpu_tot_sim_insn / difference) ); + const unsigned cycles_per_sec = (unsigned)(the_gpgpusim->g_the_gpu->gpu_tot_sim_cycle / difference); printf("gpgpu_simulation_rate = %u (cycle/sec)\n", cycles_per_sec ); - printf("gpgpu_silicon_slowdown = %ux\n", GPGPUsim_ctx_ptr()->g_the_gpu->shader_clock() * 1000 / cycles_per_sec); + printf("gpgpu_silicon_slowdown = %ux\n", the_gpgpusim->g_the_gpu->shader_clock() * 1000 / cycles_per_sec); fflush(stdout); } -int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ) +int gpgpu_context::gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ) { - GPGPUsim_ctx_ptr()->g_the_gpu->launch(grid); - sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_start)); - sem_wait(&(GPGPUsim_ctx_ptr()->g_sim_signal_finish)); + the_gpgpusim->g_the_gpu->launch(grid); + sem_post(&(the_gpgpusim->g_sim_signal_start)); + sem_wait(&(the_gpgpusim->g_sim_signal_finish)); return 0; } diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index b2e22e6..9f408df 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -76,10 +76,4 @@ class GPGPUsim_ctx { }; -void start_sim_thread(int api); - -struct GPGPUsim_ctx* GPGPUsim_ctx_ptr(); - -int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ); - #endif -- cgit v1.3