当前位置：首页 > news >正文

GPU编程实战指南01：CUDA编程极简手册

news 2025/9/2 4:57:38

1. CUDA基础概念

1.1 线程层次结构

CUDA采用层次化的线程组织结构，从小到大依次为：

线程（Thread）：最基本的执行单元
- 每个线程执行相同的核函数代码
- 通过threadIdx变量识别线程ID
线程块（Block）：多个线程的组合
- 同一块内的线程可以同步和共享内存
- 通过blockIdx变量识别块ID
- 最多可包含1024个线程（具体取决于硬件）
网格（Grid）：多个线程块的组合
- 通过gridDim变量确定网格维度
- 块之间相互独立执行

// 线程索引计算示例
int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;

1.2 内存层次结构

CUDA设备具有多层次的内存结构：

全局内存（Global Memory）
- 所有线程都可访问
- 延迟较高，带宽较大
- 通过cudaMalloc()和cudaFree()管理
共享内存（Shared Memory）
- 同一线程块内的线程共享
- 延迟低，带宽高
- 使用__shared__关键字声明
寄存器（Registers）
- 每个线程私有
- 访问速度最快
- 数量有限

// 共享内存使用示例
__shared__ float sharedData[256];

2. CUDA编程核心要素

2.1 核函数

核函数是在GPU上并行执行的函数：

使用__global__关键字声明
返回类型必须是void
可以从CPU调用，在GPU上执行

__global__ void vectorAdd(float* a, float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

2.2 内存管理

CUDA程序中的内存管理涉及以下操作：

内存分配

float *d_data;
cudaMalloc((void**)&d_data, size * sizeof(float));

内存传输

cudaMemcpy(d_data, h_data, size * sizeof(float), cudaMemcpyHostToDevice);

内存释放
```
cudaFree(d_data);
```

2.3 同步机制

CUDA提供多种同步机制：

线程同步：__syncthreads()
事件同步：cudaEventSynchronize()
流同步：cudaStreamSynchronize()

// 线程同步示例
__global__ void syncExample() {
    __shared__ float data[256];
    // ... 一些操作 ...
    __syncthreads();  // 确保所有线程完成操作
}

3. CUDA优化技巧

3.1 内存访问优化

合并访问
- 相邻线程访问相邻内存位置
- 减少内存事务数量

// 好的访问模式
int idx = blockIdx.x * blockDim.x + threadIdx.x;
data[idx] = value;  // 连续访问

// 不好的访问模式
int idx = threadIdx.x * gridDim.x + blockIdx.x;
data[idx] = value;  // 跨步访问

3.2 共享内存使用

用于频繁访问的数据
减少全局内存访问
注意bank冲突

__global__ void sharedMemExample() {
    __shared__ float sData[BLOCK_SIZE];
    int tid = threadIdx.x;
    sData[tid] = globalData[tid];  // 加载到共享内存
    __syncthreads();
    // 使用共享内存中的数据
}

3.3 线程分配优化

线程块大小是32的倍数（warp大小）
保持足够的并行度
考虑资源使用情况

// 启动核函数时的线程配置
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

4. 常见问题和解决方案

内存相关问题

内存泄漏
- 问题：未正确释放CUDA内存
- 解决：使用RAII模式或确保配对的cudaMalloc()和cudaFree()
内存访问越界
- 问题：访问超出分配范围的内存
- 解决：添加边界检查，确保正确的内存分配大小

性能问题

低设备利用率
- 问题：GPU资源未充分利用
- 解决：调整线程块大小，增加并行度
内存带宽瓶颈
- 问题：频繁的全局内存访问
- 解决：使用共享内存，优化内存访问模式

5. 实际案例分析

向量加法示例

// 核函数定义
__global__ void vectorAdd(float* a, 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() {
    int n = 1<<20;  // 1M elements
    size_t size = n * sizeof(float);

    // 分配主机内存
    float *h_a = (float*)malloc(size);
    float *h_b = (float*)malloc(size);
    float *h_c = (float*)malloc(size);

    // 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaMalloc((void**)&d_a, size);
    cudaMalloc((void**)&d_b, size);
    cudaMalloc((void**)&d_c, size);

    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_a[i] = rand()/(float)RAND_MAX;
        h_b[i] = rand()/(float)RAND_MAX;
    }

    // 复制数据到设备
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    // 启动核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

    // 复制结果回主机
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    // 清理内存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

这个示例展示了：