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 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