From 69af057daf1999f17b81d761c7a616f67f56adbb Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Tue, 16 Apr 2019 23:36:02 -0500 Subject: Adding additional fixes for the texture bugs --- libcuda/cuda_runtime_api.cc | 22 ++++++++++++---------- 1 file changed, 12 insertions(+), 10 deletions(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 7569b20..f6e5e6f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1121,10 +1121,10 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic *value= 0; break; case 75: - *value= 9 ; + *value= 7 ; break; case 76: - *value= 3 ; + *value= 0 ; break; case 78: *value= 0 ; //TODO: as of now, we dont support stream priorities. @@ -1199,22 +1199,23 @@ __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device) *device = g_active_device; return g_last_cudaError = cudaSuccess; } + __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit ) { if(g_debug_execution >= 3){ announce_call(__my_func__); } switch(limit) { - case 0: + case 0: // cudaLimitStackSize *pValue=1024; break; - case 2: + case 2: // cudaLimitMallocHeapSize *pValue=8388608; break; - case 3: + case 3: // cudaLimitDevRuntimeSyncDepth *pValue=2; break; - case 4: + case 4: // cudaLimitDevRuntimePendingLaunchCount *pValue=2048; break; default: @@ -1225,7 +1226,6 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li } - __host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority ) { if(g_debug_execution >= 3){ @@ -1585,7 +1585,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) announce_call(__my_func__); } #if (CUDART_VERSION >= 3000) - //synchronization required for application using external libraries without explicit synchronization in the code to + //per-stream synchronization required for application using external libraries without explicit synchronization in the code to //avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. stream->synchronize(); g_stream_manager->destroy_stream(stream); @@ -2714,7 +2714,6 @@ cudaError_t cudaDeviceReset ( void ) { return g_last_cudaError = cudaSuccess; } cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ - // I don't know what this should do if(g_debug_execution >= 3){ announce_call(__my_func__); } @@ -2722,7 +2721,6 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ return g_last_cudaError = cudaSuccess; } - void CUDARTAPI __cudaRegisterFunction( void **fatCubinHandle, const char *hostFun, @@ -3333,6 +3331,10 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, } function_info *entry = context->get_kernel(hostFun); gpgpu_t* gpu= context->get_device()->get_gpgpu(); + /* + Passing a snapshot of the GPU's current texture mapping to the kernel's info + as kernels should use texture bindings present at the time of their launch. + */ kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping()); if( entry == NULL ) { printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun); -- 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 'libcuda/cuda_runtime_api.cc') 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 74289a429633176e37116353347f60c75f2b644a Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Wed, 17 Apr 2019 23:30:20 -0500 Subject: Modified config files to add device limit configurations --- configs/tested-cfgs/SM2_GTX480/gpgpusim.config | 3 +++ configs/tested-cfgs/SM6_TITANX/gpgpusim.config | 6 +++++ configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 6 +++++ libcuda/cuda_runtime_api.cc | 35 ++++++++++++++++++-------- 4 files changed, 40 insertions(+), 10 deletions(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config index d71b2fd..8d54d8b 100644 --- a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config +++ b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config @@ -3,6 +3,9 @@ -gpgpu_ptx_sim_mode 0 -gpgpu_ptx_force_max_capability 20 +# Device Limits +-gpgpu_stack_size_limit 1024 +-gpgpu_heap_size_limit 8388608 # SASS execution (only supported with CUDA >= 4.0) -gpgpu_ptx_convert_to_ptxplus 0 diff --git a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config index cb23ab3..3371186 100644 --- a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config +++ b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config @@ -8,6 +8,12 @@ -gpgpu_ptx_force_max_capability 61 -gpgpu_ignore_resources_limitation 1 +# Device Limits +-gpgpu_stack_size_limit 1024 +-gpgpu_heap_size_limit 8388608 +-gpgpu_runtime_sync_depth_limit 2 +-gpgpu_runtime_pending_launch_count_limit 2048 + # 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 8ed4cd0..1af64ac 100644 --- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config +++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config @@ -13,6 +13,12 @@ -gpgpu_ptx_sim_mode 0 -gpgpu_ptx_force_max_capability 70 +# Device Limits +-gpgpu_stack_size_limit 1024 +-gpgpu_heap_size_limit 8388608 +-gpgpu_runtime_sync_depth_limit 2 +-gpgpu_runtime_pending_launch_count_limit 2048 + # 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 86435b9..6f1ab08 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1121,10 +1121,10 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic *value= 0; break; case 75: - *value= 7 ; + *value= 7 ; //cudaDevAttrComputeCapabilityMajor for Volta architecture break; case 76: - *value= 0 ; + *value= 0 ; //cudaDevAttrComputeCapabilityMinor for Volta architecture break; case 78: *value= 0 ; //TODO: as of now, we dont support stream priorities. @@ -1206,20 +1206,35 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li announce_call(__my_func__); } _cuda_device_id *dev = GPGPUSim_Init(); + const struct cudaDeviceProp *prop = dev->get_prop(); const gpgpu_sim_config& config=dev->get_gpgpu()->get_config(); switch(limit) { case 0: // cudaLimitStackSize - *pValue=config.stack_limit(); - break; + *pValue=config.stack_limit(); + break; case 2: // cudaLimitMallocHeapSize - *pValue=config.heap_limit(); - break; + *pValue=config.heap_limit(); + break; +#if (CUDART_VERSION > 5050) case 3: // cudaLimitDevRuntimeSyncDepth - *pValue=config.sync_depth_limit(); - break; + if(prop->major > 2){ + *pValue=config.sync_depth_limit(); + break; + } + else{ + printf("ERROR:Limit %s is not supported on this architecture \n",limit); + abort(); + } case 4: // cudaLimitDevRuntimePendingLaunchCount - *pValue=config.pending_launch_count_limit(); - break; + if(prop->major > 2){ + *pValue=config.pending_launch_count_limit(); + break; + } + else{ + printf("ERROR:Limit %s is not supported on this architecture \n",limit); + abort(); + } +#endif default: printf("ERROR:Limit %s unimplemented \n",limit); abort(); -- cgit v1.3