CUDA 编程基础教程（面向 Java/Python 开发者）

nvcc --version

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2024 NVIDIA Corporation
Built on ...
Cuda compilation tools, release 12.x, V12.x.x

nvidia-smi

int[] arr = new int[5];  // arr 是一个"引用"，指向堆里的数组对象

arr = [0] * 5  # arr 是一个变量名，指向内存中的列表对象

int a = 10;        // 在栈上分配一个整数，值是 10
int* p = &a;       // p 是一个"指向 int 的指针"，&a 表示取 a 的地址

printf("%d\n", *p);  // *p 表示"解引用"：去 p 指向的地址里取值，输出 10
*p = 20;             // 通过指针修改原变量的值，现在 a 变成 20 了

#include <stdlib.h>

// 在堆上申请能放 5 个 int 的内存空间（5 * 4 = 20 字节）
int* arr = (int*)malloc(5 * sizeof(int));

if (arr == NULL) {
    // 申请失败的处理
}

// 使用这块内存...
arr[0] = 100;
arr[1] = 200;

// 用完了必须手动释放！否则内存泄漏。
free(arr);

#include <stdio.h>

// __global__ 表示这个函数在 GPU 上执行，但由 CPU 调用
__global__ void helloFromGPU() {
    // threadIdx.x 是当前线程在 Block 中的编号
    printf("Hello from GPU thread %d!\n", threadIdx.x);
}

int main() {
    printf("Hello from CPU!\n");

    // <<<1, 4>>> 是 CUDA 的"启动配置"：
    // 1 个 Block，每个 Block 里有 4 个线程
    helloFromGPU<<<1, 4>>>();

    // cudaDeviceSynchronize() 等待 GPU 执行完毕，
    // 否则 CPU 可能不等 GPU 跑完就退出了
    cudaDeviceSynchronize();

    printf("All GPU threads finished.\n");
    return 0;
}

nvcc hello_cuda.cu -o hello_cuda
hello_cuda.exe        # Windows
./hello_cuda          # Linux

Hello from CPU!
Hello from GPU thread 0!
Hello from GPU thread 1!
Hello from GPU thread 2!
Hello from GPU thread 3!
All GPU threads finished.

Grid（整个工厂）
├── Block 0（车间 0）
│   ├── Thread 0（工人 0）
│   ├── Thread 1（工人 1）
│   ├── Thread 2（工人 2）
│   └── ...
├── Block 1（车间 1）
│   ├── Thread 0
│   ├── Thread 1
│   └── ...
└── Block 2（车间 2）
    └── ...

global_id = blockIdx.x * blockDim.x + threadIdx.x

// 例 1：一维，处理 1024 个元素
// 每个 Block 256 个线程，需要 4 个 Block
myKernel<<<4, 256>>>();
// 线程全局 ID：blockIdx.x * 256 + threadIdx.x，范围 [0, 1023]

// 例 2：二维，处理图像像素（宽高各 16）
dim3 blockSize(4, 4);   // 每个 Block 4x4 = 16 个线程
dim3 gridSize(4, 4);    // Grid 4x4 = 16 个 Block，总共 256 个线程
myKernel<<<gridSize, blockSize>>>();
// 线程二维坐标：(blockIdx.x * 4 + threadIdx.x, blockIdx.y * 4 + threadIdx.y)

CPU 内存（Host Memory）          GPU 显存（Device Memory）
     ↑                                ↑
   malloc                          cudaMalloc
     |                                |
     |-------- cudaMemcpy H2D -------->|
     |     (Host to Device：CPU → GPU)  |
     |                                |
     |     Kernel 在 GPU 上执行计算      |
     |                                |
     |<------- cudaMemcpy D2H ---------|
     |     (Device to Host：GPU → CPU)  |
     |                                |
   free                            cudaFree

// 1. 在 GPU 上分配内存
// 参数 1：void** devPtr（指针的指针，类似 Java 里的 AtomicReference）
// 参数 2：size_t size（要分配的字节数）
cudaError_t cudaMalloc(void** devPtr, size_t size);

// 2. 在 CPU 和 GPU 之间拷贝数据
// 参数 1：目标地址
// 参数 2：源地址
// 参数 3：拷贝字节数
// 参数 4：方向（cudaMemcpyHostToDevice / cudaMemcpyDeviceToHost / cudaMemcpyDeviceToDevice）
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);

// 3. 释放 GPU 内存
cudaError_t cudaFree(void* devPtr);

#include <stdio.h>
#include <stdlib.h>

// ---------------------- Kernel 函数 ----------------------
// __global__ 表示这是 GPU 上执行的函数，由 CPU 调用
__global__ void vectorAdd(const float* A, const float* B, float* C, int numElements) {
    // 计算当前线程的全局唯一 ID
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    // 防止越界（因为线程数通常会向上取整到 Block 大小的整数倍）
    if (i < numElements) {
        C[i] = A[i] + B[i];
    }
}

// ---------------------- 主函数 ----------------------
int main() {
    int numElements = 50000;           // 向量长度
    size_t size = numElements * sizeof(float);  // 总字节数

    // ========== 1. 在 CPU 上分配内存（Host Memory） ==========
    float* h_A = (float*)malloc(size);
    float* h_B = (float*)malloc(size);
    float* h_C = (float*)malloc(size);

    // 初始化数据
    for (int i = 0; i < numElements; ++i) {
        h_A[i] = rand() / (float)RAND_MAX;
        h_B[i] = rand() / (float)RAND_MAX;
    }

    // ========== 2. 在 GPU 上分配内存（Device Memory） ==========
    float* d_A = NULL;
    float* d_B = NULL;
    float* d_C = NULL;

    // cudaMalloc 的第一个参数是 void**（指针的指针），所以要传 &d_A
    cudaMalloc((void**)&d_A, size);
    cudaMalloc((void**)&d_B, size);
    cudaMalloc((void**)&d_C, size);

    // ========== 3. 将数据从 CPU 拷贝到 GPU ==========
    // cudaMemcpyHostToDevice：Host(CPU) → Device(GPU)
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // ========== 4. 启动 Kernel ==========
    // 每个 Block 256 个线程
    int threadsPerBlock = 256;
    // Block 数量：向上取整，确保所有元素都被处理
    int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;

    printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid, threadsPerBlock);

    // 启动！<<<blocksPerGrid, threadsPerBlock>>>
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);

    // 检查 Kernel 是否启动成功
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        printf("Kernel launch failed: %s\n", cudaGetErrorString(err));
        return -1;
    }

    // ========== 5. 将结果从 GPU 拷贝回 CPU ==========
    // cudaMemcpyDeviceToHost：Device(GPU) → Host(CPU)
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // ========== 6. 验证结果 ==========
    for (int i = 0; i < numElements; ++i) {
        if (fabs(h_A[i] + h_B[i] - h_C[i]) > 1e-5) {
            printf("Verification failed at element %d!\n", i);
            return -1;
        }
    }
    printf("Test PASSED!\n");

    // ========== 7. 释放内存 ==========
    // GPU 内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // CPU 内存
    free(h_A);
    free(h_B);
    free(h_C);

    printf("Done!\n");
    return 0;
}

nvcc vector_add.cu -o vector_add
vector_add.exe    # Windows
./vector_add      # Linux

// Java 写法：一个 for 循环搞定
for (int i = 0; i < 50000; i++) {
    C[i] = A[i] + B[i];
}

// CUDA 写法：没有 for 循环！
// 每个线程负责一个 i，50000 个线程同时执行
int i = blockIdx.x * blockDim.x + threadIdx.x;
C[i] = A[i] + B[i];

// 假设有个方法：void allocate(GpuMemoryRef ref)
// 你需要传一个"能被修改的引用"进去
AtomicReference<float[]> d_A = new AtomicReference<>();
cudaMalloc(d_A);  // 方法内部会把新分配的 GPU 地址写回 d_A

#include <stdio.h>
#include <stdlib.h>

#define WIDTH 16  // 矩阵宽高都是 16

// 2D Kernel：每个线程负责计算结果矩阵的一个元素
__global__ void matrixAdd(const float* A, const float* B, float* C, int width) {
    // 计算当前线程负责的二维坐标
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < width && col < width) {
        int idx = row * width + col;  // 二维坐标转一维索引
        C[idx] = A[idx] + B[idx];
    }
}

int main() {
    int size = WIDTH * WIDTH * sizeof(float);

    // CPU 内存
    float* h_A = (float*)malloc(size);
    float* h_B = (float*)malloc(size);
    float* h_C = (float*)malloc(size);

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

    // GPU 内存
    float* d_A; cudaMalloc(&d_A, size);
    float* d_B; cudaMalloc(&d_B, size);
    float* d_C; cudaMalloc(&d_C, size);

    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // 2D 启动配置：
    // 每个 Block 是 4x4 = 16 个线程
    dim3 blockSize(4, 4);
    // Grid 是 4x4 = 16 个 Block，总共 256 个线程处理 16x16=256 个元素
    dim3 gridSize((WIDTH + 3) / 4, (WIDTH + 3) / 4);

    matrixAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, WIDTH);

    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // 验证：1.0 + 2.0 = 3.0
    printf("C[0] = %f, C[%d] = %f\n", h_C[0], WIDTH * WIDTH - 1, h_C[WIDTH * WIDTH - 1]);

    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);

    return 0;
}

矩阵 C（16x16）：

(0,0) (0,1) ... (0,15)
(1,0) (1,1) ... (1,15)
 ...
(15,0)        (15,15)

线程 (blockIdx.x=0, blockIdx.y=0, threadIdx.x=0, threadIdx.y=0) 计算 (0,0)
线程 (blockIdx.x=0, blockIdx.y=0, threadIdx.x=1, threadIdx.y=0) 计算 (0,1)
...
每个线程各算一个"格子"，全部并行！

#include <stdio.h>

// 封装错误检查宏
#define CHECK_CUDA(call)                                                      \
    do {                                                                      \
        cudaError_t err = call;                                               \
        if (err != cudaSuccess) {                                             \
            printf("CUDA error at %s:%d - %s\n", __FILE__, __LINE__,          \
                   cudaGetErrorString(err));                                  \
            exit(EXIT_FAILURE);                                               \
        }                                                                     \
    } while (0)

// 使用示例：
int main() {
    float* d_data;
    CHECK_CUDA(cudaMalloc(&d_data, 1024));
    CHECK_CUDA(cudaMemcpy(d_data, h_data, 1024, cudaMemcpyHostToDevice));

    myKernel<<<blocks, threads>>>(d_data);
    // Kernel 启动本身不返回错误，需要用 cudaGetLastError 捕获
    CHECK_CUDA(cudaGetLastError());
    CHECK_CUDA(cudaDeviceSynchronize());

    CHECK_CUDA(cudaFree(d_data));
    return 0;
}

Phase 1（现在）：理解本教程的所有代码，能手写向量加法和矩阵加法
        ↓
Phase 2：学习 Shared Memory，实现 Tiled Matrix Multiplication（性能提升 10x+）
        ↓
Phase 3：学习 CUDA 流（Stream）和事件（Event），理解异步执行与并发
        ↓
Phase 4：学习 cuBLAS / cuDNN / Thrust 等高级库，不要重复造轮子
        ↓
Phase 5：学习 TensorRT / Triton / CUTLASS，进入 AI 推理/训练优化深水区