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