从零入门 cuda 编程？🦴 看这篇就够了！

实录：从零开始接触 cuda 编程 🐟

CUDA（全称 Compute Unified Device Architecture，中文译为“统一计算设备架构”），由 nvidia 在 2007 年发布，提供用于编程和管理 GPU 的 C/C++ 语言扩展和 API，服务于 GPU 通用计算，时至今日已历经十几年演进。本博客记录对 CUDA 的学习过程，重点内容包括：

  1. CUDA 的架构详述，包括：硬件、软件和各种常见术语。
  2. CUDA 的安装、编译工具和编辑器的配置。
  3. CUDA 的基础语法：函数标识符、变量标识符、内置向量类型和内置变量、执行配置。
  4. CUDA 的 Runtime API：内存和线程管理、设备管理、事件管理、流管理。
  5. CUDA 编程实例：测试 GPU 的乘加性能。

详解 CUDA 架构：软件层和硬件层相结合

在 CUDA 中，GPU 作为 CPU 的协处理器工作。通常，我们将 CPU 和 GPU 系统分别称为 主机(host) 和 设备(device) ，它们是具有各自内存空间的独立平台 📚。一般地，我们在 CPU 上运行串行工作负载，并将并行计算卸载到 GPU 上。

CUDA 架构：软件层和硬件层

CUDA 的架构包括软件层和硬件层，如上图所示。软件层包括：线程 (thread)，线程块 (block) 和 网格 (grid) 。硬件层包括：CUDA 核心 (或称流处理器，也即 SP)，流式多处理器 (SM, streaming multiprocessor) 和 GPU。除这些以外，还有 线程束 (warp) ，它充当沟通硬件层和软件层的桥梁。

线程和 CUDA 核心：最小的基本运算单位

线程 (thread) 🧶

线程是对计算机中 程序执行流 的形象比喻，它由 指令流 和 数据流 交织构成，是 GPU/CPU 通用的原子级概念。在 CUDA 中，线程按束调度： 同束线程共享相同的程序计数器 ，它们同步地执行相同的指令，作用于各自寄存器存储的数据。

CUDA 核心 (SP)

CUDA 核心是 CUDA 中执行标量运算指令的基本单元，其核心组件是 整数运算单元 (INT) 和 浮点数运算单元 (FP) 。CUDA 核心和线程相对应：单个 CUDA 核心执行来自单个线程的指令 。与线程相同，CUDA 核心同样按束调度： 同束核心在同一时刻执行相同的指令 ，但作业于不同的寄存器。

CUDA 核心

单指令多线程 (SIMT, Single Instruction MultiThreads)

🪶 这里，我们特别指出：上述线程和 CUDA 核心所采用的“ 多个线程在同一时刻执行相同的指令，但作用于不同的数据 ”的并行方式，即所谓的 单指令多线程 ，这也是 CUDA 所使用的方式。与 SIMT 相对的是 SIMD (Single Instruction MultiData, 单指令多数据)，它更像是 vector 架构。

线程块、网格和 Kernel：CUDA 的概念模型

线程块 (block) 🧊 和 网格 (grid) 🥅

线程块是对许多线程的概念抽象，它将数量繁多的线程按照 1 维、2 维或 3 维的方式组织，为开发者遍历、索引线程提供了极大方便。比线程块更高级的概念是网格，它将许多个线程块按照 1 维、2 维或 3 维的方式排列。由此，我们得到了 网格-线程块-线程 的线程层级概念划分。

线程坐标与索引的转换

譬如下图左半部分，对于大小为 $(D_x, D_y)$ 二维线程块，其中索引为 $(x, y)$ 的线程的 块内线程 ID 应当是 $x + yD_x$ 。然而，由于线程块的维度和网格的维度能以各种方式组合，在不同组合下获取线程的全局 ID 或者块内 ID 仍然是一个复杂的问题。

CUDA 内存的读写规则

CUDA 允许设备上的线程以 寻址、共享、缓冲 的方式读取 DRAM 和 On-Chip 内存。更具体的，如下图右半部分，执行在设备上的线程，只允许按如下方式读写内存：1. 读写每条线程的寄存器和本地内存。 2. 读写每个块的共享内存。 3. 读写每个网格的全局内存。 4. 只读每个网格的常量内存和纹理内存。

Kernel 批处理（左）与 CUDA 内存模型（右）

Kernel (核函数) 🥜

特指由主机调用，在设备上运行的函数，它指示了网格内所有线程的行为。 程序运行时，主机向设备连续地发送 kernel 调用的请求，每个 kernel 就作为一个由线程块组成的线程批处理来执行，如上图左半部分所示。单个 kernel 可能由多个线程块执行，线程块内的线程将 共享某块内存，并在必要时同步。

流式多处理器：并行计算的幕后功臣

SM (流式多处理器，Streaming Multiprocessor)

以 Fermi 为例，如下图所示，单个完整的 SM 除了若干个 SP 还应该包括：指令缓存、线程束调度器和分派单元、寄存器文件、加载/存储单元队列、特殊功能单元队列、共享内存/L1 缓存、统一缓存等。它们的作用听名字都能猜出来，感兴趣可以自己去查一下，这里不再赘述 🐈🐕。

流式多处理器 SM 架构图

线程束：连接软件层和硬件层的纽带

线程束 (warp) 🧵

线程束是真正执行并行操作的单位，通常以 32 个线程为一束。 线程束是连接软、硬件层的纽带，软件层面开发者定义的线程块，通过线程束这一执行单元映射到硬件层面的 SM 上。具体来说，每个线程块在加载到 SM 时，会被划分为若干束，而这些束正是 SM 的实际调度单位。

线程束的并行与同步

同束线程共享内存，在运行时保持同步。束与束之间则保持并行，仅在必要时同步。 束内线程仅 32 个，因此同步的时间开销是可接受的。 线程束的共享内存实际上作为缓存，用于缓存本束的计算结果，然后再上传。 对所有束的计算结果同步汇总的时间也是可接受的。由此，也能知道 CUDA 计算的时间：

$$总时间成本 = 核心计算的时间 + 束内线程同步的时间 + 对各个束进行同步的时间$$

线程束的分化

同束线程执行不同指令就叫做线程束分化。 同束线程执行相同的指令，但处理各自的数据。若它们在执行时遇不同的控制条件，就会进行不同的选择，导致线程束分化。分化期间，条件为真的线程将执行指令，为假的线程则空等（stall execution）。线程束分化严重影响性能。条件分支越多，并行性削弱越严重。因此，应尽量避免同束内线程分化，确保线程分配到线程束是有规律的。

线程束分化示意图

正式进入 CUDA 编程：安装、编译器、编辑器

我们花了大量的时间和篇幅详述 CUDA 的软件层面和硬件层面的架构，这有助于我们理解接下来的内容。CUDA 的编程主要围绕 Kernel 展开，我们将会逐步深入对它的理解。在此之前，先下载好 CUDA 的编译器 nvcc。

下载 CUDA 库及编译工具 nvcc

1
2
3
4
5
6
# 安装 cuda
sudo pacman -S cuda

# 验证安装
nvcc --version # 检查 CUDA 编译器
nvidia-smi     # 检查 GPU 驱动状态

配置环境变量

将以下内容添加到 ~/.bashrc 或 ~/.zshrc：

1
2
export PATH="/opt/cuda/bin:$PATH"
export LD_LIBRARY_PATH="/opt/cuda/lib64:$LD_LIBRARY_PATH"

然后执行 source ~/.bashrc 或重新登录。

Neovim 启用 CUDA 语法检查

1
2
3
4
5
6
7
8
9
10
11
12
13
require("lspconfig").clangd.setup({
    capabilities = require("cmp_nvim_lsp").default_capabilities(),
    cmd = {
        "clangd",
        "--background-index",
        "--clang-tidy",
        "--header-insertion=iwyu",
        "--completion-style=detailed",
        "--query-driver=/usr/local/cuda/bin/nvcc", -- 指定 CUDA 编译器路径
        "--offset-encoding=utf-16",
    },
    filetypes = { "c", "cpp", "cuda", "objc", "objcpp" }, -- 添加 cuda 类型
})

基础的 CUDA 语法：C 语言的迷你扩展集

让我们从 Hello World 程序开始，简单比较一下 C 和 CUDA，理解下 CUDA 在实际编程上的特点。

C/C++

1
2
3
4
5
6
7
8
void c_hello(){
    printf("Hello World!\n");
}

int main() {
    c_hello();
    return 0;
}

CUDA

1
2
3
4
5
6
7
8
__global__ void cuda_hello(){
    printf("Hello World from GPU!\n");
}

int main() {
    cuda_hello<<<1,1>>>(); 
    return 0;
}

可以看到主要区别有两点：一是函数声明/定义时使用 __global__，我们将其称作 函数执行空间标识符。二是调用函数时使用 <<...>>，我们将其称作 执行配置。实际上，CUDA 对 C 的扩展可归纳为四点：1. 函数执行空间标识符；2. 变量内存空间标识符；3. 内置向量类型和内置变量；4. 执行配置。

函数执行空间标识符 (Function Execution Space Specifier)

__global__

  1. 声明一个函数为 kernel，它在设备上执行，从主机调用。
  2. __global__ 函数必须返回 void 类型，且不能作为类成员方法。
  3. 调用 __global__ 函数必须给定执行配置。
  4. __global__ 的调用是异步的，这意味着它会在设备执行完毕前返回。

__device__

  1. 声明一个函数，在设备上执行，也仅可以从设备上调用。
  2. 它不能与 __global__ 标识符同时使用。

__host__

  1. 声明一个函数，在主机上执行，也仅可以从主机上调用。
  2. 通常该标识符可以省略不写，默认是主机调用并执行函数。
  3. 不能与 __global__ 标识符同时使用。
  4. 当与 __device__ 标识符同时使用时，会同时为主机和设备编译该函数。

Undefined behavior (未定义行为)
以下函数调用方式会被编译器解释为未定义行为：

  1. 在 __global__, __device__ 或 __host__ __device__ 内调用 __host__ 函数；
  2. 在 __host__ 函数内调用 __device__ 函数。

其他注意事项⚠️

  1. __device__ 和 __global__ 函数不支持递归。
  2. 不能在 __device__ 和 __global__ 函数中声明静态变量。
  3. __device__ 和 __global__ 函数不能有自变量的一个变量数字。
  4. 无法获得 __device__ 函数的地址，但指向 __global__ 是合法的。

变量内存空间标识符 (Variable Memory Space Specifiers)

默认情况下，未指定标识符的变量会驻留在寄存器上。然而，部分编译器会将其置于本地内存，这会影响并行性能。因此，应当尽可能指定变量的内存空间标识符。

__device__

  1. 声明一个驻留在设备内存的变量。
  2. 可与至多 1 个其他的变量内存标识符组合使用。
  3. 默认行为 (当它单独使用时)：
    • 位于全局内存空间。
    • 生命周期与创建它的 CUDA 上下文相同。
    • 每个设备有其独立的副本。
    • 可通过运行时 API 被网格内所有线程和主机访问。
  4. 适合存储需要全局访问的大规模数据。

__constant__

  1. 必须与 __device__ 组合使用。
  2. 驻留在常量内存空间。
  3. 生命周期与创建它的 CUDA 上下文相同。
  4. 每个设备有其独立的副本。
  5. 可通过运行时 API 被网格内所有线程和主机访问。
  6. 主机在并发内核访问时修改常量会导致未定义行为。
  7. 适合存储只读数据。

__shared__

  1. 必须与 __device__ 组合使用。
  2. 驻留在线程块的共享内存。
  3. 生命周期与线程块相同。
  4. 每个线程块有独立副本。
  5. 仅块内线程可访问，地址不固定。

在 CUDA 中，通过 extern 关键字，可使数组大小在内核启动时由执行配置 <<<...>>> 动态指定。以这种方式声明的所有变量都从内存中的相同地址开始，因此必须通过偏移量显式地管理数组中变量的布局。例如，在动态分配的共享内存中，可以用以下方式声明和初始化数组：

1
2
3
4
5
6
extern __shared__ float array[];
__device__ void func() {     // __device__ or __global__ function
    short* array0 = (short*)array;     // 存储 short
    float* array1 = (float*)&array0[128]; // 接在 128 个 short 之后，存储 float
    int* array2 = (int*)&array1[64];    // 接在 64 个 float 之后，存储 int
}

注意，指针需与其所指向类型对齐，故而以下代码不能工作，因为 array1 没有对齐到声明 array 时所指定的 float 的 4 字节。

1
2
3
4
5
extern __shared__ float array[];
__device__ void func() {     // __device__ or __global__ function
    short* array0 = (short*)array;
    float* array1 = (float*)&array0[127];   // 错误！127*2 字节不是 4 字节对齐
}

内置向量类型和内置变量

内置向量类型 (Built-in Vector Types)

这些是从基本整数和浮点类型派生出来的向量类型。它们是结构体，其第 1、2、3 和 4 个分量分别可以通过字段 x、y、z 和 w 来访问。它们都带有一个构造函数，形式为 make_<type name>；例如，

1
int2 make_int2(int x, int y);

该构造函数会创建一个值为 (x, y) 的 int2 类型的向量。向量类型的对齐要求在右表中有详细说明。

类型对齐
char1, uchar11
char2, uchar22
char3, uchar31
char4, uchar44
short1, ushort12
short2, ushort24
short3, ushort32
short4, ushort48
int1, uint14
int2, uint28
int3, uint34
int4, uint416
long2, ulong28 or 16
long3, ulong34 or 8
long4, ulong416
longlong1, ulonglong18
longlong2, ulonglong216
longlong3, ulonglong38
longlong4, ulonglong416
float14
float28
float34
float416
double18
double216
double38
double416
long1, ulong14 or 8

除上面常用的向量类型，CUDA 还提供了名为 dim3 的向量类型。该类型是一个基于 uint3 的整数向量类型，用于指定线程块和网格的尺寸，在后面讲执行配置时我们会看到它的作用。在定义类型为 dim3 的变量时，任何未指定的组件都初始化为 1。

内置变量 (Built-in Variables)

  • gridDim ：dim3，表示网格的维度 (gridDim.x, gridDim.y, gridDim.z)。
  • blockIdx ：uint3，表示网格中块的坐标 (blockIdx.x, blockIdx.y, blockIdx.z)。
  • blockDim ：dim3，表示线程块的维度 (blockDim.x, blockDim.y, blockDim.z)。
  • threadIdx：uint3，表示线程在块中的坐标 (threadIdx.x, threadIdx.y, threadIdx.z)。
  • warpSize ：int ，包含了以线程为单位的线程数的大小。

执行配置 (Execution Configuration)

调用 __global__ 函数必须指定执行配置。执行配置定义了将在设备上执行该函数的网格和块的维度，以及相关的 stream。执行配置通过在函数名和括号参数列表之间插入 <<< Dg, Db, Ns, S >>> 形式的表达式来指定，其中：

  • Dg 是 dim3 类型，指定网格的维度大小，使得启动的线程块的块数为 Dg.x * Dg.y * Dg.z。
  • Db 同样是 dim3 类型，指定每个块的维度大小，使得每个块具有 Db.x * Db.y * Db.z 个线程。
  • Ns 是 size_t 类型，指定每个块为此调用 动态分配的共享内存字节数 。此动态分配的内存用于任何声明为 extern 数组的变量（如 __shared__ 中所述）；Ns 是可选参数，默认为 0。
  • S 是 cudaStream_t 类型，指定相关的流；S 是一个可选参数，默认值为 0。

例如，声明为以下形式的函数，必须像这样调用：

1
2
3
__global__ void Func(float* parameter); // __global__ 声明其为 kernel 函数

Func<<< Dg, Db, Ns >>>(parameter);  // 调用时必须指定执行配置
  • 执行配置的参数在实际函数参数之前被评估，以下两种情况将导致函数调用失败：
    1. Dg 或 Db 大于设备允许的最大大小（如计算能力中所指定）；
    2. Ns 大于设备上可用的最大共享内存量减去静态分配所需的共享内存量。

进阶：部分常用的 Runtime API

内存与线程管理

cudaMalloc()

cudaError_t cudaMalloc(void** devPtr, size_t count);

  • 在设备上分配 count 字节的线性内存，并将指向该部分内存的指针交给 devPtr。
  • 分配的内存适合任何类型的变量。
  • 如果分配失败，返回 cudaErrorAlloction。

cudaFree()

cudaError_t cudaFree(void* devPtr);

  • 释放被 devPtr 指向的内存空间。
  • devPtr 必须被 cudaMalloc() 赋值过，否则会引发错误。
  • cudaFree() 可以重复调用，这将不会执行任何操作。

cudaMemcpy()

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

cudaError_t cudaMemcpyAsync(void* dst,constvoid*src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream);

  • 拷贝 count 字节，从 src 指向的内存区域到 dst 指向的内存区域。
  • kind 可以是 cudaMemcpyHostToHost， cudaMemcpyHostToDevice，cudaMemcpyDeviceToHost，或 cudaMemcpyDeviceToDevice 的拷贝方向。
  • cudaMemcpyAsync() 是异步的，并且可以作为一个可选参数通过流使用。

cudaDeviceSynchronize()

cudaError_t cudaDeviceSynchronize(void);

  • 阻塞，直到设备上所有请求的任务执行完毕。
  • 若任何一个任务失败，cudaThreadSynchronize() 都将返回一个错误。

示例：管理设备内存和等待设备执行

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
# include <stdio.h>

int main() {
    // 声明 host data 和 device data
    // 大小为 10 个 float
    float *h_data, *d_data;
    size_t size = 10 * sizeof(float);

    // 初始化 host data
    h_data = (float*)malloc(size)
    for (int i = 0; i < size; i++) {
        h_data[i] = 3.14 + i % 10;
    }

    cudaMalloc(&d_data, size); // 1. 为 device data 开辟空间
    cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice); // 2. 将 host data 拷贝至 device data

    myKernel<<<...>>>(); // 3. 执行核函数

    cudaDeviceSynchronize(); // 4. 阻塞，直到设备执行完毕
}

设备管理 (Device)

cudaGetDeviceCount()

cudaError_t cudaGetDeviceCount(int* count);

  • 将计算兼容性大于等于 1.0 的设备数量赋值到 count 所指向的 int 变量。
  • 如果没有相关设备，返回 1。

cudaSetDevice()

cudaError_t cudaSetDevice(int dev);

  • 指示当前所处的主机线程，使用代码为 dev 的设备。

cudaGetDevice()

cudaError_t cudaGetDevice(int* dev);

  • 将当前主机线程所用设备的代码赋值到 dev 所指向的 int 变量。

cudaGetDeviceProperties()

cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp* prop, int dev);

  • 将代码为 dev 的设备的属性赋值到 prop 所指向的 cudaDeviceProp 结构体。
  • cudaDeviceProp 结构体的定义可参阅 nvidia 文档

示例：获悉设备属性

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
#include <stdio.h>

int main() {
  int device_count;
  cudaGetDeviceCount(&device_count);  // 1. 获取设备数量
  printf("Found %d CUDA devices\n", device_count);

  cudaSetDevice(0);  // 2. 选择 0 号设备

  int dev;
  cudaGetDevice(&int);  // 3. 获取当前的设备编号

  cudaDeviceProp prop;
  cudaGetDeviceProperties(&prop, 0);  // 4. 获取设备属性
  printf("Device Name: %s\n", prop.name);
  printf("Compute Capability: %d.%d\n", prop.major, prop.minor);
  
  return 0;
}

事件管理 (Event)

cudaEventCreate()

cudaError_t cudaEventCreate(cudaEvent_t* eventPtr);

  • 创建一个事件对象，它具有 cudaEvent_t 类型，被 eventPtr 指针所指向。
  • cudaEvent_t 类型的定义可参考 nvidia 文档

cudaEventRecord()

cudaError_t cudaEventRecord(cudaEvent_t event, CUstream stream);

  • 记录事件 event，注意此处的 event 并非指针，而是 cudaEvent_t 类型。
  • 如果stream 是非零的，当流中所有的操作完毕，事件被记录。
  • 若 stream 为空，当CUDA context 中所有的操作完毕，事件被记录。
  • 该操作是异步的，必须使用 cudaEventSyncronize() 来确保事件真的被记录。

cudaEventSyncronize()

cudaError_t cudaEventSyncronize(cudaEvent_t event);

  • 阻塞线程，直到事件 event 真的被记录。
  • 如果 cudaEventRecord() 从未被调用过，则返回 cudaErrorInvalidValue。

cudaEventElapsedTime()

cudaError_t cudaEventElapsedTime(float* time, cudaEvent_t start, cudaEvent_t end);

  • 计算两个事件 start 和 end 之间所花费的时间，并赋值到 time 所指的 float 变量。
  • start 和 end 必须都被调用过 cudaEventRecord()，并且已经真的被记录了。
  • 否则，将返回 cudaErrorInvalidValue 错误。
  • 此外，若 start 和 end 中有被带非零 stream 记录的事件，结果将会是 undefined。

cudaEventDestroy()

cudaError_t cudaEventDestroy(cudaEvent_t event);

  • 销毁事件对象 event。

示例：测量内核执行时间

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
// 1. 创建事件对象
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start); // 2. 记录起始时间
myKernel<<<...>>>();  // 执行内核
cudaEventRecord(stop); // 3. 记录终止时间

cudaEventSynchronize(stop); // 4. 阻塞，直到 stop 被记录
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop); // 5. 计算时间

// 6. 销毁事件对象
cudaEventDestroy(start);
cudaEventDestroy(stop);

流管理 (stream)

cudaStreamCreate()

cudaError_t cudaStreamCreate(cudaStream_t* streamPtr);

  • 创建一个流对象，它是 cudaStream_t 类型，被 streamPtr 所指向。
  • cudaStream_t 类型的定义可参考 nvidia 文档

cudaStreamSyncronize()

cudaError_t cudaStreamSyncronize(cudaStream_t stream);

  • 阻塞当前进程，直到设备完成流 stream 中的所有操作。

cudaStreamDestroy()

cudaError_t cudaStreamDestroy(cudaStream_t stream);

  • 销毁一个流对象。

示例：异步数据传输

1
2
3
4
5
6
7
8
9
10
11
12
// 1. 创建流对象
cudaStream_t stream;
cudaStreamCreate(&stream);

// 异步内存拷贝
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream);

// 2. 异步执行内核
myKernel<<<grid, block, 0, stream>>>();

cudaStreamSynchronize(stream);  // 3. 阻塞，直到流完成
cudaStreamDestroy(stream); // 4. 销毁流对象

CUDA 编程实例：测试 GPU 的乘加性能

利用 cuda 编程测试 GPU 进行乘加浮点运算 c+=a*b 的性能，将其量化成 GFLOPS 指标 (GFLOPS, Giga FLoating-point Operations Per Second，每秒 10 亿次的浮点运算数) ，并呈现 GPU 的部分属性。拟对 $2^{24}$ 个元素进行计算，迭代 100 次。

Cuda 源代码

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
#include <stdio.h>
#include <cuda_runtime.h>

#define CHECK(cmd) { \
	cudaError_t error = cmd; \
	if (error != cudaSuccess) { \
		printf("Error: %s:%d, ", __FILE__, __LINE__); \
		printf("code:%d, reason:%s\n", error, cudaGetErrorString(error)); \
		exit(1); \
	} \
}

__global__ void multiplyAddKernel(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]; // 乘加操作：c += a * b
	}
}

int main() {
	// 获取GPU属性
	cudaDeviceProp prop;
	CHECK(cudaGetDeviceProperties(&prop, 0));
	
	// 打印GPU属性
	printf("GPU 属性:\n");
	printf("设备名称: %s\n", prop.name);
	printf("计算能力: %d.%d\n", prop.major, prop.minor);
	printf("SM 数量: %d\n", prop.multiProcessorCount);
	printf("全局内存: %.2f GB\n", prop.totalGlobalMem / (1024.0 * 1024.0 * 1024.0));
	printf("每块最大线程数: %d\n", prop.maxThreadsPerBlock);
	printf("每线程块共享内存: %zu KB\n\n", prop.sharedMemPerBlock / 1024);

	// 设置数据量
	const int N = 1 << 24; // 16,777,216 个元素
	const size_t size = N * sizeof(float);
	const int blockSize = 256;
	const int gridSize = (N + blockSize - 1) / blockSize;
	const int iterations = 100; // 重复执行次数

	// 分配主机内存并初始化
	float *h_a = (float*)malloc(size);
	float *h_b = (float*)malloc(size);
	float *h_c = (float*)malloc(size);
	for (int i = 0; i < N; ++i) {
		h_a[i] = 2.0f;
		h_b[i] = 0.5f;
		h_c[i] = 0.0f;
	}

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

	// 拷贝数据到设备
	CHECK(cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice));
	CHECK(cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice));
	CHECK(cudaMemcpy(d_c, h_c, size, cudaMemcpyHostToDevice));

	// 预热运行
	for (int i = 0; i < 5; ++i) {
		multiplyAddKernel<<<gridSize, blockSize>>>(d_a, d_b, d_c, N);
		CHECK(cudaGetLastError());
	}
	CHECK(cudaDeviceSynchronize());

	// 创建计时事件
	cudaEvent_t start, stop;
	CHECK(cudaEventCreate(&start));
	CHECK(cudaEventCreate(&stop));

	// 重置结果为初始值
	CHECK(cudaMemcpy(d_c, h_c, size, cudaMemcpyHostToDevice));

	// 执行并计时
	CHECK(cudaEventRecord(start));
	for (int i = 0; i < iterations; ++i) {
		multiplyAddKernel<<<gridSize, blockSize>>>(d_a, d_b, d_c, N);
		CHECK(cudaGetLastError());
	}
	CHECK(cudaEventRecord(stop));
	CHECK(cudaEventSynchronize(stop));

	// 计算时间
	float milliseconds;
	CHECK(cudaEventElapsedTime(&milliseconds, start, stop));
	double seconds = milliseconds / 1000.0;

	// 计算性能指标
	double totalFlops = 2.0 * N * iterations; // 每次迭代每个元素2次浮点操作
	double gflops = totalFlops / seconds / 1e9;

	printf("性能指标:\n");
	printf("数据量: %d 个元素\n", N);
	printf("总运算量: %.2f GFLOP\n", totalFlops / 1e9);
	printf("总耗时: %.3f ms\n", milliseconds);
	printf("平均性能: %.2f GFLOPS\n\n", gflops);

	// 验证结果正确性
	CHECK(cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost));
	printf("验证前5个结果:\n");
	for (int i = 0; i < 5; ++i) {
		printf("c[%d] = %.1f (预期值: %d)\n", i, h_c[i], iterations);
	}

	// 释放资源
	free(h_a);
	free(h_b);
	free(h_c);
	CHECK(cudaFree(d_a));
	CHECK(cudaFree(d_b));
	CHECK(cudaFree(d_c));
	CHECK(cudaEventDestroy(start));
	CHECK(cudaEventDestroy(stop));

	return 0;
}

编译及输出

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
$ nvcc -ccbin g++ -m64 -gencode arch=compute_35,code=sm_35 a.cu

$ ./a.out
GPU 属性:
设备名称: Tesla K20c
计算能力: 3.5
SM 数量: 13
全局内存: 4.63 GB
每块最大线程数: 1024
每线程块共享内存: 48 KB

性能指标:
数据量: 16777216 个元素
总运算量: 3.36 GFLOP
总耗时: 181.626 ms
平均性能: 18.47 GFLOPS

验证前 5 个结果:
c[0] = 100.0 (预期值: 100)
c[1] = 100.0 (预期值: 100)
c[2] = 100.0 (预期值: 100)
c[3] = 100.0 (预期值: 100)
c[4] = 100.0 (预期值: 100)

可以看到，Tesla K20c 有 13 块 SM，每块最大线程数为 1024，显存 4.63 GB，每秒可执行 184.7 亿次浮点运算。到此，我们已经掌握了入门 CUDA 所需的全部基础知识了！鼓掌👏👏👏

💻 抠点代码
#cuda 编程 #并行程序
从零入门 cuda 编程？🦴 看这篇就够了！
http://example.com/2025/02/23/cuda-programming/
Author
LazyPool
Posted on
February 23, 2025
Licensed under