clock64()

__device__ unsigned long long int clock64();

• 返回 当前 GPU 线程所处 SM（Streaming Multiprocessor）上的时钟计数器值。

• 单位是 GPU 核心的时钟周期（不是秒）。

• 这个计数器从 GPU 启动后一直累加，直到溢出。

• 对比 clock()：

  • clock() 返回 32 位计数器（容易溢出）。

  • clock64() 返回 64 位计数器，更适合长期计时或者需要高精度的测量。

使用场景：线程内微基准

__global__ void kernel() {
    unsigned long long start = clock64();
    // 你的代码
    unsigned long long end = clock64();
    printf("Thread %d took %llu cycles\n", threadIdx.x, end - start);
}

• • 可以统计单个线程执行某段代码的 GPU 时钟周期。

• 测量 warp 内/ block 内的执行时间

  • 注意，clock64() 返回的是每个线程独立的时钟计数。

  • 不同 SM 上线程的 clock64() 可能不是同步的，所以跨 SM 的比较不精确。

• 高精度调试

  • 比如调试原子操作、共享内存延迟、warp 内 shuffle 等操作的开销。

  • 适合微调 kernel 性能。

注意事项

• 返回值类型是 unsigned long long。

• 仅在 device 上可用，不能在 host 上直接调用。

• 由于是 SM 局部计数器，跨 SM 比较时间可能不准确。

• 对比 host 的 cudaEvent 或 chrono，它提供的是硬件周期级的细粒度测量，适合分析单线程或 warp 的性能瓶颈。

__isGlobal / __isShared

__isGlobal和 __isShared 是 CUDA 提供的 内置设备函数，主要用于 在编译期检查指针指向的内存类型，方便做模板编程或泛型核函数优化。它们本身不会改变程序行为，只是返回布尔值，告诉你这个指针是指向 全局内存（global memory） 还是 共享内存（shared memory）。

__device__ __device_builtin__ bool __isGlobal(const void* ptr);

作用：

• 返回 true 表示指针 ptr 指向的是 全局内存。

• 返回 false 表示不是全局内存（可能是共享内存、寄存器或者常量内存等）。

__device__ __device_builtin__ bool __isShared(const void* ptr);

作用：

• 返回 true 表示指针 ptr 指向的是 共享内存（__shared__）。

• 返回 false 表示不是共享内存。

模板泛型核函数：当你写一个模板核函数，需要在不同内存类型上执行相同逻辑，但访问方式不同时，可以在编译期判断内存类型：

template<typename T>
__device__ void load_data(T* ptr) {
    if (__isGlobal(ptr)) {
        // 使用 global memory 访问优化
    } else if (__isShared(ptr)) {
        // 使用 shared memory 优化
    }
}

__trap()

__trap() 是 CUDA 提供的 设备端调试函数，主要用在 GPU 代码中主动触发异常或中断，用于调试目的。它类似 CPU 上的 assert(0) 或者 __builtin_trap()，会让 GPU 停止执行并生成 device-side trap。

__device__ __device_builtin__ void __trap(void);

• 只能在 __device__ 或 __global__ 函数里调用。

• 不返回值。

• 在 运行时会中断当前线程的执行。

• 通常会导致 CUDA kernel 异常退出，可以配合 cuda-memcheck 或 Nsight 进行调试。

__global__ void exampleKernel(int* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx >= N) {
        // 出现非法访问时，主动触发 GPU trap
        __trap();
    }

    data[idx] = idx;
}

• 如果 idx >= N，GPU 会立即触发中断，停止该线程的执行，并在调试工具中捕获。

• 对应 host 端，可能看到类似：

cudaErrorIllegalAddress: an illegal memory access was encountered

典型用途

1. 调试非法访问

   在 warp/block 内检测索引越界、非法指针等。

2. 模板/分支调试

   在条件分支中验证不同代码路径是否被执行。

3. 结合 cuda-memcheck

   GPU kernel 异常捕获、定位问题线程。

注意事项

• 生产环境 不要使用，会导致 kernel 异常退出。

• 调试时，最好配合 小数据量 或 单 block 测试，否则大量线程触发 __trap() 会产生海量报错。

• 对比 CPU 上 assert()，它是线程级的中断，不会影响其他线程继续执行（但整个 kernel 可能失败）。