CUDA Streams 与 CPU 回调实现多线程异构计算cuda-samples simpleCallback 示例深度解析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读simpleCallback 是 NVIDIA CUDA Samples 中位于 cpp/0_Introduction/simpleCallback 的入门级示例它演示了如何利用 CUDA 5.0 引入的流与事件 CPU 回调cudaStreamAddCallback配合多线程编程构建CPU 预处理 → GPU 计算 → CPU 后处理的异构计算流水线。读完本文你将掌握回调函数的注册方式与触发时机、如何用多线程 多流让 8 个异构工作负载并发运行在多块 GPU 上以及 CUDA 运行时 API 的线程安全用法并可在本仓库中直接构建运行验证。一、示例定位什么是 CPU Callback在 CUDA 5.0 之前CPU 线程若要等待流中 GPU 工作完成通常只能阻塞式同步如cudaStreamSynchronize或轮询事件这会打断 CPU 与 GPU 的流水线节奏。CUDA 5.0 新增的cudaStreamAddCallback允许开发者把一个主机端回调函数挂到指定流上当该流中此前所有排队的操作全部完成后回调函数会在 CPU 侧被调用。这正是本示例名字 simpleCallback 的由来。回调让GPU 干完活自动通知 CPU 接手成为可能CPU 线程不必空等GPU 也能持续流水线化执行。README 中将其关键概念归纳为三点CUDA Streams流、Callback Functions回调函数、Multithreading多线程。示例声明的目标很明确This sample implements multi-threaded heterogeneous computing workloads with the new CPU callbacks for CUDA streams and events introduced with CUDA 5.0.结合 CUDA API 本身的线程安全特性实现在 CPU 线程与 GPU 之间自由流动的异构工作负载变得简单高效。二、工作负载架构CPU preprocess → GPU process → CPU postprocess整个示例的运行模型是三段式流水线见 simpleCallback.cu 文件头部注释CPU 预处理preprocess每个工作负载由专属 CPU 线程负责在线程内生成输入数据h_dataGPU 计算processCPU 线程通过异步 APIcudaMemcpyAsync、内核启动、cudaMemcpyAsync回拷把工作排入专属流然后线程退出GPU 继续在后台执行CPU 后处理postprocess流的 CPU 回调触发后重新起一个 CPU 线程消费结果并释放资源。各 CPU 处理步骤由各自专用线程承担GPU 工作则被分发到系统中所有可用 GPU 上——这是一种典型的多线程异构计算流水线。2.1 参数与数据结构工作负载的数量与规模由两个常量控制simpleCallback.cu#L51-L52const int N_workloads 8; // 异构工作负载个数 const int N_elements_per_workload 100000; // 每个负载的元素个数每个工作负载的状态封装在heterogeneous_workload结构体中simpleCallback.cu#L58-L68struct heterogeneous_workload { int id; // 工作负载编号 int cudaDeviceID; // 指派到的 GPU 编号 int *h_data; // 主机端可页面锁定数据 int *d_data; // 设备端数据 cudaStream_t stream; // 专属 CUDA 流 bool success; // 后处理验证结果 };GPU 侧的内核极其简单——对每个元素加 1simpleCallback.cu#L70-L76__global__ void incKernel(int *data, int N) { int i blockIdx.x * blockDim.x threadIdx.x; if (i N) data[i]; }它承担GPU process环节的代表性计算任务便于验证数据流正确性。2.2 核心 API 清单README 列出的 CUDA 运行时 API 涉及资源分配、流管理、异步传输与回调注册API作用cudaGetDeviceCount/cudaGetDeviceProperties/cudaSetDevice枚举 GPU、查询属性、为当前线程选定设备cudaStreamCreate/cudaStreamDestroy创建 / 销毁专属流cudaMalloc/cudaFree设备内存分配 / 释放cudaHostAlloc/cudaFreeHost页面锁定主机内存分配 / 释放cudaHostAllocPortable标志cudaMemcpyAsync异步主机↔设备拷贝需流与页面锁定内存配合cudaStreamAddCallbackCUDA 5.0 新增在流上注册 CPU 回调三、核心机制逐段解析从 launch 线程到回调再到 postprocess3.1 launch 线程CPU 预处理 排入 GPU 流水线每个工作负载对应一个 launch 线程simpleCallback.cu#L78-L118。其执行顺序如下选定 GPUcudaSetDevice(workload-cudaDeviceID)—— CUDA 运行时 API 是线程安全的每个 CPU 线程都可独立选卡分配资源cudaStreamCreate创建专属流cudaMalloc分配设备端d_datacudaHostAlloc(..., cudaHostAllocPortable)分配页面锁定主机内存h_data——cudaHostAllocPortable标志使该内存对所有设备可用这是异步拷贝的前提CPU 生成数据h_data[i] workload-id i;排入 GPU 流水线不阻塞 CPU 线程dim3 block(512); dim3 grid((N_elements_per_workload block.x - 1) / block.x); checkCudaErrors(cudaMemcpyAsync(workload-d_data, workload-h_data, N_elements_per_workload * sizeof(int), cudaMemcpyHostToDevice, workload-stream)); incKernelgrid, block, 0, workload-stream(workload-d_data, N_elements_per_workload); checkCudaErrors(cudaMemcpyAsync(workload-h_data, workload-d_data, N_elements_per_workload * sizeof(int), cudaMemcpyDeviceToHost, workload-stream));注意三个异步操作H2D 拷贝、内核、D2H 拷贝共享同一个流因此它们在 GPU 上严格按序执行不同工作负载拥有各自的流从而可以并发执行。注册 CPU 回调并结束线程checkCudaErrors(cudaStreamAddCallback(workload-stream, myStreamCallback, workload, 0)); CUT_THREADEND; // CPU 线程结束GPU 继续处理数据...launch 线程到此退出GPU 仍在后台处理该负载——这正是异步 回调带来的流水线效果。3.2 回调函数GPU 完成后的接力棒回调函数声明与实现如下simpleCallback.cu#L56、simpleCallback.cu#L146-L153void CUDART_CB myStreamCallback(cudaStream_t stream, cudaError_t status, void *data); void CUDART_CB myStreamCallback(cudaStream_t stream, cudaError_t status, void *data) { // 检查流中全部操作完成后的 GPU 状态 checkCudaErrors(status); // 启动新的 CPU 工作线程继续在 CPU 上处理 cutStartThread(postprocess, data); }回调签名包含三个参数流句柄、流执行完毕后的错误状态、cudaStreamAddCallback传入的用户数据指针此处即heterogeneous_workload*。函数体内首先用checkCudaErrors校验 GPU 执行状态——这意味着回调也是集中处理流异步错误的天然位置随后通过cutStartThread派生 postprocess 线程完成CPU 后处理。3.3 postprocess 线程CPU 后处理 资源回收 结果汇合postprocess 线程simpleCallback.cu#L120-L144负责再次选卡cudaSetDevice新线程必须重新声明其使用的设备验证结果h_data[i] i workload-id 1——因为h_data[i]初值为id i经过内核1后应为id i 1示例取前N_workloads个元素做正确性断言并把结果累计到workload-success释放资源cudaFree设备内存、cudaFreeHost页面锁定内存、cudaStreamDestroy销毁流汇合主线程cutIncrementBarrier(thread_barrier)递增屏障计数通知主线程本负载已结束。3.4 main 线程GPU 枚举、负载分发与汇合主线程simpleCallback.cu#L155-L215的执行流程枚举 GPUcudaGetDeviceCount获取 GPU 数量示例最多支持 32 块逐卡查询能力通过cudaGetDeviceProperties获取deviceProp.major/deviceProp.minor计算 SM 版本号并判断是否支持回调SMversion deviceProp.major 4 deviceProp.minor; printf(GPU[%d] %s supports SM %d.%d, devid, deviceProp.name, deviceProp.major, deviceProp.minor); printf(, %s GPU Callback Functions\n, (SMversion 0x11) ? capable : NOT capable); if (SMversion 0x11) { gpuInfo[max_gpus] devid; }SM 版本不低于 1.10x11的 GPU 才被纳入回调可用集合gpuInfo[] 3.创建屏障cutCreateBarrier(N_workloads)等待 8 个负载全部完成 4.分发负载为每个i启动一个 launch 线程设备分配采用轮询策略workloads[i].cudaDeviceID gpuInfo[i % max_gpus];——多块 GPU 时工作负载会被均匀分摊 5.汇合与判定cutWaitForBarrier(thread_barrier)阻塞主线程直到所有负载完成最后汇总success并输出Success/Failure以相应退出码结束进程。四、配套的可移植多线程库multithreading.h / .cpp示例没有直接使用裸pthread/Win32API而是封装了跨平台线程库 multithreading.h 与 multithreading.cpp这正是它能在 Linux 与 Windows 上运行的基础README 的 Supported OSes 也标注了这两个平台。4.1 线程与屏障抽象在 multithreading.h 中通过条件编译区分平台WindowsCUTThread为HANDLE线程例程类型CUT_THREADROUTINE为unsigned(WINAPI*)(void*)屏障基于CRITICAL_SECTION 事件对象实现POSIXLinux 等CUTThread为pthread_t例程为void*(*)(void*)屏障基于pthread_mutex_tpthread_cond_t实现。对外暴露的接口包括CUTThread cutStartThread(CUT_THREADROUTINE, void *data); // 创建线程 void cutEndThread(CUTThread thread); // 等待单个线程结束 void cutWaitForThreads(const CUTThread *threads, int num); // 等待多个线程 CUTBarrier cutCreateBarrier(int releaseCount); // 创建屏障 void cutIncrementBarrier(CUTBarrier *barrier); // 计数 1到达 releaseCount 时放行 void cutWaitForBarrier(CUTBarrier *barrier); // 等待屏障放行 void cutDestroyBarrier(CUTBarrier *barrier); // 销毁屏障4.2 屏障的 POSIX 实现要点在 multithreading.cpp#L120-L143 中cutIncrementBarrier以互斥锁保护计数自增当计数达到releaseCount时用条件变量唤醒等待者cutWaitForBarrier则在条件变量上循环等待直到count releaseCount。主线程正是依靠这个屏障机制一次性等待 8 个由回调派生的 postprocess 线程全部完成。这个专用 CPU 线程 屏障汇合的框架与本示例的回调机制天然契合回调派生的线程最终都通过屏障向主线程汇报避免了主线程对每个工作负载做忙等轮询。五、错误检查checkCudaErrors 的使用惯例示例中所有 CUDA API 调用都被checkCudaErrors包裹如cudaSetDevice、cudaMemcpyAsync、cudaStreamAddCallback这是 CUDA Samples 的通用惯例。其宏定义在 Common/helper_cuda.h#L585-L598template typename T void check(T result, char const *const func, const char *const file, int const line) { if (result) { fprintf(stderr, CUDA error at %s:%d code%d(%s) \%s\ \n, file, line, static_castunsigned int(result), _cudaGetErrorEnum(result), func); exit(EXIT_FAILURE); } } #define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__)当任一调用返回非cudaSuccess时程序会输出出错文件、行号、错误码对应的可读字符串与出错 API 名称后退出极大方便了多线程场景下的问题定位。六、构建与运行6.1 构建方式本示例使用 CMake 构建其 CMakeLists.txt 关键点如下要求 CMake ≥ 3.20启用C CXX CUDA三种语言find_package(CUDAToolkit REQUIRED)通过set(CMAKE_CUDA_ARCHITECTURES 75 80 86 87 89 90 100 110 120)预置编译目标架构对应 SM 7.5 至 SM 12.0目标simpleCallback由simpleCallback.cu与multithreading.cpp共同编译启用 C17/CUDA 17 标准与CUDA_SEPARABLE_COMPILATION可分离编译支持ENABLE_CUDA_DEBUG选项对应 nvcc-G调试模式通过 cmake/InstallSamples.cmake 接入统一的安装配置。按仓库根目录 README.md 的说明可任选以下方式构建# 方式一在仓库根目录统一构建 mkdir build cd build cmake .. make simpleCallback -j$(nproc) # 方式二直接进入示例目录单独构建 cd cpp/0_Introduction/simpleCallback mkdir build cd build cmake .. makeWindows 下可用 Visual Studio 的 CMake 语言服务直接导入该目录或用x64 Native Tools Command Prompt for VS执行cmake .. -G Visual Studio 16 2019 -A x64后打开生成的解决方案构建。6.2 运行与预期输出构建产物直接运行即可无需任何输入数据文件./simpleCallback预期输出大致如下Starting simpleCallback Found N CUDA capable GPUs GPU[0] GPU 名称 supports SM x.y, capable GPU Callback Functions ... M GPUs available to run Callback Functions Starting 8 heterogeneous computing workloads Total of 8 workloads finished: Success由于 8 个负载被分发到gpuInfo[i % max_gpus]选定的 GPU 上单卡环境会全部落在同一设备多卡环境则被自动分摊——程序最终依据全部 8 个负载的验证结果输出Success或Failure。七、关键概念小结概念在本示例中的体现CUDA Streams每个工作负载一个专属流流内异步操作有序执行流间并发执行Callback FunctionscudaStreamAddCallback注册主机回调流内操作全部完成后自动触发Multithreadinglaunch 线程负责预处理与排程postprocess 线程由回调派生主线程用屏障汇合异步数据搬运cudaHostAlloc(cudaHostAllocPortable)cudaMemcpyAsync实现不阻塞 CPU 的 H2D/D2H 传输运行时 API 线程安全各 CPU 线程独立cudaSetDevice选择设备配合流实现异构流水线延伸提示本示例采用的回调机制是 CUDA 5.0 时代的经典做法在更新的 CUDA 版本中CUDA 还提供了主机函数Host Function等机制来实现类似GPU 完成后在 CPU 侧执行的编排读者可结合 CUDA Runtime API 官方文档进一步了解演进。八、参考与延伸阅读示例 READMEcpp/0_Introduction/simpleCallback/README.md含支持的 SM 架构 5.0–9.0、支持的操作系统 Linux/Windows、支持的 CPU 架构 x86_64/armv7l 等清单示例源码simpleCallback.cu多线程封装multithreading.h、multithreading.cpp构建配置cpp/0_Introduction/simpleCallback/CMakeLists.txt通用错误检查宏Common/helper_cuda.h仓库统一构建说明与前置条件安装 CUDA ToolkitREADME.md同一目录下其他入门示例如 simpleStreams、simpleHyperQ、simpleMultiGPU可帮助对比理解流、多队列与多 GPU 的编程模型前置条件方面与仓库其他示例一致安装与你的平台对应的 CUDA Toolkit 即可开始构建运行本示例。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考