diff options
Diffstat (limited to 'libcuda')
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 11 |
1 files changed, 11 insertions, 0 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 97396ac..c70a570 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1501,6 +1501,17 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) sscanf(mode,"%u", &g_ptx_sim_mode); gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); kernel_config config = g_cuda_launch_stack.back(); + { + 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) + { + //can't launch + 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()) stream = 0; |