CUDA

一、CUDA简介

1.0 Clion远程配置

新建项目，左边选择CUDA Executable，C++标准选择11，服务器cmake版本应为3.28.x。
在设置Build, Execution, Deployment中：
- Toolchains：增加一个Remote Host，填写右边的Credentials。
- Deployment：在Mappings中添加服务器端对应的部署路径。
- CMake：选择刚才配置的远程主机的Toolchain，CMake Options添加
```
-DCMAKE_CUDA_COMPILER:PATH=/usr/local/cuda/bin/nvcc
```

1.1 主机与设备相关函数

在CUDA程序中，主机（如CPU）称为Host，设备（如GPU）称为Device。下面先介绍一些基本的函数：

在设备上分配内存空间。
```
cudaError_t cudaMalloc(void **devPtr, size_t size)
```
- devPtr：指向设备内存的指针。该指针主机不能解析，只能在设备中使用。
- size：以字节为单位分配的内存大小。
在设备上释放内存空间。参数意义同上。
```
cudaError_t cudaFree(void *devPtr)
```
在主机和设备之间复制内存数据。
```
cudaError_t cudaMemcpy(
    void *dst, 
    const void *src, 
    size_t count, 
    cudaMemcpyKind kind)
```
- dst：目标内存地址。
- src：源内存地址。
- count：复制的字节数。
- kind：传输类型，可选值有：cudaMemcpyHostToHost，cudaMemcpyHostToDevice，cudaMemcpyDeviceToHost，cudaMemcpyDeviceToDevice。

上面的函数是有返回值的，完整的写法应该加上错误处理：

cudaError_t err = cudaMalloc((void**)&d_A, sizeof(float) * N);
if (err != cudaSuccess) {
    printf("%s in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__);
    exit(EXIT_FAILURE);
}

1.2 kernel函数

指在设备上执行的函数。声明时需要在前面加上__global__关键字。其他CUDA扩展关键字有：

关键字	执行位置	调用位置
`__device__`	设备	设备
`__global__`	设备	主机
`__host__`（默认）	主机	主机

启动kernel函数时需要使用<<<grid_size, block_size>>>指出配置参数，第1个参数表示网格中线程块的数量，第2个参数表示线程块中的线程数。如：

vecAdd<<<1, 512>>>(d_A, d_B, d_C, N);

1.3 向量加法完整示例

#include <cstdio>
#include "cuda_runtime.h"

__global__
void vecAdd(float *A, float *B, float *C, int n) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

int main() {
    const int N = 100;
    float h_A[N], h_B[N], h_C[N];
    float *d_A, *d_B, *d_C;

    for (int i = 0; i < N; i++) {
        h_A[i] = i;
        h_B[i] = N - i;
    }

    cudaMalloc((void**)&d_A, sizeof(float) * N);
    cudaMalloc((void**)&d_B, sizeof(float) * N);
    cudaMalloc((void**)&d_C, sizeof(float) * N);
    cudaMemcpy(d_A, h_A, sizeof(float) * N, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, sizeof(float) * N, cudaMemcpyHostToDevice);
    cudaMemcpy(d_C, h_C, sizeof(float) * N, cudaMemcpyHostToDevice);

    vecAdd<<<1, 512>>>(d_A, d_B, d_C, N);

    cudaMemcpy(h_C, d_C, sizeof(float) * N, cudaMemcpyDeviceToHost);

    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    for (int i = 0; i < N; i++) {
        printf("%f\n", h_C[i]);
    }
    return 0;
}

二、数据并行执行模型

2.1 线程组织

一个kernel函数会创建一个网格（Grid），grid中的所有线程都执行这个kernel函数。一个grid包含多个线程块（Block），每个block又包含多个线程（Thread）。

下表是kernel函数预设的一些内置变量，它们都有x,y,z三个属性，x,y,z应当理解为元素在x方向、y方向、z方向上的个数（而不是x行y列）。

内置变量	含义
`blockIdx`	当前块的索引
`threadIdx`	当前线程的索引
`gridDim`	网格中线程块的维度
`blockDim`	线程块中线程的维度

grid是由block组成的三维数组，block是由thread组成的三维数组，用不到的维度置1即可。dim3是结构体，结构体名字随意，然后传入<<<,>>>中。

dim3 gridDim(2, 2, 1);
dim3 blockDim(256, 1, 1);
vecAdd<<<gridDim, blockDim>>>(d_A, d_B, d_C, N);

// 直接写的话相当于gridDim和blockDim的yz维度都置1了
vecAdd<<<4, 256>>>(d_A, d_B, d_C, N);

2.2 矩阵乘法示例

在向kernel函数传递参数时不能直接传数组，对于多维数组应当进行扁平化处理（对于C来说，二维数组的存储是行优先的）。

__global__
void matrix_mult(float *A, float *B, float *C, int W) {
    // y方向是第几行
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    // x方向是第几列
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < W && col < W) {
        float sum = 0.0;
        for (int k = 0; k < W; k++) {
            // C[row][col] = A[row][k] * B[k][col]
            sum += A[row * W + k] * B[k * W + col];
        }
        C[row * W + col] = sum;
    }
}

2.3 资源分配与调度

GPU划分成很多SM（Streaming Multiprocessor，流式多处理器），SM内部又有很多SP（streaming processor）或称作CUDA Core，每个SM中还有warp scheduler、shared memory、register等。SM能分配的线程块数量有限制，能调度的线程数也有限制。

warp是调度和运行的基本单元，Nvidia把32个线程组成一个warp，warp中所有的线程并行执行相同的指令。warp由SM的硬件warp scheduler负责调度，一个SM同一个时刻可以执行多个warp，取决于warp scheduler的数量。warp调度可以隐藏线程切换的开销，当一个warp阻塞时，调度器会调取其他的warp来执行，达到“零开销线程调度”。

CUDA编程模型与GPU的映射关系如上图：

一个Grid可以包括多个SM，也可以访问Global Memory和Constant Memory。
一个Block只能在一个SM中，一个SM包含多个Block，Block可以访问Shared Memory。
一个Block中有多个Thread，Thread只能访问Registers。

2.4 查询设备属性

int dev_count;
cudaDeviceProp dev_prop;
cudaGetDeviceCount(&dev_count);
for (int i = 0; i < dev_count; i++) {
    cudaGetDeviceProperties(&dev_prop, i);
    printf("Device: %d, %s\n", i, dev_prop.name);
    printf("MultiProcessor Count: %d\n", dev_prop.multiProcessorCount);
    printf("Max Blocks Per MultiProcessor: %d\n", dev_prop.maxBlocksPerMultiProcessor);
    printf("Max Threads Per MultiProcessor: %d\n", dev_prop.maxThreadsPerMultiProcessor);
    printf("Regs Per Multiprocessor: %d\n", dev_prop.regsPerMultiprocessor);
    printf("SharedMem Per Multiprocessor: %ld\n", dev_prop.sharedMemPerMultiprocessor);
    printf("Max Threads Per Block: %d\n", dev_prop.maxThreadsPerBlock);
    printf("Max Threads Dim: [%d, %d, %d]\n", dev_prop.maxThreadsDim[0], dev_prop.maxThreadsDim[1], dev_prop.maxThreadsDim[2]);
    printf("Warp Size: %d\n", dev_prop.warpSize);
}

示例结果如下：

Device: 0, Tesla M40
MultiProcessor Count: 24
Max Blocks Per MultiProcessor: 32
Max Threads Per MultiProcessor: 2048
Regs Per Multiprocessor: 65536
SharedMem Per Multiprocessor: 98304
Max Threads Per Block: 1024
Max Threads Dim: [1024, 1024, 64]
Warp Size: 32

程序设计

#CUDA

CUDA

https://shuusui.site/blog/2025/03/15/cuda/

作者

Shuusui

发布于

2025年3月15日

更新于

2025年3月16日

许可协议

近世代数（一）基本概念上一篇

数字电子技术基础下一篇