📚 课程目录第3章 CUDA核函数编写基础

3. CUDA核函数编写基础

CUDA程序的基本骨架可以概括为“主机准备数据—设备分配内存—数据拷入设备—启动核函数—结果拷回主机—释放资源”。这一节围绕这条主线，把核函数的声明与调用语法、启动配置、线程索引计算、内存管理与错误检查串成一个可以真正跑通的最小闭环。理解这些语法细节，是后续编写数值核函数的前提。

一、global函数的声明与调用规则

核函数用__global__修饰，它在设备上执行，但只能由主机代码调用。声明时通常写成__global__ void kernelName(参数列表)，返回类型必须是void，因为主机无法直接接收核函数的返回值。核函数内部可以读取threadIdx、blockIdx、blockDim、gridDim等内置变量来确定自己负责哪一部分数据。

调用核函数时使用三重尖括号语法：kernelName<<<grid, block>>>(实参)。这里的grid和block既可以是整数，也可以是dim3类型。例如dim3 block(256)表示每个线程块256个线程，dim3 grid((N + 255) / 256)表示按数据量向上取整得到线程块数量。核函数启动是异步的，主机调用后立即继续执行，因此后续若要把结果拷回主机，必须通过同步或拷贝操作建立依赖。

二、线程唯一标识的计算

每个线程需要知道自己处理哪个元素，最常用的全局索引公式是：

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

其中blockIdx.x是当前线程块在网格中的编号，blockDim.x是每块线程数，threadIdx.x是线程在块内的编号。三者相乘相加得到全局唯一的一维索引。由于grid是按(N + blockSize - 1) / blockSize向上取整分配的，最后一个线程块中可能存在idx >= N的线程，因此核函数开头必须加边界判断if (idx < N)，否则会越界访问。对于二维或三维数据，可以分别用blockIdx.y、threadIdx.y等组合出二维索引，但一维数组是最常见的入门场景。

三、主机与设备内存管理

主机内存用malloc或new分配，设备内存必须用cudaMalloc分配。它的第一个参数是void**类型的设备指针地址，第二个参数是字节数：

float *d_a = nullptr;
cudaMalloc((void**)&d_a, N * sizeof(float));

使用完毕后必须调用cudaFree(d_a)释放，否则会造成设备内存泄漏。主机指针不能直接传给核函数解引用，设备指针也不能在主机代码中直接读写，二者之间必须通过cudaMemcpy搬运数据。常用的拷贝方向有cudaMemcpyHostToDevice和cudaMemcpyDeviceToHost，参数顺序为“目标指针、源指针、字节数、方向”。

四、错误检查与最小可运行示例

CUDA的许多API返回cudaError_t，核函数启动本身不返回错误码，需要用cudaGetLastError()或cudaDeviceSynchronize()捕获启动和执行阶段的错误。一个实用的检查宏如下：

#define CHECK(call) do { \
    cudaError_t err = (call); \
    if (err != cudaSuccess) { \
        printf("CUDA错误: %s\n", cudaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    } \
} while(0)

下面是一个完整的向量加法示例，N取1024，每块256线程，因此grid为4：

__global__ void vecAdd(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 = 1024;
    size_t bytes = N * sizeof(float);
    float *h_a = (float*)malloc(bytes), *h_b = (float*)malloc(bytes), *h_c = (float*)malloc(bytes);
    for (int i = 0; i < N; ++i) { h_a[i] = 1.0f; h_b[i] = 2.0f; }

    float *d_a, *d_b, *d_c;
    CHECK(cudaMalloc((void**)&d_a, bytes));
    CHECK(cudaMalloc((void**)&d_b, bytes));
    CHECK(cudaMalloc((void**)&d_c, bytes));
    CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));

    vecAdd<<<(N + 255) / 256, 256>>>(d_a, d_b, d_c, N);
    CHECK(cudaGetLastError());
    CHECK(cudaDeviceSynchronize());

    CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));
    printf("c[0]=%f, c[1023]=%f\n", h_c[0], h_c[1023]);

    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b); free(h_c);
    return 0;
}

若程序正确执行，输出应为c[0]=3.000000, c[1023]=3.000000。如果出现invalid configuration argument，通常说明block大小超过1024或为0；如果出现an illegal memory access was encountered，优先检查边界判断是否遗漏、指针是否误用主机地址。下图把这条从主机到设备再回到主机的数据通路与关键检查点画在一起，便于对照代码定位问题。

五、常见错误与排查要点

入门阶段最容易踩的坑集中在三处。第一是忘记边界判断，当N不是block大小的整数倍时，最后一个块的多余线程会读写越界，表现为随机崩溃或结果错误。第二是混淆主机指针与设备指针，例如把h_a直接传给核函数，编译能通过但运行时报非法访问。第三是忽略异步性，核函数启动后立刻用主机读取h_c，此时拷贝尚未发生，读到的是旧值；正确做法是先cudaMemcpy回主机，或显式cudaDeviceSynchronize()。

建议把CHECK宏贯穿所有CUDA API调用，并在核函数启动后立即调用cudaGetLastError()，这样能把“配置错误”和“执行错误”区分开。对于数值核函数，还可以在小规模数据上先用主机端参考实现对比结果，确认索引映射与边界处理正确后，再扩大到大规模数据并关注性能。掌握这套最小闭环后，后续的共享内存、归约与矩阵运算都可以在此基础上逐步叠加。

📚 课程目录