1. 我们和CUDA基本上只有两种交互：调用CUDA API和调用kernel。
2. API有两种：同步和异步；kernel只有异步执行，但是launch有两种：立即launch和添加到某个stream等待launch。
   1. 同步API（cudaMemcpy）：有返回值，我们可以了解到其是否执行成功。
   2. 异步API（cudaMemcpyAsync）：也会返回error吗？即使返回了success也无法确定它可以顺利执行结束。【AI：异步 API（如 cudaMemcpyAsync）提交时返回的错误，只检查参数是否合法，不能说明执行是否真的成功】结论：异步API和kernel的检错方式一样
3. kernel没有返回值，我们如何知道其是否成功执行了呢？首先kernel应该有两种错误：launch error和runtime error。
4. 关于kernel的launch error：可能的原因包括不恰当的gridSize、blockSize、stream等
   1. 立即launch：在调用kernel后执行cudaGetLastError【捕获 内核启动错误（launch errors）】（启动失败是cudaErrorLaunchFailure）
   2. 添加到stream中：如何确保它可以成功launch呢？其实不用担心这种问题：launch error 是在你调用 kernel 的那一刻就被检测的，而不是等到它在 stream 中真正开始执行的时候。所以【立即launch】和【添加到stream中】是不用区分的（对于检测 launch error）
5. 关于kernel的runtime error：可能的原因包括访问越界等
   1. 在kernel执行结束后（注意这个时机），可以获取错误信息【捕获 内核执行错误（runtime execution errors）】，但是这里不一定使用cudaGetLastError。为什么呢？因为我们不进行明确的同步操作的话，是不能确保一个kernel执行完成的；如果执行Sync的话，这个Sync的API就直接可以返回error了。
   2. 但是如果是多个kernel在一个stream中执行，我不想每个kernel结束后都和host同步一下check error应该怎么做呢？实际上没有太好的办法，可以使用CUDA Event检测进行到哪一步了，但是CUDA Event不会记录错误信息。另外，runtime error会导致整个stream停止，那其实使用Event检测停在哪一步了也没啥太大用处？

1. 在kernel内部可以使用printf但是非常不推荐
2. 在kernel内部不可以使用std::cout
3. export CUDA_LAUNCH_BLOCKING=1
   1. 通常 CUDA 的 kernel 和异步 API（如 cudaMemcpyAsync、kernel launch）是异步的。
   2. 作用：强制所有 CUDA API 变成同步。每个操作执行完毕才返回。等效于在每个 kernel / memcpy 后加了cudaDeviceSynchronize。
4. 一个kernel的错误，在下一次调用CUDA API的时候会被返回
5. 直接调用cudaGetLastError不太方便，建议包装一下，cudaCheckError / cudaSafeCall之类的
6. cudaGetLastError：获取最近一次 CUDA 错误并清除该状态
7. cudaPeekAtLastError：查询最近一次 CUDA 错误但不清除状态
8. 注意：LastError是全局的（跨stream），会被一次查询操作清空，或被另一个error覆盖
9. CUDA API 返回错误，不会自动清除 CUDA 内部的错误状态
10. CUDA 支持 assert 调试
11. cudaGetLastError是否可以用于kernel runtime error的检测？
12. 如果一个 kernel 在执行时发生了 runtime error（比如非法内存访问、断言失败等），它所在的 stream 会立刻停止，后续的任务不会继续执行。

从网上抄来的几个处理错误的宏

// 应该是因为定义在了头文件里，所以用了static
static void HandleError( cudaError_t err,
                         const char *file,
                         int line ) {
    if (err != cudaSuccess) {
        printf( "%s in %s at line %d\n", cudaGetErrorString( err ),
                file, line );
        exit( EXIT_FAILURE );
    }
}
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
 
 
#define HANDLE_NULL( a ) {if (a == NULL) { \
                            printf( "Host memory failed in %s at line %d\n", \
                                    __FILE__, __LINE__ ); \
                            exit( EXIT_FAILURE );}}

// check cuda error
inline void check(cudaError_t call, const char* file, const int line)
{
    if (call != cudaSuccess)
    {
        std::cout << "cuda error: " << cudaGetErrorName(call) << std::endl;
        std::cout << "at file: " << file << ", line: " << line << std::endl;
        std::cout << cudaGetErrorString(call) << std::endl;
    }
}
 
#define CHECK(call) (check(call, __FILE__, __LINE__))

// Generated by AI
// 最简单的情况：可以接受Device和Host一直同步，也就是不需要多个kernel并行
// 每个kernel的launch error和runtime error都可以很方便地捕获
 
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
 
// Kernel A on stream1
kernelA<<<grid, block, 0, stream1>>>(...);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    fprintf(stderr, "[Kernel A] Launch Error: %s\n", cudaGetErrorString(err));
    exit(EXIT_FAILURE);
}
err = cudaStreamSynchronize(stream1); // 同步
if (err != cudaSuccess) {
    fprintf(stderr, "[Kernel A] Execution Error: %s\n", cudaGetErrorString(err));
    exit(EXIT_FAILURE);
}
 
// Kernel B on stream2
kernelB<<<grid, block, 0, stream2>>>(...);
err = cudaGetLastError();
if (err != cudaSuccess) {
    fprintf(stderr, "[Kernel B] Launch Error: %s\n", cudaGetErrorString(err));
    exit(EXIT_FAILURE);
}
err = cudaStreamSynchronize(stream2);
if (err != cudaSuccess) {
    fprintf(stderr, "[Kernel B] Execution Error: %s\n", cudaGetErrorString(err));
    exit(EXIT_FAILURE);
}

Nsight和NVTX