threadMigration 深度解析用 CUDA Driver API 实现多线程多 GPU 的 CUDA Context 创建、迁移与管理【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples本篇技术指南围绕 CUDA Samples 仓库中的 threadMigration 示例展开讲解如何使用 CUDA Driver API 的 Context Management 接口在多个 CPU 线程上分别创建、挂载并迁移 CUDA Context实现多 GPU、多线程并行提交 Kernel。读完本文你将掌握cuCtxCreate/cuCtxPushCurrent/cuCtxPopCurrent/cuCtxDestroy的完整使用链路、Driver API 下通过 fatbin 模块加载 Kernel 的流程以及 CUDA 4.0 引入的两种 Kernel 参数传递与启动方式并能在自己的多线程 CUDA 程序中安全地管理 Context 的迁移。示例概述一个线程一套 ContextthreadMigration 是一个演示型程序其核心目标见 README是展示CUDA Context Management API的用法演示CUDA 4.0 参数传递与 Kernel 启动 API即cuLaunchKernel风格验证CUDA Context 可以被独立创建并分别挂载attach到不同线程上。程序的运行模型非常清晰主线程负责初始化、枚举设备、为每块 GPU 创建独立的 CUDA Context随后为每个 Context 派生NumThreads个工作线程每个工作线程在自己的栈上调用cuCtxPushCurrent把对应 Context 设为当前然后独立完成分配显存 → 加载模块 → 启动 Kernel → 拷回结果 → 校验数据的全过程。从源码结构看threadMigration.cpp线程总数等于设备数 × 每设备线程数默认每设备 2 个线程。整个示例完全基于CUDA Driver API关键概念见 README不依赖运行时 API 的隐式上下文管理因此是理解Context 与线程关系的绝佳教材。CUDA Context 管理模型为什么需要迁移在 CUDA 中Context 是资源所有权与状态的核心载体它持有设备内存分配、加载的模块Module、流Stream以及 Kernel 参数等状态。Driver API 的 Context 管理遵循如下规则源码注释中亦有明确说明见 threadMigration.cpp一个 CUDA Context 与一个CPU 进程相关联可视为浮动的floating对象一个宿主线程在同一时刻只能有一个当前 Contextcurrent contextContext 可以在线程之间迁移创建它的线程可以弹出pop另一个线程再推入push使其成为自己的当前 Context。threadMigration 将这一模型落到了实处对应的关键 API 调用点如下阶段API作用源码位置初始化cuInit初始化 Driver API 运行环境threadMigration.cpp设备枚举cuDeviceGetCount/cuDeviceGet/cuDeviceGetName获取设备数量、句柄与名称threadMigration.cpp设备属性cuDeviceGetAttribute查询算力、共享内存、常量内存、寄存器、时钟等threadMigration.cpp创建 ContextcuCtxCreate在指定设备上创建一个新 ContextthreadMigration.cpp模块加载cuModuleLoadData/cuModuleGetFunction从 fatbin 二进制加载模块并获取 Kernel 句柄threadMigration.cpp弹出/推入cuCtxPopCurrent/cuCtxPushCurrent实现 Context 在创建线程与工作线程之间的迁移threadMigration.cpp 与 L195显存操作cuMemAlloc/cuMemcpyDtoH/cuMemFree分配、回拷、释放设备内存threadMigration.cppKernel 启动cuLaunchKernel以 CUDA 4.0 风格启动 KernelthreadMigration.cpp清理cuModuleUnload/cuCtxDestroy卸载模块、销毁 ContextthreadMigration.cpp说明完整 API 清单可参见 README 的 CUDA APIs involved 一节。关键数据结构每个线程一份上下文档案为了让多个线程共享同一套初始化结果示例定义了一个名为CUDAContext的结构体threadMigration.cpp把一个 Context 关联的全部资源打包在一起typedef struct _CUDAContext_st { CUcontext hcuContext; // CUDA Context 句柄 CUmodule hcuModule; // 从 fatbin 加载的模块 CUfunction hcuFunction; // 模块中的 Kernel 函数句柄 CUdeviceptr dptr; // 设备端内存指针 int deviceID; // 关联的设备编号 int threadNum; // 使用该 Context 的线程编号 } CUDAContext;全局数组g_ThreadParams[MAXTHREADS]MAXTHREADS定义为 256见 threadMigration.cpp保存每个线程的参数。注意这里每个设备只创建一份 Context但会有多个线程共享它——这是本示例与每线程独立 Context方案的关键区别也正因如此示例才需要在创建后先cuCtxPopCurrent把 Context 从主线程释放再让多个工作线程轮流cuCtxPushCurrent使用它。初始化链路从 cuInit 到模块加载1. 初始化与设备枚举runTest首先调用cuInit(0)初始化驱动随后cuDeviceGetCount获取设备数若为 0 则直接返回失败threadMigration.cpp。对每块设备依次执行cuDeviceGet取得设备句柄cuDeviceGetName打印设备名称一组cuDeviceGetAttribute查询并打印计算能力major/minor、每块共享内存、总常量内存、每块寄存器数与时钟频率threadMigration.cpp。2. 创建 Context新版 cuCtxCreate 签名示例在InitCUDAContext中调用threadMigration.cppCUctxCreateParams ctxCreateParams {}; CUresult status cuCtxCreate(hcuContext, ctxCreateParams, 0, hcuDevice);注意这里的签名采用了带CUctxCreateParams的新版 Driver API 形式ctxCreateParams以零初始化、标志位传 0表示创建默认配置的 Context。创建成功后该 Context 会立即成为当前线程的当前 Context。3. 从 fatbin 加载模块由于本示例完全走 Driver APIKernel 必须以可加载的二进制形式存在。构建系统通过 CMake 自定义命令把 threadMigration_kernel.cu 编译为threadMigration_kernel64.fatbin见 CMakeLists.txt。运行时流程findFatbinPath定义在 helper_cuda_drvapi.h根据可执行文件路径定位 fatbin 文件并以二进制流读入内存cuModuleLoadData(hcuModule, fatbinData)直接从内存中的二进制创建模块threadMigration.cppcuModuleGetFunction(hcuFunction, hcuModule, kernelFunction)按名称取出 Kernel 句柄threadMigration.cpp。4. 关键一步把 Context 从主线程弹出模块加载完成后示例立即调用cuCtxPopCurrent(NULL)threadMigration.cpp。源码注释写得很清楚Here we must release the CUDA context from the thread context——把 Context 从创建它的线程上解绑使其成为浮动状态等待工作线程接管。这正体现了Context 在线程间的迁移。线程内执行链路Push → 分配 → 启动 → 回拷 → Pop每个工作线程执行ThreadProcthreadMigration.cpp流程与初始化阶段严格对称cuCtxPushCurrent(pParams-hcuContext)把该线程对应的 Context 设为当前完成迁移cuMemAlloc分配NUM_INTS * sizeof(int)32 个 int的设备内存cuLaunchKernel以单 block1×1×1、每 block 32 线程的配置启动kernelFunctioncuMemcpyDtoH把结果拷回主机端pInt缓冲区校验期望值应为32 - ii为元素下标不匹配即报错计数cuMemFree释放显存cuCtxPopCurrent(NULL)把 Context 再次弹出结束本次迁移。工作线程通过宏在错误路径上打印Error并提前返回THREAD_QUIT见 threadMigration.cpp正常路径则在退出前打印设备号、Context 指针与线程号便于观察多线程并行执行的交错输出。Kernel 实现Kernel 本身极其简单threadMigration_kernel.cuextern C __global__ void kernelFunction(int *input) { input[threadIdx.x] 32 - threadIdx.x; }每个线程把自己的threadIdx.x写入input[threadIdx.x] 32 - threadIdx.x因此 32 个线程恰好填满 32 个 int主机端校验逻辑与之一一对应——它存在的意义不是计算而是验证每个工作线程都能在自己的 Context 上正确启动 Kernel 并取回结果。两种 Kernel 参数传递方式CUDA 4.0 API示例在同一段代码中演示了cuLaunchKernel的两种参数传入方式默认启用第一种见 threadMigration.cpp方式一简单参数数组默认启用void *args[5] {pParams-dptr}; status cuLaunchKernel(pParams-hcuFunction, 1, 1, 1, 32, 1, 1, 0, NULL, args, NULL);args数组按顺序存放各 Kernel 参数的指针启动配置为grid 维度 1×1×1、block 维度 32×1×1、共享内存 0 字节、默认流NULL。方式二参数缓冲区高级方法代码中以else分支保留int offset 0; char argBuffer[256]; *((CUdeviceptr *)argBuffer[offset]) pParams-dptr; offset sizeof(CUdeviceptr); void *kernel_launch_config[5] { CU_LAUNCH_PARAM_BUFFER_POINTER, argBuffer, CU_LAUNCH_PARAM_BUFFER_SIZE, offset, CU_LAUNCH_PARAM_END}; status cuLaunchKernel(pParams-hcuFunction, 1, 1, 1, 32, 1, 1, 0, 0, NULL, (void **)kernel_launch_config);该方式把所有参数按 ABI 布局连续打包进argBuffer再通过CU_LAUNCH_PARAM_BUFFER_POINTER/CU_LAUNCH_PARAM_BUFFER_SIZE/CU_LAUNCH_PARAM_END三个标记描述缓冲区的指针、大小与结束位置。两种方式结果等价方式二更贴近参数即二进制字节流的底层语义适合参数数量多、需动态拼装的场景。并发与同步线程创建、临界区与结果汇总线程创建与等待主线程对每块设备循环派生NumThreads个线程threadMigration.cppLinux 下使用POSIX 线程pthread_create 后续pthread_joinWindows 下使用Win32 线程CreateThreadWaitForMultipleObjects通过#if defined(WIN32) || defined(_WIN32) || ...条件编译切换见 threadMigration.cpp。ThreadLaunchCount是全局计数器用于统计成功完成校验的线程数多个线程同时更新它因此更新操作被EnterCriticalSection/LeaveCriticalSectionWindows或pthread_mutex_lock/pthread_mutex_unlockLinux保护的临界区包裹threadMigration.cpp避免数据竞争。最终校验所有线程join完成后FinalErrorCheckthreadMigration.cpp断言ThreadLaunchCount NumThreads * deviceCount即每个设备上的每个线程都必须成功完成一次 Kernel 启动与结果校验随后对所有 Context 执行cuCtxDestroy并释放设备资源。满足条件时main以EXIT_SUCCESS退出否则以EXIT_FAILURE退出threadMigration.cpp因此该示例可被自动化测试直接用作通过/失败的判定程序。命令行参数与运行方式支持的参数在runTest中解析threadMigration.cpp参数说明默认值 / 约束-qatest或-noprompt自动化测试模式跳过交互提示自动退出无-nthreads/-numthreadsthreads每块 GPU 上派生的工作线程数默认 2取值范围 115越界时打印用法提示并返回失败命令行解析依赖 helper_cuda_drvapi.h 中提供的checkCmdLineFlag/getCmdLineArgumentInt工具函数。构建与运行按仓库根 README 的标准 CMake 流程构建mkdir build cd build cmake .. make -j$(nproc)构建产物含threadMigration_kernel64.fatbin位于对应平台的build/bin目录下在build目录内找到 sample 目录后直接运行可执行文件即可例如./bin/x64/linux/release/threadMigration # 默认每设备 2 线程 ./bin/x64/linux/release/threadMigration -n4 # 每设备 4 线程 ./bin/x64/linux/release/threadMigration -qatest # 自动化测试模式构建细节值得注意Kernel 由 CMake 自定义命令以-fatbin选项直接编译为 fatbin 文件见 CMakeLists.txt目标threadMigration只链接CUDA::cuda_driverLinux 上额外链接pthread见 CMakeLists.txt——这是纯 Driver API 自编译 fatbin示例的典型构建模式。若需用 cuda-gdb 调试可在顶层 CMake 配置时传入-DENABLE_CUDA_DEBUGTrue。平台与架构支持范围按 README 的官方声明支持的 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覆盖 Maxwell 到 Hopper/Blackwell 主流算力当前 CMakeLists.txt 默认面向75 80 86 87 89 90 100 110 120生成 fatbin支持的操作系统Linux、Windows支持的 CPU 架构x86_64、armv7l前置条件安装对应平台的 CUDA ToolkitREADME 的 Prerequisites 一节说明见 README。小结threadMigration 虽然只有几十行有效逻辑却浓缩了 Driver API 多线程编程的完整要点Context 生命周期cuCtxCreate创建 → 创建线程cuCtxPopCurrent释放 → 工作线程cuCtxPushCurrent接管 → 用完cuCtxPopCurrent→ 最后cuCtxDestroy销毁形成可复用的迁移闭环模块化加载Kernel 以 fatbin 形式随构建生成运行时经cuModuleLoadDatacuModuleGetFunction解析两代参数传递既演示了简单args数组也保留了CU_LAUNCH_PARAM_BUFFER_*的高级打包方式线程安全跨线程共享的计数器以临界区/互斥锁保护最终以线程数 × 设备数的精确断言完成自检。在多 GPU 或多线程的 CUDA 应用中这套一 Context 多线程迁移模式尤其适合线程池场景——线程不必各自持有 Context而是按需 Push/Pop 复用浮动 Context。若需进一步研究可对照参考 CUDA 编程指南中关于 Context Management 的章节并结合本仓库同类的 Driver API 示例如 vectorAddDrv、matrixMulDrv交叉阅读。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考