尧图网站建设  精准测量内核块级耗时)
cuda-samples 之 clock_nvrtc用 libNVRTC 运行时编译与 clock() 精准测量内核块级耗时【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读本文围绕 NVIDIA cuda-samples 仓库中的clock_nvrtc示例展开讲解如何借助clock()设备端时钟函数对内核中每个线程块的执行耗时进行精确测量并同时演示 CUDA Runtime CompilationNVRTC运行时编译这一关键能力。读完本文你将掌握块级计时的原理与实现套路计时采样写入显存、主机端回读求差、NVRTC 从 CUDA 源码到 CUBIN 再到模块加载的完整调用链以及通过调整网格/块规模观察硬件占用与延迟隐藏的实验方法。示例定位与核心概念clock_nvrtc位于 cpp/0_Introduction/clock_nvrtc/与同目录的clock示例见 cpp/0_Introduction/clock/clock.cu解决同一问题——使用clock()函数测量内核中线程块block of threads的执行性能二者的区别在于本示例将设备端内核源码交由libNVRTC在运行时动态编译而非在构建期由 NVCC 预先编译。根据 README.md 的声明本示例涉及两个关键概念Performance Strategies性能策略通过设备端计时定位内核热点、验证硬件占用与并行度Runtime Compilation运行时编译使用 NVRTC 在程序运行期把 CUDA C 设备代码编译成 GPU 二进制CUBIN实现写源码即所得、免去构建期依赖的灵活分发。关于 NVRTC 的官方定义仓库顶层 README.md 的依赖章节也做了说明NVRTCCUDA RunTime Compilation是一个面向 CUDA C 的运行时编译库本示例正是依赖该库构建与运行的典型样例。支持环境与依赖矩阵按 README.md 的说明本示例的兼容性要求如下维度支持范围SM 架构SM 5.0、5.2、5.3、6.0、6.1、7.0、7.2、7.5、8.0、8.6、8.7、8.9、9.0操作系统Linux、Windows、QNXCPU 架构x86_64、aarch64CUDA APIDriver APIcuMemcpyDtoH、cuLaunchKernel、cuMemcpyHtoD、cuCtxSynchronize、cuMemAlloc、cuMemFree、cuModuleGetFunction Runtime APIcudaBlockSize、cudaGridSize运行依赖NVRTC 库运行前提是下载并安装与平台匹配的 CUDA Toolkit并确保依赖章节提到的 NVRTC 已随 Toolkit 一并安装。值得留意的是与 clock/README.md 相比本示例额外支持 QNX 系统与 aarch64 CPU 架构而从源码看虽然主程序 clock.cpp 以 Driver API 为主但仍通过dim3与cudaBlockSize/cudaGridSize这类 Runtime 风格结构组织网格与块维度因此 README 将两类 API 一并列入。块级计时原理为什么要在设备端打点GPU 内核中的线程块block是并行、乱序执行的——不同块可能由不同 SM流多处理器调度块之间没有任何同步机制。因此不能在主机端用clock()或gettimeofday()整体计时来区分每个块究竟花了多少时钟周期。正确做法是在每个块内、由tid 0的线程在进入核心计算前读取一次clock()将起始时间戳写入显存完成计算后再由tid 0读取一次clock()把结束时间戳写入显存另一区域内核结束后主机端把两组时间戳回读按块配对做差即可得到每个块的实际执行时钟周期数。由于clock()返回的是设备端高精度时钟周期计数clock_t且时间戳随每个块独立采样这种块内打点、显存暂存、主机求差的机制能够规避块间无同步带来的计时误差这正是该示例命名为 Clock 的原因也是其核心计时思路。内核源码剖析共享内存归约 双时间戳设备端内核定义在 clock_kernel.cu 中通过extern C导出符号便于运行时模块查找extern C __global__ void timedReduction(const float *input, float *output, clock_t *timer) { extern __shared__ float shared[]; const int tid threadIdx.x; const int bid blockIdx.x; if (tid 0) timer[bid] clock(); // 起始时间戳 shared[tid] input[tid]; shared[tid blockDim.x] input[tid blockDim.x]; 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; } } if (tid 0) output[bid] shared[0]; __syncthreads(); if (tid 0) timer[bid gridDim.x] clock(); // 结束时间戳 }几个实现细节值得展开动态共享内存extern __shared__ float shared[]声明动态共享内存实际大小由启动配置在运行时指定本示例为sizeof(float) * 2 * NUM_THREADS字节内核把input[tid]与input[tid blockDim.x]两份数据载入对应每块 512 个float树形归约for (int d blockDim.x; d 0; d / 2)是经典的分治归约循环配合__syncthreads()保证跨线程可见性最终shared[0]保存该块输入的最小值双时间戳布局起始戳写入timer[bid]结束戳写入timer[bid gridDim.x]偏移一个网格宽度两块区域连续存放主机端只需一次cuMemcpyDtoH即可全部回读再做timer[i NUM_BLOCKS] - timer[i]求差该内核刻意让每个块都执行相同的工作量同一份输入数据因此耗时差异可直接反映调度与硬件占用状况。主机端流程NVRTC 编译与 Driver API 启动主程序 clock.cpp 完整演示了NVRTC 编译 → 模块加载 → 内核启动 → 结果回读的运行时编译链路1. 编译期参数与输入初始化#define NUM_BLOCKS 64 #define NUM_THREADS 256示例默认以 64 个块、每块 256 线程运行输入为 512 个floatNUM_THREADS * 2。源码注释明确提示调整块数与线程数来观察硬件占用情况是本示例最有价值的实验。2. NVRTC 运行时编译内核kernel_file sdkFindFilePath(clock_kernel.cu, argv[0]); compileFileToCUBIN(kernel_file, argc, argv, cubin, cubinSize, 0);compileFileToCUBIN与loadCUBIN均定义在通用辅助头 Common/nvrtc_helper.h 中其内部完成了标准 NVRTC 调用序列nvrtcCreateProgram从文件内容创建 NVRTC 程序对象nvrtcCompileProgram执行编译编译选项为--gpu-architecturesm_majorminor——即针对当前运行设备的具体计算能力SM 版本现场编译这使同一份二进制可在不同 GPU 上自动适配是运行时编译相对预编译的优势nvrtcGetProgramLog/nvrtcGetCUBINSize/nvrtcGetCUBIN取回编译日志与最终 CUBIN 二进制loadCUBINcuInit→cuCtxCreate创建上下文 →cuModuleLoadData把 CUBIN 载入为 CUmodule并打印设备的 SM 计算能力。整个流程中所有 NVRTC 调用都被NVRTC_SAFE_CALL宏包裹同样定义于 nvrtc_helper.h失败时输出nvrtcGetErrorString错误描述并退出便于定位编译问题。3. 符号查找与内核启动CUfunction kernel_addr; checkCudaErrors(cuModuleGetFunction(kernel_addr, module, timedReduction));通过cuModuleGetFunction按名称从模块中取得内核句柄这也是内核必须extern C导出、避免名字改编的原因随后分配显存并把输入拷贝到设备端checkCudaErrors(cuMemAlloc(dinput, sizeof(float) * NUM_THREADS * 2)); checkCudaErrors(cuMemAlloc(doutput, sizeof(float) * NUM_BLOCKS)); checkCudaErrors(cuMemAlloc(dtimer, sizeof(clock_t) * NUM_BLOCKS * 2)); checkCudaErrors(cuMemcpyHtoD(dinput, input, sizeof(float) * NUM_THREADS * 2));启动内核时使用cuLaunchKernel注意它需要显式传递网格/块各维度、动态共享内存大小、可选流与参数数组void *arr[] {(void *)dinput, (void *)doutput, (void *)dtimer}; checkCudaErrors(cuLaunchKernel(kernel_addr, cudaGridSize.x, cudaGridSize.y, cudaGridSize.z, /* grid dim */ cudaBlockSize.x, cudaBlockSize.y, cudaBlockSize.z, /* block dim */ sizeof(float) * 2 * NUM_THREADS, /* 动态共享内存字节数 */ 0, /* stream */ arr[0], /* kernel 参数数组 */ 0)); /* 可选额外参数 */其中cudaBlockSize与cudaGridSize即 README 中列出的两个 Runtime API 相关符号dim3类型负责把 256×1×1 与 64×1×1 的启动配置传入 Driver API。4. 同步、回读与平均耗时统计checkCudaErrors(cuCtxSynchronize()); checkCudaErrors(cuMemcpyDtoH(timer, dtimer, sizeof(clock_t) * NUM_BLOCKS * 2));cuCtxSynchronize等待内核执行完毕随后cuMemcpyDtoH一次性回读 128 个时间戳主机端按块配对求差并取平均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);程序最终输出Average clocks/block 数值即每个块完成归约平均消耗的时钟周期数——这就是块级性能的直接量化指标。与 clock 示例的对照同一内核的 Runtime/Driver 两种形态仓库中同时提供了不使用 NVRTC 的 clock/clock.cu 版本两者内核timedReduction完全一致差异集中在主机端对比项clockclock.cuclock_nvrtcclock.cpp内核编译方式构建期由 NVCC 编译运行期由 libNVRTC 编译为 CUBIN设备端内存cudaMalloc/cudaFreecuMemAlloc/cuMemFree数据搬运cudaMemcpyHostToDevice/DeviceToHostcuMemcpyHtoD/cuMemcpyDtoH内核启动timedReductionNUM_BLOCKS, NUM_THREADS, sharedBytes(...)cuModuleGetFunctioncuLaunchKernel依赖库仅 Runtime APINVRTC Driver API见 CMakeLists.txt 中CUDA::nvrtc与CUDA::cuda_driver对照阅读这两个示例可以直观理解同一段 CUDA 内核代码如何在预编译NVCC与运行时编译NVRTC两种模式下组织主机端代码是理解 Driver API 与 Runtime API 差异的绝佳入门素材。构建与运行方式本示例采用 CMake 构建配置见 CMakeLists.txt。其中几个值得注意的构建细节通过find_package(CUDAToolkit REQUIRED)定位 Toolkit并将CMAKE_CUDA_ARCHITECTURES设为75 80 86 87 89 90 100 110 120即为主机端代码预编译的架构列表设备端内核实际架构由 NVRTC 在运行时按当前 GPU 决定默认编译带-lineinfo行号信息便于调试工具使用开启ENABLE_CUDA_DEBUG时切换为-G以支持 cuda-gdb链接目标为CUDA::nvrtc与CUDA::cuda_driver与 README 声明的依赖一致构建后通过add_custom_command将 clock_kernel.cu 复制到构建输出目录——因为该文件需在程序运行期被读取编译必须随可执行文件一起分发。典型构建流程假设已安装 CMake 3.20 与 CUDA Toolkit# 在 cuda-samples 仓库根目录下 mkdir -p build cd build cmake ../cpp/0_Introduction/clock_nvrtc make ./clock_nvrtc程序运行后会先打印CUDA Clock sample随后输出 GPU Device has SM x.y compute capability来自 nvrtc_helper.h 的loadCUBIN最后给出Average clocks/block ...的平均块耗时结果。实验视角块数、线程数与硬件占用clock.cpp 源码中保留了开发者针对早期 GPU 的实测数据该注释同时出现在 clock.cu 中用于指导调整网格规模blocks - clocks 1 - 3096 8 - 3232 16 - 3364 32 - 4615 64 - 9981对应的解读逻辑是块数少于 16 时部分 SM 处于空闲16 个块左右虽能占满所有 SM但每个 SM 上只有一个块无法通过块间切换隐藏访存延迟块数超过 32 后随并行度提升耗时近似线性增长。需要说明的是这些数值来自早期硬件环境在当前 GPU 上的绝对值会不同但其揭示的规律——通过调节NUM_BLOCKS与NUM_THREADS观察平均块耗时变化从而判断硬件是否被充分占用——依然有效。读者可以把修改这两个宏作为动手实验实际验证块太少 SM 空闲、块太多调度开销上升的硬件行为。小结clock_nvrtc用一个 200 行不到的示例同时串起了三条对 CUDA 开发者极具价值的知识线设备端clock()块级计时方法论块内打点、显存暂存、主机回读求差、NVRTC 运行时编译完整链路nvrtcCreateProgram → nvrtcCompileProgram → nvrtcGetCUBIN → cuModuleLoadData → cuLaunchKernel辅助实现见 nvrtc_helper.h以及网格规模与硬件占用的调优实验。它与 clock 示例的 Runtime/Driver 双版本对照也为理解两套 CUDA API 的编程差异提供了最直观的样本。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 cuda-samples 之 clock_nvrtc：用 libNVRTC 运行时编译与 clock() 精准测量内核块级耗时
cuda-samples 之 clock_nvrtc用 libNVRTC 运行时编译与 clock() 精准测量内核块级耗时【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读本文围绕 NVIDIA cuda-samples 仓库中的clock_nvrtc示例展开讲解如何借助clock()设备端时钟函数对内核中每个线程块的执行耗时进行精确测量并同时演示 CUDA Runtime CompilationNVRTC运行时编译这一关键能力。读完本文你将掌握块级计时的原理与实现套路计时采样写入显存、主机端回读求差、NVRTC 从 CUDA 源码到 CUBIN 再到模块加载的完整调用链以及通过调整网格/块规模观察硬件占用与延迟隐藏的实验方法。示例定位与核心概念clock_nvrtc位于 cpp/0_Introduction/clock_nvrtc/与同目录的clock示例见 cpp/0_Introduction/clock/clock.cu解决同一问题——使用clock()函数测量内核中线程块block of threads的执行性能二者的区别在于本示例将设备端内核源码交由libNVRTC在运行时动态编译而非在构建期由 NVCC 预先编译。根据 README.md 的声明本示例涉及两个关键概念Performance Strategies性能策略通过设备端计时定位内核热点、验证硬件占用与并行度Runtime Compilation运行时编译使用 NVRTC 在程序运行期把 CUDA C 设备代码编译成 GPU 二进制CUBIN实现写源码即所得、免去构建期依赖的灵活分发。关于 NVRTC 的官方定义仓库顶层 README.md 的依赖章节也做了说明NVRTCCUDA RunTime Compilation是一个面向 CUDA C 的运行时编译库本示例正是依赖该库构建与运行的典型样例。支持环境与依赖矩阵按 README.md 的说明本示例的兼容性要求如下维度支持范围SM 架构SM 5.0、5.2、5.3、6.0、6.1、7.0、7.2、7.5、8.0、8.6、8.7、8.9、9.0操作系统Linux、Windows、QNXCPU 架构x86_64、aarch64CUDA APIDriver APIcuMemcpyDtoH、cuLaunchKernel、cuMemcpyHtoD、cuCtxSynchronize、cuMemAlloc、cuMemFree、cuModuleGetFunction Runtime APIcudaBlockSize、cudaGridSize运行依赖NVRTC 库运行前提是下载并安装与平台匹配的 CUDA Toolkit并确保依赖章节提到的 NVRTC 已随 Toolkit 一并安装。值得留意的是与 clock/README.md 相比本示例额外支持 QNX 系统与 aarch64 CPU 架构而从源码看虽然主程序 clock.cpp 以 Driver API 为主但仍通过dim3与cudaBlockSize/cudaGridSize这类 Runtime 风格结构组织网格与块维度因此 README 将两类 API 一并列入。块级计时原理为什么要在设备端打点GPU 内核中的线程块block是并行、乱序执行的——不同块可能由不同 SM流多处理器调度块之间没有任何同步机制。因此不能在主机端用clock()或gettimeofday()整体计时来区分每个块究竟花了多少时钟周期。正确做法是在每个块内、由tid 0的线程在进入核心计算前读取一次clock()将起始时间戳写入显存完成计算后再由tid 0读取一次clock()把结束时间戳写入显存另一区域内核结束后主机端把两组时间戳回读按块配对做差即可得到每个块的实际执行时钟周期数。由于clock()返回的是设备端高精度时钟周期计数clock_t且时间戳随每个块独立采样这种块内打点、显存暂存、主机求差的机制能够规避块间无同步带来的计时误差这正是该示例命名为 Clock 的原因也是其核心计时思路。内核源码剖析共享内存归约 双时间戳设备端内核定义在 clock_kernel.cu 中通过extern C导出符号便于运行时模块查找extern C __global__ void timedReduction(const float *input, float *output, clock_t *timer) { extern __shared__ float shared[]; const int tid threadIdx.x; const int bid blockIdx.x; if (tid 0) timer[bid] clock(); // 起始时间戳 shared[tid] input[tid]; shared[tid blockDim.x] input[tid blockDim.x]; 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; } } if (tid 0) output[bid] shared[0]; __syncthreads(); if (tid 0) timer[bid gridDim.x] clock(); // 结束时间戳 }几个实现细节值得展开动态共享内存extern __shared__ float shared[]声明动态共享内存实际大小由启动配置在运行时指定本示例为sizeof(float) * 2 * NUM_THREADS字节内核把input[tid]与input[tid blockDim.x]两份数据载入对应每块 512 个float树形归约for (int d blockDim.x; d 0; d / 2)是经典的分治归约循环配合__syncthreads()保证跨线程可见性最终shared[0]保存该块输入的最小值双时间戳布局起始戳写入timer[bid]结束戳写入timer[bid gridDim.x]偏移一个网格宽度两块区域连续存放主机端只需一次cuMemcpyDtoH即可全部回读再做timer[i NUM_BLOCKS] - timer[i]求差该内核刻意让每个块都执行相同的工作量同一份输入数据因此耗时差异可直接反映调度与硬件占用状况。主机端流程NVRTC 编译与 Driver API 启动主程序 clock.cpp 完整演示了NVRTC 编译 → 模块加载 → 内核启动 → 结果回读的运行时编译链路1. 编译期参数与输入初始化#define NUM_BLOCKS 64 #define NUM_THREADS 256示例默认以 64 个块、每块 256 线程运行输入为 512 个floatNUM_THREADS * 2。源码注释明确提示调整块数与线程数来观察硬件占用情况是本示例最有价值的实验。2. NVRTC 运行时编译内核kernel_file sdkFindFilePath(clock_kernel.cu, argv[0]); compileFileToCUBIN(kernel_file, argc, argv, cubin, cubinSize, 0);compileFileToCUBIN与loadCUBIN均定义在通用辅助头 Common/nvrtc_helper.h 中其内部完成了标准 NVRTC 调用序列nvrtcCreateProgram从文件内容创建 NVRTC 程序对象nvrtcCompileProgram执行编译编译选项为--gpu-architecturesm_majorminor——即针对当前运行设备的具体计算能力SM 版本现场编译这使同一份二进制可在不同 GPU 上自动适配是运行时编译相对预编译的优势nvrtcGetProgramLog/nvrtcGetCUBINSize/nvrtcGetCUBIN取回编译日志与最终 CUBIN 二进制loadCUBINcuInit→cuCtxCreate创建上下文 →cuModuleLoadData把 CUBIN 载入为 CUmodule并打印设备的 SM 计算能力。整个流程中所有 NVRTC 调用都被NVRTC_SAFE_CALL宏包裹同样定义于 nvrtc_helper.h失败时输出nvrtcGetErrorString错误描述并退出便于定位编译问题。3. 符号查找与内核启动CUfunction kernel_addr; checkCudaErrors(cuModuleGetFunction(kernel_addr, module, timedReduction));通过cuModuleGetFunction按名称从模块中取得内核句柄这也是内核必须extern C导出、避免名字改编的原因随后分配显存并把输入拷贝到设备端checkCudaErrors(cuMemAlloc(dinput, sizeof(float) * NUM_THREADS * 2)); checkCudaErrors(cuMemAlloc(doutput, sizeof(float) * NUM_BLOCKS)); checkCudaErrors(cuMemAlloc(dtimer, sizeof(clock_t) * NUM_BLOCKS * 2)); checkCudaErrors(cuMemcpyHtoD(dinput, input, sizeof(float) * NUM_THREADS * 2));启动内核时使用cuLaunchKernel注意它需要显式传递网格/块各维度、动态共享内存大小、可选流与参数数组void *arr[] {(void *)dinput, (void *)doutput, (void *)dtimer}; checkCudaErrors(cuLaunchKernel(kernel_addr, cudaGridSize.x, cudaGridSize.y, cudaGridSize.z, /* grid dim */ cudaBlockSize.x, cudaBlockSize.y, cudaBlockSize.z, /* block dim */ sizeof(float) * 2 * NUM_THREADS, /* 动态共享内存字节数 */ 0, /* stream */ arr[0], /* kernel 参数数组 */ 0)); /* 可选额外参数 */其中cudaBlockSize与cudaGridSize即 README 中列出的两个 Runtime API 相关符号dim3类型负责把 256×1×1 与 64×1×1 的启动配置传入 Driver API。4. 同步、回读与平均耗时统计checkCudaErrors(cuCtxSynchronize()); checkCudaErrors(cuMemcpyDtoH(timer, dtimer, sizeof(clock_t) * NUM_BLOCKS * 2));cuCtxSynchronize等待内核执行完毕随后cuMemcpyDtoH一次性回读 128 个时间戳主机端按块配对求差并取平均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);程序最终输出Average clocks/block 数值即每个块完成归约平均消耗的时钟周期数——这就是块级性能的直接量化指标。与 clock 示例的对照同一内核的 Runtime/Driver 两种形态仓库中同时提供了不使用 NVRTC 的 clock/clock.cu 版本两者内核timedReduction完全一致差异集中在主机端对比项clockclock.cuclock_nvrtcclock.cpp内核编译方式构建期由 NVCC 编译运行期由 libNVRTC 编译为 CUBIN设备端内存cudaMalloc/cudaFreecuMemAlloc/cuMemFree数据搬运cudaMemcpyHostToDevice/DeviceToHostcuMemcpyHtoD/cuMemcpyDtoH内核启动timedReductionNUM_BLOCKS, NUM_THREADS, sharedBytes(...)cuModuleGetFunctioncuLaunchKernel依赖库仅 Runtime APINVRTC Driver API见 CMakeLists.txt 中CUDA::nvrtc与CUDA::cuda_driver对照阅读这两个示例可以直观理解同一段 CUDA 内核代码如何在预编译NVCC与运行时编译NVRTC两种模式下组织主机端代码是理解 Driver API 与 Runtime API 差异的绝佳入门素材。构建与运行方式本示例采用 CMake 构建配置见 CMakeLists.txt。其中几个值得注意的构建细节通过find_package(CUDAToolkit REQUIRED)定位 Toolkit并将CMAKE_CUDA_ARCHITECTURES设为75 80 86 87 89 90 100 110 120即为主机端代码预编译的架构列表设备端内核实际架构由 NVRTC 在运行时按当前 GPU 决定默认编译带-lineinfo行号信息便于调试工具使用开启ENABLE_CUDA_DEBUG时切换为-G以支持 cuda-gdb链接目标为CUDA::nvrtc与CUDA::cuda_driver与 README 声明的依赖一致构建后通过add_custom_command将 clock_kernel.cu 复制到构建输出目录——因为该文件需在程序运行期被读取编译必须随可执行文件一起分发。典型构建流程假设已安装 CMake 3.20 与 CUDA Toolkit# 在 cuda-samples 仓库根目录下 mkdir -p build cd build cmake ../cpp/0_Introduction/clock_nvrtc make ./clock_nvrtc程序运行后会先打印CUDA Clock sample随后输出 GPU Device has SM x.y compute capability来自 nvrtc_helper.h 的loadCUBIN最后给出Average clocks/block ...的平均块耗时结果。实验视角块数、线程数与硬件占用clock.cpp 源码中保留了开发者针对早期 GPU 的实测数据该注释同时出现在 clock.cu 中用于指导调整网格规模blocks - clocks 1 - 3096 8 - 3232 16 - 3364 32 - 4615 64 - 9981对应的解读逻辑是块数少于 16 时部分 SM 处于空闲16 个块左右虽能占满所有 SM但每个 SM 上只有一个块无法通过块间切换隐藏访存延迟块数超过 32 后随并行度提升耗时近似线性增长。需要说明的是这些数值来自早期硬件环境在当前 GPU 上的绝对值会不同但其揭示的规律——通过调节NUM_BLOCKS与NUM_THREADS观察平均块耗时变化从而判断硬件是否被充分占用——依然有效。读者可以把修改这两个宏作为动手实验实际验证块太少 SM 空闲、块太多调度开销上升的硬件行为。小结clock_nvrtc用一个 200 行不到的示例同时串起了三条对 CUDA 开发者极具价值的知识线设备端clock()块级计时方法论块内打点、显存暂存、主机回读求差、NVRTC 运行时编译完整链路nvrtcCreateProgram → nvrtcCompileProgram → nvrtcGetCUBIN → cuModuleLoadData → cuLaunchKernel辅助实现见 nvrtc_helper.h以及网格规模与硬件占用的调优实验。它与 clock 示例的 Runtime/Driver 双版本对照也为理解两套 CUDA API 的编程差异提供了最直观的样本。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考