CUDA Samples 之 simpleMultiGPU多 GPU 并行规约的上下文管理与异步流实战指南【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samplessimpleMultiGPU 是 CUDA Samples 仓库README.md0_Introduction目录下的入门级多 GPU 示例它演示了如何借助 CUDA 4.0 引入的每设备上下文per-device context管理接口通过cudaSetDevice与 CUDA Stream 在多个 GPU 上异步执行规约Reduction内核并最终与 CPU 结果交叉验证。读完本文你将掌握多 GPU 场景下最核心的四步套路——设备枚举、数据划分、按设备建立异步流水线、跨设备汇总结果可直接迁移到真实的多卡并行任务中。示例概览与核心概念该示例位于 cpp/0_Introduction/simpleMultiGPU/其设计目标正如源码注释所述“demonstrates how to use the CUDA API to use multiple GPUs, with an emphasis on simple illustration of the techniques (not on performance)”即重在讲清楚多 GPU 编程方法而非追求性能极致。示例覆盖的关键技术点异步数据传输Asynchronous Data Transfers使用cudaMemcpyAsync让拷贝与内核执行在流中异步排队CUDA 流与事件CUDA Streams and Events每个 GPU 各自持有独立 Stream实现设备间并行多线程/多设备Multithreading, Multi-GPU通过cudaSetDevice在单进程内切换、管理多个设备的上下文。在仓库目录体系上它与其他入门示例如 asyncAPI、simpleStreams、simpleMultiCopy同属0_Introduction但 simpleMultiGPU 是其中唯一以“多 GPU 协同”为主线的示例与其最接近的进阶参考是同目录下的 simpleP2P关注 GPU 间点对点通信和 simpleMultiGPU 对应的TGPUplan数据结构设计。环境要求与适用平台根据 simpleMultiGPU/README.md该示例的官方支持矩阵如下维度支持范围GPU 计算能力SM5.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、WindowsCPU 架构x86_64、armv7l前置条件为安装对应平台的 CUDA Toolkit。从仓库根 README.md 可知当前仓库支持 CUDA Toolkit 13.3构建工具要求CMake 3.20 及以上。需要特别强调的是源码文件头部的运行前提警告simpleMultiGPU.cu在 NVIDIA 控制面板中必须禁用 SLI否则应用只能看到一块 GPU反之即使不组建 SLI也可以将桌面扩展到两块 GPU 各自连接的显示器上。这是排查“明明有两张卡却只检测到一张”时最常遇到的坑。构建与运行simpleMultiGPU 已注册进0_Introduction的构建体系cpp/0_Introduction/CMakeLists.txt 中的add_subdirectory(simpleMultiGPU)其自身的 CMakeLists.txt 声明LANGUAGES C CXX CUDA并通过find_package(CUDAToolkit REQUIRED)定位 Toolkit。既可以跟随仓库根 README.md 的 Linux 流程整体构建mkdir build cd build cmake .. make -j$(nproc)也可以按根 README 的建议在任意单个示例目录内独立构建cd cpp/0_Introduction/simpleMultiGPU mkdir build cd build cmake .. make -j$(nproc) ./simpleMultiGPU构建配置中有几个值得注意的细节cpp/0_Introduction/simpleMultiGPU/CMakeLists.txt默认编译目标架构为75 80 86 87 89 90 100 110 120与仓库根 CMakeLists.txt 保持一致即覆盖 Volta 至 Blackwell 世代默认附加-lineinfo编译选项为性能分析工具保留行号信息若开启-DENABLE_CUDA_DEBUGTrue则改用-G以支持 cuda-gdb 调试代价是性能显著下降开启CUDA_SEPARABLE_COMPILATION可分离编译便于后续扩展为多编译单元项目。程序结构TGPUplan 数据结构多 GPU 程序的第一步是为每个 GPU 建立独立的工作计划plan。示例在 simpleMultiGPU.h 中定义了核心数据结构TGPUplantypedef struct { int dataN; // 分配给该 GPU 的数据元素个数 float *h_Data; // 主机端输入数据页锁定内存 float *h_Sum; // 该 GPU 的部分和结果主机端 float *d_Data, *d_Sum; // 设备端输入缓冲与规约结果缓冲 float *h_Sum_from_device; // 从设备拷回的部分和页锁定内存 cudaStream_t stream; // 该 GPU 专属的异步执行流 } TGPUplan;这个结构体清晰地映射了多 GPU 编程的“每设备资源独立”原则每个 GPU 拥有自己的输入/输出缓冲区、自己的 Stream甚至自己的主机端 pinned memory。注意h_Data与h_Sum_from_device使用cudaMallocHost分配这正是后续cudaMemcpyAsync能够异步执行的前提见下文。多 GPU 执行流程逐段拆解1. 设备枚举与数据划分程序在 simpleMultiGPU.cu 中首先调用cudaGetDeviceCount查询系统中可用 GPU 数量并声明上限常量MAX_GPU_COUNT 32超出则截断。输入数据总量为DATA_N 1048576 * 32即 33,554,432 个float128 MB。数据划分采用均分加余数补偿策略for (i 0; i GPU_N; i) { plan[i].dataN DATA_N / GPU_N; } // 余数逐个分摊保证所有数据都被覆盖 for (i 0; i DATA_N % GPU_N; i) { plan[i].dataN; }这种“整数除法 余数分配”的模式避免了多 GPU 任务里常见的负载不均或数据遗漏问题是通用的数据分区写法。2. 逐设备初始化上下文、流与内存核心初始化循环simpleMultiGPU.cu展示了 CUDA 多设备编程的经典三连for (i 0; i GPU_N; i) { checkCudaErrors(cudaSetDevice(i)); // ① 切换到设备 i checkCudaErrors(cudaStreamCreate(plan[i].stream)); // ② 为该设备创建专属流 checkCudaErrors(cudaMalloc((void **)plan[i].d_Data, ...)); // ③ 设备端分配 checkCudaErrors(cudaMalloc((void **)plan[i].d_Sum, ...)); checkCudaErrors(cudaMallocHost((void **)plan[i].h_Sum_from_device, ...)); checkCudaErrors(cudaMallocHost((void **)plan[i].h_Data, ...)); ... }cudaSetDevice是贯穿全程的关键接口CUDA 4.0 之后运行时 API 为每个设备维护独立上下文任何“当前设备”相关的操作内存分配、内核启动、流操作都作用于最近一次cudaSetDevice指定的设备。因此分配内存、启动内核、同步流之前都必须先切换到正确设备——这也是本示例反复调用cudaSetDevice(i)的原因。所有 CUDA 调用均包裹在checkCudaErrors宏中定义于 Common/helper_cuda.h它会在调用返回非零错误码时打印出错文件、行号、错误码及对应的可读错误字符串并终止程序是贯穿整个 CUDA Samples 仓库的标准错误处理范式。3. 异步执行Memcpy Kernel 全部入流执行阶段simpleMultiGPU.cu将每个 GPU 的 H2D 拷贝、内核启动、D2H 回拷全部放入该 GPU 的专属流中for (i 0; i GPU_N; i) { checkCudaErrors(cudaSetDevice(i)); checkCudaErrors(cudaMemcpyAsync( plan[i].d_Data, plan[i].h_Data, plan[i].dataN * sizeof(float), cudaMemcpyHostToDevice, plan[i].stream)); reduceKernelBLOCK_N, THREAD_N, 0, plan[i].stream( plan[i].d_Sum, plan[i].d_Data, plan[i].dataN); getLastCudaError(reduceKernel() execution failed.\n); checkCudaErrors(cudaMemcpyAsync( plan[i].h_Sum_from_device, plan[i].d_Sum, ACCUM_N * sizeof(float), cudaMemcpyDeviceToHost, plan[i].stream)); }这里的关键点有两个异步的前提是页锁定内存cudaMemcpyAsync只有配合cudaMallocHost分配的 pinned memoryh_Data、h_Sum_from_device才能真正异步化。如果使用普通malloc页内存运行时 API 会自动退化为同步拷贝多 GPU 并行将名存实亡。多 GPU 并行自动发生由于每块 GPU 的操作都在自己的流中而内核启动与异步拷贝均不阻塞 CPU主机端循环可以依次向各 GPU“下单”各设备随即并行执行自己的流水线。规约内核本身是一个简单的 grid-stride 累加simpleMultiGPU.cu__global__ static void reduceKernel(float *d_Result, float *d_Input, int N) { const int tid blockIdx.x * blockDim.x threadIdx.x; const int threadN gridDim.x * blockDim.x; float sum 0; for (int pos tid; pos N; pos threadN) sum d_Input[pos]; d_Result[tid] sum; }启动配置为BLOCK_N 32个 block、THREAD_N 256个线程因此每个 GPU 产生ACCUM_N 32 × 256 8192个部分和。源码注释也明确指出追求高性能规约应参考仓库中的 reduction 示例本示例的规约实现仅作多 GPU 流程演示之用。4. 结果回收与逐设备清理由于各 GPU 的 D2H 拷贝是异步的主机端在读取结果前必须逐设备同步simpleMultiGPU.cufor (i 0; i GPU_N; i) { checkCudaErrors(cudaSetDevice(i)); cudaStreamSynchronize(plan[i].stream); // 等待该流上所有操作完成 sum 0; for (j 0; j ACCUM_N; j) { sum plan[i].h_Sum_from_device[j]; // CPU 端汇总 8192 个部分和 } *(plan[i].h_Sum) (float)sum; checkCudaErrors(cudaFreeHost(plan[i].h_Sum_from_device)); checkCudaErrors(cudaFree(plan[i].d_Sum)); checkCudaErrors(cudaFree(plan[i].d_Data)); checkCudaErrors(cudaStreamDestroy(plan[i].stream)); }此处使用cudaStreamSynchronize而非cudaDeviceSynchronize语义是“等待该 GPU 上指定流的所有工作完成”——与cudaSetDevice配合后即为“等待该 GPU 完成”。随后在 CPU 上把每个 GPU 的 8192 个部分和累加为整机部分和并立即释放设备内存与销毁流保证多 GPU 程序退出时不留资源泄漏。5. CPU 参考实现与精度校验程序末尾simpleMultiGPU.cu在主机端用同样的数据逐元素累加得到 CPU 参考和然后计算相对误差作为正确性判据diff fabs(sumCPU - sumGPU) / fabs(sumCPU); printf( GPU sum: %f\n CPU sum: %f\n, sumGPU, sumCPU); printf( Relative difference: %E \n\n, diff); ... exit((diff 1e-5) ? EXIT_SUCCESS : EXIT_FAILURE);只有当相对误差小于1e-5时程序才以EXIT_SUCCESS退出否则以EXIT_FAILURE退出。这套“GPU 结果 vs CPU 黄金参考 相对误差阈值”的验证模式在 CUDA Samples 中大量复用例如 mergeSort、histogram 等可作为自研 CUDA 程序的测试范式。另外执行阶段用sdkCreateTimer/sdkStartTimer/sdkGetTimerValue来自 Common/helper_functions.h统计 GPU 处理耗时并打印便于直观对比单卡与多卡的吞吐差异。涉及 CUDA Runtime API 一览按 simpleMultiGPU/README.md 的列举本示例完整使用了以下 Runtime API其职责与出现位置汇总如下API作用源码位置cudaGetDeviceCount查询可用 GPU 数量simpleMultiGPU.cucudaSetDevice切换当前设备上下文循环初始化与执行段cudaStreamCreate/cudaStreamDestroy创建/销毁每设备专属流simpleMultiGPU.cucudaMalloc/cudaFree设备端显存分配/释放simpleMultiGPU.cucudaMallocHost/cudaFreeHost主机端页锁定内存分配/释放simpleMultiGPU.cucudaMemcpyAsync流内异步主机-设备双向拷贝simpleMultiGPU.cucudaStreamSynchronize等待指定流上的操作完成simpleMultiGPU.cu从示例到实战多 GPU 编程的通用要点从 simpleMultiGPU 可以提炼出可复用到任何多 GPU 项目的四条经验设备即上下文操作前必须切换所有设备相关 API分配、拷贝、启动、同步都受cudaSetDevice影响多 GPU 程序的常见 Bug 就是在错误的设备上执行了操作。进阶的替代方案是基于线程绑定每线程固定一个设备对应 README 提到的 Multithreading 概念或使用流序内存分配stream-ordered allocation实现更细粒度的跨设备调度相关进阶示例可参考 streamOrderedAllocation。Pinned memory 是异步化的前提cudaMemcpyAsync只有配合cudaMallocHost才能让数据传输与计算真正重叠这是多 GPU 并行的基础。每设备一条流流水线自然并行把“H2D → Kernel → D2H”放入各自 GPU 的流中主机端只需依次发起设备间即可并行执行多流之间的同步可用事件Event完成。GPU 只算部分和汇总交给 CPU规约这类“分而治之”任务天然适合多 GPU——每卡产生部分结果主机端做最终合并。若需要设备间直接交换数据而不经过主机则要进一步使用cudaMemcpyPeer/ P2P参考 simpleP2P。延伸阅读示例入口与说明cpp/0_Introduction/simpleMultiGPU/README.md完整源码simpleMultiGPU.cu、simpleMultiGPU.h构建配置cpp/0_Introduction/simpleMultiGPU/CMakeLists.txt及其父级注册项 cpp/0_Introduction/CMakeLists.txt通用工具头文件Common/helper_cuda.h错误检查与设备选择辅助函数仓库级构建与运行说明README.md【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CUDA Samples 之 simpleMultiGPU：多 GPU 并行规约的上下文管理与异步流实战指南
CUDA Samples 之 simpleMultiGPU多 GPU 并行规约的上下文管理与异步流实战指南【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samplessimpleMultiGPU 是 CUDA Samples 仓库README.md0_Introduction目录下的入门级多 GPU 示例它演示了如何借助 CUDA 4.0 引入的每设备上下文per-device context管理接口通过cudaSetDevice与 CUDA Stream 在多个 GPU 上异步执行规约Reduction内核并最终与 CPU 结果交叉验证。读完本文你将掌握多 GPU 场景下最核心的四步套路——设备枚举、数据划分、按设备建立异步流水线、跨设备汇总结果可直接迁移到真实的多卡并行任务中。示例概览与核心概念该示例位于 cpp/0_Introduction/simpleMultiGPU/其设计目标正如源码注释所述“demonstrates how to use the CUDA API to use multiple GPUs, with an emphasis on simple illustration of the techniques (not on performance)”即重在讲清楚多 GPU 编程方法而非追求性能极致。示例覆盖的关键技术点异步数据传输Asynchronous Data Transfers使用cudaMemcpyAsync让拷贝与内核执行在流中异步排队CUDA 流与事件CUDA Streams and Events每个 GPU 各自持有独立 Stream实现设备间并行多线程/多设备Multithreading, Multi-GPU通过cudaSetDevice在单进程内切换、管理多个设备的上下文。在仓库目录体系上它与其他入门示例如 asyncAPI、simpleStreams、simpleMultiCopy同属0_Introduction但 simpleMultiGPU 是其中唯一以“多 GPU 协同”为主线的示例与其最接近的进阶参考是同目录下的 simpleP2P关注 GPU 间点对点通信和 simpleMultiGPU 对应的TGPUplan数据结构设计。环境要求与适用平台根据 simpleMultiGPU/README.md该示例的官方支持矩阵如下维度支持范围GPU 计算能力SM5.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、WindowsCPU 架构x86_64、armv7l前置条件为安装对应平台的 CUDA Toolkit。从仓库根 README.md 可知当前仓库支持 CUDA Toolkit 13.3构建工具要求CMake 3.20 及以上。需要特别强调的是源码文件头部的运行前提警告simpleMultiGPU.cu在 NVIDIA 控制面板中必须禁用 SLI否则应用只能看到一块 GPU反之即使不组建 SLI也可以将桌面扩展到两块 GPU 各自连接的显示器上。这是排查“明明有两张卡却只检测到一张”时最常遇到的坑。构建与运行simpleMultiGPU 已注册进0_Introduction的构建体系cpp/0_Introduction/CMakeLists.txt 中的add_subdirectory(simpleMultiGPU)其自身的 CMakeLists.txt 声明LANGUAGES C CXX CUDA并通过find_package(CUDAToolkit REQUIRED)定位 Toolkit。既可以跟随仓库根 README.md 的 Linux 流程整体构建mkdir build cd build cmake .. make -j$(nproc)也可以按根 README 的建议在任意单个示例目录内独立构建cd cpp/0_Introduction/simpleMultiGPU mkdir build cd build cmake .. make -j$(nproc) ./simpleMultiGPU构建配置中有几个值得注意的细节cpp/0_Introduction/simpleMultiGPU/CMakeLists.txt默认编译目标架构为75 80 86 87 89 90 100 110 120与仓库根 CMakeLists.txt 保持一致即覆盖 Volta 至 Blackwell 世代默认附加-lineinfo编译选项为性能分析工具保留行号信息若开启-DENABLE_CUDA_DEBUGTrue则改用-G以支持 cuda-gdb 调试代价是性能显著下降开启CUDA_SEPARABLE_COMPILATION可分离编译便于后续扩展为多编译单元项目。程序结构TGPUplan 数据结构多 GPU 程序的第一步是为每个 GPU 建立独立的工作计划plan。示例在 simpleMultiGPU.h 中定义了核心数据结构TGPUplantypedef struct { int dataN; // 分配给该 GPU 的数据元素个数 float *h_Data; // 主机端输入数据页锁定内存 float *h_Sum; // 该 GPU 的部分和结果主机端 float *d_Data, *d_Sum; // 设备端输入缓冲与规约结果缓冲 float *h_Sum_from_device; // 从设备拷回的部分和页锁定内存 cudaStream_t stream; // 该 GPU 专属的异步执行流 } TGPUplan;这个结构体清晰地映射了多 GPU 编程的“每设备资源独立”原则每个 GPU 拥有自己的输入/输出缓冲区、自己的 Stream甚至自己的主机端 pinned memory。注意h_Data与h_Sum_from_device使用cudaMallocHost分配这正是后续cudaMemcpyAsync能够异步执行的前提见下文。多 GPU 执行流程逐段拆解1. 设备枚举与数据划分程序在 simpleMultiGPU.cu 中首先调用cudaGetDeviceCount查询系统中可用 GPU 数量并声明上限常量MAX_GPU_COUNT 32超出则截断。输入数据总量为DATA_N 1048576 * 32即 33,554,432 个float128 MB。数据划分采用均分加余数补偿策略for (i 0; i GPU_N; i) { plan[i].dataN DATA_N / GPU_N; } // 余数逐个分摊保证所有数据都被覆盖 for (i 0; i DATA_N % GPU_N; i) { plan[i].dataN; }这种“整数除法 余数分配”的模式避免了多 GPU 任务里常见的负载不均或数据遗漏问题是通用的数据分区写法。2. 逐设备初始化上下文、流与内存核心初始化循环simpleMultiGPU.cu展示了 CUDA 多设备编程的经典三连for (i 0; i GPU_N; i) { checkCudaErrors(cudaSetDevice(i)); // ① 切换到设备 i checkCudaErrors(cudaStreamCreate(plan[i].stream)); // ② 为该设备创建专属流 checkCudaErrors(cudaMalloc((void **)plan[i].d_Data, ...)); // ③ 设备端分配 checkCudaErrors(cudaMalloc((void **)plan[i].d_Sum, ...)); checkCudaErrors(cudaMallocHost((void **)plan[i].h_Sum_from_device, ...)); checkCudaErrors(cudaMallocHost((void **)plan[i].h_Data, ...)); ... }cudaSetDevice是贯穿全程的关键接口CUDA 4.0 之后运行时 API 为每个设备维护独立上下文任何“当前设备”相关的操作内存分配、内核启动、流操作都作用于最近一次cudaSetDevice指定的设备。因此分配内存、启动内核、同步流之前都必须先切换到正确设备——这也是本示例反复调用cudaSetDevice(i)的原因。所有 CUDA 调用均包裹在checkCudaErrors宏中定义于 Common/helper_cuda.h它会在调用返回非零错误码时打印出错文件、行号、错误码及对应的可读错误字符串并终止程序是贯穿整个 CUDA Samples 仓库的标准错误处理范式。3. 异步执行Memcpy Kernel 全部入流执行阶段simpleMultiGPU.cu将每个 GPU 的 H2D 拷贝、内核启动、D2H 回拷全部放入该 GPU 的专属流中for (i 0; i GPU_N; i) { checkCudaErrors(cudaSetDevice(i)); checkCudaErrors(cudaMemcpyAsync( plan[i].d_Data, plan[i].h_Data, plan[i].dataN * sizeof(float), cudaMemcpyHostToDevice, plan[i].stream)); reduceKernelBLOCK_N, THREAD_N, 0, plan[i].stream( plan[i].d_Sum, plan[i].d_Data, plan[i].dataN); getLastCudaError(reduceKernel() execution failed.\n); checkCudaErrors(cudaMemcpyAsync( plan[i].h_Sum_from_device, plan[i].d_Sum, ACCUM_N * sizeof(float), cudaMemcpyDeviceToHost, plan[i].stream)); }这里的关键点有两个异步的前提是页锁定内存cudaMemcpyAsync只有配合cudaMallocHost分配的 pinned memoryh_Data、h_Sum_from_device才能真正异步化。如果使用普通malloc页内存运行时 API 会自动退化为同步拷贝多 GPU 并行将名存实亡。多 GPU 并行自动发生由于每块 GPU 的操作都在自己的流中而内核启动与异步拷贝均不阻塞 CPU主机端循环可以依次向各 GPU“下单”各设备随即并行执行自己的流水线。规约内核本身是一个简单的 grid-stride 累加simpleMultiGPU.cu__global__ static void reduceKernel(float *d_Result, float *d_Input, int N) { const int tid blockIdx.x * blockDim.x threadIdx.x; const int threadN gridDim.x * blockDim.x; float sum 0; for (int pos tid; pos N; pos threadN) sum d_Input[pos]; d_Result[tid] sum; }启动配置为BLOCK_N 32个 block、THREAD_N 256个线程因此每个 GPU 产生ACCUM_N 32 × 256 8192个部分和。源码注释也明确指出追求高性能规约应参考仓库中的 reduction 示例本示例的规约实现仅作多 GPU 流程演示之用。4. 结果回收与逐设备清理由于各 GPU 的 D2H 拷贝是异步的主机端在读取结果前必须逐设备同步simpleMultiGPU.cufor (i 0; i GPU_N; i) { checkCudaErrors(cudaSetDevice(i)); cudaStreamSynchronize(plan[i].stream); // 等待该流上所有操作完成 sum 0; for (j 0; j ACCUM_N; j) { sum plan[i].h_Sum_from_device[j]; // CPU 端汇总 8192 个部分和 } *(plan[i].h_Sum) (float)sum; checkCudaErrors(cudaFreeHost(plan[i].h_Sum_from_device)); checkCudaErrors(cudaFree(plan[i].d_Sum)); checkCudaErrors(cudaFree(plan[i].d_Data)); checkCudaErrors(cudaStreamDestroy(plan[i].stream)); }此处使用cudaStreamSynchronize而非cudaDeviceSynchronize语义是“等待该 GPU 上指定流的所有工作完成”——与cudaSetDevice配合后即为“等待该 GPU 完成”。随后在 CPU 上把每个 GPU 的 8192 个部分和累加为整机部分和并立即释放设备内存与销毁流保证多 GPU 程序退出时不留资源泄漏。5. CPU 参考实现与精度校验程序末尾simpleMultiGPU.cu在主机端用同样的数据逐元素累加得到 CPU 参考和然后计算相对误差作为正确性判据diff fabs(sumCPU - sumGPU) / fabs(sumCPU); printf( GPU sum: %f\n CPU sum: %f\n, sumGPU, sumCPU); printf( Relative difference: %E \n\n, diff); ... exit((diff 1e-5) ? EXIT_SUCCESS : EXIT_FAILURE);只有当相对误差小于1e-5时程序才以EXIT_SUCCESS退出否则以EXIT_FAILURE退出。这套“GPU 结果 vs CPU 黄金参考 相对误差阈值”的验证模式在 CUDA Samples 中大量复用例如 mergeSort、histogram 等可作为自研 CUDA 程序的测试范式。另外执行阶段用sdkCreateTimer/sdkStartTimer/sdkGetTimerValue来自 Common/helper_functions.h统计 GPU 处理耗时并打印便于直观对比单卡与多卡的吞吐差异。涉及 CUDA Runtime API 一览按 simpleMultiGPU/README.md 的列举本示例完整使用了以下 Runtime API其职责与出现位置汇总如下API作用源码位置cudaGetDeviceCount查询可用 GPU 数量simpleMultiGPU.cucudaSetDevice切换当前设备上下文循环初始化与执行段cudaStreamCreate/cudaStreamDestroy创建/销毁每设备专属流simpleMultiGPU.cucudaMalloc/cudaFree设备端显存分配/释放simpleMultiGPU.cucudaMallocHost/cudaFreeHost主机端页锁定内存分配/释放simpleMultiGPU.cucudaMemcpyAsync流内异步主机-设备双向拷贝simpleMultiGPU.cucudaStreamSynchronize等待指定流上的操作完成simpleMultiGPU.cu从示例到实战多 GPU 编程的通用要点从 simpleMultiGPU 可以提炼出可复用到任何多 GPU 项目的四条经验设备即上下文操作前必须切换所有设备相关 API分配、拷贝、启动、同步都受cudaSetDevice影响多 GPU 程序的常见 Bug 就是在错误的设备上执行了操作。进阶的替代方案是基于线程绑定每线程固定一个设备对应 README 提到的 Multithreading 概念或使用流序内存分配stream-ordered allocation实现更细粒度的跨设备调度相关进阶示例可参考 streamOrderedAllocation。Pinned memory 是异步化的前提cudaMemcpyAsync只有配合cudaMallocHost才能让数据传输与计算真正重叠这是多 GPU 并行的基础。每设备一条流流水线自然并行把“H2D → Kernel → D2H”放入各自 GPU 的流中主机端只需依次发起设备间即可并行执行多流之间的同步可用事件Event完成。GPU 只算部分和汇总交给 CPU规约这类“分而治之”任务天然适合多 GPU——每卡产生部分结果主机端做最终合并。若需要设备间直接交换数据而不经过主机则要进一步使用cudaMemcpyPeer/ P2P参考 simpleP2P。延伸阅读示例入口与说明cpp/0_Introduction/simpleMultiGPU/README.md完整源码simpleMultiGPU.cu、simpleMultiGPU.h构建配置cpp/0_Introduction/simpleMultiGPU/CMakeLists.txt及其父级注册项 cpp/0_Introduction/CMakeLists.txt通用工具头文件Common/helper_cuda.h错误检查与设备选择辅助函数仓库级构建与运行说明README.md【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考