In this post, I will introduce how to debug async kernels or async APIs in CUDA. <\/p>\n
The async operations will not block CPU codes. When we check the return type of the functions calls, it may be SUCCESS but there are bugs like "illegal memory access". On the other hand, when we find the return type is FAILED, it may mean that previous operations failed. But we do not know which operation has bugs. <\/p>\n
Insert the following codes into the suspicious positions:<\/p>\n
{\n cudaMemoryTest();\n kernel_A<<<...>>>;\n\n ......\n\n cudaMemoryTest();\n kernel_B<<<...>>>;\n\n ......\n}<\/code><\/pre>\nThe macros are:<\/p>\n
#define cudaSafeCall(call) \\\n do {\\\n cudaError_t err = call;\\\n if (cudaSuccess != err) \\\n {\\\n std::cerr << "CUDA error in " << __FILE__ << "(" << __LINE__ << "): " \\\n << cudaGetErrorString(err);\\\n exit(EXIT_FAILURE);\\\n }\\\n } while(0)\n\nvoid cudaMemoryTest()\n{\n const unsigned int N = 1048576;\n const unsigned int bytes = N * sizeof(int);\n int *h_a = (int*)malloc(bytes);\n int *d_a;\n cudaSafeCall(cudaMalloc((int**)&d_a, bytes));\n\n memset(h_a, 0, bytes);\n cudaSafeCall(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));\n cudaSafeCall(cudaMemcpy(h_a, d_a, bytes, cudaMemcpyDeviceToHost));\n}<\/code><\/pre>\nDetails<\/h2>\n
The above codes work since there are specific behaviors for the streams. In short, the function cudaMemoryTest()<\/code> uses the default stream. The default stream will block all previous operations in all other streams (except for non-blocking streams). See my another post<\/a> for details.<\/p>\nSo when we call the function cudaMemoryTest()<\/code>, it will block until all previous async codes finish. So if there are bugs, the function cudaMemoryTest()<\/code> will fail. Thus, we can easily locate the bugs. <\/p>\nReference<\/h2>\n