From a98e7db0402054aee7332380f6679be8527fc844 Mon Sep 17 00:00:00 2001 From: Deval Shah Date: Wed, 30 Jan 2019 17:40:17 -0800 Subject: changes for Ubuntu 18.04 support --- libcuda/cuda_runtime_api.cc | 39 +++++++++++++++++++++++++++++++++++++++ 1 file changed, 39 insertions(+) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 95a3c24..042bebd 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1199,6 +1199,35 @@ __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__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaSuccess; + +} + +__host__ cudaError_t CUDARTAPI cudaCreateTextureObject ( cudaTextureObject_t* pTexObject, const cudaResourceDesc* pResDesc, const cudaTextureDesc* pTexDesc, const cudaResourceViewDesc* pResViewDesc ) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaSuccess; + +} + +__host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority ) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaSuccess; + +} __host__ cudaError_t CUDARTAPI cudaDeviceGetPCIBusId ( char *pciBusId, @@ -1235,6 +1264,16 @@ __host__ cudaError_t cudaIpcOpenMemHandle( return g_last_cudaError = cudaErrorUnknown; } +__host__ cudaError_t CUDARTAPI cudaDestroyTextureObject(cudaTextureObject_t texObject) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; +} + + /******************************************************************************* * * * * -- cgit v1.3 From fae9182e8954c7d0258abf403a6fe910b0745758 Mon Sep 17 00:00:00 2001 From: Deval Shah Date: Wed, 30 Jan 2019 22:51:12 -0800 Subject: changes for ubuntu 18.04 support --- libcuda/cuda_runtime_api.cc | 19 ++++++++++--------- 1 file changed, 10 insertions(+), 9 deletions(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 042bebd..c81c483 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1209,15 +1209,6 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li } -__host__ cudaError_t CUDARTAPI cudaCreateTextureObject ( cudaTextureObject_t* pTexObject, const cudaResourceDesc* pResDesc, const cudaTextureDesc* pTexDesc, const cudaResourceViewDesc* pResViewDesc ) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaSuccess; - -} __host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority ) { @@ -4486,6 +4477,16 @@ CUresult CUDAAPI cuPointerGetAttribute(void *data, CUpointer_attribute attribute #endif /* CUDART_VERSION >= 4000 */ #if CUDART_VERSION >= 8000 +__host__ cudaError_t CUDARTAPI cudaCreateTextureObject ( cudaTextureObject_t* pTexObject, const cudaResourceDesc* pResDesc, const cudaTextureDesc* pTexDesc, const cudaResourceViewDesc* pResViewDesc ) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaSuccess; + +} + CUresult CUDAAPI cuMemPrefetchAsync(CUdeviceptr devPtr, size_t count, CUdevice dstDevice, CUstream hStream) { if(g_debug_execution >= 3){ -- cgit v1.3 From 9fb5819697e4caa107980635887226d507f9cb36 Mon Sep 17 00:00:00 2001 From: Deval Shah Date: Thu, 14 Feb 2019 14:57:57 -0800 Subject: Skip app version check for pytorch --- libcuda/cuda_runtime_api.cc | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index c81c483..0f7aac2 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2582,8 +2582,17 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) // Making this a runtime variable based on the app, enables GPGPU-Sim compiled // with a newer version of CUDA to run apps compiled with older versions of // CUDA. This is especially useful for PTXPLUS execution. - int app_cuda_version = get_app_cuda_version(); - assert( app_cuda_version == CUDART_VERSION / 1000 && "The app must be compiled with same major version as the simulator." ); + //Skip cuda version check for pytorch application + std::string app_binary_path = "ldd " + get_app_binary(); + int pos = app_binary_path.find("python"); + if (pos==std::string::npos){ + printf("Not pytorch app : checking cuda version"); + int app_cuda_version = get_app_cuda_version(); + assert( app_cuda_version == CUDART_VERSION / 1000 && "The app must be compiled with same major version as the simulator." ); + } + + //int app_cuda_version = get_app_cuda_version(); + //assert( app_cuda_version == CUDART_VERSION / 1000 && "The app must be compiled with same major version as the simulator." ); const char* filename; #if CUDART_VERSION < 6000 // FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC) -- cgit v1.3 From 431381ece357b1bfae7171743e56ab9361b749f4 Mon Sep 17 00:00:00 2001 From: Deval Shah Date: Thu, 14 Feb 2019 15:15:29 -0800 Subject: Skip app version check for pytorch --- libcuda/cuda_runtime_api.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 0f7aac2..e6b53c3 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2583,10 +2583,10 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) // with a newer version of CUDA to run apps compiled with older versions of // CUDA. This is especially useful for PTXPLUS execution. //Skip cuda version check for pytorch application - std::string app_binary_path = "ldd " + get_app_binary(); + std::string app_binary_path = get_app_binary(); int pos = app_binary_path.find("python"); if (pos==std::string::npos){ - printf("Not pytorch app : checking cuda version"); + // Not pytorch app : checking cuda version int app_cuda_version = get_app_cuda_version(); assert( app_cuda_version == CUDART_VERSION / 1000 && "The app must be compiled with same major version as the simulator." ); } -- cgit v1.3 From 19d53139ce6676849b35266c5693c177aa1ee4e7 Mon Sep 17 00:00:00 2001 From: aamir Date: Wed, 20 Feb 2019 23:56:31 -0800 Subject: debugged the cutlass deadlock issue after pytorch and cudnn merge --- libcuda/cuda_runtime_api.cc | 42 +++++++++++++----------------------------- 1 file changed, 13 insertions(+), 29 deletions(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index e6b53c3..27644b3 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1499,42 +1499,26 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) return g_last_cudaError = cudaSuccess; } - __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream ) { - struct CUstream_st *s = (struct CUstream_st *)stream; - g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) ); - - //printf("cudaLaunchKernel:sizeof(Arg[0])=%d)\n ",sizeof(args[0])); - kernel_config &config = g_cuda_launch_stack.back(); - config.set_arg(args[0],432,0);//standard interface for cutlass library #TODO Implementing a generalized kernel - - CUctx_st* context = GPGPUSim_Context(); - char *mode = getenv("PTX_SIM_MODE_FUNC"); - if( mode ) - sscanf(mode,"%u", &g_ptx_sim_mode); - gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); - kernel_config config1 = g_cuda_launch_stack.back(); - struct CUstream_st *stream1 = config1.get_stream(); - printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun, - g_ptx_sim_mode?"functional simulation":"performance simulation", stream1?stream1->get_uid():0 ); - kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config1.get_args(),config1.grid_dim(),config1.block_dim(),context); - std::string kname = grid->name(); - dim3 gridDim1 = config1.grid_dim(); - dim3 blockDim1 = config1.block_dim(); - printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n", - kname.c_str(), stream1?stream1->get_uid():0, gridDim1.x,gridDim1.y,gridDim1.z,blockDim1.x,blockDim1.y,blockDim1.z ); - /*Kernel is hardcoded to enable the cutlass library*/ - std::string cutlass("cutlass"); - assert(kname.find(cutlass) != std::string::npos); + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + CUctx_st *context = GPGPUSim_Context(); + function_info *entry = context->get_kernel(hostFun); + + cudaConfigureCall(gridDim, blockDim, sharedMem, stream); + for(unsigned i = 0; i < entry->num_args(); i++){ + std::pair p = entry->get_param_config(i); + cudaSetupArgument(args[i], p.first, p.second); + } - stream_operation op(grid,g_ptx_sim_mode,stream1); - g_stream_manager->push(op); - g_cuda_launch_stack.pop_back(); + cudaLaunch(hostFun); return g_last_cudaError = cudaSuccess; } + /******************************************************************************* * * * * -- cgit v1.3 From f911bdf3b5bfe793ad04df5a62e425d3ed5508b5 Mon Sep 17 00:00:00 2001 From: PSuchita <31720043+PSuchita@users.noreply.github.com> Date: Fri, 15 Mar 2019 15:43:48 -0500 Subject: Update cudaStreamDestroy function Update cudaStreamDestroy's functionality to synchronize before destroying streams. --- libcuda/cuda_runtime_api.cc | 1 + 1 file changed, 1 insertion(+) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 27644b3..3ff3212 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1569,6 +1569,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) announce_call(__my_func__); } #if (CUDART_VERSION >= 3000) + synchronize(); g_stream_manager->destroy_stream(stream); #endif return g_last_cudaError = cudaSuccess; -- cgit v1.3 From 981c79b0f3e1c637755bfa8dd0b2dbac1d44a865 Mon Sep 17 00:00:00 2001 From: PSuchita <31720043+PSuchita@users.noreply.github.com> Date: Fri, 15 Mar 2019 16:26:28 -0500 Subject: Added comment Adding comment to describe the issue resolved by the previous commit related to cudaStreamDestroy. --- libcuda/cuda_runtime_api.cc | 1 + 1 file changed, 1 insertion(+) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 3ff3212..920be46 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1569,6 +1569,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) announce_call(__my_func__); } #if (CUDART_VERSION >= 3000) + //synchronization required to avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. synchronize(); g_stream_manager->destroy_stream(stream); #endif -- cgit v1.3 From d575f9464df6358d1be60953e984bcfa9257d1a3 Mon Sep 17 00:00:00 2001 From: PSuchita <31720043+PSuchita@users.noreply.github.com> Date: Fri, 15 Mar 2019 17:13:16 -0500 Subject: Updated comment. --- libcuda/cuda_runtime_api.cc | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 920be46..cf16d9b 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1569,7 +1569,8 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) announce_call(__my_func__); } #if (CUDART_VERSION >= 3000) - //synchronization required to avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. + //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. synchronize(); g_stream_manager->destroy_stream(stream); #endif -- cgit v1.3 From 73050b730b31e9f58f8c2a0bfc4b3a8cd92be7ff Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Wed, 10 Apr 2019 23:27:42 -0500 Subject: Adding texture Fix --- libcuda/cuda_runtime_api.cc | 28 ++++++++++++++++++++++++---- src/abstract_hardware_model.cc | 4 +++- src/abstract_hardware_model.h | 22 +++++++++++++++++++++- src/cuda-sim/cuda-sim.cc | 5 +++-- src/cuda-sim/cuda_device_runtime.cc | 3 ++- src/cuda-sim/instructions.cc | 5 +++-- 6 files changed, 56 insertions(+), 11 deletions(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index cf16d9b..7569b20 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1204,8 +1204,24 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li if(g_debug_execution >= 3){ announce_call(__my_func__); } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaSuccess; + switch(limit) { + case 0: + *pValue=1024; + break; + case 2: + *pValue=8388608; + break; + case 3: + *pValue=2; + break; + case 4: + *pValue=2048; + break; + default: + printf("ERROR:Limit %s unimplemented \n",limit); + abort(); + } + return g_last_cudaError = cudaSuccess; } @@ -1571,7 +1587,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) #if (CUDART_VERSION >= 3000) //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. - synchronize(); + stream->synchronize(); g_stream_manager->destroy_stream(stream); #endif return g_last_cudaError = cudaSuccess; @@ -2702,6 +2718,7 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ if(g_debug_execution >= 3){ announce_call(__my_func__); } + synchronize(); return g_last_cudaError = cudaSuccess; } @@ -3135,6 +3152,8 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v } return g_last_cudaError = cudaSuccess; } + + #endif #endif @@ -3313,7 +3332,8 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, announce_call(__my_func__); } function_info *entry = context->get_kernel(hostFun); - kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry); + gpgpu_t* gpu= context->get_device()->get_gpgpu(); + 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); abort(); diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 307df40..3db87c9 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -688,7 +688,7 @@ 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 ) +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map NameToCudaArray, std::map NameToTexureInfo) { m_kernel_entry=entry; m_grid_dim=gridDim; @@ -708,6 +708,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_launch_latency = g_kernel_launch_latency; volta_cache_config_set=false; + m_NameToCudaArray = NameToCudaArray; + m_NameToTexureInfo = NameToTexureInfo; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 28a22e0..bbc31e5 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 ); + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map m_NameToCudaArray, std::map m_NameToTexureInfo); ~kernel_info_t(); void inc_running() { m_num_cores_running++; } @@ -275,6 +275,20 @@ public: std::list &active_threads() { return m_active_threads; } class memory_space *get_param_memory() { return m_param_mem; } + const struct cudaArray* get_texarray( const std::string &texname ) const + { + std::map::const_iterator t=m_NameToCudaArray.find(texname); + assert(t != m_NameToCudaArray.end()); + return t->second; + } + + const struct textureInfo* get_texinfo( const std::string &texname ) const + { + std::map::const_iterator t=m_NameToTexureInfo.find(texname); + assert(t != m_NameToTexureInfo.end()); + return t->second; + } + private: kernel_info_t( const kernel_info_t & ); // disable copy constructor void operator=( const kernel_info_t & ); // disable copy operator @@ -284,6 +298,9 @@ private: unsigned m_uid; static unsigned m_next_uid; + std::map m_NameToCudaArray; + std::map m_NameToTexureInfo; + dim3 m_grid_dim; dim3 m_block_dim; dim3 m_next_cta; @@ -595,6 +612,9 @@ public: const gpgpu_functional_sim_config &get_config() const { return m_function_model_config; } FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; } + std::map getNameArrayMapping() {return m_NameToCudaArray;} + std::map getNameInfoMapping() {return m_NameToTexureInfo;} + protected: const gpgpu_functional_sim_config &m_function_model_config; FILE* ptx_inst_debug_file; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index ec68b5b..a51ea3e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1507,7 +1507,8 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th std::string texname = src1.name(); gpgpu_t *gpu = thread->get_gpu(); - const struct textureInfo* texInfo = gpu->get_texinfo(texname); + kernel_info_t& k = thread->get_kernel(); + const struct textureInfo* texInfo = k.get_texinfo(texname); unsigned data_size = texInfo->texel_size; return data_size; @@ -1918,7 +1919,7 @@ kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, struct dim3 blockDim, gpgpu_t *gpu ) { - kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry); + kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping()); unsigned argcount=args.size(); unsigned argn=1; for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) { diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index b399133..917e7a8 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -198,7 +198,8 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * memory_space *device_kernel_param_mem; //create child kernel_info_t and index it with parameter_buffer address - device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry); + gpgpu_t* gpu=thread->get_gpu(); + device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry, gpu->getNameArrayMapping(), gpu->getNameInfoMapping()); device_grid->launch_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; kernel_info_t & parent_grid = thread->get_kernel(); DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index c85654c..d29fc8b 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -5260,9 +5260,10 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->get_vector_operand_values(src2, ptx_tex_regs, nelem); //ptx_reg should be 4 entry vector type...coordinates into texture gpgpu_t *gpu = thread->get_gpu(); + kernel_info_t &k = thread->get_kernel(); const struct textureReference* texref = gpu->get_texref(texname); - const struct cudaArray* cuArray = gpu->get_texarray(texname); - const struct textureInfo* texInfo = gpu->get_texinfo(texname); + const struct cudaArray* cuArray = k.get_texarray(texname); + const struct textureInfo* texInfo = k.get_texinfo(texname); const struct textureReferenceAttr* texAttr = gpu->get_texattr(texname); //assume always 2D f32 input -- cgit v1.3 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 ++++++++++++---------- src/abstract_hardware_model.cc | 13 ++++++++++--- src/abstract_hardware_model.h | 36 +++++++++++++++++++++--------------- src/cuda-sim/cuda-sim.cc | 7 +++++++ src/cuda-sim/instructions.cc | 8 +++++++- 5 files changed, 57 insertions(+), 29 deletions(-) (limited to 'libcuda') 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); diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 3db87c9..820f5fd 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -192,6 +192,9 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config ) checkpoint_CTA_t = m_function_model_config.get_checkpoint_CTA_t(); checkpoint_insn_Y = m_function_model_config.get_checkpoint_insn_Y(); + // initialize texture mappings to empty + m_NameToTextureInfo.clear(); + m_NameToCudaArray.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"); @@ -688,7 +691,11 @@ 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, std::map NameToCudaArray, std::map NameToTexureInfo) +/*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) { m_kernel_entry=entry; m_grid_dim=gridDim; @@ -708,8 +715,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_launch_latency = g_kernel_launch_latency; volta_cache_config_set=false; - m_NameToCudaArray = NameToCudaArray; - m_NameToTexureInfo = NameToTexureInfo; + m_NameToCudaArray = nameToCudaArray; + m_NameToTexureInfo = nameToTexureInfo; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index bbc31e5..201b21f 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 m_NameToCudaArray, std::map m_NameToTexureInfo); + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo); ~kernel_info_t(); void inc_running() { m_num_cores_running++; } @@ -275,19 +275,22 @@ public: std::list &active_threads() { return m_active_threads; } class memory_space *get_param_memory() { return m_param_mem; } + + //The following functions access texture bindings present at the kernel's launch + const struct cudaArray* get_texarray( const std::string &texname ) const - { - std::map::const_iterator t=m_NameToCudaArray.find(texname); - assert(t != m_NameToCudaArray.end()); - return t->second; - } + { + std::map::const_iterator t=m_NameToCudaArray.find(texname); + assert(t != m_NameToCudaArray.end()); + return t->second; + } - const struct textureInfo* get_texinfo( const std::string &texname ) const - { - std::map::const_iterator t=m_NameToTexureInfo.find(texname); - assert(t != m_NameToTexureInfo.end()); - return t->second; - } + const struct textureInfo* get_texinfo( const std::string &texname ) const + { + std::map::const_iterator t=m_NameToTexureInfo.find(texname); + assert(t != m_NameToTexureInfo.end()); + return t->second; + } private: kernel_info_t( const kernel_info_t & ); // disable copy constructor @@ -297,7 +300,8 @@ private: unsigned m_uid; static unsigned m_next_uid; - + + //These maps contain the snapshot of the texture mappings at kernel launch std::map m_NameToCudaArray; std::map m_NameToTexureInfo; @@ -611,7 +615,9 @@ public: const gpgpu_functional_sim_config &get_config() const { return m_function_model_config; } 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;} @@ -624,7 +630,7 @@ protected: class memory_space *m_surf_mem; unsigned long long m_dev_malloc; - + // These maps contain the current texture mappings for the GPU at any given time. std::map > m_NameToTextureRef; std::map m_TextureRefToName; std::map m_NameToCudaArray; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index a51ea3e..9d59411 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1507,6 +1507,13 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th 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 + the state of the texture mappings when it was launched (so that we + don't try to access the incorrect texture mapping if it's been updated, + or that we don't access a mapping that has been unbound). + */ kernel_info_t& k = thread->get_kernel(); const struct textureInfo* texInfo = k.get_texinfo(texname); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index d29fc8b..11001d7 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -5258,7 +5258,13 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) if (!ptx_tex_regs) ptx_tex_regs = new ptx_reg_t[4]; unsigned nelem = src2.get_vect_nelem(); thread->get_vector_operand_values(src2, ptx_tex_regs, nelem); //ptx_reg should be 4 entry vector type...coordinates into texture - + /* + For programs with many streams, textures can be bound and unbound + asynchronously. This means we need to use the kernel's "snapshot" of + the state of the texture mappings when it was launched (so that we + don't try to access the incorrect texture mapping if it's been updated, + or that we don't access a mapping that has been unbound). + */ gpgpu_t *gpu = thread->get_gpu(); kernel_info_t &k = thread->get_kernel(); const struct textureReference* texref = gpu->get_texref(texname); -- 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') 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') 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 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 'libcuda') 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 From 40a11341f74f2043d46df00be36c6d6fb5294b71 Mon Sep 17 00:00:00 2001 From: Preyesh Dalmia Date: Fri, 19 Apr 2019 10:17:22 -0500 Subject: updating comment to reflect the correct name of attribute --- libcuda/cuda_runtime_api.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 9068700..3a9d613 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1120,10 +1120,10 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic case 41: *value= 0; break; - case 75:// Major compute capability version number + case 75://cudaDevAttrComputeCapabilityMajor *value= prop->major ; break; - case 76:// Minor compute capability version number + case 76://cudaDevAttrComputeCapabilityMinor *value= prop->minor ; break; case 78: -- cgit v1.3