CUDA 学习笔记CUDA编程学习笔记CUDA编程学习笔记CUDA编程学习笔记CUDA编程学习笔记CUDA编程学习笔记C

CUDA编程-01：搭建CUDA编程环境

nvidia-smi
nvcc --version # 租赁的服务器未安装 nvcc
apt install nvidia-cuda-toolkit
nvcc --version

CUDA编程-02：初识CUDA编程

Hello World

新建一个CUDA C源文件hello_world.cu

#include <stdio.h>

__global__ void helloFromGPU() {
    printf("Hello, World from GPU!\n");
}

int main() {
    printf("Hello from CPU!\n");
    helloFromGPU<<<1, 5>>>();
    cudaDeviceSynchronize();
    return 0;
}

用nvcc进行编译生成可执行文件：

nvcc hello_world.cu -o hello_world

root@085ffcdb1fa5:/data/xm/code/20251209# ./hello_world 
Hello from CPU!
Hello, World from GPU!
Hello, World from GPU!
Hello, World from GPU!
Hello, World from GPU!
Hello, World from GPU!

global

被__global__函数类型限定符修饰的函数被称为内核函数，该函数在host上被调用，在device上执行，只能返回void类型，不能作为类的成员函数。调用__global__修饰的函数是异步的，也就是说它未执行完就会返回。

<<<1, 5>>>

第一个数字 1：线程块（Block）的数量。这里启动 1个线程块。第二个数字 5：每个线程块内的线程（Thread）数量。这里每个线程块包含5个线程。这个配置总共启动了 1 × 5 = 5 个线程来并行执行 helloFromGPU 函数。

用CUDA实现数组相加

一个典型的CUDA程序结构一般由以下5个步骤组成：

分配GPU内存；
从CPU内存中拷贝数据到GPU内存中；
调用CUDA内核函数执行程序指定的计算任务；
从GPU内存中把数据拷贝回CPU内存中；
释放GPU内存；

CPU上实现数组相加的代码：

#include <iostream>

void VectorAddCPU(const float *const a, const float *const b, float *const c,
                  const int n) {
  for (int i = 0; i < n; ++i) {
    c[i] = a[i] + b[i];
  }
}

int main(void) {
  // alloc memory for host
  const size_t n = 2560;
  float *ha = new float[n]();
  float *hb = new float[n]();
  float *hc = new float[n]();

  for (int i = 0; i < n; ++i) {
    ha[i] = i;
    hb[i] = n - i;
  }

  VectorAddCPU(ha, hb, hc, n);
  std::cout << "Vector add result (first 10 elements):" << std::endl;
  for (int i = 0; i < 10; ++i) { // 只打印前10个，避免刷屏
      std::cout << "hc[" << i << "] = " << hc[i] << std::endl;
  } 

  delete[] ha;
  delete[] hb;
  delete[] hc;

  return 0;
}

函数VectorAddCPU用了一个for循环在CPU上实现数组相加的过程。如果要用GPU来实现该过程，则调用CUDA的API按照前面说的5个步骤编写代码：

#include <cuda_runtime.h>
#include <iostream>

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

int main(void) {
    // 分配CPU内存
    const size_t size = 10240;
    const size_t n = size / sizeof(float);
    float *ha = new float[n]();
    float *hb = new float[n]();
    float *hc = new float[n]();

    for (int i = 0; i < n; ++i) {
        ha[i] = i;
        hb[i] = n - i;
    }

    // 分配GPU内存
    float *da = nullptr;
    float *db = nullptr;
    float *dc = nullptr;
    cudaMalloc((void **)&da, size);
    cudaMalloc((void **)&db, size);
    cudaMalloc((void **)&dc, size);

    // 把数据从CPU拷贝到GPU
    cudaMemcpy(da, ha, size, cudaMemcpyHostToDevice);
    cudaMemcpy(db, hb, size, cudaMemcpyHostToDevice);
    cudaMemcpy(dc, hc, size, cudaMemcpyHostToDevice);

    const int thread_per_block = 256;
    const int block_per_grid = (n + thread_per_block - 1) / thread_per_block;

    // 调用核函数
    VectorAddGPU<<<block_per_grid, thread_per_block>>>(da, db, dc, n);

    // 把数据从GPU拷贝回CPU
    cudaMemcpy(hc, dc, size, cudaMemcpyDeviceToHost);
    std::cout << "Vector add result (first 10 elements):" << std::endl;
    for (int i = 0; i < 10; ++i) { // 只打印前10个，避免刷屏
        std::cout << "hc[" << i << "] = " << hc[i] << std::endl;
    }

    // 释放GPU内存
    cudaFree(da);
    cudaFree(db);
    cudaFree(dc);

    // 释放CPU内存
    delete[] ha;
    delete[] hb;
    delete[] hc;

    return 0;
}

cudaMalloc

函数原型：

cudaError_t cudaMalloc(void** devPtr, size_t size)

该函数用于在GPU上分配指定大小的内存空间，类似于标准C语言中的malloc函数。

cudaMemcpy

函数原型：

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind)

该函数用于内存拷贝，类似于标准C语言中的memcpy函数。数据拷贝的流向由参数kind指定，分为以下4种方式：

cudaMemcpyHostToHost
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice

从它们的字面意思就能知道，如果是从host拷贝数据到device，那么kind参数应该设置为cudaMemcpyHostToDevice；如果是从device拷贝数据到host，那么kind参数应该设置为cudaMemcpyDeviceToHost。上面的代码体现了这一点。

cudaFree

udaError_t cudaFree(void* devPtr)

该函数用于释放已分配的GPU内存空间，类似于标准C语言中的free函数。

CUDA编程-03:线程层级

blockIdx: 线程块在线程网格中的索引
threadIdx: 线程块内的线程索引
blockDim: 线程块的维度，用每个线程块中的线程数量来表示
gridDim: 线程网格的维度，用每个线程网格中的线程块数量来表示
blockIdx.x, blockIdx.y：当前线程所在的块在网格中的二维坐标（第几列块，第几行块）。
blockDim.x, blockDim.y：每个线程块的二维尺寸（一个块有多少列线程，多少行线程）。
threadIdx.x, threadIdx.y：当前线程在其所属块内的二维坐标（在块中是第几列，第几行）

const unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x; 
const unsigned int idy = blockDim.y * blockIdx.y + threadIdx.y; 
const unsigned int thread_id = (gridDim.x * blockDim.x) * idy + idx;

可以用一个生动的比喻来理解它 把GPU线程组织想象成一个巨大的“办公大楼”

整栋大楼 = 线程网格 (Grid) ：负责一个完整的计算任务。
每个楼层 = 线程块 (Block) ：楼层内的办公室可以紧密协作（共享资源）。
每个工位 = 线程 (Thread) ：真正干活的“员工”。

给这栋大楼里的每一位员工计算一个唯一的、连续的工号。

假设我们的“办公大楼”（线程网格）配置

配置项	值	说明
数据尺寸	6列 x 4行
块尺寸 (`blockDim`)	3线程宽 x 2线程高	每个“办公室”是3列2行
网格尺寸 (`gridDim`)	2块宽 x 2块高	整个“大楼”是2x2个办公室

那么，一个位于 blockIdx=(1,0), threadIdx=(0,1) 的线程（即：第1列第0行的块中，第0列第1行的线程），它的工号计算过程是：

// 已知：
// blockDim.x = 3, blockDim.y = 2
// gridDim.x = 2, gridDim.y = 2
// blockIdx.x = 1, blockIdx.y = 0
// threadIdx.x = 0, threadIdx.y = 1

const unsigned int idx = 3 * 1 + 0 = 3; // 全局第3列
const unsigned int idy = 2 * 0 + 1 = 1; // 全局第1行
const unsigned int thread_id = (2 * 3) * 1 + 3 = 6 * 1 + 3 = 9; // 唯一工号

可视化理解：整个网格被组织成如下布局，thread_id 就是每个位置的最终编号：

(网格Grid，共2x2个块，每个块3x2线程)

块(0,0)              块(1,0)
+-----------+        +-----------+
| 0  1  2   |        | 3  4  5   | <- idy=0 行
| 6  7  8   |        | 9 10 11   | <- idy=1 行 (我们的线程在这里！它是9号)
+-----------+        +-----------+

块(0,1)              块(1,1)
+-----------+        +-----------+
|12 13 14   |        |15 16 17   | <- idy=2 行
|18 19 20   |        |21 22 23   | <- idy=3 行
+-----------+        +-----------+

CUDA 学习笔记

CUDA编程-01： 搭建CUDA编程环境

CUDA编程-02： 初识CUDA编程

Hello World

用CUDA实现数组相加

CUDA编程-03:线程层级

CUDA编程-04：CUDA内存模型

CUDA编程-01：搭建CUDA编程环境

CUDA编程-02：初识CUDA编程