PMPP Chapter 02: Heterogeneous data parallel computing

以 RGB 图片 -> 灰度图片转换为例，说明 Data Parallel 如何运行并行计算：

CUDA 编程步骤：

• 分配 device global memory
• 拷贝 host memory 内容到 device memory
• launch kernel 在 device 进行计算
• 从 device 拷贝计算结果到 host
• 释放 device 内存

注：

• 实际项目中，热点数据常驻 device memory，避免频繁 malloc/copy/free 影响性能。
• launch kernel 以及部份 CUDA API 是异步操作，实际项目中可以利用此特性进行 CPU 和 GPU 计算的 overlap，同时需要注意 host 侧操作依赖 device 侧计算完成时，要做好同步，确保 device 侧计算完成，host 侧读取到正确数据。

关于 memcpy/memset API 的同步/异步行为解释，见：

• Nvidia Docs-API Synchronization behavior

实际编程时，最佳实践是尽量使用 pinned memory + device memory + 显式指定 CUDA Stream（而不是依赖默认 stream），host 侧依赖 device 计算结果时，显式调用 cudaMemcpyAsync + cudaStreamSynchronize 接口等待 device 侧计算/拷贝到 host 完成。示例代码：

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>

// 错误检查宏：任何 CUDA API 调用返回非 cudaSuccess 时打印错误并退出
#define CUDA_CALL(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 i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) C[i] = A[i] + B[i];
}

int main() {
    const int N = 1 << 20;              // 1M elements
    const size_t bytes = N * sizeof(float);

    // 1. 分配 pinned host memory（异步传输的前提）
    float *h_A, *h_B, *h_C;
    CUDA_CALL(cudaMallocHost(&h_A, bytes));
    CUDA_CALL(cudaMallocHost(&h_B, bytes));
    CUDA_CALL(cudaMallocHost(&h_C, bytes));

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

    // 2. 分配 device memory
    float *d_A, *d_B, *d_C;
    CUDA_CALL(cudaMalloc(&d_A, bytes));
    CUDA_CALL(cudaMalloc(&d_B, bytes));
    CUDA_CALL(cudaMalloc(&d_C, bytes));

    // 3. 创建显式 stream
    cudaStream_t stream;
    CUDA_CALL(cudaStreamCreate(&stream));

    // 4. 异步拷贝 H2D
    CUDA_CALL(cudaMemcpyAsync(d_A, h_A, bytes, cudaMemcpyHostToDevice, stream));
    CUDA_CALL(cudaMemcpyAsync(d_B, h_B, bytes, cudaMemcpyHostToDevice, stream));

    // 5. 在指定 stream 上 launch kernel
    int threads = 256;
    int blocks = (N + threads - 1) / threads;
    vectorAdd<<<blocks, threads, 0, stream>>>(d_A, d_B, d_C, N);
    CUDA_CALL(cudaGetLastError());      // 检查 kernel launch 是否成功

    // 6. 异步拷贝 D2H
    CUDA_CALL(cudaMemcpyAsync(h_C, d_C, bytes, cudaMemcpyDeviceToHost, stream));

    // 7. 同步 stream：确保 device 计算和拷贝全部完成
    CUDA_CALL(cudaStreamSynchronize(stream));

    printf("h_C[0] = %f, h_C[N-1] = %f\n", h_C[0], h_C[N-1]);

    // 8. 清理
    CUDA_CALL(cudaFreeHost(h_A));
    CUDA_CALL(cudaFreeHost(h_B));
    CUDA_CALL(cudaFreeHost(h_C));
    CUDA_CALL(cudaFree(d_A));
    CUDA_CALL(cudaFree(d_B));
    CUDA_CALL(cudaFree(d_C));
    CUDA_CALL(cudaStreamDestroy(stream));
    return 0;
}

CUDA 线程层次

• many blocks -> grid
• many threads -> block

编译过程

• Use the CUDA keywords to separate the host code and device code.
• Compile host code with the host’s standard C/C++ compilers.
• The NVCC compiler compiles the device code into virtual binary files called PTX files. These PTX files are further compiled by a runtime component of NVCC into the real object files and executed on a CUDA-capable GPU device.