CUDA 示例解析UnifiedMemoryStreams —— 用 OpenMP 多线程与 CUDA Streams 在统一内存上调度混合主机/设备任务【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读UnifiedMemoryStreams 是 cuda-samples 仓库cpp/0_Introduction目录下的入门级示例演示如何在单 GPU上结合OpenMP 多线程与CUDA Streams让一批数据全部驻留于Unified Memory统一内存的任务被主机线程和 GPU 协同消费。本文基于该示例的 README 与其完整 源码实现讲解任务划分、内存附着memory attach、host/device 双路径执行与 cuBLAS 调用的底层细节读完即可掌握如何在单 GPU 上编排一个任务消费者式的混合异构工作流。示例要解决什么问题该示例实现了一个简单的任务消费者task consumer程序生成一批大小随机、数据存于统一内存的任务然后通过线程与流并行地消费这些任务。关键点在于——任务既可能在主机上执行也可能在设备上执行任务规模较小时直接在 CPU 上用朴素循环完成矩阵向量乘DGEMV任务规模较大时交给 GPU通过cuBLAS的cublasDgemv在对应的 CUDA Stream 中执行。由于所有数据都是用cudaMallocManaged分配的托管内存CPU 与 GPU 无需显式拷贝即可访问同一份数据这正是 Unified Memory 的核心价值。原 README 将其核心概念归纳为CUDA Systems Integration、OpenMP、CUBLAS、Multithreading、Unified Memory、CUDA Streams and Events。支持的平台与运行前提根据 UnifiedMemoryStreams/README.md支持的操作系统Linux、Windows支持的 CPU 架构x86_64、armv7l构建依赖OpenMP、UVMUnified Virtual Memory、CUBLAS前置条件下载并安装与平台对应的 CUDA Toolkit并确保上述依赖已安装。需要注意的是原 README 的 Supported SM Architectures 一节为空因此本文不臆造任何架构列表实际可编译的架构由构建脚本中的CMAKE_CUDA_ARCHITECTURES决定见下文构建章节。涉及的 CUDA API 清单原 README 明确列出了本示例使用的 CUDA Runtime API共 9 个其作用与在源码中的位置如下CUDA API作用源码位置cudaMallocManaged分配统一内存data/result/vectorUnifiedMemoryStreams.cu#L82-L84cudaFree释放统一内存UnifiedMemoryStreams.cu#L92-L94cudaStreamCreate创建 CUDA Stream线程数 1 个UnifiedMemoryStreams.cu#L273cudaStreamDestroy销毁 CUDA StreamUnifiedMemoryStreams.cu#L330cudaStreamAttachMemAsync将托管内存异步附着到指定流host 或 device 路径UnifiedMemoryStreams.cu#L163-L165、#L179-L181cudaStreamSynchronize同步单个流确保异步 attach 完成UnifiedMemoryStreams.cu#L167cudaDeviceSynchronize同步整个设备保证内存可安全释放/访问UnifiedMemoryStreams.cu#L85、#L326cudaSetDevice在每个线程中显式绑定当前设备UnifiedMemoryStreams.cu#L290、#L320cudaGetDeviceProperties查询设备属性校验 Unified Memory 支持UnifiedMemoryStreams.cu#L241核心数据结构Task 与统一内存分配源码定义了一个模板化的Task结构体UnifiedMemoryStreams.cu#L61-L117每个任务包含size矩阵维度方阵决定任务在 host 还是 device 执行id任务唯一编号datasize×size矩阵、vectorsize向量、resultsize结果向量三者全部通过cudaMallocManaged分配为统一内存。值得注意的两个工程细节构造与分配分离构造函数中分配内存另有一个allocate()方法在外部以随机大小重新分配并填充数据便于任务列表批量初始化UnifiedMemoryStreams.cu#L97-L116析构前同步~Task()在cudaFree之前先调用cudaDeviceSynchronize()确保所有可能仍在设备上访问该内存的异步操作已经完成否则可能出现内存仍在使用的未定义行为UnifiedMemoryStreams.cu#L88-L95。任务规模在initialise_tasks()中随机生成size max(drand48() * 1000, 64)即范围落在64~1000之间UnifiedMemoryStreams.cu#L226-L234。执行策略按任务规模分流到 Host 或 Device示例的核心调度逻辑位于execute()函数判断标准简单直观任务规模 100 时在主机执行否则在设备执行UnifiedMemoryStreams.cu#L191-L222。主机路径cudaMemAttachHost小任务走主机路径// attach managed memory to a (dummy) stream to allow host access // while the device is running checkCudaErrors(cudaStreamAttachMemAsync(stream[0], t.data, 0, cudaMemAttachHost)); checkCudaErrors(cudaStreamAttachMemAsync(stream[0], t.vector, 0, cudaMemAttachHost)); checkCudaErrors(cudaStreamAttachMemAsync(stream[0], t.result, 0, cudaMemAttachHost)); checkCudaErrors(cudaStreamSynchronize(stream[0])); gemv(t.size, t.size, 1.0, t.data, t.vector, 0.0, t.result);这里的关键是cudaStreamAttachMemAsync的cudaMemAttachHost标志它将托管内存附着到一个哑流stream[0]上表示这段内存当前应被CPU 访问。注释点明了设计意图——当设备还有其他任务在运行时这样可以让主机安全地访问该内存实现 CPU 与 GPU 并发执行互不干扰。随后的cudaStreamSynchronize确保异步 attach 操作全部完成后再由 CPU 读取数据。主机侧的计算由模板函数gemv()完成它以朴素的三重循环按行主序实现矩阵向量乘UnifiedMemoryStreams.cu#L133-L143。设备路径cudaMemAttachSingle cuBLAS大任务走设备路径并把托管内存附着到该线程专属的流double one 1.0, zero 0.0; checkCudaErrors(cublasSetStream(handle[tid 1], stream[tid 1])); checkCudaErrors(cudaStreamAttachMemAsync(stream[tid 1], t.data, 0, cudaMemAttachSingle)); checkCudaErrors(cudaStreamAttachMemAsync(stream[tid 1], t.vector, 0, cudaMemAttachSingle)); checkCudaErrors(cudaStreamAttachMemAsync(stream[tid 1], t.result, 0, cudaMemAttachSingle)); checkCudaErrors(cublasDgemv( handle[tid 1], CUBLAS_OP_N, t.size, t.size, one, t.data, t.size, t.vector, 1, zero, t.result, 1));三个要点值得展开cudaMemAttachSingle语义与cudaMemAttachHost附着给主机相对cudaMemAttachSingle将内存附着到单个流使该段内存只对这一个流可见、可访问。多个线程各自用stream[tid 1]互不干扰从而支持并行流之间的内存访问隔离与并发cublasSetStream把每个线程对应的 cuBLAS handle 绑定到专属流使cublasDgemv的 kernel 提交进该流与内存附着目标流保持一致索引约定stream[0]是哑流专用于 host 路径的内存附着设备路径使用stream[tid 1]与handle[tid 1]即线程 tid 使用第 tid1 个流/handleUnifiedMemoryStreams.cu#L269-L275。线程模型OpenMP 与 pthread 双实现示例支持两种线程后端由宏USE_PTHREADS切换UnifiedMemoryStreams.cu#L38-L42默认 OpenMP 路径omp_set_num_threads(4)设定 4 个线程用#pragma omp parallel for schedule(dynamic)以动态调度方式遍历全部 40 个任务每个线程先cudaSetDevice(dev_id)再执行execute()UnifiedMemoryStreams.cu#L316-L324pthread 路径USE_PTHREADS定义时构造threadData传递 tid/streams/handles/task 列表按TaskList.size() / nthreads静态切分任务用pthread_createpthread_join并行执行UnifiedMemoryStreams.cu#L285-L315。omp_get_thread_num()返回线程编号作为tid用于索引专属流与 cuBLAS handle。运行时的设备校验逻辑main()在真正执行任务前做了两道关键校验UnifiedMemoryStreams.cu#L236-L259Unified Memory 支持检查通过cudaGetDeviceProperties读取device_prop.managedMemory若设备不支持统一内存则打印Unified Memory not supported on this device并以EXIT_WAIVED退出表示测试被放弃而非失败计算模式检查通过cudaDeviceGetAttribute(..., cudaDevAttrComputeMode, ...)判断计算模式若为cudaComputeModeProhibited禁止计算则同样以EXIT_WAIVED退出并提示需要 default 或 process exclusive 模式。设备选择本身复用Common目录下的工具函数findCudaDevice()Common/helper_cuda.h#L882-L912默认挑选算力最高Gflops 最大的设备也支持命令行-deviceid显式指定。构建与运行方式该示例通过 UnifiedMemoryStreams/CMakeLists.txt 构建并被 cpp/0_Introduction/CMakeLists.txt 以add_subdirectory纳入整体构建。构建脚本的几个关键点要求CMake 3.20及以上project(... LANGUAGES C CXX CUDA)通过find_package(CUDAToolkit REQUIRED)定位 Toolkit默认编译架构列表为75 80 86 87 89 90 100 110 120CMakeLists.txt#L11覆盖 Volta 到 Blackwell 时代的常见算力OpenMP 探测find_package(OpenMP COMPONENTS CXX)从 CMake 3.31 起裸find_package(OpenMP)可能额外探测 CUDA 语言的 OpenMP 而导致误判因此显式只请求 CXX 组件未找到 OpenMP 时仅打印提示并跳过该示例CMakeLists.txt#L28-L49链接依赖为CUDA::cublas与OpenMP::OpenMP_CXX并开启CUDA_SEPARABLE_COMPILATION可分离编译QNX 不支持检测到CMAKE_SYSTEM_NAME QNX时直接return()跳过CMakeLists.txt#L22-L26。在仓库根目录执行标准的 CMake 流程即可构建README.md#L35-L54mkdir build cd build cmake .. make -j$(nproc)构建产物默认安装到build/bin/${TARGET_ARCH}/${TARGET_OS}/${BUILD_TYPE}Linux x86_64 Release 对应build/bin/x64/linux/release。也可单独进入cpp/0_Introduction/UnifiedMemoryStreams目录对单个示例进行同样的 CMake 构建。运行示例可用-deviceid指定 GPU./UnifiedMemoryStreams程序会输出每个任务的执行位置例如Task [0], thread [0] executing on device (XXX) Task [1], thread [1] executing on host (XX) ... All Done!该示例也可通过仓库根目录的 run_tests.py 配合 test_args.json 作为快速健全性检查的一部分批量运行。小结单 GPU 上编排混合异构任务的参考范式UnifiedMemoryStreams 虽然是一个入门级示例却浓缩了单 GPU 异构编程的几个关键工程点值得在实际项目中复用统一内存免拷贝cudaMallocManaged分配的数据被 CPU 与 GPU 直接共享省去显式cudaMemcpy代价是需要用同步点cudaDeviceSynchronize/cudaStreamSynchronize管理访问顺序按规模分流以任务大小阈值本示例为 100在 host 与 device 之间分流小任务避免 kernel 启动开销大任务获得 GPU 加速内存附着策略cudaStreamAttachMemAsync配合cudaMemAttachHost允许主机在设备运行时访问与cudaMemAttachSingle附着到单条流实现流间隔离是精细控制托管内存归属的关键 API多线程 多流每个工作线程绑定独立 stream 与 cuBLAS handle通过 OpenMP或 pthread并行消费任务队列实现并发计算。理解这些模式后你可以将同样的任务消费者 流 统一内存框架扩展到更复杂的异构工作负载中。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
