From 37a709bb69fcd9f2b0fe53a189e92e548164cc4b Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Wed, 5 Jun 2019 19:18:29 -0400 Subject: adding new cuda 9 APIs to run the deepbench workloads --- src/gpgpu-sim/gpu-sim.cc | 10 ++++++++++ 1 file changed, 10 insertions(+) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index c1ba934..6de5845 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -806,6 +806,16 @@ int gpgpu_sim::shader_clock() const return m_config.core_freq/1000; } +int gpgpu_sim::max_cta_per_core() const +{ + return m_shader_config->max_cta_per_core; +} + +int gpgpu_sim::get_max_cta( const kernel_info_t &k ) const +{ + return m_shader_config->max_cta(k); +} + void gpgpu_sim::set_prop( cudaDeviceProp *prop ) { m_cuda_properties = prop; -- cgit v1.3 From 17815dfdd7b6e9b558997bc5a9ba157a0493da16 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Fri, 12 Jul 2019 21:31:54 -0400 Subject: fixing device and function parameters config --- configs/tested-cfgs/SM2_GTX480/gpgpusim.config | 1 + configs/tested-cfgs/SM6_TITANX/gpgpusim.config | 1 + configs/tested-cfgs/SM7_QV100/gpgpusim.config | 11 +++++++---- configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 3 +++ libcuda/cuda_runtime_api.cc | 21 ++++++++++++--------- src/gpgpu-sim/gpu-sim.cc | 4 ++-- 6 files changed, 26 insertions(+), 15 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config index cf3627b..4a7a3c3 100644 --- a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config +++ b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config @@ -61,6 +61,7 @@ # Note: Hashing set index function (H) only applies to a set size of 32 or 64. -gpgpu_cache:dl1 N:32:128:4,L:L:m:N:H,S:64:8,8 -gpgpu_shmem_size 49152 +-gpgpu_shmem_sizeDefault 49152 -icnt_flit_size 40 -gmem_skip_L1D 0 -gpgpu_n_cluster_ejection_buffer_size 32 diff --git a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config index 2fe898a..e6d8f1d 100644 --- a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config +++ b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config @@ -81,6 +81,7 @@ -gpgpu_cache:dl1PrefL1 S:4:128:48,L:L:s:N:L,A:256:8,16:0,32 -gpgpu_cache:dl1PrefShared S:4:128:48,L:L:s:N:L,A:256:8,16:0,32 -gpgpu_shmem_size 49152 +-gpgpu_shmem_sizeDefault 49152 -gpgpu_shmem_size_PrefL1 49152 -gpgpu_shmem_size_PrefShared 49152 # By default, L1 cache is disabled in Pascal P102. diff --git a/configs/tested-cfgs/SM7_QV100/gpgpusim.config b/configs/tested-cfgs/SM7_QV100/gpgpusim.config index 1a34d0f..5f64908 100644 --- a/configs/tested-cfgs/SM7_QV100/gpgpusim.config +++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config @@ -11,7 +11,7 @@ # functional simulator specification -gpgpu_ptx_instruction_classification 0 -gpgpu_ptx_sim_mode 0 --gpgpu_ptx_force_max_capability 60 +-gpgpu_ptx_force_max_capability 70 # Device Limits @@ -21,7 +21,7 @@ -gpgpu_runtime_pending_launch_count_limit 2048 # Compute Capability --gpgpu_compute_capability_major 6 +-gpgpu_compute_capability_major 7 -gpgpu_compute_capability_minor 0 # SASS execution (only supported with CUDA >= 4.0) @@ -44,12 +44,13 @@ # shader core pipeline config -gpgpu_shader_registers 65536 --gpgpu_occupancy_sm_number 60 +-gpgpu_registers_per_block 65536 +-gpgpu_occupancy_sm_number 70 # This implies a maximum of 64 warps/SM -gpgpu_shader_core_pipeline 2048:32 -gpgpu_shader_cta 32 --gpgpu_simd_model 1 +-gpgpu_simd_model 1 # Pipeline widths and number of FUs # ID_OC_SP,ID_OC_DP,ID_OC_INT,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_DP,OC_EX_INT,OC_EX_SFU,OC_EX_MEM,EX_WB,ID_OC_TENSOR_CORE,OC_EX_TENSOR_CORE @@ -91,6 +92,8 @@ -mem_unit_ports 4 -gpgpu_cache:dl1 S:4:128:64,L:L:s:N:L,A:256:8,16:0,32 -gpgpu_shmem_size 98304 +-gpgpu_shmem_sizeDefault 98304 +-gpgpu_shmem_per_block 65536 -gmem_skip_L1D 0 -icnt_flit_size 40 -gpgpu_n_cluster_ejection_buffer_size 32 diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config index b06f048..6c21dcb 100644 --- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config +++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config @@ -44,6 +44,7 @@ # shader core pipeline config -gpgpu_shader_registers 65536 +-gpgpu_registers_per_block 65536 -gpgpu_occupancy_sm_number 70 # This implies a maximum of 64 warps/SM @@ -91,6 +92,8 @@ -mem_unit_ports 4 -gpgpu_cache:dl1 S:4:128:64,L:L:s:N:L,A:256:8,16:0,32 -gpgpu_shmem_size 98304 +-gpgpu_shmem_sizeDefault 98304 +-gpgpu_shmem_per_block 65536 -gmem_skip_L1D 0 -icnt_flit_size 40 -gpgpu_n_cluster_ejection_buffer_size 32 diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index c70a570..45511d4 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -383,7 +383,7 @@ struct _cuda_device_id *GPGPUSim_Init() prop->sharedMemPerMultiprocessor = the_gpu->shared_mem_size(); #endif prop->sharedMemPerBlock = the_gpu->shared_mem_per_block(); - prop->regsPerBlock = the_gpu->num_registers_per_core(); + prop->regsPerBlock = the_gpu->num_registers_per_block(); prop->warpSize = the_gpu->wrp_size(); prop->clockRate = the_gpu->shader_clock(); #if (CUDART_VERSION >= 2010) @@ -1014,7 +1014,7 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic prop = dev->get_prop(); switch (attr) { case 1: - *value= prop->maxThreadsDim[0] * prop->maxThreadsDim[1] * prop->maxThreadsDim[2] * prop->maxGridSize[0] * prop->maxGridSize[1] * prop->maxGridSize[2]; + *value= prop->maxThreadsPerBlock; break; case 2: *value= prop->maxThreadsDim[0]; @@ -1504,13 +1504,13 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) { dim3 gridDim = config.grid_dim(); dim3 blockDim = config.block_dim(); - if (gridDim.x * gridDim.y * gridDim.z == 0 || blockDim.x * blockDim.y * blockDim.z == 0) - { + //if (gridDim.x * gridDim.y * gridDim.z == 0 || blockDim.x * blockDim.y * blockDim.z == 0) + //{ //can't launch - printf("can't launch a empty kernel\n"); - g_cuda_launch_stack.pop_back(); - return g_last_cudaError = cudaErrorInvalidConfiguration; - } + // printf("can't launch a empty kernel\n"); + // g_cuda_launch_stack.pop_back(); + // return g_last_cudaError = cudaErrorInvalidConfiguration; + //} } struct CUstream_st *stream = config.get_stream(); if(g_stream_manager->is_blocking()) @@ -3151,9 +3151,12 @@ size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr) { size_t max = prop.maxThreadsPerBlock; - if ((prop.regsPerBlock / attr->numRegs) < max) + if (attr->numRegs && (prop.regsPerBlock / attr->numRegs) < max) max = prop.regsPerBlock / attr->numRegs; + if (attr->sharedSizeBytes && (prop.sharedMemPerBlock / attr->sharedSizeBytes) < max) + max = prop.sharedMemPerBlock / attr->sharedSizeBytes; + return max; } diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 6de5845..72cb32b 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -320,7 +320,7 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-adaptive_volta_cache_config", OPT_BOOL, &adaptive_volta_cache_config, "adaptive_volta_cache_config", "0"); - option_parser_register(opp, "-gpgpu_shmem_size", OPT_UINT32, &gpgpu_shmem_sizeDefault, + option_parser_register(opp, "-gpgpu_shmem_sizeDefault", OPT_UINT32, &gpgpu_shmem_sizeDefault, "Size of shared memory per shader core (default 16kB)", "16384"); option_parser_register(opp, "-gpgpu_shmem_size_PrefL1", OPT_UINT32, &gpgpu_shmem_sizePrefL1, @@ -1065,7 +1065,7 @@ void gpgpu_sim::change_cache_config(FuncCache cache_config) if(cache_config != m_shader_config->m_L1D_config.get_cache_status()){ printf("FLUSH L1 Cache at configuration change between kernels\n"); for (unsigned i=0;in_simt_clusters;i++) { - m_cluster[i]->cache_flush(); + m_cluster[i]->cache_invalidate(); } } -- cgit v1.3 From 7e65afd701b62111b0c28a98fabad0319993cb1d Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Mon, 15 Jul 2019 17:15:06 -0400 Subject: adding the new elapsed_cycles_sm_tot stats --- src/gpgpu-sim/gpu-sim.cc | 2 ++ src/gpgpu-sim/gpu-sim.h | 1 + src/gpgpu-sim/shader.cc | 1 + 3 files changed, 4 insertions(+) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 72cb32b..343ff86 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -85,6 +85,7 @@ bool g_interactive_debugger_enabled=false; unsigned long long gpu_sim_cycle = 0; unsigned long long gpu_tot_sim_cycle = 0; +unsigned long long elapsed_cycles_sm_tot = 0; //this is a equivalent metric generated as nvprof. that only counts when SM is active // performance counter for stalls due to congestion. @@ -1119,6 +1120,7 @@ void gpgpu_sim::gpu_print_stat() printf("gpu_sim_insn = %lld\n", gpu_sim_insn); printf("gpu_ipc = %12.4f\n", (float)gpu_sim_insn / gpu_sim_cycle); printf("gpu_tot_sim_cycle = %lld\n", gpu_tot_sim_cycle+gpu_sim_cycle); + printf("elapsed_cycles_sm_tot = %lld\n", elapsed_cycles_sm_tot); 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); diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 7336cac..e98e499 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -295,6 +295,7 @@ struct memory_config { // global counters and flags (please try not to add to this list!!!) extern unsigned long long gpu_sim_cycle; extern unsigned long long gpu_tot_sim_cycle; +extern unsigned long long elapsed_cycles_sm_tot; extern bool g_interactive_debugger_enabled; class gpgpu_sim_config : public power_config, public gpgpu_functional_sim_config { diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 007ad42..e38eefd 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -3045,6 +3045,7 @@ void shader_core_ctx::cycle() if(!isactive() && get_not_completed() == 0) return; + elapsed_cycles_sm_tot++; m_stats->shader_cycles[m_sid]++; writeback(); execute(); -- cgit v1.3 From 21d937256fbca004c926531cfef1adefcedeef91 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Mon, 29 Jul 2019 21:15:59 -0400 Subject: adding simple dram model --- src/gpgpu-sim/gpu-sim.cc | 7 ++++- src/gpgpu-sim/gpu-sim.h | 1 + src/gpgpu-sim/l2cache.cc | 77 ++++++++++++++++++++++++++++++++++++++++++------ src/gpgpu-sim/l2cache.h | 1 + 4 files changed, 76 insertions(+), 10 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 343ff86..92d5366 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -156,6 +156,8 @@ void memory_config::reg_options(class OptionParser * opp) { option_parser_register(opp, "-perf_sim_memcpy", OPT_BOOL, &m_perf_sim_memcpy, "Fill the L2 cache on memcpy", "1"); + option_parser_register(opp, "-simple_dram_model", OPT_BOOL, &simple_dram_model, + "simple_dram_model with fixed latency and BW", "0"); option_parser_register(opp, "-gpgpu_dram_scheduler", OPT_INT32, &scheduler_type, "0 = fifo, 1 = FR-FCFS (defaul)", "1"); option_parser_register(opp, "-gpgpu_dram_partition_queues", OPT_CSTR, &gpgpu_L2_queue_config, @@ -1606,7 +1608,10 @@ void gpgpu_sim::cycle() if (clock_mask & DRAM) { for (unsigned i=0;im_n_mem;i++){ - m_memory_partition_unit[i]->dram_cycle(); // Issue the dram command (scheduler + delay model) + if(m_memory_config->simple_dram_model) + m_memory_partition_unit[i]->simple_dram_model_cycle(); + else + m_memory_partition_unit[i]->dram_cycle(); // Issue the dram command (scheduler + delay model) // Update performance counters for DRAM m_memory_partition_unit[i]->set_dram_power_stats(m_power_stats->pwr_mem_stat->n_cmd[CURRENT_STAT_IDX][i], m_power_stats->pwr_mem_stat->n_activity[CURRENT_STAT_IDX][i], m_power_stats->pwr_mem_stat->n_nop[CURRENT_STAT_IDX][i], m_power_stats->pwr_mem_stat->n_act[CURRENT_STAT_IDX][i], m_power_stats->pwr_mem_stat->n_pre[CURRENT_STAT_IDX][i], diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index e98e499..c14d0a7 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -290,6 +290,7 @@ struct memory_config { unsigned write_high_watermark; unsigned write_low_watermark; bool m_perf_sim_memcpy; + bool simple_dram_model; }; // global counters and flags (please try not to add to this list!!!) diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc index 0edc3b7..f24a596 100644 --- a/src/gpgpu-sim/l2cache.cc +++ b/src/gpgpu-sim/l2cache.cc @@ -203,7 +203,67 @@ int memory_partition_unit::global_sub_partition_id_to_local_id(int global_sub_pa return (global_sub_partition_id - m_id * m_config->m_n_sub_partition_per_memory_channel); } -void memory_partition_unit::dram_cycle() +void memory_partition_unit::simple_dram_model_cycle() +{ + + // pop completed memory request from dram and push it to dram-to-L2 queue + // of the original sub partition + if (!m_dram_latency_queue.empty() && ( (gpu_sim_cycle+gpu_tot_sim_cycle) >= m_dram_latency_queue.front().ready_cycle )) { + mem_fetch* mf_return = m_dram_latency_queue.front().req; + if( mf_return->get_access_type() != L1_WRBK_ACC && mf_return->get_access_type() != L2_WRBK_ACC ) { + mf_return->set_reply(); + + unsigned dest_global_spid = mf_return->get_sub_partition_id(); + int dest_spid = global_sub_partition_id_to_local_id(dest_global_spid); + assert(m_sub_partition[dest_spid]->get_id() == dest_global_spid); + if (!m_sub_partition[dest_spid]->dram_L2_queue_full()) { + if( mf_return->get_access_type() == L1_WRBK_ACC ) { + m_sub_partition[dest_spid]->set_done(mf_return); + delete mf_return; + } else { + m_sub_partition[dest_spid]->dram_L2_queue_push(mf_return); + mf_return->set_status(IN_PARTITION_DRAM_TO_L2_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + m_arbitration_metadata.return_credit(dest_spid); + MEMPART_DPRINTF("mem_fetch request %p return from dram to sub partition %d\n", mf_return, dest_spid); + } + m_dram_latency_queue.pop_front(); + } + + } else { + this->set_done(mf_return); + delete mf_return; + m_dram_latency_queue.pop_front(); + } + } + + // mem_fetch *mf = m_sub_partition[spid]->L2_dram_queue_top(); + //if( !m_dram->full(mf->is_write()) ) { + // L2->DRAM queue to DRAM latency queue + // Arbitrate among multiple L2 subpartitions + int last_issued_partition = m_arbitration_metadata.last_borrower(); + for (unsigned p = 0; p < m_config->m_n_sub_partition_per_memory_channel; p++) { + int spid = (p + last_issued_partition + 1) % m_config->m_n_sub_partition_per_memory_channel; + if (!m_sub_partition[spid]->L2_dram_queue_empty() && can_issue_to_dram(spid)) { + mem_fetch *mf = m_sub_partition[spid]->L2_dram_queue_top(); + if(m_dram->full(mf->is_write()) ) + break; + + m_sub_partition[spid]->L2_dram_queue_pop(); + MEMPART_DPRINTF("Issue mem_fetch request %p from sub partition %d to dram\n", mf, spid); + dram_delay_t d; + d.req = mf; + d.ready_cycle = gpu_sim_cycle+gpu_tot_sim_cycle + m_config->dram_latency; + m_dram_latency_queue.push_back(d); + mf->set_status(IN_PARTITION_DRAM_LATENCY_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + m_arbitration_metadata.borrow_credit(spid); + break; // the DRAM should only accept one request per cycle + } + } + //} + +} + +void memory_partition_unit::dram_cycle() { // pop completed memory request from dram and push it to dram-to-L2 queue // of the original sub partition @@ -228,8 +288,8 @@ void memory_partition_unit::dram_cycle() m_dram->return_queue_pop(); } - m_dram->cycle(); - m_dram->dram_log(SAMPLELOG); + m_dram->cycle(); + m_dram->dram_log(SAMPLELOG); // mem_fetch *mf = m_sub_partition[spid]->L2_dram_queue_top(); //if( !m_dram->full(mf->is_write()) ) { @@ -257,12 +317,11 @@ void memory_partition_unit::dram_cycle() //} // DRAM latency queue - - if( !m_dram_latency_queue.empty() && ( (gpu_sim_cycle+gpu_tot_sim_cycle) >= m_dram_latency_queue.front().ready_cycle ) && !m_dram->full(m_dram_latency_queue.front().req->is_write()) ) { - mem_fetch* mf = m_dram_latency_queue.front().req; - m_dram_latency_queue.pop_front(); - m_dram->push(mf); - } + if( !m_dram_latency_queue.empty() && ( (gpu_sim_cycle+gpu_tot_sim_cycle) >= m_dram_latency_queue.front().ready_cycle ) && !m_dram->full(m_dram_latency_queue.front().req->is_write()) ) { + mem_fetch* mf = m_dram_latency_queue.front().req; + m_dram_latency_queue.pop_front(); + m_dram->push(mf); + } } void memory_partition_unit::set_done( mem_fetch *mf ) diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h index beafdd3..9a51c0e 100644 --- a/src/gpgpu-sim/l2cache.h +++ b/src/gpgpu-sim/l2cache.h @@ -65,6 +65,7 @@ public: void cache_cycle( unsigned cycle ); void dram_cycle(); + void simple_dram_model_cycle(); void set_done( mem_fetch *mf ); -- cgit v1.3 From da0c8dff3b4e89acaf2f2dd31bf8940ab4a1e71c Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Mon, 26 Aug 2019 12:23:43 -0400 Subject: missing some files to add --- src/gpgpu-sim/gpu-sim.cc | 9 ++++++--- src/gpgpu-sim/local_interconnect.h | 30 +++++++++++++++++++++--------- 2 files changed, 27 insertions(+), 12 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 92d5366..4f071c7 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -257,9 +257,12 @@ void shader_core_config::reg_options(class OptionParser * opp) "per-shader L1 data cache config " " {::,:::,::, | none}", "none" ); + option_parser_register(opp, "-l1_banks", OPT_UINT32, &m_L1D_config.l1_banks, + "The number of L1 cache banks", + "1"); option_parser_register(opp, "-l1_latency", OPT_UINT32, &m_L1D_config.l1_latency, "L1 Hit Latency", - "0"); + "1"); option_parser_register(opp, "-smem_latency", OPT_UINT32, &smem_latency, "smem Latency", "3"); @@ -320,8 +323,8 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-gpgpu_shmem_size", OPT_UINT32, &gpgpu_shmem_size, "Size of shared memory per shader core (default 16kB)", "16384"); - option_parser_register(opp, "-adaptive_volta_cache_config", OPT_BOOL, &adaptive_volta_cache_config, - "adaptive_volta_cache_config", + option_parser_register(opp, "-adaptive_cache_config", OPT_BOOL, &adaptive_volta_cache_config, + "adaptive_cache_config", "0"); option_parser_register(opp, "-gpgpu_shmem_sizeDefault", OPT_UINT32, &gpgpu_shmem_sizeDefault, "Size of shared memory per shader core (default 16kB)", diff --git a/src/gpgpu-sim/local_interconnect.h b/src/gpgpu-sim/local_interconnect.h index f4a2af1..a784da8 100644 --- a/src/gpgpu-sim/local_interconnect.h +++ b/src/gpgpu-sim/local_interconnect.h @@ -35,27 +35,35 @@ using namespace std; +enum Interconnect_type { + REQ_NET=0, + REPLY_NET=1 +}; + +enum Arbiteration_type { + NAIVE_RR=0, + iSLIP=1 +}; + struct inct_config { - //config for local interconnect unsigned in_buffer_limit; unsigned out_buffer_limit; unsigned subnets; + Arbiteration_type arbiter_algo; }; -enum Interconnect_type { - REQ_NET=0, - REPLY_NET=1 -}; class xbar_router { public: - xbar_router(unsigned router_id, enum Interconnect_type m_type, unsigned n_shader, unsigned n_mem, unsigned m_in_buffer_limit, unsigned m_out_buffer_limit); + xbar_router(unsigned router_id, enum Interconnect_type m_type, unsigned n_shader, unsigned n_mem, unsigned m_in_buffer_limit, unsigned m_out_buffer_limit, enum Arbiteration_type m_arbit_type); ~xbar_router(); void Push(unsigned input_deviceID, unsigned output_deviceID, void* data, unsigned int size); void* Pop(unsigned ouput_deviceID); - void Advance(); + void Advance( ); + + bool Busy() const; bool Has_Buffer_In(unsigned input_deviceID, unsigned size, bool update_counter=false); bool Has_Buffer_Out(unsigned output_deviceID, unsigned size); @@ -70,6 +78,9 @@ public: unsigned long long packets_num; private: + void iSLIP_Advance(); + void RR_Advance(); + struct Packet{ Packet(void* m_data, unsigned m_output_deviceID) { data = m_data; @@ -82,11 +93,12 @@ private: vector > out_buffers; unsigned _n_shader, _n_mem, total_nodes; unsigned in_buffer_limit, out_buffer_limit; - vector next_node; -// unsigned next_node; + vector next_node; //used for iSLIP arbit + unsigned next_node_id; //used for RR arbit unsigned m_id; enum Interconnect_type router_type; unsigned active_in_buffers,active_out_buffers; + Arbiteration_type arbit_type; friend class LocalInterconnect; -- cgit v1.3 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 --- libcuda/cuda_runtime_api.cc | 10 ++--- src/abstract_hardware_model.h | 18 ++++----- src/cuda-sim/cuda-sim.cc | 18 ++------- src/cuda-sim/instructions.cc | 74 ++++++++++++++++++------------------- src/cuda-sim/ptx_parser.cc | 2 +- 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 +-- 15 files changed, 82 insertions(+), 101 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 10a651a..43a5864 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -548,7 +548,7 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( size_t* pValue, cudaL break; } else{ - printf("ERROR:Limit %s is not supported on this architecture \n",limit); + printf("ERROR:Limit %d is not supported on this architecture \n", limit); abort(); } case 4: // cudaLimitDevRuntimePendingLaunchCount @@ -557,12 +557,12 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( size_t* pValue, cudaL break; } else{ - printf("ERROR:Limit %s is not supported on this architecture \n",limit); + printf("ERROR:Limit %d is not supported on this architecture \n",limit); abort(); } #endif default: - printf("ERROR:Limit %s unimplemented \n",limit); + printf("ERROR:Limit %d unimplemented \n",limit); abort(); } return g_last_cudaError = cudaSuccess; @@ -2471,7 +2471,6 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){ while (std::getline(infile, line)) { //int pos = line.find(std::string(get_app_binary_name(app_binary))); - const char *ptx_file = line.c_str(); int pos1 = line.find("sm_"); int pos2 = line.find_last_of("."); if (pos1==std::string::npos&&pos2==std::string::npos){ @@ -2499,11 +2498,10 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){ * */ void cuda_runtime_api::extract_code_using_cuobjdump(){ CUctx_st *context = GPGPUSim_Context(); - unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); //prevent the dumping by cuobjdump everytime we execute the code! const char *override_cuobjdump = getenv("CUOBJDUMP_SIM_FILE"); - char command[1000], ptx_file[1000]; + char command[1000]; std::string app_binary = get_app_binary(); //Running cuobjdump using dynamic link to current process snprintf(command,1000,"md5sum %s ", app_binary.c_str()); diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index d13b8c6..1982e04 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -519,10 +519,10 @@ private: int checkpoint_option; int checkpoint_kernel; int checkpoint_CTA; - int resume_option; - int resume_kernel; - int resume_CTA; - int checkpoint_CTA_t; + unsigned resume_option; + unsigned resume_kernel; + unsigned resume_CTA; + unsigned checkpoint_CTA_t; int checkpoint_insn_Y; int g_ptx_inst_debug_to_file; char* g_ptx_inst_debug_file; @@ -540,10 +540,10 @@ public: int checkpoint_option; int checkpoint_kernel; int checkpoint_CTA; - int resume_option; - int resume_kernel; - int resume_CTA; - int checkpoint_CTA_t; + unsigned resume_option; + unsigned resume_kernel; + unsigned resume_CTA; + unsigned checkpoint_CTA_t; int checkpoint_insn_Y; //Move some cycle core stats here instead of being global @@ -992,7 +992,7 @@ public: printf("Printing mem access generated\n"); std::list::iterator it; for (it = m_accessq.begin(); it != m_accessq.end(); ++it){ - printf("MEM_TXN_GEN:%s:%x, Size:%d \n",mem_access_type_str(it->get_type()), it->get_addr(),it->get_size()); + printf("MEM_TXN_GEN:%s:%llx, Size:%d \n",mem_access_type_str(it->get_type()), it->get_addr(),it->get_size()); } } } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index b9e6552..f8d0b3e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1491,7 +1491,6 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th const operand_info &src1 = pI->src1(); //the name of the texture std::string texname = src1.name(); - gpgpu_t *gpu = thread->get_gpu(); /* For programs with many streams, textures can be bound and unbound asynchronously. This means we need to use the kernel's "snapshot" of @@ -1577,7 +1576,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) } //Tensorcore is warp synchronous operation. So these instructions needs to be executed only once. To make the simulation faster removing the redundant tensorcore operation - if(!tensorcore_op(inst_opcode)||(tensorcore_op(inst_opcode))&&(lane_id==0)){ + if(!tensorcore_op(inst_opcode)||((tensorcore_op(inst_opcode))&&(lane_id==0))){ switch ( inst_opcode ) { #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break; #define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,get_core(),inst); op_classification = CLASSIFICATION; break; @@ -2141,13 +2140,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL unsigned max_cta_tot = max_cta(kernel_info,kernel.threads_per_cta(), g_the_gpu()->getShaderCoreConfig()->warp_size, g_the_gpu()->getShaderCoreConfig()->n_thread_per_shader, g_the_gpu()->getShaderCoreConfig()->gpgpu_shmem_size, g_the_gpu()->getShaderCoreConfig()->gpgpu_shader_registers, g_the_gpu()->getShaderCoreConfig()->max_cta_per_core); printf("Max CTA : %d\n",max_cta_tot); - - - - - int inst_count=50; int cp_op= g_the_gpu()->checkpoint_option; - int cp_CTA = g_the_gpu()->checkpoint_CTA; int cp_kernel= g_the_gpu()->checkpoint_kernel; cp_count= g_the_gpu()->checkpoint_insn_Y; cp_cta_resume= g_the_gpu()->checkpoint_CTA_t; @@ -2184,7 +2177,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL { char f1name[2048]; snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", kernel.get_uid() ); - g_checkpoint->store_global_mem(g_the_gpu()->get_global_memory(), f1name , "%08x"); + g_checkpoint->store_global_mem(g_the_gpu()->get_global_memory(), f1name , (char *)"%08x"); } @@ -2312,18 +2305,15 @@ void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp) checkpoint *g_checkpoint; g_checkpoint = new checkpoint(); - symbol * sym; ptx_reg_t regval; regval.u64= 123; - symbol_table * symtab= m_kernel->entry()->get_symtab(); - unsigned ctaid =m_kernel->get_next_cta_id_single(); if(m_gpu->checkpoint_option==1 && (m_kernel->get_uid()==m_gpu->checkpoint_kernel) && (ctaid_cp>=m_gpu->checkpoint_CTA) && (ctaid_cpcheckpoint_CTA_t)) { char fname[2048]; snprintf(fname,2048,"checkpoint_files/shared_mem_%d.txt",ctaid-1 ); - g_checkpoint->store_global_mem(m_thread[0]->m_shared_mem, fname , "%08x"); + g_checkpoint->store_global_mem(m_thread[0]->m_shared_mem, fname , (char *)"%08x"); for(int i=0; i<32*m_warp_count;i++) { char fname[2048]; @@ -2331,7 +2321,7 @@ void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp) m_thread[i]->print_reg_thread(fname); char f1name[2048]; snprintf(f1name,2048,"checkpoint_files/local_mem_thread_%d_%d_reg.txt",i,ctaid-1 ); - g_checkpoint->store_global_mem(m_thread[i]->m_local_mem, f1name , "%08x"); + g_checkpoint->store_global_mem(m_thread[i]->m_local_mem, f1name , (char *)"%08x"); m_thread[i]->set_done(); m_thread[i]->exitCore(); m_thread[i]->registerExit(); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 58a077e..a44b03f 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -203,7 +203,7 @@ void ptx_thread_info::print_reg_thread(char * fname) const std::string &name = it->first->name(); const std::string &dec= it->first->decl_location(); unsigned size = it->first->get_size_in_bytes(); - fprintf(fp,"%s %llu %s %d\n",name.c_str(),it->second, dec.c_str(),size ); + fprintf(fp,"%s %llu %s %d\n", name.c_str(), it->second, dec.c_str(), size); } //m_regs.pop_back(); @@ -232,7 +232,6 @@ void ptx_thread_info::resume_reg_thread(char * fname, symbol_table * symtab) pch = strtok (NULL," "); data = atoi(pch); pch = strtok (NULL," "); - char * decl= pch; pch = strtok (NULL," "); size = atoi(pch); @@ -1819,9 +1818,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) ptx_reg_t matrix_d[16][16]; ptx_reg_t src_data; ptx_thread_info *thread; - int stride; - unsigned wmma_type = pI->get_wmma_type(); unsigned a_layout = pI->get_wmma_layout(0); unsigned b_layout = pI->get_wmma_layout(1); unsigned type = pI->get_type(); @@ -1833,7 +1830,6 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) tid= inst.warp_id_func()*core->get_warp_size(); else tid= inst.warp_id()*core->get_warp_size(); - unsigned thread_group_index; float temp; half temp2; @@ -1847,9 +1843,9 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) ptx_reg_t v[8]; thread->get_vector_operand_values( src_a, v, nelem ); if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){ - printf("Thread%d_Iteration=%d\n:",thrd,operand_num); - for(k=0;kset_vector_operand_values(dst,nw_data1,nw_data2,nw_data3,nw_data4); if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) - printf("thread%d=%x,%x,%x,%x",thrd,nw_data1.s64,nw_data2.s64,nw_data3.s64,nw_data4.s64); + printf("thread%d=%llx,%llx,%llx,%llx", thrd, nw_data1.s64, nw_data2.s64, nw_data3.s64, nw_data4.s64); } else{ @@ -2298,9 +2294,8 @@ unsigned int saturatei(unsigned int a, unsigned int max) ptx_reg_t f2x( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, int rounding_mode, int saturation_mode ) { - half mytemp; - float myfloat; - half_float::half tmp_h; + half mytemp; + half_float::half tmp_h; //assert( from_width == 32); enum cudaRoundMode mode = cudaRoundZero; @@ -3085,7 +3080,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) size_t size; unsigned smid; int t; - int thrd,odd,inx,k; + int thrd, odd, inx, k; ptx_thread_info *thread; const operand_info &src = pI->operand_lookup(1); @@ -3105,15 +3100,15 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) _memory_op_t insn_memory_op = pI->has_memory_read() ? memory_load : memory_store; for (thrd=0; thrd < core->get_warp_size(); thrd++) { thread = core->get_thread_info()[tid+thrd]; - odd=thrd%2; - inx=thrd/2; - ptx_reg_t addr_reg = thread->get_operand_value(src1, src, type, thread, 1); + odd= thrd % 2; + inx= thrd / 2; + ptx_reg_t addr_reg = thread->get_operand_value(src1, src, type, thread, 1); ptx_reg_t src2_data = thread->get_operand_value(src2, src, type, thread, 1); const operand_info &src_a= pI->operand_lookup(1); unsigned nelem = src_a.get_vect_nelem(); ptx_reg_t* v= new ptx_reg_t[8]; thread->get_vector_operand_values( src_a, v, nelem ); - stride=src2_data.u32; + stride = src2_data.u32; memory_space_t space = pI->get_space(); @@ -3130,9 +3125,9 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) } decode_space(space,thread,src1,mem,addr); - type_info_key::type_decode(type,size,t); + type_info_key::type_decode(type, size, t); if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) - printf("mma_st: thrd=%d,addr=%x, fp(size=%d), stride=%d\n",thrd,addr_reg.u32,size,src2_data.u32); + printf("mma_st: thrd=%d, addr=%x, fp(size=%zu), stride=%d\n", thrd, addr_reg.u32, size, src2_data.u32); addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8; addr_t push_addr; @@ -3152,7 +3147,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) mem_txn_addr[num_mem_txn++]=push_addr; if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){ - printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,v[0].s64,v[1].s64,v[2].s64,v[3].s64,v[4].s64,v[5].s64,v[6].s64,v[7].s64); + printf("wmma:store:thread%d=%llx,%llx,%llx,%llx,%llx,%llx,%llx,%llx\n",thrd,v[0].s64,v[1].s64,v[2].s64,v[3].s64,v[4].s64,v[5].s64,v[6].s64,v[7].s64); float temp; int l; printf("thread=%d:",thrd); @@ -3179,7 +3174,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) } if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) - printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,nw_v[0].s64,nw_v[1].s64,nw_v[2].s64,nw_v[3].s64,nw_v[4].s64,nw_v[5].s64,nw_v[6].s64,nw_v[7].s64); + printf("wmma:store:thread%d=%llx,%llx,%llx,%llx,%llx,%llx,%llx,%llx\n",thrd,nw_v[0].s64,nw_v[1].s64,nw_v[2].s64,nw_v[3].s64,nw_v[4].s64,nw_v[5].s64,nw_v[6].s64,nw_v[7].s64); } } @@ -3238,11 +3233,11 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) } decode_space(space,thread,src1,mem,addr); - type_info_key::type_decode(type,size,t); + type_info_key::type_decode(type, size, t); ptx_reg_t data[16]; if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) - printf("mma_ld: thrd=%d,addr=%x, fpsize=%d, stride=%d\n",thrd,src1_data.u32,size,src2_data.u32); + printf("mma_ld: thrd=%d,addr=%x, fpsize=%zu, stride=%d\n", thrd, src1_data.u32, size, src2_data.u32); addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8; addr_t fetch_addr; @@ -3341,7 +3336,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) if(type==F16_TYPE){ printf("\nmma_ld:thread%d= ",thrd); for(i=0;i<16;i++){ - printf("%x ",data[i].u64); + printf("%llx ",data[i].u64); } printf("\n"); @@ -3361,7 +3356,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) printf("\n"); printf("\nmma_ld:thread%d= ",thrd); for(i=0;i<8;i++){ - printf("%x ",data[i].u64); + printf("%llx ",data[i].u64); } printf("\n"); } @@ -3388,15 +3383,15 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) else thread->set_wmma_vector_operand_values(dst,nw_data[0],nw_data[1],nw_data[2],nw_data[3],nw_data[4],nw_data[5],nw_data[6],nw_data[7]); if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){ - printf("mma_ld:data[0].s64=%x,data[1].s64=%x,new_data[0].s64=%x\n",data[0].u64,data[1].u64,nw_data[0].u64); - printf("mma_ld:data[2].s64=%x,data[3].s64=%x,new_data[1].s64=%x\n",data[2].u64,data[3].u64,nw_data[1].u64); - printf("mma_ld:data[4].s64=%x,data[5].s64=%x,new_data[2].s64=%x\n",data[4].u64,data[5].u64,nw_data[2].u64); - printf("mma_ld:data[6].s64=%x,data[7].s64=%x,new_data[3].s64=%x\n",data[6].u64,data[7].u64,nw_data[3].u64); + printf("mma_ld:data[0].s64=%llx,data[1].s64=%llx,new_data[0].s64=%llx\n",data[0].u64,data[1].u64,nw_data[0].u64); + printf("mma_ld:data[2].s64=%llx,data[3].s64=%llx,new_data[1].s64=%llx\n",data[2].u64,data[3].u64,nw_data[1].u64); + printf("mma_ld:data[4].s64=%llx,data[5].s64=%llx,new_data[2].s64=%llx\n",data[4].u64,data[5].u64,nw_data[2].u64); + printf("mma_ld:data[6].s64=%llx,data[7].s64=%llx,new_data[3].s64=%llx\n",data[6].u64,data[7].u64,nw_data[3].u64); if(wmma_type!=LOAD_C){ - printf("mma_ld:data[8].s64=%x,data[9].s64=%x,new_data[4].s64=%x\n",data[8].u64,data[9].u64,nw_data[4].s64); - printf("mma_ld:data[10].s64=%x,data[11].s64=%x,new_data[5].s64=%x\n",data[10].u64,data[11].u64,nw_data[5].u64); - printf("mma_ld:data[12].s64=%x,data[13].s64=%x,new_data[6].s64=%x\n",data[12].u64,data[13].u64,nw_data[6].u64); - printf("mma_ld:data[14].s64=%x,data[15].s64=%x,new_data[7].s64=%x\n",data[14].u64,data[15].u64,nw_data[3].u64); + printf("mma_ld:data[8].s64=%llx,data[9].s64=%llx,new_data[4].s64=%llx\n",data[8].u64,data[9].u64,nw_data[4].s64); + printf("mma_ld:data[10].s64=%llx,data[11].s64=%llx,new_data[5].s64=%llx\n",data[10].u64,data[11].u64,nw_data[5].u64); + printf("mma_ld:data[12].s64=%llx,data[13].s64=%llx,new_data[6].s64=%llx\n",data[12].u64,data[13].u64,nw_data[6].u64); + printf("mma_ld:data[14].s64=%llx,data[15].s64=%llx,new_data[7].s64=%llx\n",data[14].u64,data[15].u64,nw_data[3].u64); } } } @@ -4132,9 +4127,9 @@ int prmt_mode_present(int mode) } return returnval; } -int read_byte(int mode,int control,int d_sel_index,signed long long value){ +int read_byte(int mode, int control, int d_sel_index, signed long long value){ - int returnval; + int returnval = 0; int prmt_f4e_mode[4][4]={{0,1,2,3},{1,2,3,4},{2,3,4,5},{3,4,5,6}}; int prmt_b4e_mode[4][4]={{0,7,6,5},{1,0,7,6},{2,1,0,7},{3,2,1,0}}; int prmt_rc8_mode[4][4]={{0,0,0,0},{1,1,1,1},{2,2,2,2},{3,3,3,3}}; @@ -4157,11 +4152,12 @@ int read_byte(int mode,int control,int d_sel_index,signed long long value){ case PRMT_RC8_MODE: returnval=prmt_rc8_mode[control][d_sel_index];break; case PRMT_ECL_MODE: returnval=prmt_ecl_mode[control][d_sel_index];break; case PRMT_ECR_MODE: returnval=prmt_ecr_mode[control][d_sel_index];break; - case PRMT_RC16_MODE: returnval=prmt_rc16_mode[control][d_sel_index];break; - default: printf("ERROR\n");break; + case PRMT_RC16_MODE: returnval=prmt_rc16_mode[control][d_sel_index];break; + // Change the default from printing "ERROR" to just asserting + default: assert(false); } } - return (returnval<<8*d_sel_index); + return (returnval << 8 * d_sel_index); } void prmt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 81b70af..a4f4a0c 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -421,7 +421,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi assert( (num_bits%8) == 0 ); addr = g_current_symbol_table->get_sstarr_next(); addr_pad = pad_address(addr, num_bits/8, 128); - printf("from 0x%x to 0x%lx (sstarr memory space)\n", + printf("from 0x%llx to 0x%llx (sstarr memory space)\n", addr+addr_pad, addr+addr_pad + num_bits/8); fflush(stdout); 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/gpu-sim.cc') 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