CUDA编程 - 测量每个block内线程块的执行时间 - 如何应用到自己的项目中 - clock()

完整代码与例程目的

代码地址：https://github.com/NVIDIA/cuda-samples/tree/v11.8/Samples/0_Introduction/clock

clock() 直接使用GPU硬件的时钟计数器，精度更高（时钟周期级别）。
如果使用 cudaEvent 或 nsys工具更侧重于测量 kernel 整体耗时

本示例演示了如何使用时钟函数精确测量内核中线程块的执行性能。由于线程块是并行且无序执行的，且块间缺乏同步机制，我们通过为每个块单独测量时钟值的方式实现性能监控。所有时钟采样值将被写入设备内存。

关键要点说明：
测量对象：单个线程块(block)的执行周期数
并行特性：块间并行执行且无固定顺序 实现限制：块间无同步机制 → 独立测量每个块
数据存储：时钟采样值直接写入显存(device memory)

完整代码：
clock.cu

// System includes
#include <assert.h>
#include <stdint.h>
#include <stdio.h>// CUDA runtime
#include <cuda_runtime.h>// helper functions and utilities to work with CUDA
#include <helper_cuda.h>
#include <helper_functions.h>// This kernel computes a standard parallel reduction and evaluates the
// time it takes to do that for each block. The timing results are stored
// in device memory.
__global__ static void timedReduction(const float *input, float *output,clock_t *timer) {// __shared__ float shared[2 * blockDim.x];extern __shared__ float shared[];const int tid = threadIdx.x;const int bid = blockIdx.x;if (tid == 0) timer[bid] = clock();// Copy input.shared[tid] = input[tid];shared[tid + blockDim.x] = input[tid + blockDim.x];// Perform reduction to find minimum.for (int d = blockDim.x; d > 0; d /= 2) {__syncthreads();if (tid < d) {float f0 = shared[tid];float f1 = shared[tid + d];if (f1 < f0) {shared[tid] = f1;}}}// Write result.if (tid == 0) output[bid] = shared[0];__syncthreads();if (tid == 0) timer[bid + gridDim.x] = clock();
}#define NUM_BLOCKS 64
#define NUM_THREADS 256// It's interesting to change the number of blocks and the number of threads to
// understand how to keep the hardware busy.
//
// Here are some numbers I get on my G80:
//    blocks - clocks
//    1 - 3096
//    8 - 3232
//    16 - 3364
//    32 - 4615
//    64 - 9981
//
// With less than 16 blocks some of the multiprocessors of the device are idle.
// With more than 16 you are using all the multiprocessors, but there's only one
// block per multiprocessor and that doesn't allow you to hide the latency of
// the memory. With more than 32 the speed scales linearly.// Start the main CUDA Sample here
int main(int argc, char **argv) {printf("CUDA Clock sample\n");// This will pick the best possible CUDA capable deviceint dev = findCudaDevice(argc, (const char **)argv);float *dinput = NULL;float *doutput = NULL;clock_t *dtimer = NULL;clock_t timer[NUM_BLOCKS * 2];float input[NUM_THREADS * 2];for (int i = 0; i < NUM_THREADS * 2; i++) {input[i] = (float)i;// std::cout << input[i] << std::endl;}checkCudaErrors(cudaMalloc((void **)&dinput, sizeof(float) * NUM_THREADS * 2));checkCudaErrors(cudaMalloc((void **)&doutput, sizeof(float) * NUM_BLOCKS));checkCudaErrors(cudaMalloc((void **)&dtimer, sizeof(clock_t) * NUM_BLOCKS * 2));checkCudaErrors(cudaMemcpy(dinput, input, sizeof(float) * NUM_THREADS * 2,cudaMemcpyHostToDevice));timedReduction<<<NUM_BLOCKS, NUM_THREADS, sizeof(float) * 2 * NUM_THREADS>>>(dinput, doutput, dtimer);checkCudaErrors(cudaMemcpy(timer, dtimer, sizeof(clock_t) * NUM_BLOCKS * 2,cudaMemcpyDeviceToHost));checkCudaErrors(cudaFree(dinput));checkCudaErrors(cudaFree(doutput));checkCudaErrors(cudaFree(dtimer));long double avgElapsedClocks = 0;for (int i = 0; i < NUM_BLOCKS; i++) {avgElapsedClocks += (long double)(timer[i + NUM_BLOCKS] - timer[i]);}avgElapsedClocks = avgElapsedClocks / NUM_BLOCKS;printf("Average clocks/block = %Lf\n", avgElapsedClocks);return EXIT_SUCCESS;
}

代码拆解与复用

一、计时机制设计原理（块级独立计时）（应用到自己的项目中）

每个线程块独立记录起始/结束时钟值：

__global__ void timedReduction(...) {if (tid == 0) timer[bid] = clock();       // 块开始时间// ... 计算逻辑if (tid == 0) timer[bid + gridDim.x] = clock(); // 块结束时间
}

这种设计避免了块间同步问题，因为GPU的SM（流处理器簇）会并行执行多个块，无法保证全局同步

所以可以直接参考这种方式，应用到自己的项目中进行计时。

二、关键实现细节

2.1、​共享内存优化

通过extern __shared__ float shared[] 声明动态共享内存：

__global__ static void timedReduction(...) {extern __shared__ float shared[];// 加载数据到共享内存shared[tid] = input[tid];shared[tid + blockDim.x] = input[...];
}

确保线程块内数据访问的高效性，避免全局内存延迟对计时的影响

2.2、​​同步控制机制

使用__syncthreads()保证块内线程同步：

for (int d = blockDim.x; d > 0; d /= 2) {__syncthreads();  // 同步所有线程// 归约计算
}

2.3、​统计处理策略

主机端计算每个块的时钟周期差：

long double avgElapsedClocks = 0;
for (int i = 0; i < NUM_BLOCKS; i++) {avgElapsedClocks += (timer[i + NUM_BLOCKS] - timer[i]);
}

通过平均多个块的执行时间，消除硬件调度波动的影响。可以调整 block 和 thread 数量进行测试。

三、优势和劣势

优势：

• 避免全局同步开销，适应GPU并行执行特性
• 块级细粒度测量，定位性能瓶颈更精确
• 无需额外硬件支持（如CUDA事件需要特定计算能力）

局限：

• 不同SM时钟域可能存在微小偏差
• 无法测量内核启动/数据传输时间
• 需手动处理线程束发散（Warp Divergence）的影响