CUDA 编程入门

CUDA（Compute Unified Device Architecture）是 NVIDIA 推出的并行计算平台和编程模型。本文将介绍 CUDA 编程的基础知识，包括环境搭建、基本概念和常用编程技巧。

CUDA 简介

CUDA 是一种通用的并行计算架构，允许开发者使用 C/C++ 语言编写在 GPU 上运行的程序。CUDA 的主要特点包括：

• 易于使用：基于 C/C++，学习曲线平缓
• 高性能：充分利用 GPU 的并行计算能力
• 丰富的库：提供 cuBLAS、cuDNN、cuFFT 等高性能库
• 跨平台：支持 Windows、Linux、macOS

环境搭建

1. 安装 NVIDIA 驱动

首先需要安装 NVIDIA 显卡驱动：

# Ubuntu/Debian
sudo apt-get update
sudo apt-get install nvidia-driver-535

# 检查驱动是否安装成功
nvidia-smi

2. 安装 CUDA Toolkit

从 NVIDIA 官网下载并安装 CUDA Toolkit：

# Ubuntu/Debian
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-ubuntu2204.pin
sudo mv cuda-ubuntu2204.pin /etc/apt/preferences.d/cuda-repository-pin-600
wget https://developer.download.nvidia.com/compute/cuda/12.2.0/local_installers/cuda-repo-ubuntu2204-12-2-local_12.2.0-535.54.03-1_amd64.deb
sudo dpkg -i cuda-repo-ubuntu2204-12-2-local_12.2.0-535.54.03-1_amd64.deb
sudo cp /var/cuda-repo-ubuntu2204-12-2-local/cuda-*-keyring.gpg /usr/share/keyrings/
sudo apt-get update
sudo apt-get install cuda

# 配置环境变量
echo 'export PATH=/usr/local/cuda/bin:$PATH' >> ~/.bashrc
echo 'export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc
source ~/.bashrc

# 验证安装
nvcc --version

3. 安装 cuDNN

cuDNN 是 NVIDIA 的深度学习加速库：

# 下载 cuDNN（需要 NVIDIA 账号）
tar -xvf cudnn-linux-x86_64-8.9.7.29_cuda12-archive.tar.xz
sudo cp cudnn-linux-x86_64-8.9.7.29_cuda12-archive/include/cudnn*.h /usr/local/cuda/include
sudo cp -P cudnn-linux-x86_64-8.9.7.29_cuda12-archive/lib/libcudnn* /usr/local/cuda/lib64
sudo chmod a+r /usr/local/cuda/include/cudnn*.h /usr/local/cuda/lib64/libcudnn*

CUDA 编程基础

Hello World

第一个 CUDA 程序：

#include <stdio.h>

// CUDA Kernel 函数
__global__ void helloCUDA() {
    printf("Hello from GPU thread %d\n", threadIdx.x);
}

int main() {
    // 启动 Kernel，使用 1 个 Block，每个 Block 10 个线程
    helloCUDA<<<1, 10>>>();

    // 等待 GPU 执行完成
    cudaDeviceSynchronize();

    return 0;
}

编译并运行：

nvcc hello.cu -o hello
./hello

CUDA 关键字

关键字	执行位置	调用位置
__global__	Device（GPU）	Host（CPU）
__device__	Device（GPU）	Device（GPU）
__host__	Host（CPU）	Host（CPU）

// __global__: Kernel 函数，从 Host 调用，在 Device 执行
__global__ void kernel(float *data) {
    // ...
}

// __device__: 设备函数，从 Device 调用，在 Device 执行
__device__ float deviceFunction(float x) {
    return x * 2.0f;
}

// __host__: 主机函数，从 Host 调用，在 Host 执行（默认）
__host__ float hostFunction(float x) {
    return x * 2.0f;
}

// __host__ __device__: 既可以在 Host 也可以在 Device 调用
__host__ __device__ float sharedFunction(float x) {
    return x * 2.0f;
}

线程层次结构

CUDA 的线程层次结构包括：

• Thread（线程）：最基本的执行单元
• Block（线程块）：一组线程，可以访问共享内存和同步
• Grid（网格）：一组线程块，构成一个完整的 kernel

__global__ void kernel(float *data) {
    // 线程 ID
    int tid = threadIdx.x;

    // Block ID
    int bid = blockIdx.x;

    // 全局线程 ID
    int globalTid = threadIdx.x + blockIdx.x * blockDim.x;

    data[globalTid] = globalTid;
}

int main() {
    int N = 1024;
    float *d_data;

    cudaMalloc(&d_data, N * sizeof(float));

    // 启动 Kernel：4 个 Block，每个 Block 256 个线程
    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;
    kernel<<<numBlocks, blockSize>>>(d_data);

    cudaFree(d_data);
    return 0;
}

内存管理

设备内存分配

int main() {
    int N = 1024;
    int *h_data;  // 主机内存
    int *d_data;  // 设备内存

    // 分配主机内存
    h_data = (int*)malloc(N * sizeof(int));

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

    // 初始化主机数据
    for (int i = 0; i < N; i++) {
        h_data[i] = i;
    }

    // 主机到设备拷贝
    cudaMemcpy(d_data, h_data, N * sizeof(int), cudaMemcpyHostToDevice);

    // 执行 Kernel
    kernel<<<1, 256>>>(d_data);

    // 设备到主机拷贝
    cudaMemcpy(h_data, d_data, N * sizeof(int), cudaMemcpyDeviceToHost);

    // 释放内存
    free(h_data);
    cudaFree(d_data);

    return 0;
}

统一内存（Managed Memory）

CUDA 6.0 引入了统一内存，简化了内存管理：

int main() {
    int N = 1024;
    int *data;  // 统一内存

    // 分配统一内存
    cudaMallocManaged(&data, N * sizeof(int));

    // 初始化数据
    for (int i = 0; i < N; i++) {
        data[i] = i;
    }

    // 执行 Kernel
    kernel<<<1, 256>>>(data);

    // 同步
    cudaDeviceSynchronize();

    // 释放内存
    cudaFree(data);

    return 0;
}

常用编程模式

向量加法

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

int main() {
    int N = 1024 * 1024;
    float *a, *b, *c;

    cudaMallocManaged(&a, N * sizeof(float));
    cudaMallocManaged(&b, N * sizeof(float));
    cudaMallocManaged(&c, N * sizeof(float));

    // 初始化数据
    for (int i = 0; i < N; i++) {
        a[i] = i;
        b[i] = i * 2;
    }

    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;
    vectorAdd<<<numBlocks, blockSize>>>(a, b, c, N);

    cudaDeviceSynchronize();

    // 释放内存
    cudaFree(a);
    cudaFree(b);
    cudaFree(c);

    return 0;
}

矩阵乘法

#define TILE_SIZE 16

__global__ void matrixMul(float *A, float *B, float *C, int N) {
    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
    __shared__ float tileB[TILE_SIZE][TILE_SIZE];

    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;

    int row = by * TILE_SIZE + ty;
    int col = bx * TILE_SIZE + tx;

    float sum = 0.0f;

    for (int t = 0; t < N / TILE_SIZE; ++t) {
        tileA[ty][tx] = A[row * N + t * TILE_SIZE + tx];
        tileB[ty][tx] = B[(t * TILE_SIZE + ty) * N + col];
        __syncthreads();

        for (int k = 0; k < TILE_SIZE; ++k) {
            sum += tileA[ty][k] * tileB[k][tx];
        }
        __syncthreads();
    }

    C[row * N + col] = sum;
}

int main() {
    int N = 1024;
    float *A, *B, *C;

    cudaMallocManaged(&A, N * N * sizeof(float));
    cudaMallocManaged(&B, N * N * sizeof(float));
    cudaMallocManaged(&C, N * N * sizeof(float));

    // 初始化数据...

    dim3 blockSize(TILE_SIZE, TILE_SIZE);
    dim3 gridSize(N / TILE_SIZE, N / TILE_SIZE);
    matrixMul<<<gridSize, blockSize>>>(A, B, C, N);

    cudaDeviceSynchronize();

    cudaFree(A);
    cudaFree(B);
    cudaFree(C);

    return 0;
}

错误处理

#define CUDA_CHECK(call) \
    do { \
        cudaError_t error = call; \
        if (error != cudaSuccess) { \
            printf("CUDA error: %s:%d, code: %d, reason: %s\n", \
                   __FILE__, __LINE__, error, cudaGetErrorString(error)); \
            exit(1); \
        } \
    } while(0)

int main() {
    int *d_data;
    int N = 1024;

    // 使用错误检查宏
    CUDA_CHECK(cudaMalloc(&d_data, N * sizeof(int)));

    // 检查 Kernel 执行错误
    kernel<<<1, 256>>>(d_data);
    CUDA_CHECK(cudaGetLastError());  // 检查 Kernel 启动错误
    CUDA_CHECK(cudaDeviceSynchronize());  // 检查执行错误

    CUDA_CHECK(cudaFree(d_data));

    return 0;
}

性能分析

使用 nvprof

# 编译程序
nvcc -O3 program.cu -o program

# 使用 nvprof 分析
nvprof ./program

# 详细分析
nvprof --print-gpu-trace ./program

使用 Nsight Compute

# 安装 Nsight Compute
# 从 NVIDIA 官网下载并安装

# 分析 Kernel
ncu --set full ./program

# 查看报告
ncu-ui report.ncu-rep

最佳实践

• 内存合并访问：确保相邻线程访问连续的内存地址
• 避免线程分歧：Warp 内的线程应执行相同的代码路径
• 使用共享内存：利用共享内存减少全局内存访问
• 最大化占用率：保持足够的活跃线程
• 异步执行：使用流和异步内存拷贝
• 错误处理：始终检查 CUDA API 调用的返回值

总结

CUDA 是一个强大的并行计算平台，通过简单的 C/C++ 扩展，开发者可以充分利用 GPU 的计算能力。本文介绍了 CUDA 编程的基础知识，包括环境搭建、基本概念、内存管理和常用编程模式。掌握这些基础知识后，可以进一步学习更高级的优化技巧和 CUDA 库的使用。