From ec2ada3b5f2ef91341d30e8f56f3c4d9ecafb2a4 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Thu, 11 Apr 2019 22:32:08 -0400 Subject: adding some stats and random hashing --- src/gpgpu-sim/gpu-sim.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.h') diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 6ce5524..efda65a 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -62,8 +62,7 @@ #define SAMPLELOG 222 #define DUMPLOG 333 - - +extern tr1_hash_map address_random_interleaving; enum dram_ctrl_t { -- cgit v1.3 From 49c6eeff7752b5816fb08c0d6baf2b47c789bc18 Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Wed, 17 Apr 2019 22:09:17 -0500 Subject: Adding configs for device limits and fixed other misc bugs --- libcuda/cuda_runtime_api.cc | 11 +++++++---- src/abstract_hardware_model.cc | 9 ++++----- src/abstract_hardware_model.h | 17 ++++++++--------- src/cuda-sim/cuda-sim.cc | 4 ++-- src/gpgpu-sim/gpu-sim.cc | 8 ++++++++ src/gpgpu-sim/gpu-sim.h | 11 ++++++++++- 6 files changed, 39 insertions(+), 21 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index f6e5e6f..86435b9 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1205,18 +1205,20 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li if(g_debug_execution >= 3){ announce_call(__my_func__); } + _cuda_device_id *dev = GPGPUSim_Init(); + const gpgpu_sim_config& config=dev->get_gpgpu()->get_config(); switch(limit) { case 0: // cudaLimitStackSize - *pValue=1024; + *pValue=config.stack_limit(); break; case 2: // cudaLimitMallocHeapSize - *pValue=8388608; + *pValue=config.heap_limit(); break; case 3: // cudaLimitDevRuntimeSyncDepth - *pValue=2; + *pValue=config.sync_depth_limit(); break; case 4: // cudaLimitDevRuntimePendingLaunchCount - *pValue=2048; + *pValue=config.pending_launch_count_limit(); break; default: printf("ERROR:Limit %s unimplemented \n",limit); @@ -2717,6 +2719,7 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ if(g_debug_execution >= 3){ announce_call(__my_func__); } + //Blocks until the device has completed all preceding requested tasks synchronize(); return g_last_cudaError = cudaSuccess; } diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 820f5fd..cebdb25 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -88,8 +88,6 @@ void checkpoint::load_global_mem(class memory_space *temp_mem, char * f1name) fclose ( fp2 ); } - - void checkpoint::store_global_mem(class memory_space * mem, char *fname, char * format) { @@ -195,6 +193,8 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config ) // initialize texture mappings to empty m_NameToTextureInfo.clear(); m_NameToCudaArray.clear(); + m_TextureRefToName.clear(); + m_NameToAttribute.clear(); if(m_function_model_config.get_ptx_inst_debug_to_file() != 0) ptx_inst_debug_file = fopen(m_function_model_config.get_ptx_inst_debug_file(), "w"); @@ -694,8 +694,7 @@ unsigned kernel_info_t::m_next_uid = 1; /*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. */ - -kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo) +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTextureInfo) { m_kernel_entry=entry; m_grid_dim=gridDim; @@ -716,7 +715,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * volta_cache_config_set=false; m_NameToCudaArray = nameToCudaArray; - m_NameToTexureInfo = nameToTexureInfo; + m_NameToTextureInfo = nameToTextureInfo; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 201b21f..e8716ab 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -212,7 +212,7 @@ public: // m_num_cores_running=0; // m_param_mem=NULL; // } - kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo); + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTextureInfo); ~kernel_info_t(); void inc_running() { m_num_cores_running++; } @@ -287,8 +287,8 @@ public: const struct textureInfo* get_texinfo( const std::string &texname ) const { - std::map::const_iterator t=m_NameToTexureInfo.find(texname); - assert(t != m_NameToTexureInfo.end()); + std::map::const_iterator t=m_NameToTextureInfo.find(texname); + assert(t != m_NameToTextureInfo.end()); return t->second; } @@ -303,7 +303,7 @@ private: //These maps contain the snapshot of the texture mappings at kernel launch std::map m_NameToCudaArray; - std::map m_NameToTexureInfo; + std::map m_NameToTextureInfo; dim3 m_grid_dim; dim3 m_block_dim; @@ -601,8 +601,8 @@ public: const struct textureInfo* get_texinfo( const std::string &texname ) const { - std::map::const_iterator t=m_NameToTexureInfo.find(texname); - assert(t != m_NameToTexureInfo.end()); + std::map::const_iterator t=m_NameToTextureInfo.find(texname); + assert(t != m_NameToTextureInfo.end()); return t->second; } @@ -617,9 +617,8 @@ public: FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; } // These maps return the current texture mappings for the GPU at any given time. - std::map getNameArrayMapping() {return m_NameToCudaArray;} - std::map getNameInfoMapping() {return m_NameToTexureInfo;} + std::map getNameInfoMapping() {return m_NameToTextureInfo;} protected: const gpgpu_functional_sim_config &m_function_model_config; @@ -634,7 +633,7 @@ protected: std::map > m_NameToTextureRef; std::map m_TextureRefToName; std::map m_NameToCudaArray; - std::map m_NameToTexureInfo; + std::map m_NameToTextureInfo; std::map m_NameToAttribute; }; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 9d59411..f7bb9cc 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -218,7 +218,7 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te texInfo->Ty_numbits = intLOGB2(Ty); texInfo->texel_size = texel_size; texInfo->texel_size_numbits = intLOGB2(texel_size); - m_NameToTexureInfo[texname] = texInfo; + m_NameToTextureInfo[texname] = texInfo; } void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) @@ -226,7 +226,7 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) //assumes bind-use-unbind-bind-use-unbind pattern std::string texname = gpgpu_ptx_sim_findNamefromTexture(texref); m_NameToCudaArray.erase(texname); - m_NameToTexureInfo.erase(texname); + m_NameToTextureInfo.erase(texname); } unsigned g_assemble_code_next_pc=0; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index ec570bf..2a862a4 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -532,6 +532,14 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-visualizer_zlevel", OPT_INT32, &g_visualizer_zlevel, "Compression level of the visualizer output log (0=no comp, 9=highest)", "6"); + option_parser_register(opp, "-gpgpu_stack_size_limit", OPT_INT32, &stack_size_limit, + "GPU thread stack size", "1024" ); + option_parser_register(opp, "-gpgpu_heap_size_limit", OPT_INT32, &heap_size_limit, + "GPU malloc heap size ", "8388608" ); + option_parser_register(opp, "-gpgpu_runtime_sync_depth_limit", OPT_INT32, &runtime_sync_depth_limit, + "GPU device runtime synchronize depth", "2" ); + option_parser_register(opp, "-gpgpu_runtime_pending_launch_count_limit", OPT_INT32, &runtime_pending_launch_count_limit, + "GPU device runtime pending launch count", "2048" ); option_parser_register(opp, "-trace_enabled", OPT_BOOL, &Trace::enabled, "Turn on traces", "0"); diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 6ce5524..11ffe49 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -337,6 +337,11 @@ public: unsigned get_max_concurrent_kernel() const { return max_concurrent_kernel; } unsigned checkpoint_option; + size_t stack_limit() const {return stack_size_limit; } + size_t heap_limit() const {return heap_size_limit; } + size_t sync_depth_limit() const {return runtime_sync_depth_limit; } + size_t pending_launch_count_limit() const {return runtime_pending_launch_count_limit;} + private: void init_clock_domains(void ); @@ -377,7 +382,11 @@ private: int gpu_stat_sample_freq; int gpu_runtime_stat_flag; - + // Device Limits + size_t stack_size_limit; + size_t heap_size_limit; + size_t runtime_sync_depth_limit; + size_t runtime_pending_launch_count_limit; unsigned long long liveness_message_freq; -- cgit v1.3 From 73f1a8a424cff7d689c138eaa8a59780b86a76d1 Mon Sep 17 00:00:00 2001 From: Preyesh Dalmia Date: Thu, 18 Apr 2019 17:26:40 -0500 Subject: changes for compute capabilty --- configs/tested-cfgs/SM2_GTX480/gpgpusim.config | 4 ++++ configs/tested-cfgs/SM6_TITANX/gpgpusim.config | 4 ++++ configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 4 ++++ libcuda/cuda_runtime_api.cc | 12 ++++++------ src/gpgpu-sim/gpu-sim.cc | 16 ++++++++++++++++ src/gpgpu-sim/gpu-sim.h | 5 +++++ 6 files changed, 39 insertions(+), 6 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.h') diff --git a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config index 8d54d8b..cf3627b 100644 --- a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config +++ b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config @@ -7,6 +7,10 @@ -gpgpu_stack_size_limit 1024 -gpgpu_heap_size_limit 8388608 +# Compute Capability +-gpgpu_compute_capability_major 2 +-gpgpu_compute_capability_minor 0 + # SASS execution (only supported with CUDA >= 4.0) -gpgpu_ptx_convert_to_ptxplus 0 -gpgpu_ptx_save_converted_ptxplus 0 diff --git a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config index 3371186..23d044c 100644 --- a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config +++ b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config @@ -14,6 +14,10 @@ -gpgpu_runtime_sync_depth_limit 2 -gpgpu_runtime_pending_launch_count_limit 2048 +# Compute Capability +-gpgpu_compute_capability_major 6 +-gpgpu_compute_capability_minor 1 + # SASS execution (only supported with CUDA >= 4.0) -gpgpu_ptx_convert_to_ptxplus 0 -gpgpu_ptx_save_converted_ptxplus 0 diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config index 1af64ac..f3a5f65 100644 --- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config +++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config @@ -19,6 +19,10 @@ -gpgpu_runtime_sync_depth_limit 2 -gpgpu_runtime_pending_launch_count_limit 2048 +# Compute Capability +-gpgpu_compute_capability_major 7 +-gpgpu_compute_capability_minor 0 + # SASS execution (only supported with CUDA >= 4.0) -gpgpu_ptx_convert_to_ptxplus 0 -gpgpu_ptx_save_converted_ptxplus 0 diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 6f1ab08..9068700 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -354,8 +354,8 @@ struct _cuda_device_id *GPGPUSim_Init() cudaDeviceProp *prop = (cudaDeviceProp *) calloc(sizeof(cudaDeviceProp),1); snprintf(prop->name,256,"GPGPU-Sim_v%s", g_gpgpusim_version_string ); - prop->major = 5; - prop->minor = 2; + prop->major = the_gpu->compute_capability_major(); + prop->minor = the_gpu->compute_capability_minor(); prop->totalGlobalMem = 0x80000000 /* 2 GB */; prop->memPitch = 0; if(prop->major >= 2) { @@ -1120,11 +1120,11 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic case 41: *value= 0; break; - case 75: - *value= 7 ; //cudaDevAttrComputeCapabilityMajor for Volta architecture + case 75:// Major compute capability version number + *value= prop->major ; break; - case 76: - *value= 0 ; //cudaDevAttrComputeCapabilityMinor for Volta architecture + case 76:// Minor compute capability version number + *value= prop->minor ; break; case 78: *value= 0 ; //TODO: as of now, we dont support stream priorities. diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 2a862a4..3687448 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -499,6 +499,12 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-liveness_message_freq", OPT_INT64, &liveness_message_freq, "Minimum number of seconds between simulation liveness messages (0 = always print)", "1"); + option_parser_register(opp, "-gpgpu_compute_capability_major", OPT_UINT32, &gpgpu_compute_capability_major, + "Major compute capability version number", + "7"); + option_parser_register(opp, "-gpgpu_compute_capability_minor", OPT_UINT32, &gpgpu_compute_capability_minor, + "Minor compute capability version number", + "0"); option_parser_register(opp, "-gpgpu_flush_l1_cache", OPT_BOOL, &gpgpu_flush_l1_cache, "Flush L1 cache at the end of each kernel call", "0"); @@ -802,6 +808,16 @@ void gpgpu_sim::set_prop( cudaDeviceProp *prop ) m_cuda_properties = prop; } +int gpgpu_sim::compute_capability_major() const +{ + return m_config.gpgpu_compute_capability_major; +} + +int gpgpu_sim::compute_capability_minor() const +{ + return m_config.gpgpu_compute_capability_minor; +} + const struct cudaDeviceProp *gpgpu_sim::get_prop() const { return m_cuda_properties; diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 11ffe49..d4cc18f 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -388,6 +388,9 @@ private: size_t runtime_sync_depth_limit; size_t runtime_pending_launch_count_limit; + //gpu compute capability options + unsigned int gpgpu_compute_capability_major; + unsigned int gpgpu_compute_capability_minor; unsigned long long liveness_message_freq; friend class gpgpu_sim; @@ -447,6 +450,8 @@ public: int shared_mem_size() const; int shared_mem_per_block() const; + int compute_capability_major() const; + int compute_capability_minor() const; int num_registers_per_core() const; int num_registers_per_block() const; int wrp_size() const; -- cgit v1.3