3、CUDA编程模型：主机与设备、内核函数（kernel）定义与调用、<<<>>>语法、线程索引与块索引的使用

CUDA（Compute Unified Device Architecture）编程模型的核心思想是将一个计算任务分解为两个执行域：主机（Host）与设备（Device）。主机通常指CPU及其内存系统，负责逻辑控制、数据准备和任务调度；设备则指GPU及其显存，负责大规模并行计算。理解这一模型是掌握CUDA编程的基石，而内核函数（Kernel）的编写与调用方式、线程索引与块索引的映射关系，则是决定并行程序性能的关键。本章将从编程模型的基本概念出发，深入剖析内核函数的定义与调用语法，并通过实际案例演示如何正确使用线程和块索引来组织并行计算。

3.1 主机与设备：异构计算的分工协作

核心原理：CPU与GPU的职责边界

在CUDA架构中，CPU（主机）与GPU（设备）通过PCIe总线或NVLink互联，形成异构计算系统。主机负责执行串行逻辑、管理内存分配、发起数据传输以及启动内核；设备则专注于执行高度并行的计算任务。这种分工源于两者硬件设计的根本差异：CPU拥有少量但强大的核心，适合处理复杂分支和依赖关系；GPU拥有数千个简单核心，适合处理数据并行任务。

主机与设备各自拥有独立的内存空间：主机内存（DRAM）和设备显存（VRAM）。CUDA提供了显式的内存管理API，开发者必须手动完成数据在两者之间的搬运。这一设计虽然增加了编程复杂度，但也为性能优化提供了精细控制的可能性。

操作流程：典型CUDA程序执行步骤

1. 初始化：检查GPU设备是否存在，获取设备属性（如计算能力、显存大小）。
2. 内存分配：使用 cudaMalloc 在设备端分配显存，使用 malloc 或 new 在主机端分配内存。
3. 数据准备：在主机端初始化输入数据。
4. 数据传输：使用 cudaMemcpy 将数据从主机内存拷贝到设备显存（H2D）。
5. 内核启动：通过 kernel<<<grid, block>>>(args) 语法在设备上启动并行计算。
6. 结果回传：使用 cudaMemcpy 将计算结果从设备显存拷贝回主机内存（D2H）。
7. 资源释放：使用 cudaFree 释放设备显存，使用 free 释放主机内存。

案例拆解：向量加法的完整流程

以向量加法为例，假设需要计算 C[i] = A[i] + B[i]，其中向量长度为N=1024。主机端代码负责分配内存、初始化数据、调用内核并回收结果；设备端内核负责并行计算每个元素的加法。

// 主机代码
int N = 1024;
size_t bytes = N * sizeof(float);
float *h_A, *h_B, *h_C;
float *d_A, *d_B, *d_C;

// 1. 分配主机内存
h_A = (float*)malloc(bytes);
h_B = (float*)malloc(bytes);
h_C = (float*)malloc(bytes);

// 2. 初始化数据
for (int i = 0; i < N; i++) { h_A[i] = i * 1.0f; h_B[i] = i * 2.0f; }

// 3. 分配设备显存
cudaMalloc(&d_A, bytes);
cudaMalloc(&d_B, bytes);
cudaMalloc(&d_C, bytes);

// 4. 数据拷贝 H2D
cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

// 5. 启动内核
int threads = 256;
int blocks = (N + threads - 1) / threads;
vecAddKernel<<<blocks, threads>>>(d_A, d_B, d_C, N);

// 6. 同步并拷贝结果 D2H
cudaDeviceSynchronize();
cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

// 7. 释放资源
cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
free(h_A); free(h_B); free(h_C);

常见误区与总结

• 误区一：忘记调用 cudaDeviceSynchronize() 就读取主机端结果，导致数据未就绪。
• 误区二：在主机代码中直接访问设备指针，引发段错误。
• 误区三：忽略错误检查。每个CUDA API调用都应检查返回值（如 cudaError_t），否则错误难以定位。

主机与设备的分工是CUDA性能的基石。合理规划数据流、减少不必要的H2D/D2H拷贝，是优化CUDA程序的首要原则。

3.2 内核函数（Kernel）的定义与调用

核心原理：内核函数的语法约束

内核函数是在设备端执行的函数，必须使用 __global__ 修饰符声明。其返回值必须为 void，因为GPU无法直接向主机返回复杂类型。内核函数可以接受任意参数，但参数传递方式与普通C函数不同——参数会先被复制到设备的常量内存或共享内存中，因此应避免传递大型结构体。

内核函数的调用语法为：kernel_name<<<gridDim, blockDim, sharedMemSize, stream>>>(args)。其中 gridDim 和 blockDim 可以是整数或 dim3 类型，分别指定网格和块的维度；sharedMemSize 为可选参数，指定动态共享内存大小；stream 为可选参数，用于指定执行流。

判断方法：如何设计网格与块

参数	类型	取值范围	说明
gridDim.x	uint32_t	1 ~ 2^31-1	网格在x维度的块数量
gridDim.y	uint32_t	1 ~ 65535	网格在y维度的块数量
gridDim.z	uint32_t	1 ~ 65535	网格在z维度的块数量
blockDim.x	uint32_t	1 ~ 1024	块在x维度的线程数量
blockDim.y	uint32_t	1 ~ 1024	块在y维度的线程数量
blockDim.z	uint32_t	1 ~ 64	块在z维度的线程数量

块内线程总数（blockDim.x × blockDim.y × blockDim.z）不得超过1024。网格和块的维度选择直接影响线程的组织方式和内存访问模式。

案例拆解：二维矩阵乘法中的内核设计

假设需要计算两个32×32矩阵的乘法。我们可以将每个输出元素分配给一个线程，使用二维块（如8×8）和二维网格（如4×4）来组织线程。

__global__ void matMulKernel(float *A, float *B, float *C, int n) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    float sum = 0.0f;
    for (int k = 0; k < n; k++) {
        sum += A[row * n + k] * B[k * n + col];
    }
    C[row * n + col] = sum;
}

// 启动配置
dim3 block(8, 8);
dim3 grid((32 + 7) / 8, (32 + 7) / 8);
matMulKernel<<<grid, block>>>(d_A, d_B, d_C, 32);

这种设计使得每个线程负责一个输出元素，线程间无数据依赖，实现了完全并行。但需要注意，这种朴素实现存在大量重复的全局内存访问，后续章节将介绍共享内存优化。

常见误区与总结

• 误区一：在 __global__ 函数中使用 printf 调试（虽然CUDA支持，但输出顺序不确定且影响性能）。
• 误区二：内核函数中调用主机函数或使用 malloc（设备端不支持标准库）。
• 误区三：网格或块维度设置过大，超出硬件限制导致启动失败。

内核函数的定义与调用是CUDA编程的核心语法。掌握 __global__ 的约束、<<<>>> 的配置参数，以及网格/块的维度设计原则，是编写高效并行程序的前提。

3.3 线程索引与块索引的使用

核心原理：内置变量与索引计算

CUDA为每个线程提供了内置变量：threadIdx（线程在块内的索引）、blockIdx（块在网格内的索引）、blockDim（块的维度）、gridDim（网格的维度）。这些变量是 dim3 类型，包含 x、y、z 三个分量。

全局线程ID的计算公式为：globalThreadId = blockIdx.x * blockDim.x + threadIdx.x（一维情况）。对于二维或三维情况，需要分别计算每个维度的全局索引，再映射到线性内存地址。

操作流程：从线程索引到数据映射

1. 确定数据维度：判断数据是一维数组、二维矩阵还是三维体。
2. 选择线程组织方式：根据数据维度选择一维、二维或三维的块和网格。
3. 计算全局索引：使用 blockIdx、blockDim、threadIdx 计算每个线程对应的数据位置。
4. 边界检查：确保索引不越界，特别是当数据大小不是块大小的整数倍时。

案例拆解：图像灰度化中的索引映射

假设有一张1920×1080的RGB图像，需要转换为灰度图。每个像素由3个字节（R、G、B）组成，输出为1个字节。我们可以将每个像素分配给一个线程，使用二维块（16×16）和二维网格。

__global__ void rgbToGrayKernel(const unsigned char *rgb, unsigned char *gray, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (x < width && y < height) {
        int idx = (y * width + x) * 3;
        unsigned char r = rgb[idx];
        unsigned char g = rgb[idx + 1];
        unsigned char b = rgb[idx + 2];
        gray[y * width + x] = (unsigned char)(0.299f * r + 0.587f * g + 0.114f * b);
    }
}

// 启动配置
dim3 block(16, 16);
dim3 grid((1920 + 15) / 16, (1080 + 15) / 16);
rgbToGrayKernel<<<grid, block>>>(d_rgb, d_gray, 1920, 1080);

这里的关键是边界检查 if (x < width && y < height)，因为1920不能被16整除，网格边缘的线程会超出图像范围。这种检查虽然增加了一点开销，但保证了程序的正确性。

常见误区与总结

• 误区一：忘记边界检查，导致越界访问显存，产生未定义行为。
• 误区二：混淆 blockIdx.x 和 threadIdx.x 的用途，导致索引计算错误。
• 误区三：使用一维索引访问二维数据，忽略了行优先存储的偏移量计算。

线程索引与块索引的正确使用是CUDA编程中最容易出错的部分。建议在编写内核时，先在纸上画出线程到数据的映射关系，再编写代码。同时，善用边界检查可以避免大量难以调试的内存错误。

通过本章的学习，读者应掌握CUDA编程模型的基本框架：主机与设备的分工协作、内核函数的定义与调用语法、以及线程索引与块索引的映射方法。这些知识是后续学习内存管理、性能优化和高级并行模式的基础。在实际编程中，建议从简单的向量运算开始，逐步过渡到矩阵运算和图像处理等复杂场景，通过不断实践加深对CUDA编程模型的理解。