CUDA Samples 入门：使用 cuda.core 与运行时编译在 Python 中实现 GPU 向量加法（vectorAdd）

CUDA Samples 入门：使用 cuda.core 与运行时编译在 Python 中实现 GPU 向量加法（vectorAdd）

【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples

导读

本篇技术指南以 CUDA Samples 仓库中的 Python vectorAdd 示例 为主体，完整讲解如何在 Python 中编写、编译并启动你的第一个 CUDA 内核：通过cuda.core高层 API 完成设备管理、内核源码的运行时编译（Runtime Compilation）与启动，借助 CuPy 管理 GPU 显存，并实现C = A + B的逐元素向量加法。阅读完本文，你将掌握 C++ 模板内核的编写、Device/Program/ProgramOptions/LaunchConfig/launch五类核心 API 的用法、网格（grid）与线程块（block）的配置方法，以及一套可复用的 GPU 结果验证流程。

示例概览：从 C++ 到 Python 的一站式 CUDA 编程

vectorAdd 是 CUDA 编程的经典入门样例。仓库中同时提供了 C++ 运行时 API 版本（使用cudaMalloc、cudaMemcpy等底层 API）和本示例的 Python 版本。Python 版的最大特点是不需要编写独立的.cu文件再离线编译，而是把内核源码作为字符串直接嵌入 Python 脚本，运行时再编译为 cubin 并加载执行。

本示例的核心数据流如下：

1. 初始化 CUDA 设备并创建流（stream）；
2. 将 C++ 模板内核源码经Program提交给 NVRTC 风格的运行时编译器；
3. 编译得到 cubin 模块，取出vectorAdd<float>内核句柄；
4. 用 CuPy 在 GPU 上分配向量A、B、C的显存并初始化；
5. 按grid/block尺寸构造LaunchConfig，通过launch在指定流上启动内核；
6. 同步后调用仓库工具函数verify_array_result与 CPU 计算结果比对，输出Test PASSED。

运行环境与依赖

硬件要求

• NVIDIA GPU，Compute Capability 7.0 及以上；
• 最小 GPU 显存：512 MB。

软件要求

依赖清单由示例目录下的 requirements.txt 维护，内容与仓库根级 python/requirements.txt 保持一致（后者额外包含numpy>=2.3.2，作为 Python 样例的公共底座）。

安装依赖

cd /path/to/cuda-samples/python/1_GettingStarted/vectorAdd pip install -r requirements.txt

核心库与 API 速览

cuda.core：Pythonic 的 CUDA 运行时封装

• Device—— 初始化并管理 CUDA 设备，提供set_current()、create_stream()、sync()等设备级操作；
• Program—— 由内核源码字符串创建编译程序对象；
• ProgramOptions—— 设置编译选项（C++ 标准、目标架构等）；
• LaunchConfig—— 配置内核启动参数（grid 与 block 维度）；
• launch—— 在指定流上执行内核。

按官方建议，这些稳定符号应从顶层包导入（即from cuda.core import ...），而不是从cuda.core.experimental导入，以保证 API 的稳定性。

CuPy：GPU 数组库

• cp.random.rand()—— 在 GPU 上生成随机数组；
• cp.empty()—— 分配未初始化的 GPU 数组；
• cp.allclose()—— 带容差地验证结果一致性。

仓库公共工具：cuda_samples_utils

• verify_array_result()—— 验证计算结果，源码见 python/Utilities/cuda_samples_utils.py。该函数会自动检测传入数组是 NumPy 还是 CuPy 类型，并选用对应的allclose（默认rtol=1e-5, atol=1e-8），从而避免不必要的跨设备数据拷贝；比对失败时还会打印Test FAILED - Max error: ...辅助排查。

内核源码：模板化的一维网格跨步写法

内核以字符串常量形式定义在 vectorAdd.py 中：

template<typename T> __global__ void vectorAdd(const T *A, const T *B, T *C, int numElements) { int i = blockDim.x * blockIdx.x + threadIdx.x; if (i < numElements) { C[i] = A[i] + B[i]; } }

这段内核体现了三类关键技巧：

• 1D Grid-Stride 思路的雏形：每个线程通过blockDim.x * blockIdx.x + threadIdx.x计算全局线性索引。虽然本示例每个线程只处理一个元素，但这一索引公式是后续网格跨步循环（Grid-Stride Loop）的基础，可让固定大小的网格处理任意规模的数组。
• 模板编程：template<typename T>使同一份内核源码可为float、double等不同类型实例化，编译期通过name_expressions指定实例化哪个特化版本。
• 边界检查：if (i < numElements)防止越界访问，确保当numElements不能被blockDim.x * gridDim.x整除时行为正确。

从源码到内核：设备、编译与启动全流程

以下代码取自 vectorAdd.py 的vector_add_cuda_core()函数主体。

第一步：初始化设备并创建流

device = Device(device_id) device.set_current() print(f"Device: {device.name}") print(f"Compute Capability: sm_{device.arch}") stream = device.create_stream()

device.arch会返回形如90的计算能力值，用于拼装编译目标架构。创建独立的流（stream）可以让内核启动、同步等操作与默认流解耦。

第二步：运行时编译内核

program_options = ProgramOptions(std="c++17", arch=f"sm_{device.arch}") program = Program(VECTOR_ADD_KERNEL, code_type="c++", options=program_options) module = program.compile("cubin", name_expressions=("vectorAdd<float>",)) kernel = module.get_kernel("vectorAdd<float>")

• ProgramOptions(std="c++17", arch="sm_XX")指定使用 C++17 标准，并针对当前设备架构（如sm_90）编译；
• program.compile("cubin", name_expressions=("vectorAdd<float>",))输出 cubin 格式，并通过name_expressions显式实例化模板特化版本vectorAdd<float>；
• module.get_kernel("vectorAdd<float>")取出可启动的内核句柄。

这一"字符串源码 → 运行时编译 → 句柄获取"的链路，正是cuda.core相对传统编译-链接-加载流程的核心优势：内核与 Python 代码同源、同文件维护，迭代时无需重启工具链。

第三步：用 CuPy 管理 GPU 显存

dtype = cp.float32 a = cp.random.rand(num_elements).astype(dtype) b = cp.random.rand(num_elements).astype(dtype) c = cp.empty(num_elements, dtype=dtype) # 内核启动前先同步设备 device.sync()

a、b在 GPU 上直接生成随机数，c作为输出容器。启动前调用device.sync()可确保数据初始化操作全部完成，避免与后续内核启动产生数据竞争。

第四步：配置并启动内核

threads_per_block = 256 blocks_per_grid = (num_elements + threads_per_block - 1) // threads_per_block config = LaunchConfig(grid=blocks_per_grid, block=threads_per_block) launch( stream, config, kernel, a.data.ptr, b.data.ptr, c.data.ptr, cp.int32(num_elements), ) stream.sync()

• 网格尺寸采用"向上取整"公式(num_elements + threads_per_block - 1) // threads_per_block，与内核内的边界检查配合，可安全处理任意元素个数；
• 默认num_elements = 50000时，threads_per_block = 256，则blocks_per_grid = 196（50000 / 256 向上取整）；
• launch以a.data.ptr传入 CuPy 数组的裸设备指针，cp.int32(num_elements)保证内核签名中的int numElements参数类型匹配；
• 启动后stream.sync()等待内核执行完成，随后才能安全读取结果。

第五步：验证结果

expected = a + b if not verify_array_result(c, expected): return False

利用 CuPy 的广播加法在 GPU 上直接算出期望结果，再调用 cuda_samples_utils.verify_array_result 在容差范围内比对，输出Test PASSED。注意verify_array_result要求传入的两个数组同为 NumPy 或同为 CuPy 类型，否则会抛出TypeError。

命令行运行方式

基本用法

cd samples/python/1_GettingStarted/vectorAdd python vectorAdd.py

自定义参数

# 自定义向量规模 python vectorAdd.py --elements 1000000 # 指定使用某块 GPU python vectorAdd.py --device 1 # 跳过验证（适合只做性能基准测试） python vectorAdd.py --no-verify

三个命令行参数的解析与校验位于 vectorAdd.py 的main()中：--elements默认 50000，--device默认 0，--no-verify为布尔开关。脚本会对--elements <= 0的非法输入报错退出，并在成功时打印Done、以退出码 0 结束。

预期输出

[Vector addition using CUDA Core API] Device: <Your GPU Name> Compute Capability: sm_<XX> Compiling kernel 'vectorAdd<float>'... Kernel compiled successfully [Vector addition of 50000 elements] CUDA kernel launch with 196 blocks of 256 threads Verifying result... Test PASSED Done

注意：设备名称与计算能力会随实际 GPU 不同而变化；196 blocks of 256 threads对应默认 50000 个元素的网格配置。

示例文件清单

工具模块的导入方式值得留意：脚本通过sys.path.insert(0, str(Path(__file__).parent.parent.parent / "Utilities"))将python/Utilities目录加入搜索路径，再from cuda_samples_utils import verify_array_result，这种轻量路径式导入避免了复杂的包结构（详见 python/Utilities/README.md）。

进阶方向

• 多类型实例化：借助模板与name_expressions，同一内核可继续实例化vectorAdd<double>等特化版本，验证模板编程在 GPU 内核中的通用性；
• 网格跨步循环：将本示例的"每线程单元素"模式改写为循环跨步模式，让固定网格处理远超其规模的数组，同时减少线程启动开销；
• 与 C++ 版本对照：对比 cpp/0_Introduction/vectorAdd 中cudaMalloc/cudaMemcpy/cudaGetLastError的显式错误检查流程，可以直观体会cuda.core封装后的简洁性；
• 性能测试：使用--no-verify跳过验证环节，结合 CUDA 工具链（如 Nsight Systems）对内核启动与执行进行剖析。

总结

Python 版 vectorAdd 示例完整展示了cuda.core驱动的"设备管理 → 运行时编译 → 内核启动 → 结果验证"闭环，是进入 CUDA Python 编程最直接、依赖最少的起点。它既保留了 CUDA 内核编写（模板、索引计算、边界检查）的全部精髓，又将繁琐的编译加载过程收敛为Program → ProgramOptions → LaunchConfig → launch几步调用，配合 CuPy 的显存管理与仓库内置的验证工具，可在几分钟内跑通你的第一个 GPU 程序。

【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples