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/abstract_hardware_model.cc | 22 ++++++++++++++++++++++ src/abstract_hardware_model.h | 1 + src/gpgpu-sim/gpu-sim.cc | 10 ++++++++++ src/gpgpu-sim/gpu-sim.h | 2 ++ 4 files changed, 35 insertions(+) (limited to 'src') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index cebdb25..023f51b 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -691,6 +691,28 @@ unsigned g_kernel_launch_latency; unsigned kernel_info_t::m_next_uid = 1; +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry) +{ + m_kernel_entry=entry; + m_grid_dim=gridDim; + m_block_dim=blockDim; + m_next_cta.x=0; + m_next_cta.y=0; + m_next_cta.z=0; + m_next_tid=m_next_cta; + m_num_cores_running=0; + m_uid = m_next_uid++; + m_param_mem = new memory_space_impl<8192>("param",64*1024); + + //Jin: parent and child kernel management for CDP + m_parent_kernel = NULL; + + //Jin: launch latency management + m_launch_latency = g_kernel_launch_latency; + + volta_cache_config_set=false; +} + /*A snapshot of the texture mappings needs to be stored in the kernel's info as kernels should use the texture bindings seen at the time of launch and textures can be bound/unbound asynchronously with respect to streams. */ diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 22ef509..64bbaa2 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -212,6 +212,7 @@ public: // m_num_cores_running=0; // m_param_mem=NULL; // } + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry); kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTextureInfo); ~kernel_info_t(); 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; diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index c8dad89..7336cac 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -455,6 +455,8 @@ public: int num_registers_per_block() const; int wrp_size() const; int shader_clock() const; + int max_cta_per_core() const; + int get_max_cta( const kernel_info_t &k ) const; const struct cudaDeviceProp *get_prop() const; enum divergence_support_t simd_model() const; -- cgit v1.3 From 556191bcbe6cfb2a8a1008166aa51b7a8567be82 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Mon, 10 Jun 2019 11:29:38 -0400 Subject: adding new wmma instruction cogif --- src/cuda-sim/opcodes.h | 6 +++++- src/cuda-sim/ptx.l | 2 ++ src/cuda-sim/ptx_ir.cc | 2 ++ 3 files changed, 9 insertions(+), 1 deletion(-) (limited to 'src') diff --git a/src/cuda-sim/opcodes.h b/src/cuda-sim/opcodes.h index b91d92f..86d3b99 100644 --- a/src/cuda-sim/opcodes.h +++ b/src/cuda-sim/opcodes.h @@ -68,6 +68,10 @@ enum wmma_type{ MMA, ROW, COL, - M16N16K16 + M16N16K16, + M32N8K16, + M8N32K16 + + }; #endif diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 3232361..36a8a4d 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -169,6 +169,8 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; \.row TC; ptx_lval.int_value = ROW; return LAYOUT; \.col TC; ptx_lval.int_value = COL; return LAYOUT; \.m16n16k16 TC; ptx_lval.int_value = M16N16K16; return CONFIGURATION; +\.m32n8k16 TC; ptx_lval.int_value = M32N8K16; return CONFIGURATION; +\.m8n32k16 TC; ptx_lval.int_value = M8N32K16; return CONFIGURATION; \.f4e TC; return PRMT_F4E_MODE; \.b4e TC; return PRMT_B4E_MODE; \.rc8 TC; return PRMT_RC8_MODE; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 73db24d..4ad9ddf 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1144,6 +1144,8 @@ ptx_instruction::ptx_instruction( int opcode, m_wmma_layout[rr++]=last_ptx_inst_option; break; case M16N16K16: + case M32N8K16: + case M8N32K16: break; default: assert(0); -- cgit v1.3 From fe58efe9c8ca38f7d0f3781e54b04bc526bdfd07 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Mon, 10 Jun 2019 18:06:42 -0400 Subject: fixing thrust error --- libcuda/cuda_runtime_api.cc | 3 +++ src/stream_manager.h | 2 ++ 2 files changed, 5 insertions(+) (limited to 'src') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index fb3e07a..17a5c96 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1498,6 +1498,9 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); kernel_config config = g_cuda_launch_stack.back(); struct CUstream_st *stream = config.get_stream(); + if(g_stream_manager->is_blocking()) + stream = 0; + printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun, g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); diff --git a/src/stream_manager.h b/src/stream_manager.h index 91d1b36..3fbdbaf 100644 --- a/src/stream_manager.h +++ b/src/stream_manager.h @@ -258,6 +258,8 @@ public: void pushCudaStreamWaitEventToAllStreams( CUevent_st *e, unsigned int flags ); bool operation(bool * sim); void stop_all_running_kernels(); + unsigned size() {return m_streams.size(); }; + bool is_blocking() {return m_cuda_launch_blocking; }; private: void print_impl( FILE *fp); -- cgit v1.3 From 8b276989af2f82ffc51306016409e01156678273 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Tue, 18 Jun 2019 19:23:06 -0400 Subject: fixing shuffle inst bug --- src/cuda-sim/cuda-sim.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index f7bb9cc..4c0fc58 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -873,7 +873,7 @@ void ptx_instruction::set_opcode_and_latency() op=TENSOR_CORE_OP; break; case SHFL_OP: - latency = 32; + latency = 4; initiation_interval = 4; break; default: -- cgit v1.3 From f23a5de2c7eec680fc8f5c6ba45c64fcd9544e65 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Sun, 23 Jun 2019 15:22:54 -0400 Subject: fixing the buffer limit for function names --- libcuda/cuda_runtime_api.cc | 4 ++++ src/cuda-sim/ptx_ir.cc | 4 ++-- 2 files changed, 6 insertions(+), 2 deletions(-) (limited to 'src') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 17a5c96..97396ac 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -987,6 +987,10 @@ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int * context->get_device()->get_gpgpu()->get_config().num_shader()); dim3 blockDim(blockSize); kernel_info_t result(gridDim, blockDim, entry); + //if(entry == NULL){ + // *numBlocks = 1; + // return g_last_cudaError = cudaErrorUnknown; + //} *numBlocks = context->get_device()->get_gpgpu()->get_max_cta(result); printf("Maximum size is %d with gridDim %d and blockDim %d\n", *numBlocks, gridDim.x, blockDim.x); return g_last_cudaError = cudaSuccess; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 4ad9ddf..c4d5a6c 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -177,8 +177,8 @@ void symbol_table::add_function( function_info *func, const char *filename, unsi //Jin: handle instruction group for cdp symbol_table* symbol_table::start_inst_group() { - char inst_group_name[1024]; - snprintf(inst_group_name, 1024, "%s_inst_group_%u", m_scope_name.c_str(), m_inst_group_id); + char inst_group_name[2048]; + snprintf(inst_group_name, 2048, "%s_inst_group_%u", m_scope_name.c_str(), m_inst_group_id); //previous added assert(m_inst_group_symtab.find(std::string(inst_group_name)) == m_inst_group_symtab.end()); -- cgit v1.3 From b40c38d7e4d711b818718ae6e291ce6c108392ec Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Thu, 27 Jun 2019 18:15:28 -0400 Subject: fixing the booksim id buffer flip assertion --- src/intersim2/flit.hpp | 4 ++-- src/intersim2/gputrafficmanager.cpp | 6 +++--- src/intersim2/routers/iq_router.cpp | 2 +- src/intersim2/stats.hpp | 3 +++ src/intersim2/trafficmanager.cpp | 8 ++++---- src/intersim2/trafficmanager.hpp | 14 +++++++------- 6 files changed, 20 insertions(+), 17 deletions(-) (limited to 'src') diff --git a/src/intersim2/flit.hpp b/src/intersim2/flit.hpp index fd48306..1c58c68 100644 --- a/src/intersim2/flit.hpp +++ b/src/intersim2/flit.hpp @@ -57,8 +57,8 @@ public: int itime; int atime; - int id; - int pid; + unsigned long long id; + unsigned long long pid; bool record; diff --git a/src/intersim2/gputrafficmanager.cpp b/src/intersim2/gputrafficmanager.cpp index bf422d6..6897a22 100644 --- a/src/intersim2/gputrafficmanager.cpp +++ b/src/intersim2/gputrafficmanager.cpp @@ -105,7 +105,7 @@ void GPUTrafficManager::_RetireFlit( Flit *f, int dest ) if(f->head) { head = f; } else { - map::iterator iter = _retired_packets[f->cl].find(f->pid); + map::iterator iter = _retired_packets[f->cl].find(f->pid); assert(iter != _retired_packets[f->cl].end()); head = iter->second; _retired_packets[f->cl].erase(iter); @@ -195,8 +195,8 @@ void GPUTrafficManager::_GeneratePacket(int source, int stype, int cl, int time, // Flit::FlitType packet_type = Flit::ANY_TYPE; int size = packet_size; //input size - int pid = _cur_pid++; - assert(_cur_pid); + unsigned long long pid = _cur_pid++; + assert(_cur_pid > 0); int packet_destination = dest; bool record = false; bool watch = gWatchOut && (_packets_to_watch.count(pid) > 0); diff --git a/src/intersim2/routers/iq_router.cpp b/src/intersim2/routers/iq_router.cpp index d97f485..7dffb3a 100644 --- a/src/intersim2/routers/iq_router.cpp +++ b/src/intersim2/routers/iq_router.cpp @@ -306,7 +306,7 @@ bool IQRouter::_ReceiveFlits( ) if(f->watch) { *gWatchOut << GetSimTime() << " | " << FullName() << " | " - << "Received flit " << f->id + << "Received flit " << (unsigned) f->id << " from channel at input " << input << "." << endl; } diff --git a/src/intersim2/stats.hpp b/src/intersim2/stats.hpp index 1aaf013..e186f4d 100644 --- a/src/intersim2/stats.hpp +++ b/src/intersim2/stats.hpp @@ -62,6 +62,9 @@ public: inline void AddSample( int val ) { AddSample( (double)val ); } + inline void AddSample( unsigned long long val ) { + AddSample( (double)val ); + } int GetBin(int b){ return _hist[b];} diff --git a/src/intersim2/trafficmanager.cpp b/src/intersim2/trafficmanager.cpp index 8a015bb..7a20d07 100644 --- a/src/intersim2/trafficmanager.cpp +++ b/src/intersim2/trafficmanager.cpp @@ -679,7 +679,7 @@ void TrafficManager::_RetireFlit( Flit *f, int dest ) if(f->head) { head = f; } else { - map::iterator iter = _retired_packets[f->cl].find(f->pid); + map::iterator iter = _retired_packets[f->cl].find(f->pid); assert(iter != _retired_packets[f->cl].end()); head = iter->second; _retired_packets[f->cl].erase(iter); @@ -1380,7 +1380,7 @@ void TrafficManager::_DisplayRemaining( ostream & os ) const { for(int c = 0; c < _classes; ++c) { - map::const_iterator iter; + map::const_iterator iter; int i; os << "Class " << c << ":" << endl; @@ -1463,7 +1463,7 @@ bool TrafficManager::_SingleSim( ) double latency = (double)_plat_stats[c]->Sum(); double count = (double)_plat_stats[c]->NumSamples(); - map::const_iterator iter; + map::const_iterator iter; for(iter = _total_in_flight_flits[c].begin(); iter != _total_in_flight_flits[c].end(); iter++) { @@ -1568,7 +1568,7 @@ bool TrafficManager::_SingleSim( ) double acc_latency = _plat_stats[c]->Sum(); double acc_count = (double)_plat_stats[c]->NumSamples(); - map::const_iterator iter; + map::const_iterator iter; for(iter = _total_in_flight_flits[c].begin(); iter != _total_in_flight_flits[c].end(); iter++) { diff --git a/src/intersim2/trafficmanager.hpp b/src/intersim2/trafficmanager.hpp index 9694df4..97564ea 100644 --- a/src/intersim2/trafficmanager.hpp +++ b/src/intersim2/trafficmanager.hpp @@ -113,9 +113,9 @@ protected: vector > _qdrained; vector > > _partial_packets; - vector > _total_in_flight_flits; - vector > _measured_in_flight_flits; - vector > _retired_packets; + vector > _total_in_flight_flits; + vector > _measured_in_flight_flits; + vector > _retired_packets; bool _empty_network; bool _hold_switch_for_packet; @@ -229,12 +229,12 @@ protected: vector _warmup_threshold; vector _acc_warmup_threshold; - int _cur_id; - int _cur_pid; + unsigned long long _cur_id; + unsigned long long _cur_pid; int _time; - set _flits_to_watch; - set _packets_to_watch; + set _flits_to_watch; + set _packets_to_watch; bool _print_csv_results; -- cgit v1.3 From c9fb34c320a5adbd9649a410b417be70937345ad Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Thu, 27 Jun 2019 18:40:57 -0400 Subject: increase the function buffer again --- src/cuda-sim/ptx_ir.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'src') diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index c4d5a6c..e4e0c09 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -177,8 +177,8 @@ void symbol_table::add_function( function_info *func, const char *filename, unsi //Jin: handle instruction group for cdp symbol_table* symbol_table::start_inst_group() { - char inst_group_name[2048]; - snprintf(inst_group_name, 2048, "%s_inst_group_%u", m_scope_name.c_str(), m_inst_group_id); + char inst_group_name[4096]; + snprintf(inst_group_name, 4096, "%s_inst_group_%u", m_scope_name.c_str(), m_inst_group_id); //previous added assert(m_inst_group_symtab.find(std::string(inst_group_name)) == m_inst_group_symtab.end()); -- 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') 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 f9ca59092e3270106f7fa9e2b4273a9cce30610b Mon Sep 17 00:00:00 2001 From: "Mahmoud Khairy A. Abdallah" Date: Fri, 12 Jul 2019 21:35:18 -0400 Subject: fixing hguh gcc iostream exception --- src/option_parser.cc | 1 + 1 file changed, 1 insertion(+) (limited to 'src') diff --git a/src/option_parser.cc b/src/option_parser.cc index 4fa2343..497316f 100644 --- a/src/option_parser.cc +++ b/src/option_parser.cc @@ -101,6 +101,7 @@ public: try { ss >> m_variable; } catch (stringstream::failure &e) { + } catch (exception &e) { return false; } m_isParsed = true; -- cgit v1.3 From ba7145d18dd65391e9510664d26c963fcaab20ee Mon Sep 17 00:00:00 2001 From: "Mahmoud Khairy A. Abdallah" Date: Sun, 14 Jul 2019 10:58:39 -0400 Subject: fixing the volta local memory size per thread --- src/abstract_hardware_model.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) (limited to 'src') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 64bbaa2..27a1ba6 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -438,9 +438,9 @@ protected: #define GLOBAL_HEAP_START 0xC0000000 // start allocating from this address (lower values used for allocating globals in .ptx file) -#define SHARED_MEM_SIZE_MAX (64*1024) -#define LOCAL_MEM_SIZE_MAX (8*1024) -#define MAX_STREAMING_MULTIPROCESSORS 64 +#define SHARED_MEM_SIZE_MAX (96*1024) +#define LOCAL_MEM_SIZE_MAX (16*1024) +#define MAX_STREAMING_MULTIPROCESSORS 80 //scale it to Volta #define MAX_THREAD_PER_SM 2048 #define MAX_WARP_PER_SM 64 #define TOTAL_LOCAL_MEM_PER_SM (MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX) -- 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') 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') 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 2f5b3332c9b9b3fa9fea43d61276bddb24aa7df2 Mon Sep 17 00:00:00 2001 From: shen203 <52721142+shen203@users.noreply.github.com> Date: Tue, 30 Jul 2019 15:48:32 -0400 Subject: changing x-bar algorithm to islip --- src/gpgpu-sim/local_interconnect.cc | 58 +++++++++++++++++++++++++++++++++++-- src/gpgpu-sim/local_interconnect.h | 3 +- 2 files changed, 58 insertions(+), 3 deletions(-) (limited to 'src') diff --git a/src/gpgpu-sim/local_interconnect.cc b/src/gpgpu-sim/local_interconnect.cc index 66d6648..1416b2c 100644 --- a/src/gpgpu-sim/local_interconnect.cc +++ b/src/gpgpu-sim/local_interconnect.cc @@ -45,7 +45,8 @@ xbar_router::xbar_router(unsigned router_id, enum Interconnect_type m_type, unsi total_nodes = n_shader+n_mem; in_buffers.resize(total_nodes); out_buffers.resize(total_nodes); - next_node=0; + next_node.resize(total_nodes,0); +// next_node = 0; in_buffer_limit = m_in_buffer_limit; out_buffer_limit = m_out_buffer_limit; if(m_type == REQ_NET) { @@ -108,7 +109,7 @@ bool xbar_router::Has_Buffer_In(unsigned input_deviceID, unsigned size, bool upd bool xbar_router::Has_Buffer_Out(unsigned output_deviceID, unsigned size){ return (out_buffers[output_deviceID].size() + size <= out_buffer_limit); } - +/* void xbar_router::Advance() { cycles++; @@ -142,6 +143,59 @@ void xbar_router::Advance() { out_buffer_util+=out_buffers[i].size(); } } +*/ + +void xbar_router::Advance() { + cycles++; + + vector node_tmp; + + for (unsigned i=0; i > out_buffers; unsigned _n_shader, _n_mem, total_nodes; unsigned in_buffer_limit, out_buffer_limit; - unsigned next_node; + vector next_node; +// unsigned next_node; unsigned m_id; enum Interconnect_type router_type; unsigned active_in_buffers,active_out_buffers; -- cgit v1.3 From 56c52cf6c4b369e9fd05759e9b16ea37ff6e332c Mon Sep 17 00:00:00 2001 From: shen203 <52721142+shen203@users.noreply.github.com> Date: Tue, 30 Jul 2019 15:48:32 -0400 Subject: changing x-bar algorithm to islip --- src/gpgpu-sim/local_interconnect.cc | 58 +++++++++++++++++++++++++++++++++++-- src/gpgpu-sim/local_interconnect.h | 3 +- 2 files changed, 58 insertions(+), 3 deletions(-) (limited to 'src') diff --git a/src/gpgpu-sim/local_interconnect.cc b/src/gpgpu-sim/local_interconnect.cc index 66d6648..1416b2c 100644 --- a/src/gpgpu-sim/local_interconnect.cc +++ b/src/gpgpu-sim/local_interconnect.cc @@ -45,7 +45,8 @@ xbar_router::xbar_router(unsigned router_id, enum Interconnect_type m_type, unsi total_nodes = n_shader+n_mem; in_buffers.resize(total_nodes); out_buffers.resize(total_nodes); - next_node=0; + next_node.resize(total_nodes,0); +// next_node = 0; in_buffer_limit = m_in_buffer_limit; out_buffer_limit = m_out_buffer_limit; if(m_type == REQ_NET) { @@ -108,7 +109,7 @@ bool xbar_router::Has_Buffer_In(unsigned input_deviceID, unsigned size, bool upd bool xbar_router::Has_Buffer_Out(unsigned output_deviceID, unsigned size){ return (out_buffers[output_deviceID].size() + size <= out_buffer_limit); } - +/* void xbar_router::Advance() { cycles++; @@ -142,6 +143,59 @@ void xbar_router::Advance() { out_buffer_util+=out_buffers[i].size(); } } +*/ + +void xbar_router::Advance() { + cycles++; + + vector node_tmp; + + for (unsigned i=0; i > out_buffers; unsigned _n_shader, _n_mem, total_nodes; unsigned in_buffer_limit, out_buffer_limit; - unsigned next_node; + vector next_node; +// unsigned next_node; unsigned m_id; enum Interconnect_type router_type; unsigned active_in_buffers,active_out_buffers; -- cgit v1.3 From bd14ce38470dfc54c690db09f00ee5c18b577575 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Fri, 23 Aug 2019 12:26:46 -0400 Subject: fixing CUDA 10 fail --- src/gpgpu-sim/gpu-sim.h | 3 --- src/gpgpu-sim/l2cache.cc | 2 +- 2 files changed, 1 insertion(+), 4 deletions(-) (limited to 'src') diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index f841bf9..76c7a06 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -298,9 +298,6 @@ class memory_config { gpgpu_context* gpgpu_ctx; }; -// 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 bool g_interactive_debugger_enabled; diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc index f1672f9..39a5812 100644 --- a/src/gpgpu-sim/l2cache.cc +++ b/src/gpgpu-sim/l2cache.cc @@ -257,7 +257,7 @@ void memory_partition_unit::simple_dram_model_cycle() d.req = mf; d.ready_cycle = m_gpu->gpu_sim_cycle+m_gpu->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); + mf->set_status(IN_PARTITION_DRAM_LATENCY_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); m_arbitration_metadata.borrow_credit(spid); break; // the DRAM should only accept one request per cycle } -- cgit v1.3 From c2a1e3a668f9a88239184e13460f7e1725b15c90 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Mon, 26 Aug 2019 12:17:52 -0400 Subject: Banked L1, adding iSLIP and RR arbiteratio and adding some comments --- configs/tested-cfgs/SM7_QV100/gpgpusim.config | 11 +- configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 13 +- src/gpgpu-sim/addrdec.cc | 2 + src/gpgpu-sim/gpu-cache.cc | 8 ++ src/gpgpu-sim/gpu-cache.h | 4 + src/gpgpu-sim/icnt_wrapper.cc | 2 + src/gpgpu-sim/local_interconnect.cc | 40 ++++-- src/gpgpu-sim/shader.cc | 168 +++++++++++++------------ src/gpgpu-sim/shader.h | 2 +- 9 files changed, 152 insertions(+), 98 deletions(-) (limited to 'src') diff --git a/configs/tested-cfgs/SM7_QV100/gpgpusim.config b/configs/tested-cfgs/SM7_QV100/gpgpusim.config index f70ad55..f807e11 100644 --- a/configs/tested-cfgs/SM7_QV100/gpgpusim.config +++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config @@ -87,9 +87,10 @@ # if the assigned shd mem = 0, then L1 cache = 128KB # For more info, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x # disable this mode in case of multi kernels/apps execution --adaptive_volta_cache_config 1 -# Volta unified cache has four ports --mem_unit_ports 4 +-adaptive_cache_config 1 +# Volta unified cache has four banks +-l1_banks 4 +#-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 @@ -97,8 +98,8 @@ -gmem_skip_L1D 0 -icnt_flit_size 40 -gpgpu_n_cluster_ejection_buffer_size 32 --l1_latency 80 --smem_latency 19 +-l1_latency 20 +-smem_latency 20 -gpgpu_flush_l1_cache 1 # 32 sets, each 128 bytes 24-way for each memory sub partition (96 KB per memory sub partition). This gives us 6MB L2 cache diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config index 786e75e..f8e4afe 100644 --- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config +++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config @@ -88,8 +88,9 @@ # For more info, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x # disable this mode in case of multi kernels/apps execution -adaptive_volta_cache_config 1 -# Volta unified cache has four ports --mem_unit_ports 4 +# Volta unified cache has four banks +-l1_banks 4 +#-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 @@ -97,8 +98,8 @@ -gmem_skip_L1D 0 -icnt_flit_size 40 -gpgpu_n_cluster_ejection_buffer_size 32 --l1_latency 28 --smem_latency 19 +-l1_latency 20 +-smem_latency 20 -gpgpu_flush_l1_cache 1 # 32 sets, each 128 bytes 24-way for each memory sub partition (96 KB per memory sub partition). This gives us 4.5MB L2 cache @@ -106,7 +107,7 @@ -gpgpu_cache:dl2_texture_only 0 -gpgpu_dram_partition_queues 64:64:64:64 -perf_sim_memcpy 1 --memory_partition_indexing 0 +-memory_partition_indexing 4 # 128 KB Inst. -gpgpu_cache:il1 N:64:128:16,L:R:f:N:L,S:2:48,4 @@ -144,7 +145,7 @@ # "-network_mode 2 -inct_in_buffer_limit 512 -inct_out_buffer_limit 512 -inct_subnets 2" # memory partition latency config --rop_latency 120 +-rop_latency 160 -dram_latency 100 # dram model config diff --git a/src/gpgpu-sim/addrdec.cc b/src/gpgpu-sim/addrdec.cc index ca88ec9..b0db034 100644 --- a/src/gpgpu-sim/addrdec.cc +++ b/src/gpgpu-sim/addrdec.cc @@ -167,6 +167,8 @@ void linear_to_raw_address_translation::addrdec_tlx(new_addr_type addr, addrdec_ } case RANDOM: { + //This is an unrealistic hashing using software hashtable + //we generate a random set for each memory address and save the value in a big hashtable for future reuse new_addr_type chip_address = (addr>>ADDR_CHIP_S); tr1_hash_map::const_iterator got = address_random_interleaving.find (chip_address); if ( got == address_random_interleaving.end() ) { diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc index 370f6e6..db9701d 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -63,6 +63,14 @@ const char * cache_fail_status_str(enum cache_reservation_fail_reason status) return static_cache_reservation_fail_reason_str[status]; } +unsigned l1d_cache_config::set_bank(new_addr_type addr) const{ + + if(m_cache_type == SECTOR) + return (addr >> m_sector_sz_log2) & (l1_banks-1); + else + return (addr >> m_line_sz_log2) & (l1_banks-1); +} + unsigned l1d_cache_config::set_index(new_addr_type addr) const{ unsigned set_index = m_nset; // Default to linear set index function unsigned lower_xor = 0; diff --git a/src/gpgpu-sim/gpu-cache.h b/src/gpgpu-sim/gpu-cache.h index 337f710..90adbb5 100644 --- a/src/gpgpu-sim/gpu-cache.h +++ b/src/gpgpu-sim/gpu-cache.h @@ -584,6 +584,7 @@ public: m_nset_log2 = LOGB2(m_nset); m_valid = true; m_atom_sz = (m_cache_type == SECTOR)? SECTOR_SIZE : m_line_sz; + m_sector_sz_log2 = LOGB2(SECTOR_SIZE); original_m_assoc = m_assoc; //For more details about difference between FETCH_ON_WRITE and WRITE VALIDAE policies @@ -734,6 +735,7 @@ protected: unsigned m_nset_log2; unsigned m_assoc; unsigned m_atom_sz; + unsigned m_sector_sz_log2; unsigned original_m_assoc; bool m_is_streaming; @@ -775,7 +777,9 @@ class l1d_cache_config : public cache_config{ public: l1d_cache_config() : cache_config(){} virtual unsigned set_index(new_addr_type addr) const; + unsigned set_bank(new_addr_type addr) const; unsigned l1_latency; + unsigned l1_banks; }; class l2_cache_config : public cache_config { diff --git a/src/gpgpu-sim/icnt_wrapper.cc b/src/gpgpu-sim/icnt_wrapper.cc index 6e0950c..67724d0 100644 --- a/src/gpgpu-sim/icnt_wrapper.cc +++ b/src/gpgpu-sim/icnt_wrapper.cc @@ -182,6 +182,8 @@ void icnt_reg_options( class OptionParser * opp ) option_parser_register(opp, "-inct_in_buffer_limit", OPT_UINT32, &g_inct_config.in_buffer_limit, "in_buffer_limit", "64"); option_parser_register(opp, "-inct_out_buffer_limit", OPT_UINT32, &g_inct_config.out_buffer_limit, "out_buffer_limit", "64"); option_parser_register(opp, "-inct_subnets", OPT_UINT32, &g_inct_config.subnets, "subnets", "2"); + option_parser_register(opp, "-arbiter_algo", OPT_UINT32, &g_inct_config.arbiter_algo, "arbiter_algo", "1"); + } diff --git a/src/gpgpu-sim/local_interconnect.cc b/src/gpgpu-sim/local_interconnect.cc index 1416b2c..da8a65c 100644 --- a/src/gpgpu-sim/local_interconnect.cc +++ b/src/gpgpu-sim/local_interconnect.cc @@ -36,7 +36,7 @@ #include "local_interconnect.h" #include "mem_fetch.h" -xbar_router::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::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) { m_id=router_id; router_type=m_type; @@ -46,9 +46,10 @@ xbar_router::xbar_router(unsigned router_id, enum Interconnect_type m_type, unsi in_buffers.resize(total_nodes); out_buffers.resize(total_nodes); next_node.resize(total_nodes,0); -// next_node = 0; in_buffer_limit = m_in_buffer_limit; out_buffer_limit = m_out_buffer_limit; + arbit_type = m_arbit_type; + next_node_id=0; if(m_type == REQ_NET) { active_in_buffers=n_shader; active_out_buffers=n_mem; @@ -109,14 +110,25 @@ bool xbar_router::Has_Buffer_In(unsigned input_deviceID, unsigned size, bool upd bool xbar_router::Has_Buffer_Out(unsigned output_deviceID, unsigned size){ return (out_buffers[output_deviceID].size() + size <= out_buffer_limit); } -/* + void xbar_router::Advance() { + + if(arbit_type == NAIVE_RR) + RR_Advance(); + else if(arbit_type == iSLIP) + iSLIP_Advance(); + else + assert(0); + +} + +void xbar_router::RR_Advance() { cycles++; vector issued(total_nodes, false); for(unsigned i=0; i node_tmp; + + //calcaulte how many conflicts are there for stats for (unsigned i=0; i(i), m_n_shader, m_n_mem, m_inct_config.in_buffer_limit, m_inct_config.out_buffer_limit ); + net[i] = new xbar_router( i, static_cast(i), m_n_shader, m_n_mem, m_inct_config.in_buffer_limit, m_inct_config.out_buffer_limit,m_inct_config.arbiter_algo); } } diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index e38eefd..ffd3035 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1645,35 +1645,46 @@ mem_stage_stall_type ldst_unit::process_memory_access_queue_l1cache( l1_cache *c if( inst.accessq_empty() ) return result; - mem_fetch *mf = m_mf_allocator->alloc(inst,inst.accessq_back()); - if(m_config->m_L1D_config.l1_latency > 0) { - if((l1_latency_queue[m_config->m_L1D_config.l1_latency-1]) == NULL) - { - l1_latency_queue[m_config->m_L1D_config.l1_latency-1] = mf; + for(int j=0; jm_L1D_config.l1_banks; j++) { //We can handle at max l1_banks reqs per cycle - if( mf->get_inst().is_store() ) { - unsigned inc_ack = (m_config->m_L1D_config.get_mshr_type() == SECTOR_ASSOC)? - (mf->get_data_size()/SECTOR_SIZE) : 1; + if( inst.accessq_empty() ) + return result; - for(unsigned i=0; i< inc_ack; ++i) - m_core->inc_store_req( inst.warp_id() ); - } + mem_fetch *mf = m_mf_allocator->alloc(inst,inst.accessq_back()); + unsigned bank_id = m_config->m_L1D_config.set_bank(mf->get_addr()); + assert(bank_id < m_config->m_L1D_config.l1_banks); + + if((l1_latency_queue[bank_id][m_config->m_L1D_config.l1_latency-1]) == NULL) + { + l1_latency_queue[bank_id][m_config->m_L1D_config.l1_latency-1] = mf; - inst.accessq_pop_back(); + if( mf->get_inst().is_store() ) { + unsigned inc_ack = (m_config->m_L1D_config.get_mshr_type() == SECTOR_ASSOC)? + (mf->get_data_size()/SECTOR_SIZE) : 1; + + for(unsigned i=0; i< inc_ack; ++i) + m_core->inc_store_req( inst.warp_id() ); + } + + inst.accessq_pop_back(); + } + else + { + result = BK_CONF; + delete mf; + break; //do not try again, just break from the loop and try the next cycle + } } - else - { - result = BK_CONF; - delete mf; - } - if( !inst.accessq_empty() && result !=BK_CONF) + if( !inst.accessq_empty() && result !=BK_CONF) result = COAL_STALL; - return result; + + return result; } else { + mem_fetch *mf = m_mf_allocator->alloc(inst,inst.accessq_back()); std::list events; enum cache_request_status status = cache->access(mf->get_addr(),mf,gpu_sim_cycle+gpu_tot_sim_cycle,events); return process_cache_access( cache, mf->get_addr(), inst, events, mf, status ); @@ -1683,62 +1694,64 @@ mem_stage_stall_type ldst_unit::process_memory_access_queue_l1cache( l1_cache *c void ldst_unit::L1_latency_queue_cycle() { //std::deque< std::pair >::iterator it = m_latency_queue.begin(); - if((l1_latency_queue[0]) != NULL) - { - mem_fetch* mf_next = l1_latency_queue[0]; - std::list events; - enum cache_request_status status = m_L1D->access(mf_next->get_addr(),mf_next,gpu_sim_cycle+gpu_tot_sim_cycle,events); - - bool write_sent = was_write_sent(events); - bool read_sent = was_read_sent(events); - - if ( status == HIT ) { - assert( !read_sent ); - l1_latency_queue[0] = NULL; - if ( mf_next->get_inst().is_load() ) { - for ( unsigned r=0; r < MAX_OUTPUT_VALUES; r++) - if (mf_next->get_inst().out[r] > 0) - { - assert(m_pending_writes[mf_next->get_inst().warp_id()][mf_next->get_inst().out[r]]>0); - unsigned still_pending = --m_pending_writes[mf_next->get_inst().warp_id()][mf_next->get_inst().out[r]]; - if(!still_pending) + for(int j=0; jm_L1D_config.l1_banks; j++) { + if((l1_latency_queue[j][0]) != NULL) + { + mem_fetch* mf_next = l1_latency_queue[j][0]; + std::list events; + enum cache_request_status status = m_L1D->access(mf_next->get_addr(),mf_next,gpu_sim_cycle+gpu_tot_sim_cycle,events); + + bool write_sent = was_write_sent(events); + bool read_sent = was_read_sent(events); + + if ( status == HIT ) { + assert( !read_sent ); + l1_latency_queue[j][0] = NULL; + if ( mf_next->get_inst().is_load() ) { + for ( unsigned r=0; r < MAX_OUTPUT_VALUES; r++) + if (mf_next->get_inst().out[r] > 0) { - m_pending_writes[mf_next->get_inst().warp_id()].erase(mf_next->get_inst().out[r]); - m_scoreboard->releaseRegister(mf_next->get_inst().warp_id(),mf_next->get_inst().out[r]); - m_core->warp_inst_complete(mf_next->get_inst()); + assert(m_pending_writes[mf_next->get_inst().warp_id()][mf_next->get_inst().out[r]]>0); + unsigned still_pending = --m_pending_writes[mf_next->get_inst().warp_id()][mf_next->get_inst().out[r]]; + if(!still_pending) + { + m_pending_writes[mf_next->get_inst().warp_id()].erase(mf_next->get_inst().out[r]); + m_scoreboard->releaseRegister(mf_next->get_inst().warp_id(),mf_next->get_inst().out[r]); + m_core->warp_inst_complete(mf_next->get_inst()); + } } - } - } - - //For write hit in WB policy - if(mf_next->get_inst().is_store() && !write_sent) - { - unsigned dec_ack = (m_config->m_L1D_config.get_mshr_type() == SECTOR_ASSOC)? - (mf_next->get_data_size()/SECTOR_SIZE) : 1; - - mf_next->set_reply(); - - for(unsigned i=0; i< dec_ack; ++i) - m_core->store_ack(mf_next); - } - - if( !write_sent ) - delete mf_next; - - } else if ( status == RESERVATION_FAIL ) { - assert( !read_sent ); - assert( !write_sent ); - } else { - assert( status == MISS || status == HIT_RESERVED ); - l1_latency_queue[0] = NULL; - } - } + } - for( unsigned stage = 0; stagem_L1D_config.l1_latency-1; ++stage) - if( l1_latency_queue[stage] == NULL) { - l1_latency_queue[stage] = l1_latency_queue[stage+1] ; - l1_latency_queue[stage+1] = NULL; - } + //For write hit in WB policy + if(mf_next->get_inst().is_store() && !write_sent) + { + unsigned dec_ack = (m_config->m_L1D_config.get_mshr_type() == SECTOR_ASSOC)? + (mf_next->get_data_size()/SECTOR_SIZE) : 1; + + mf_next->set_reply(); + + for(unsigned i=0; i< dec_ack; ++i) + m_core->store_ack(mf_next); + } + + if( !write_sent ) + delete mf_next; + + } else if ( status == RESERVATION_FAIL ) { + assert( !read_sent ); + assert( !write_sent ); + } else { + assert( status == MISS || status == HIT_RESERVED ); + l1_latency_queue[j][0] = NULL; + } + } + + for( unsigned stage = 0; stagem_L1D_config.l1_latency-1; ++stage) + if( l1_latency_queue[j][stage] == NULL) { + l1_latency_queue[j][stage] = l1_latency_queue[j][stage+1] ; + l1_latency_queue[j][stage+1] = NULL; + } + } } @@ -2117,11 +2130,12 @@ ldst_unit::ldst_unit( mem_fetch_interface *icnt, m_mf_allocator, IN_L1D_MISS_QUEUE ); - if(m_config->m_L1D_config.l1_latency > 0) - { - for(int i=0; im_L1D_config.l1_latency; i++ ) - l1_latency_queue.push_back((mem_fetch*)NULL); - } + l1_latency_queue.resize(m_config->m_L1D_config.l1_banks); + assert(m_config->m_L1D_config.l1_latency > 0); + + for(int j=0; jm_L1D_config.l1_banks; j++ ) + l1_latency_queue[j].resize(m_config->m_L1D_config.l1_latency,(mem_fetch*)NULL); + } m_name = "MEM "; } diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index a0c2b63..05476f4 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1325,7 +1325,7 @@ protected: unsigned long long m_last_inst_gpu_sim_cycle; unsigned long long m_last_inst_gpu_tot_sim_cycle; - std::deque l1_latency_queue; + std::vector> l1_latency_queue; void L1_latency_queue_cycle(); }; -- 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') 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') 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: Thu, 29 Aug 2019 14:18:26 -0400 Subject: Fix a few ambiguity warning --- src/cuda-sim/half.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) (limited to 'src') diff --git a/src/cuda-sim/half.h b/src/cuda-sim/half.h index 8f1a8eb..9f74bb7 100644 --- a/src/cuda-sim/half.h +++ b/src/cuda-sim/half.h @@ -642,10 +642,10 @@ namespace half_float if(exp > 16) { if(R == std::round_toward_infinity) - return hbits | 0x7C00 - (hbits>>15); + return hbits | (0x7C00 - (hbits>>15)); else if(R == std::round_toward_neg_infinity) - return hbits | 0x7BFF + (hbits>>15); - return hbits | 0x7BFF + (R!=std::round_toward_zero); + return hbits | (0x7BFF + (hbits>>15)); + return hbits | (0x7BFF + (R!=std::round_toward_zero)); } if(exp < -13) value = std::ldexp(value, 24); -- cgit v1.3 From 7fbb56d988998d7dc421d949886b5edbcc2fd961 Mon Sep 17 00:00:00 2001 From: tgrogers 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') 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 6f2be094e18a705b320889c7e196609dcccef884 Mon Sep 17 00:00:00 2001 From: tgrogers Date: Sat, 31 Aug 2019 09:17:53 -0400 Subject: Adding a stat for slowdown versus silicon --- src/gpgpusim_entrypoint.cc | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) (limited to 'src') diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 816159f..d3deb24 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -269,7 +269,9 @@ void print_simulation_time() 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) ); - printf("gpgpu_simulation_rate = %u (cycle/sec)\n", (unsigned)(GPGPUsim_ctx_ptr()->g_the_gpu->gpu_tot_sim_cycle / difference) ); + const unsigned cycles_per_sec = (unsigned)(GPGPUsim_ctx_ptr()->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); fflush(stdout); } -- cgit v1.3 From 7bc0007c0dac14528409d33ca1a28491a2568d12 Mon Sep 17 00:00:00 2001 From: Nick Date: Tue, 3 Sep 2019 10:21:51 -0400 Subject: get rid of more UB --- src/abstract_hardware_model.h | 32 +++++++++++++++++++------------- src/cuda-sim/ptx_loader.cc | 3 ++- 2 files changed, 21 insertions(+), 14 deletions(-) (limited to 'src') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index fdc4dc3..29e4a30 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -442,19 +442,25 @@ protected: class gpgpu_sim * m_gpu; }; -#define GLOBAL_HEAP_START 0xC0000000 - // start allocating from this address (lower values used for allocating globals in .ptx file) -#define SHARED_MEM_SIZE_MAX (96*1024) -#define LOCAL_MEM_SIZE_MAX (16*1024) -#define MAX_STREAMING_MULTIPROCESSORS 80 //scale it to Volta -#define MAX_THREAD_PER_SM 2048 -#define MAX_WARP_PER_SM 64 -#define TOTAL_LOCAL_MEM_PER_SM (MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX) -#define TOTAL_SHARED_MEM (MAX_STREAMING_MULTIPROCESSORS*SHARED_MEM_SIZE_MAX) -#define TOTAL_LOCAL_MEM (MAX_STREAMING_MULTIPROCESSORS*MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX) -#define SHARED_GENERIC_START (GLOBAL_HEAP_START-TOTAL_SHARED_MEM) -#define LOCAL_GENERIC_START (SHARED_GENERIC_START-TOTAL_LOCAL_MEM) -#define STATIC_ALLOC_LIMIT (GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM+TOTAL_SHARED_MEM)) +// Let's just upgrade to C++11 so we can use constexpr here... +// start allocating from this address (lower values used for allocating globals in .ptx file) +const unsigned long long GLOBAL_HEAP_START = 0xC0000000; +// Volta max shmem size is 96kB +const unsigned long long SHARED_MEM_SIZE_MAX = 96 * (1 << 10); +// Volta max local mem is 16kB +const unsigned long long LOCAL_MEM_SIZE_MAX = 1 << 14; +// Volta Titan V has 80 SMs +const unsigned MAX_STREAMING_MULTIPROCESSORS = 80; +// Max 2048 threads / SM +const unsigned MAX_THREAD_PER_SM = 1 << 11; +// MAX 64 warps / SM +const unsigned MAX_WARP_PER_SM = 1 << 6; +const unsigned long long TOTAL_LOCAL_MEM_PER_SM = MAX_THREAD_PER_SM * LOCAL_MEM_SIZE_MAX; +const unsigned long long TOTAL_SHARED_MEM = MAX_STREAMING_MULTIPROCESSORS * SHARED_MEM_SIZE_MAX; +const unsigned long long TOTAL_LOCAL_MEM = MAX_STREAMING_MULTIPROCESSORS * MAX_THREAD_PER_SM * LOCAL_MEM_SIZE_MAX; +const unsigned long long SHARED_GENERIC_START = GLOBAL_HEAP_START - TOTAL_SHARED_MEM; +const unsigned long long LOCAL_GENERIC_START = SHARED_GENERIC_START - TOTAL_LOCAL_MEM; +const unsigned long long STATIC_ALLOC_LIMIT = GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM + TOTAL_SHARED_MEM); #if !defined(__CUDA_RUNTIME_API_H__) diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index dca3cec..33bcf45 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -214,7 +214,8 @@ void fix_duplicate_errors(char fname2[1024]) { long filesize = ftell(ptxsource); rewind(ptxsource); char *ptxdata = (char*)malloc((filesize+1)*sizeof(char)); - fread(ptxdata, filesize, 1, ptxsource); + // Fail if we do not read the file + assert(fread(ptxdata, filesize, 1, ptxsource) == 1); fclose(ptxsource); FILE *ptxdest = fopen(fname2,"w"); -- cgit v1.3 From ddab5e0282e6004c6674338f2b35b86c627219e1 Mon Sep 17 00:00:00 2001 From: Nick Date: Tue, 3 Sep 2019 10:33:18 -0400 Subject: remove more unused variables and added asserts to function return values --- src/cuda-sim/cuda-sim.cc | 6 +++--- src/cuda-sim/instructions.cc | 8 +------- src/cuda-sim/ptx_ir.cc | 2 +- 3 files changed, 5 insertions(+), 11 deletions(-) (limited to 'src') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index fbce75b..28b4bf9 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1309,7 +1309,7 @@ void function_info::ptx_jit_config(std::map mallocPt char buff[1024]; std::string filename_c(filename+"_c"); snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename_c.c_str()); - system(buff); + assert(system(buff) != NULL); FILE *fp = fopen(filename_c.c_str(), "r"); fgets(buff, 1024, fp); fclose(fp); @@ -1432,13 +1432,13 @@ void function_info::ptx_jit_config(std::map mallocPt fout = fopen(ptx_config_fn.c_str(), "a"); assert(fout!=NULL); for (unsigned i = 0; ilookup(name); @@ -233,9 +232,6 @@ void ptx_thread_info::resume_reg_thread(char * fname, symbol_table * symtab) data = atoi(pch); pch = strtok (NULL," "); pch = strtok (NULL," "); - size = atoi(pch); - - m_regs.back()[reg] = data; } fclose ( fp2 ); @@ -3080,7 +3076,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, k; ptx_thread_info *thread; const operand_info &src = pI->operand_lookup(1); @@ -3100,8 +3096,6 @@ 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); ptx_reg_t src2_data = thread->get_operand_value(src2, src, type, thread, 1); const operand_info &src_a= pI->operand_lookup(1); diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 6978cc1..d8943d2 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1415,7 +1415,7 @@ unsigned function_info::print_insn( unsigned pc, FILE * fp ) const snprintf(command,1024,"c++filt -p %s",m_name.c_str()); FILE *p = popen(command,"r"); buffer[0]=0; - fgets(buffer, 1023, p); + assert(fgets(buffer, 1023, p) != NULL); // Remove trailing "\n" in buffer char *c; if ((c=strchr(buffer, '\n')) != NULL) *c = '\0'; -- cgit v1.3