CUDA C/C++ 编程基础

CUDA C/C++ 是在标准 C/C++ 基础上扩展的并行编程语言，通过 nvcc 编译器将 GPU 代码编译�?PTX �?SASS 指令�?

开发环境与编译

nvcc 编译�?

# 编译单个文件
nvcc -o my_program my_kernel.cu

# 指定 GPU 架构（推荐显式指定）
nvcc -arch=sm_80 -o my_program my_kernel.cu   # Ampere A100
nvcc -arch=sm_90 -o my_program my_kernel.cu   # Hopper H100

# 生成多架�?fat binary
nvcc -gencode arch=compute_80,code=sm_80 \
     -gencode arch=compute_90,code=sm_90 \
     -o my_program my_kernel.cu

# 开启优�?
nvcc -O3 -arch=sm_80 -o my_program my_kernel.cu

# 生成 PTX（中间表示，可读�?
nvcc -ptx my_kernel.cu

架构代号对照

GPU	架构	sm_xx
V100	Volta	sm_70
T4	Turing	sm_75
A100	Ampere	sm_80
A10/A30	Ampere	sm_86
H100	Hopper	sm_90
RTX 4090	Ada	sm_89

内存管理

基础 API

// 分配设备内存
float* d_data;
cudaMalloc(&d_data, N * sizeof(float));

// 主机 �?设备
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);

// 设备 �?主机
cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);

// 设备内存初始化（类似 memset�?
cudaMemset(d_data, 0, N * sizeof(float));

// 释放
cudaFree(d_data);

统一内存（Unified Memory�?

// 分配统一内存（CPU �?GPU 都可直接访问�?
float* data;
cudaMallocManaged(&data, N * sizeof(float));

// CPU 端初始化
for (int i = 0; i < N; i++) data[i] = i;

// 直接传给 Kernel，无需显式拷贝
kernel<<<grid, block>>>(data, N);
cudaDeviceSynchronize();

// CPU 端直接读取结�?
printf("%f\n", data[0]);

cudaFree(data);

统一内存的代�?

统一内存通过页面迁移实现，首次访问会触发 Page Fault，有额外开销。生产环境中显式内存管理通常性能更好�?

固定内存（Pinned Memory�?

// 分配页锁定内存（不会被换出到磁盘�?
float* h_pinned;
cudaMallocHost(&h_pinned, N * sizeof(float));

// 固定内存�?H2D/D2H 传输速度更快（可�?PCIe 峰值带宽）
// 普�?malloc 内存：~6 GB/s
// 固定内存：~12 GB/s（PCIe 4.0 x16�?

cudaFreeHost(h_pinned);

错误处理

// 检�?CUDA API 调用
#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, N * sizeof(float)));
CUDA_CHECK(cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice));

// Kernel 启动后检查错�?
kernel<<<grid, block>>>(args);
CUDA_CHECK(cudaGetLastError());       // 检查启动错�?
CUDA_CHECK(cudaDeviceSynchronize());  // 等待完成并检查运行时错误

设备查询

int deviceCount;
cudaGetDeviceCount(&deviceCount);

for (int i = 0; i < deviceCount; i++) {
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, i);
    
    printf("Device %d: %s\n", i, prop.name);
    printf("  Compute Capability: %d.%d\n", prop.major, prop.minor);
    printf("  SM Count: %d\n", prop.multiProcessorCount);
    printf("  Global Memory: %.1f GB\n", prop.totalGlobalMem / 1e9);
    printf("  Shared Mem per Block: %zu KB\n", prop.sharedMemPerBlock / 1024);
    printf("  Max Threads per Block: %d\n", prop.maxThreadsPerBlock);
    printf("  Warp Size: %d\n", prop.warpSize);
    printf("  Memory Bandwidth: %.1f GB/s\n",
           2.0 * prop.memoryClockRate * (prop.memoryBusWidth / 8) / 1e6);
}

原子操作

__global__ void histogram(int* data, int* hist, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        // 原子加法，避免多线程竞争
        atomicAdd(&hist[data[idx]], 1);
    }
}

常用原子操作

函数	说明
atomicAdd(addr, val)	原子加法，返回旧�?
atomicSub(addr, val)	原子减法
atomicMin/Max(addr, val)	原子最�?最大�?
atomicExch(addr, val)	原子交换
atomicCAS(addr, compare, val)	比较并交换（CAS�?
atomicAnd/Or/Xor(addr, val)	原子位运�?

原子操作的性能

原子操作会导致序列化，高竞争场景下性能很差。优化方案：先在共享内存做局部归约，再原子写回全局内存�?

动态并行（Dynamic Parallelism�?

�?Kepler 架构开始，Kernel 内部可以启动新的 Kernel�?

__global__ void parentKernel(float* data, int N) {
    if (threadIdx.x == 0 && blockIdx.x == 0) {
        // �?GPU 上动态启动子 Kernel
        childKernel<<<N/256, 256>>>(data, N);
        cudaDeviceSynchronize();  // 等待�?Kernel 完成
    }
}

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

完整程序示例

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

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

__global__ void saxpy(int N, float alpha, float* x, float* y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) y[i] = alpha * x[i] + y[i];
}

int main() {
    const int N = 1 << 20;  // 1M 元素
    const float alpha = 2.0f;
    
    // 分配主机内存
    float *h_x = new float[N];
    float *h_y = new float[N];
    for (int i = 0; i < N; i++) { h_x[i] = 1.0f; h_y[i] = 2.0f; }
    
    // 分配设备内存
    float *d_x, *d_y;
    CUDA_CHECK(cudaMalloc(&d_x, N * sizeof(float)));
    CUDA_CHECK(cudaMalloc(&d_y, N * sizeof(float)));
    
    // 传输数据
    CUDA_CHECK(cudaMemcpy(d_x, h_x, N * sizeof(float), cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_y, h_y, N * sizeof(float), cudaMemcpyHostToDevice));
    
    // 启动 Kernel
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    saxpy<<<gridSize, blockSize>>>(N, alpha, d_x, d_y);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());
    
    // 取回结果
    CUDA_CHECK(cudaMemcpy(h_y, d_y, N * sizeof(float), cudaMemcpyDeviceToHost));
    
    printf("y[0] = %f (expected 4.0)\n", h_y[0]);
    
    // 清理
    cudaFree(d_x); cudaFree(d_y);
    delete[] h_x; delete[] h_y;
    return 0;
}

---

下一篇：内存层次结构 →