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 015fbc0..61f8415 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -775,6 +775,17 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = sscanf(mode,"%u", &g_ptx_sim_mode); gpgpusim_ptx_assert( !ctx->api->g_cuda_launch_stack.empty(), "empty launch stack" ); kernel_config config = ctx->api->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"); + ctx->api->g_cuda_launch_stack.pop_back(); + return g_last_cudaError = cudaErrorInvalidConfiguration; + } + } struct CUstream_st *stream = config.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", stream?stream->get_uid():0 ); |