cuda入门（一）

一、内核（Kernel）、线程块（Block）和线程网格（Grid）

1.1 Kernel

Kernel 函数必须使用 global 限定符来声明，这是 CUDA 编译器识别的特殊标记。global 限定符表明这个函数可以从 CPU 端（主机端）调用，在 GPU 端（设备端）执行。Kernel 函数的返回类型必须是 void，因为从 CPU 的角度来看，Kernel 调用是异步的，CPU 不会等待 GPU 完成所有线程的执行就继续执行后续代码。这种设计允许 CPU 和 GPU 同时工作，最大化系统整体的计算效率。

启动一个 Kernel 需要使用三重尖括号 <<< >>> 语法，这组尖括号被称为执行配置（Execution Configuration）。执行配置中需要指定线程网格和线程块的维度信息，这些参数决定了启动多少个线程以及如何组织这些线程。例如，<<<dimGrid, dimBlock>>> 表示启动一个包含 dimGrid 个线程块的网格，每个线程块包含 dimBlock 个线程。通过合理设置这些参数，可以控制并行计算的规模和粒度，以达到最佳的性能表现。

Kernel 启动的第三个可选参数是一个指向共享内存的指针，用于在同一个线程块内的线程之间传递数据。这个参数在需要线程协作的场景中非常有用，例如规约操作（Reduction）或卷积计算中的窗口数据共享。Kernel 调用的实际语法形式为 kernel_name<<<grid, block, shared_mem_size, stream>>>(args)，其中 shared_mem_size 指定了为每个线程块分配的动态共享内存大小，stream 则指定了执行流，允许多个 Kernel 在同一个 GPU 上重叠执行。

__global__ void add_vectors(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() {
    int n = 1024;
    float *a, *b, *c;
    // 内存分配和数据初始化代码...
    
    // 启动 Kernel：1 个线程块，1024 个线程
    add_vectors<<<1, 1024>>>(a, b, c, n);
    
    // 等待 GPU 完成计算
    cudaDeviceSynchronize();
    
    return 0;
}

几个重要的内存分配和初始化函数：

// 内存分配函数
cudaMalloc((void**)&d_A, sizeA); 
// 将数据从 Host 拷贝到 Device
cudaMemcpy(d_A, h_A, sizeA, cudaMemcpyHostToDevice);
// 配置线程块和网格
dim3 threadsPerBlock(16, 16);  // 16x16的线程块
dim3 blocksPerGrid((N + threadsPerBlock.x - 1) / threadsPerBlock.x, (M + threadsPerBlock.y - 1) / threadsPerBlock.y);

这段代码用于检查 kernel 启动错误，并在执行完成后同步，确保结果可用。

    // 检查错误
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("CUDA Kernel Launch Error: %s\n", cudaGetErrorString(err));
    return EXIT_FAILURE;
}
cudaDeviceSynchronize();

1.2 Thread（线程）：并行的基本单元

CUDA 中的线程执行相同的 Kernel 代码，但处理不同的数据元素，这种执行模式被称为"单指令多数据"（SIMT）。每个线程都有一个唯一的全局索引，用于确定它应该处理数据集中的哪个部分。

每个 CUDA 线程都可以通过内置变量获取自己在整个线程层次结构中的位置。这些内置变量包括 threadIdx、blockIdx、blockDim 和 gridDim，它们分别表示当前线程在所属线程块中的索引、当前线程块在整个线程网格中的索引、线程块的维度以及线程网格的维度。这些变量都是预定义的结构体，通常包含 x、y（可能还有 z）三个分量，用于支持一维、二维或三维的线程组织方式。

计算线程全局索引的公式为：global_idx = (blockIdx.x * blockDim.x) + threadIdx.x。这个公式假设使用一维的线程组织和索引计算。对于二维场景（如处理图像矩阵），可以使用 (blockIdx.y * blockDim.y + threadIdx.y) * width + (blockIdx.x * blockDim.x + threadIdx.x) 这样的公式将二维坐标转换为一维全局索引。正确计算线程索引是编写 CUDA 程序的基本功，因为只有每个线程都知道自己应该处理哪个数据元素，整个并行计算才能正确进行。