CUDA 编程指南 第4章：应用编程接口

本文整理自 NVIDIA《CUDA Programming Guide Version 1.1》第4章，系统介绍 CUDA 的编程接口，包括 C 语言扩展、Runtime 组件、以及 Driver API 的使用方式。

4.1 C 编程语言扩展

CUDA 编程接口的目标是为熟悉 C 语言的开发者提供简单的路径来编写设备端程序。它包括：

4.2 语言扩展

C 语言的扩展包含四个部分：

1. 函数类型限定符 — 指定函数执行位置和调用来源
2. 变量类型限定符 — 指定变量在设备上的内存位置
3. 执行配置 — 指定如何从主机启动设备内核
4. 内置变量 — 指定网格和块的维度及索引

包含这些扩展的源文件必须使用 nvcc 编译。违反限制时 nvcc 会给出错误或警告。

4.2.1 函数类型限定符

__device__

• 在设备上执行
• 只能从设备中调用

__global__（内核函数）

• 在设备上执行
• 只能从主机中调用
• 调用是异步的（设备完成前即返回）
• 必须返回 void
• 参数通过共享内存传递，限制为 256 字节

__host__

• 在主机上执行
• 只能从主机中调用
• 可与 __device__ 结合使用，同时为主机和设备编译

限制

• __device__ 和 __global__ 函数不支持迭代
• 不能声明静态变量
• 不能有可变参数
• __device__ 函数不能取地址；__global__ 函数指针受支持
• __global__ 和 __host__ 不能一起使用

4.2.2 变量类型限定符

关键规则：

• __shared__ 变量：线程间完全顺序一致性，但仅 __syncthreads() 执行后，其他线程的写入才可见
• __shared__ 不能在声明中初始化
• __constant__ 不能从设备赋值，只能从主机通过 runtime 函数赋值
• 不能用于 struct/union 成员、形参
• 不能使用 extern（但共享内存可声明为外部以支持动态大小）
• 自动变量（无限定符）通常驻留在寄存器中；若过大则可能溢出到本地内存（local memory）

4.2.3 执行配置

任何 __global__ 函数调用必须指定执行配置，语法为：

kernel_function<<< Dg, Db, Ns, S >>>(args);

Dg 或 Db 超过设备最大限制，或 Ns 超过可用共享内存时，调用失败。

4.2.4 内置变量

不允许取内置变量的地址，也不允许为其赋值。

4.2.5 使用 NVCC 编译

nvcc 是 CUDA 编译器驱动程序，工作流：

1. 分离设备代码与主机代码
2. 设备代码 → 编译为二进制（cubin）对象
3. 主机代码 → 输出为 C 代码或目标代码

• 主机代码完全支持 C++
• 设备代码只完全支持 C++ 的 C 子集（类、继承等不支持）

__noinline__

__device__ 函数默认始终内联。使用 __noinline__ 作为编译器提示，尽量不内联。

#pragma unroll

控制循环展开：

#pragma unroll 5    // 展开 5 次
#pragma unroll 1    // 禁止展开
#pragma unroll      // 若循环计数为常量则完全展开

4.3 公共 Runtime 组件

公共 runtime 组件可供主机和设备函数共同使用。

4.3.1 内置向量类型

CUDA 提供了从基本类型派生的向量类型：

• 组件通过字段 x、y、z、w 访问
• 使用 make_<type_name> 构造函数创建，如 make_int2(x, y)

dim3 类型

基于 uint3，用于指定维度。未指定的组件默认初始化为 1。

4.3.2 数学函数

支持 C/C++ 标准库数学函数（参见附录 B-1），在设备和主机上均可使用。

4.3.3 时间函数

clock_t clock();

返回每个时钟周期递增的计数器值。在内核开始和结束时取样，求差可计量线程执行所用时钟周期数。

注意：计量的是线程完全执行所用周期数（含分时等待），非实际执行指令的周期数。

4.3.4 纹理类型

CUDA 支持 GPU 纹理硬件来访问纹理内存。从纹理内存读取数据具有性能优势（参见第 5.4 节）。

关键概念：

• 纹理内存通过纹理拾取（texture fetch）在内核中读取
• 纹理拾取的第一个参数是纹理参考（texture reference）
• 纹理参考必须先绑定到内存区域才能使用

纹理参考声明：

texture<Type, Dim, ReadMode> texRef;

纹理参考属性（运行时可变）：

• 规格化坐标：坐标范围 [0.0, 1.0) 而非 [0, N)
• 寻址模式：Clamp（固定到边界）或 Wrap（重复/包装）
• 线性筛选：纹理元素间的低精度插值（仅浮点返回值）

线性内存中的纹理限制：

• 维度只能为 1
• 不支持纹理筛选
• 仅非规格化整数坐标
• 超出范围的访问返回零

4.4 设备 Runtime 组件

设备 runtime 组件只能在设备函数中使用。

4.4.1 数学函数（快速版）

设备端提供快速但精度略低的数学函数，名称加前缀 __：

__sin(x), __cos(x), __exp(x), __log(x) ...

编译器选项 -use_fast_math 可强制全部使用快速版本。

4.4.2 同步函数

void __syncthreads();

同步块中所有线程。所有线程到达此点时，执行继续。

• 用于协调块内线程通信，消除读后写/写后读/写后写数据冒险
• 允许出现在条件代码中，但仅当条件在整个线程块中求值相同时才允许

4.4.3 类型转换函数

IEEE-754 取整模式标识：

• rn：取整到最近偶数
• rz：向零取整
• ru：向上取整（到 +∞）
• rd：向下取整（到 -∞）

int   __float2int_[rn|rz|ru|rd](float);
uint  __float2uint_[rn|rz|ru|rd](float);
float __int2float_[rn|rz|ru|rd](int);
float __uint2float_[rn|rz|ru|rd](uint);

4.4.4 类型重解释函数

保留位模式不变，在不同类型之间重解释：

float __int_as_float(int);    // 如 0xC0000000 → -2.0
int   __float_as_int(float);  // 如 1.0f → 0x3f800000

4.4.5 纹理函数

从设备内存取纹理 (tex1Dfetch 族)：

float tex1Dfetch(texture<...> texRef, int x);

不支持纹理筛选和寻址模式。

从 CUDA 数组取纹理 (tex1D / tex2D)：

float tex1D(texture<...> texRef, float x);
float tex2D(texture<...> texRef, float x, float y);

纹理参考的编译时属性和运行时属性共同决定坐标解释和返回值处理。

4.4.6 原子函数

仅适用于计算能力 1.1 的设备。

原子函数在全局内存中的 32 位字上执行读-改-写原子操作：

int atomicAdd(int* address, int val);
int atomicSub(int* address, int val);
int atomicExch(int* address, int val);
int atomicCAS(int* address, int compare, int val);
// ... 等（参见附录 C）

• 保证不受其他线程干扰
• 仅适用于 32 位有符号/无符号整数

4.5 主机 Runtime 组件

主机 runtime 组件只能由主机函数使用，提供以下功能：

• 设备管理
• 上下文管理
• 内存管理
• 代码模块管理
• 执行控制
• OpenGL / Direct3D 互操作性

它由两层 API 组成：

两个 API 互斥：应用程序应使用其中之一。

4.5.1 常用概念

设备管理

• 两种 API 都支持列举设备、查询属性、选择设备
• 一个主机线程只能在一个设备上执行设备代码
• 多设备需要多主机线程

内存模型

异步并发执行

以下函数是异步的：设备完成任务前即返回控制权：

• 内核启动
• 以 Async 为后缀的内存复制
• 设备↔设备内存复制
• 内存设置函数

流（Stream）：顺序执行的操作序列。不同流之间可以不按顺序或并发执行。

cudaStream_t stream;
cudaStreamCreate(&stream);
kernel<<<grid, block, 0, stream>>>(args);
cudaStreamDestroy(stream);

零流参数的操作：所有先前操作完成后才开始，完成后才能开始后续操作。

事件（Event）：用于监控设备进度和执行精准定时。

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
// ... 内核执行 ...
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float time;
cudaEventElapsedTime(&time, start, stop);

设置 CUDA_LAUNCH_BLOCKING=1 环境变量可全局禁用异步执行（仅用于调试）。

OpenGL / Direct3D 互操作性

• OpenGL 缓冲对象可映射到 CUDA 地址空间
• Direct3D 9.0 顶点缓冲可映射到 CUDA
• CUDA 上下文和 Direct3D 设备必须位于同一 GPU
• 仅支持 Direct3D 9.0 和顶点缓冲

4.5.2 Runtime API

初始化

无显式初始化函数；第一次调用 runtime 函数时自动初始化。

设备管理

cudaGetDeviceCount(&count);
cudaGetDeviceProperties(&prop, device);
cudaSetDevice(device);

内存管理

线性内存：

cudaMalloc(&devPtr, size);          // 分配
cudaMallocPitch(&devPtr, &pitch, width, height);  // 2D 分配（推荐）
cudaFree(devPtr);                   // 释放

CUDA 数组：

cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
cudaMallocArray(&cuArray, &desc, width, height);
cudaFreeArray(cuArray);

数据复制：

cudaMemcpy(dst, src, size, cudaMemcpyHostToDevice);
cudaMemcpy(dst, src, size, cudaMemcpyDeviceToHost);
cudaMemcpy2D(dst, dpitch, src, spitch, width, height, kind);

符号地址：

cudaGetSymbolAddress(&devPtr, symbol_name);
cudaGetSymbolSize(&size, symbol_name);

纹理参考管理

texture<float, 2, cudaReadModeElementType> texRef;
// 绑定到线性内存
cudaBindTexture(0, texRef, devPtr, size);
// 或绑定到 CUDA 数组
cudaBindTextureToArray(texRef, cuArray);
// 解除绑定
cudaUnbindTexture(texRef);

纹理参考属性（可直接修改）：

• normalized — 坐标是否规格化
• filterMode — cudaFilterModePoint 或 cudaFilterModeLinear
• addressMode[2] — cudaAddressModeClamp 或 cudaAddressModeWrap

设备仿真模式

编译选项 -deviceemu 使设备代码在主机上编译运行，可利用主机原生调试器。

• 定义宏 __DEVICE_EMULATION__
• 每个线程在主机上创建真正的线程
• 限制：主机最大线程数 = 每块最大线程数 + 1，每线程 256KB 堆栈

仿真模式用于查找算法错误，但时序、并发行为、浮点精度可能与实际设备不同。

4.5.3 驱动程序 API

Driver API 是基于句柄的命令式 API：

核心流程：

1. 初始化：cuInit(0)
2. 设备管理：cuDeviceGetCount(), cuDeviceGet()
3. 上下文管理：cuCtxCreate(), cuCtxAttach(), cuCtxDetach()
   • 上下文类似于 CPU 进程，具有独立的地址空间
   • 上下文与主机线程一对一对应
4. 模块管理：cuModuleLoad(), cuModuleGetFunction()
5. 内存管理：cuMemAlloc(), cuMemAllocPitch(), cuMemFree(), cuArrayCreate()
6. 执行控制：cuFuncSetBlockShape(), cuFuncSetSharedSize(), cuParam*(), cuLaunchGrid()
7. 流/事件管理：cuStreamCreate(), cuEventCreate(), cuEventElapsedTime()

Runtime API vs Driver API 对比

小结

第 4 章系统介绍了 CUDA 的编程接口，核心要点：

1. 函数限定符（__device__ / __global__ / __host__）划分代码执行位置
2. 变量限定符（__device__ / __constant__ / __shared__）控制内存位置和访问范围
3. 执行配置 <<<Dg, Db, Ns, S>>> 定义内核启动的并行度
4. 内置变量（gridDim / blockIdx / blockDim / threadIdx）提供线程索引
5. 公共组件提供向量类型、数学函数、时钟和纹理支持
6. 设备组件提供快速数学、同步、原子操作等设备端功能
7. 主机组件分 Runtime API（高层）和 Driver API（低层），提供设备管理、内存管理、流/事件、互操作等完整功能

理解这些 API 是编写高质量 CUDA 程序的基础。