CUDA 编程模型

CUDA 编程模型是理解 GPU 并行计算的基础。本章将详细介绍 CUDA 的异构计算模型、核函数的概念以及基本的编程流程。

异构计算模型

CUDA 采用异构计算模型，程序在 CPU（Host，主机）和 GPU（Device，设备）上协同执行。这种设计充分利用了 CPU 擅长逻辑控制和串行计算、GPU 擅长大规模并行计算的特点。

Host 与 Device

概念	说明	角色
Host	CPU 及其内存	负责逻辑控制、数据准备、结果处理
Device	GPU 及其显存	负责大规模并行计算
Kernel	核函数	在 Device 上并行执行的函数

执行流程

一个典型的 CUDA 程序执行流程如下：

代码示例

下面是一个完整的 CUDA 程序示例，展示向量加法的实现：

#include <cuda_runtime.h>
#include <stdio.h>

#define N 1024

__global__ void vectorAdd(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    float *h_a, *h_b, *h_c;
    float *d_a, *d_b, *d_c;
    size_t size = N * sizeof(float);

    h_a = (float*)malloc(size);
    h_b = (float*)malloc(size);
    h_c = (float*)malloc(size);

    for (int i = 0; i < N; i++) {
        h_a[i] = 1.0f;
        h_b[i] = 2.0f;
    }

    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, N);

    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

核函数（Kernel）

核函数是在 GPU 上并行执行的函数，是 CUDA 编程的核心概念。

核函数定义

使用 __global__ 修饰符定义核函数：

__global__ void kernelFunction(parameters) {
}

核函数的特点：

返回类型必须是 void
在 Device 上执行，由 Host 调用
每个线程独立执行一份副本
不能是类的成员函数（C++11 前）

函数修饰符

CUDA 提供了多种函数修饰符：

修饰符	执行位置	调用位置	说明
`__global__`	Device	Host 或 Device	核函数，返回 void
`__device__`	Device	Device	设备函数，只能被设备代码调用
`__host__`	Host	Host	主机函数（默认）
`__host__ __device__`	Host/Device	Host/Device	同时在 Host 和 Device 上可用

__device__ float deviceFunction(float x) {
    return x * 2.0f;
}

__global__ void kernel(float *data) {
    int idx = threadIdx.x;
    data[idx] = deviceFunction(data[idx]);
}

__host__ __device__ int commonFunction(int x) {
    return x + 1;
}

核函数调用

核函数使用特殊的执行配置语法调用：

kernel<<<gridSize, blockSize>>>(arguments);

其中：

gridSize：Grid 中的 Block 数量
blockSize：每个 Block 中的 Thread 数量

完整形式：

kernel<<<gridDim, blockDim, sharedMemSize, stream>>>(args);

参数	类型	说明
gridDim	dim3	Grid 维度（Block 数量）
blockDim	dim3	Block 维度（Thread 数量）
sharedMemSize	size_t	动态共享内存大小（字节）
stream	cudaStream_t	执行流

dim3 类型

dim3 是 CUDA 定义的三维向量类型，用于指定 Grid 和 Block 的维度：

dim3 grid(2, 2, 1);
dim3 block(16, 16, 1);

kernel<<<grid, block>>>(args);

一维示例：

int n = 1024;
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;

kernel<<<gridSize, blockSize>>>(data, n);

二维示例（适合图像处理）：

dim3 blockSize(16, 16);
dim3 gridSize((width + 15) / 16, (height + 15) / 16);

imageKernel<<<gridSize, blockSize>>>(image, width, height);

内置变量

CUDA 提供了一组内置变量，用于在核函数中获取线程信息：

线程索引变量

变量	类型	说明
`threadIdx`	uint3	线程在 Block 中的索引
`blockIdx`	uint3	Block 在 Grid 中的索引
`blockDim`	dim3	Block 的维度（线程数量）
`gridDim`	dim3	Grid 的维度（Block 数量）

计算全局索引

一维 Grid 和 Block：

int idx = blockIdx.x * blockDim.x + threadIdx.x;

二维 Grid 和 Block：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;

三维 Grid 和 Block：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
int idx = z * width * height + y * width + x;

边界检查

由于 Grid 大小可能不是 Block 大小的整数倍，需要进行边界检查：

__global__ void kernel(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        data[idx] = data[idx] * 2.0f;
    }
}

内存管理函数

CUDA 提供了一组内存管理函数，用于在 Host 和 Device 之间分配、拷贝和释放内存。

显存分配

cudaError_t cudaMalloc(void **devPtr, size_t size);

示例：

float *d_data;
cudaMalloc(&d_data, N * sizeof(float));

显存释放

cudaError_t cudaFree(void *devPtr);

示例：

cudaFree(d_data);

内存拷贝

cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, cudaMemcpyKind kind);

拷贝方向：

枚举值	说明
`cudaMemcpyHostToDevice`	Host → Device
`cudaMemcpyDeviceToHost`	Device → Host
`cudaMemcpyDeviceToDevice`	Device → Device
`cudaMemcpyHostToHost`	Host → Host

示例：

cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
cudaMemcpy(h_result, d_data, size, cudaMemcpyDeviceToHost);

错误处理

所有 CUDA 运行时 API 函数都返回 cudaError_t 类型：

cudaError_t err = cudaMalloc(&d_data, size);
if (err != cudaSuccess) {
    printf("CUDA error: %s\n", cudaGetErrorString(err));
    exit(EXIT_FAILURE);
}

推荐使用宏简化错误处理：

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            fprintf(stderr, "CUDA error at %s:%d: %s\n", \
                    __FILE__, __LINE__, cudaGetErrorString(err)); \
            exit(EXIT_FAILURE); \
        } \
    } while(0)

CUDA_CHECK(cudaMalloc(&d_data, size));

核函数错误检查

核函数不直接返回错误，需要使用 cudaGetLastError() 检查：

kernel<<<grid, block>>>(args);

cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("Kernel launch error: %s\n", cudaGetErrorString(err));
}

cudaDeviceSynchronize();
err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("Kernel execution error: %s\n", cudaGetErrorString(err));
}

完整示例：向量加法

下面是一个完整的向量加法示例，包含错误处理：

#include <cuda_runtime.h>
#include <stdio.h>

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            fprintf(stderr, "CUDA error at %s:%d: %s\n", \
                    __FILE__, __LINE__, cudaGetErrorString(err)); \
            exit(EXIT_FAILURE); \
        } \
    } while(0)

__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    const int N = 1 << 20;
    const size_t size = N * sizeof(float);
    
    float *h_a = (float*)malloc(size);
    float *h_b = (float*)malloc(size);
    float *h_c = (float*)malloc(size);
    
    for (int i = 0; i < N; i++) {
        h_a[i] = 1.0f;
        h_b[i] = 2.0f;
    }
    
    float *d_a, *d_b, *d_c;
    CUDA_CHECK(cudaMalloc(&d_a, size));
    CUDA_CHECK(cudaMalloc(&d_b, size));
    CUDA_CHECK(cudaMalloc(&d_c, size));
    
    CUDA_CHECK(cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice));
    
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    
    vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, N);
    
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());
    
    CUDA_CHECK(cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost));
    
    bool success = true;
    for (int i = 0; i < N; i++) {
        if (fabs(h_c[i] - 3.0f) > 1e-5) {
            success = false;
            break;
        }
    }
    printf("Result: %s\n", success ? "PASS" : "FAIL");
    
    CUDA_CHECK(cudaFree(d_a));
    CUDA_CHECK(cudaFree(d_b));
    CUDA_CHECK(cudaFree(d_c));
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

小结

本章介绍了 CUDA 编程模型的核心概念：

异构计算模型：Host 与 Device 协同执行
核函数：使用 __global__ 定义，在 GPU 上并行执行
执行配置：通过 <<<grid, block>>> 指定线程组织
内置变量：threadIdx、blockIdx、blockDim、gridDim
内存管理：cudaMalloc、cudaMemcpy、cudaFree

下一章将详细介绍线程层次结构，深入理解 Grid、Block、Thread 的组织方式。

异构计算模型​

Host 与 Device​

执行流程​

代码示例​

核函数（Kernel）​

核函数定义​

函数修饰符​

核函数调用​

dim3 类型​

内置变量​

线程索引变量​

计算全局索引​

边界检查​

内存管理函数​

显存分配​

显存释放​

内存拷贝​

错误处理​

核函数错误检查​

完整示例：向量加法​

小结​