Zhangzhe's Blog

The projection of my life.

0%

CUDA 学习笔记 01 —— CUDA 基础

Posted on Edited on In Valine:

0. 学习 CUDA 的目的

  • 作为一个算法工程师，平时接触 HPC (High Performance Computing) 的机会并不多，那为什么还要学习 CUDA 呢？
  • 学习 CUDA 的目的不是为了用 CUDA 做模型加速，而是从 CUDA 角度理解目前较新的大模型设计理念，这些高性能模型是如何从原理上做到又快又好的。
  • 例如火出圈的 DeepSeek 系列模型，在模型设计角度做了较多创新，并开源了部分 CUDA 代码，对于不了解 CUDA 的工程师，很难 get 到算法设计的优雅之处。
  • 反观某家大模型基座公司，曾开源某个模型结构，论文中一通自夸，分析理论计算量有多低。但很多人实测发现速度并没有很快，究其原因，实际上是这家公司还用的小模型时代的旧思维，即：一个模型理论计算量低，那就是快。
  • 大模型时代不了解硬件，不尊重硬件，在算法创新上不太可能走的远。

1. Hello World

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
#include <cuda_runtime_api.h>
#include <stdio.h>

// cuda 中 host 表示 cpu 端，device 表示 gpu 端
// __device__ 是设备函数的声明符号，表明该函数在 device 执行，且只能在 device
// 中调用
__device__ const char *device_hello_world(void) {
  return "GPU: Hello world!\n";
}

// __host__ 是主机函数的声明符号，表明该函数在 host 执行，且只能在 host 中调用
__host__ const char *host_hello_world(void) { return "CPU: Hello world!\n"; }

// __global__ 是核函数的声明符号，表明该函数在 device 执行，且只能在 host 中调用
__global__ void hello_world(void) {
  const char *str = device_hello_world();
  printf("%s", str);
}

int main(int argc, char **argv) {
  printf("%s", host_hello_world());
  // <<<grid_dim, block_dim>>> 是核函数的调用符号，表示启动 grid_dim 个 block，
  // 每个 block 有 block_dim 个线程
  hello_world<<<1, 10>>>();
  cudaDeviceReset();
  return 0;
}
  • cuda 的三个函数声明符号：
    • __host__：主机函数，表示该函数在 CPU 上执行，且只能在 CPU 中调用
    • __device__：设备函数，表示该函数在 GPU 上执行，且只能在 GPU 中调用
    • __global__：核函数，表示该函数在 GPU 上执行，且只能在 CPU 中调用
  • 其中 __global__ 声明的函数类型被称为 核函数，是 CUDA 中最重要的函数类型
    • 核函数通过 <<<grid_dim, block_dim>>> 的方式调用，其中 <<<>>> 是 cuda 扩展关键字
    • grid_dim 表示启动的 block 数量，block_dim 表示每个 block 中的线程数量
    • grid_dim 和 block_dim 都是 dim3 类型的变量，表示三维数组，如果使用整形则模型 y 和 z 维度都为 1

编译

1
nvcc hello_world.cu -g -o hello_world

运行

1
2
3
4
5
6
7
8
9
10
11
12
./hello_world
# CPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!

2. dimension 测试

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
#include <cuda_runtime.h>
#include <stdio.h>
__global__ void checkIndex(void) {
  printf("threadIdx:(%d,%d,%d) blockIdx:(%d,%d,%d) blockDim:(%d,%d,%d)\
  gridDim(%d,%d,%d)\n",
         threadIdx.x, threadIdx.y, threadIdx.z, blockIdx.x, blockIdx.y,
         blockIdx.z, blockDim.x, blockDim.y, blockDim.z, gridDim.x, gridDim.y,
         gridDim.z);
}
int main(int argc, char **argv) {
  int nElem = 6;  // number of elements
  dim3 block(3);  // block size
  int nBlock = (nElem + block.x - 1) / block.x; // number of blocks
  dim3 grid(nBlock);  // grid size
  printf("grid.x %d grid.y %d grid.z %d\n", grid.x, grid.y, grid.z);  
  printf("block.x %d block.y %d block.z %d\n", block.x, block.y, block.z);
  checkIndex<<<grid, block>>>();
  cudaDeviceReset();
  return 0;
}
  • 执行结果：
1
2
3
4
5
6
7
8
9
./check_dimension
# grid.x 2 grid.y 1 grid.z 1
# block.x 3 block.y 1 block.z 1
# threadIdx:(0,0,0) blockIdx:(1,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(1,0,0) blockIdx:(1,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(2,0,0) blockIdx:(1,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(0,0,0) blockIdx:(0,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(1,0,0) blockIdx:(0,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(2,0,0) blockIdx:(0,0,0) blockDim:(3,1,1)  gridDim(2,1,1)

3. 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
#include "freshman.h"
#include <cuda_runtime.h>
#include <stdio.h>

__host__ void sumArrays(float *a, float *b, float *res, const int size) {
  for (int i = 0; i < size; i++) {
    res[i] = a[i] + b[i];
  }
}
__global__ void sumArraysGPU(float *a, float *b, float *res, const int size) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < size) // 线程索引越界检查
    res[i] = a[i] + b[i];
}
int main(int argc, char **argv) {
  // set up device
  initDevice(0);

  // allocate host memory
  int nElem = 1 << 24;
  printf("Vector size:%d\n", nElem);
  int nByte = sizeof(float) * nElem;
  float *a_h = (float *)malloc(nByte);
  float *b_h = (float *)malloc(nByte);
  float *res_h = (float *)malloc(nByte);
  float *res_from_gpu_h = (float *)malloc(nByte);
  memset(res_h, 0, nByte);
  memset(res_from_gpu_h, 0, nByte);

  // allocate device memory
  float *a_d, *b_d, *res_d;
  CHECK(cudaMalloc((float **)&a_d, nByte));
  CHECK(cudaMalloc((float **)&b_d, nByte));
  CHECK(cudaMalloc((float **)&res_d, nByte));

  // randomly initialize the input data
  initialData(a_h, nElem);
  initialData(b_h, nElem);
  CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
  CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));

  // set up execution configuration
  // 1. 计算最佳块大小
  int minGridSize, bestBlockSize;
  cudaOccupancyMaxPotentialBlockSize(&minGridSize, &bestBlockSize,
                                     (void *)sumArraysGPU,
                                     0, // 动态共享内存大小
                                     0  // 无块大小限制
  );

  // 2. 设置网格和块维度
  dim3 block(bestBlockSize);
  dim3 grid((nElem + bestBlockSize - 1) / bestBlockSize);

  // 3. 设备执行并统计耗时
  double iStart, iElaps;
  iStart = cpuSecond();
  sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d, nElem);
  iElaps = cpuSecond() - iStart;
  CHECK(cudaGetLastError());
  CHECK(cudaDeviceSynchronize());
  CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));
  printf("Execution configuration<<<%d,%d>>> Time elapsed %f sec\n", grid.x,
         block.x, iElaps);

  // 4. CPU执行并统计耗时
  iStart = cpuSecond();
  sumArrays(a_h, b_h, res_h, nElem);
  iElaps = cpuSecond() - iStart;
  printf("CPU Time elapsed %f sec\n", iElaps);

  // 5. 检查结果
  checkResult(res_h, res_from_gpu_h, nElem);

  // 6. 释放内存
  cudaFree(a_d);
  cudaFree(b_d);
  cudaFree(res_d);

  free(a_h);
  free(b_h);
  free(res_h);
  free(res_from_gpu_h);

  return 0;
}
  • cudaOccupancyMaxPotentialBlockSize 函数用于计算最佳块大小
  • max thread per block 是 cuda 中的一个限制，表示每个块中最多可以有多少个线程，一般为 1024，当超过这个限制时，CHECK(cudaGetLastError()); 会报错

运行结果

1
2
3
4
5
6
./sum_arrays
# Using device 0: NVIDIA GeForce RTX 4060 Laptop GPU
# Vector size:16777216
# Execution configuration<<<21846,768>>> Time elapsed 0.000030 sec
# CPU Time elapsed 0.076604 sec
# Check result success!
  • 可以看出，cuda 的执行时间远远小于 cpu 的执行时间，相差了 2553 倍

总体流程

graph TD
    A[Host 端申请内存] --> B[Host 端初始化输入数据]
    B --> C[Device 端申请内存]
    C --> D[拷贝 Host 输入数据到 Device 端]
    D --> E[Device 端执行核函数]
    E --> F[拷贝 Device 端输出数据到 Host 端]
    F --> G[Host 端检查结果]
    G --> H[释放 Host 端和 Device 端全部内存]
    B --> J[Host 端执行普通函数] --> G

细节分析

1. cuda 内存分配是怎么做的？

1
2
float *a_d; // 空指针
cudaMalloc((float **)&a_d, nByte);  // 将指针的地址转成二级指针（指针的指针）传入内存分配函数
  • 这里的二级指针应用很巧妙，由于 c++ 中的指针是值传递，所以如果是一级指针传入 cudaMalloc 函数时，指针 a_d 的值不会改变，因此只能将指针的地址转成二级指针传入内存分配函数

2. cuda 内存拷贝是怎么做的？

1
cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice);
  • cudaMemcpy 函数的四个参数分别是：
    • dst：目标地址
    • src：源地址
    • size：拷贝的字节数
    • kind：拷贝的类型，属于 cudaMemcpyKind 枚举类型
      • cudaMemcpyHostToHost = 0, /**< Host -> Host */
      • cudaMemcpyHostToDevice = 1, /**< Host -> Device */
      • cudaMemcpyDeviceToHost = 2, /**< Device -> Host */
      • cudaMemcpyDeviceToDevice = 3, /**< Device -> Device */
      • cudaMemcpyDefault = 4 /**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */

3. cuda 核函数如何解析线程索引？

1
2
3
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < size) // 线程索引越界检查
  res[i] = a[i] + b[i];
  • blockIdx：表示当前块的索引
  • blockDim：表示当前块的维度（每个块中的线程数）
  • threadIdx：表示当前线程的索引
  • 每个线程中计算两个标量的和
  • 由于 gridDim * blockDim 可能大于 size，所以需要判断线程索引是否越界

4. 如何计算最佳块大小？

1
2
3
4
5
6
int minGridSize, bestBlockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &bestBlockSize,
                                     (void *)sumArraysGPU,
                                     0, // 动态共享内存大小
                                     0  // 无块大小限制
);
  • cudaOccupancyMaxPotentialBlockSize 函数用于计算最佳块大小，五个参数分别是：
    • minGridSize：最小网格大小变量地址
    • bestBlockSize：最佳块大小变量地址
    • kernel：核函数指针
    • dynamicSMemSize：动态共享内存大小
    • blockSizeLimit：块大小限制
  • 函数名就是函数地址，可强转为 void * 函数指针（也可以写成：(void *)&sumArraysGPU）

4. CUDA 编程模型

线程块

  • 线程块 block 是 CUDA 中的逻辑执行单元，是一个三维逻辑结构：
    • block.x：表示块的 x 维度大小
    • block.y：表示块的 y 维度大小
    • block.z：表示块的 z 维度大小
    • 其中 block.x 是最内层的循环，block.y 是第二层循环，block.z 是最外层的循环
    • 用三维数组可以表示为：tread[z][y][x]，即 tid = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y

线程束

  • 线程束 warp 是 CUDA 基本调度执行单元，一个 warp 由 32 个线程组成
    • 一个 warp 中的线程在一个时钟周期内执行同一条指令（单指令多线程，SIMT）
    • 一个 warp 中的线程可以共享指令指针和执行资源（如寄存器、缓存等）
    • Warp 调度器（warp scheduler）负责将 warp 分配到物理执行单元上执行
  • 线程块会被划分为多个 warp，WarpsPerBlock=ceil(ThreadsPerBlock32)WarpsPerBlock=ceil(\frac{ThreadsPerBlock}{32})

CUDA core

  • CUDA core 是 CUDA 物理执行单元，负责实际的计算任务
  • 一个 CUDA core 一个时钟周期只能计算一个线程的指令

StreamMultiprocessor

  • StreamMultiprocessor 流式多处理器（简称 SM），负责执行 CUDA 线程块中的并行计算任务
  • 每个 GPU 包含多个 SM，每个 SM 包含多个 CUDA core，例如：RTX 4060 有 24 个 SM，每个 SM 有 128 个 CUDA core

5. Reduce

  • 规约（Reduce）是 CUDA 编程中常见的操作，主要用于将多个数据元素规约为一个数据元素
  • 规约操作通常是一个二元操作，例如：sum、mul、max、min 等，简单的规约可以合并成强大的算子，甚至可以说规约算子是神经网络的基础

规约求和

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
// CPU 规约求和
int recursiveReduce(int *data, int const size) {
  // terminate check
  if (size == 1)
    return data[0];
  // renew the stride
  int const stride = size / 2;
  if (size % 2 == 1) {
    for (int i = 0; i < stride; i++) {
      data[i] += data[i + stride];
    }
    data[0] += data[size - 1];
  } else {
    for (int i = 0; i < stride; i++) {
      data[i] += data[i + stride];
    }
  }
  // call
  return recursiveReduce(data, stride);
}
// GPU 规约相邻求和
__global__ void reduceNeighbored(int *g_idata, int *g_odata, unsigned int n) {
  // set thread ID
  unsigned int tid = threadIdx.x;
  // boundary check
  if (tid >= n)
    return;
  // convert global data pointer to local point of this block
  int *idata = g_idata + blockIdx.x * blockDim.x;
  // in-place reduction in global memory
  for (int stride = 1; stride < blockDim.x; stride *= 2) {
    if ((tid % (2 * stride)) == 0) {
      idata[tid] += idata[tid + stride];
    }
    // synchronize within block
    __syncthreads();
  }
  // write result for this block to global mem
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}
// GPU 规约相邻求和（简化版）
__global__ void reduceNeighboredLess(int *g_idata, int *g_odata,
                                     unsigned int n) {
  unsigned int tid = threadIdx.x;
  unsigned idx = blockIdx.x * blockDim.x + threadIdx.x;
  // convert global data pointer to the local point of this block
  int *idata = g_idata + blockIdx.x * blockDim.x;
  if (idx > n)
    return;
  // in-place reduction in global memory
  for (int stride = 1; stride < blockDim.x; stride *= 2) {
    // convert tid into local array index
    int index = 2 * stride * tid;
    if (index < blockDim.x) {
      idata[index] += idata[index + stride];
    }
    __syncthreads();
  }
  // write result for this block to global men
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}
// GPU 规约交错求和，主要是 stride 的计算方式不同
__global__ void reduceInterleaved(int *g_idata, int *g_odata, unsigned int n) {
  unsigned int tid = threadIdx.x;
  unsigned idx = blockIdx.x * blockDim.x + threadIdx.x;
  // convert global data pointer to the local point of this block
  int *idata = g_idata + blockIdx.x * blockDim.x;
  if (idx >= n)
    return;
  // in-place reduction in global memory
  for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {

    if (tid < stride) {
      idata[tid] += idata[tid + stride];
    }
    __syncthreads();
  }
  // write result for this block to global men
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}

效率分析

  • 本代码中使用了三种核函数实现方式做同一个规约操作，分别是：
    • reduceNeighbored：相邻线程规约
    • reduceNeighboredLess：相邻线程规约（简化版）
    • reduceInterleaved：交错线程规约
  • 三者效率从高到低依次是：
    • reduceInterleaved > reduceNeighboredLess > reduceNeighbored
  • 三者的示意图分别如下：
    • reduceNeighbored：相邻线程规约的实现
    • reduceNeighboredLess：相邻线程规约的简化版实现（注意圆圈中的符号已和上图不一致）
    • reduceInterleaved：交错线程规约的实现
  • 三者效率差异主要来自于 线程分支分化，后续会详细介绍

6. 循环展开

  • 循环展开（Loop Unrolling）是 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
// 数据总长是 8 * blockDim.x * gridDim.x
// 线程数是 blockDim.x * gridDim.x
// 每个线程计算 8 个数据
__global__ void reduceUnroll8(int *g_idata, int *g_odata, unsigned int n) {
  unsigned int tid = threadIdx.x;
  unsigned int idx = blockDim.x * blockIdx.x * 8 + threadIdx.x;
  if (tid >= n)
    return;
  int *idata = g_idata + blockIdx.x * blockDim.x * 8;
  // 循环展开，每个线程计算 8 个数据
  // 直接把 8 * blockDim.x * gridDim.x 的数据总长
  // 聚合到了 blockDim.x * gridDim.x 的线程数上
  if (idx + 7 * blockDim.x < n) {
    g_idata[idx] += g_idata[idx + blockDim.x];
    g_idata[idx] += g_idata[idx + blockDim.x * 2];
    g_idata[idx] += g_idata[idx + blockDim.x * 3];
    g_idata[idx] += g_idata[idx + blockDim.x * 4];
    g_idata[idx] += g_idata[idx + blockDim.x * 5];
    g_idata[idx] += g_idata[idx + blockDim.x * 6];
    g_idata[idx] += g_idata[idx + blockDim.x * 7];
  }
  // 这里需要同步，也就是线程阻塞直到所有线程都执行完
  __syncthreads();
  // 然后就是一个最简单的规约操作了，和上面一样
  for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
    if (tid < stride) {
      idata[tid] += idata[tid + stride];
    }
    __syncthreads();
  }
  // write result for this block to global mem
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}

7. 核函数递归调用

  • 和 CPU 一样，CUDA 也支持核函数递归调用，调用方式和普通递归函数一样
  • 需要注意的是在编译的时候需要加上 -rdc=true 选项

代码

1
2
3
4
5
6
7
8
9
10
11
12
13
__global__ void nesthelloworld(int iSize, int iDepth) {
  unsigned int tid = threadIdx.x;
  printf("depth : %d blockIdx: %d,threadIdx: %d\n", iDepth, blockIdx.x,
         threadIdx.x);
  if (iSize == 1)
    return;
  int nthread = (iSize >> 1);
  if (tid == 0 && nthread > 0) {
    // 递归调用核函数
    nesthelloworld<<<1, nthread>>>(nthread, ++iDepth);
    printf("-----------> nested execution depth: %d\n", iDepth);
  }
}

8. 固定内存

  • Pinned Memory 是 CUDA 中的一种特殊内存类型（不是显存，是内存），主要用于提高数据传输效率
  • 普通内存是分页管理，分页管理存在两个问题：
    1. 一页内存逻辑上连续，但物理上不连续
    2. 操作系统可能会将内存页交换到磁盘上，导致数据不在物理内存中
  • Pinned Memory 就是解决了这两个问题，分配了一块连续物理地址且固定的主机内存（host 内存），方便整块拷贝数据到显存（DMA）

代码

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
#include "freshman.h"
#include <cuda_runtime.h>

int main(int argc, char **argv) {
  int dev = 0;
  cudaSetDevice(dev);

  int nElem = 1 << 14;
  int nByte = sizeof(float) * nElem;
  float *a_h, *b_h, *res_h, *res_from_gpu_h;
  // 注意这里的 cudaMallocHost 和 cudaMalloc 是不同的
  // 前者申请的是 host 固定内存，后者申请的是 device 显存
  // cudaMallocHost 是 malloc 的一个平替
  CHECK(cudaMallocHost((float **)&a_h, nByte));
  CHECK(cudaMallocHost((float **)&b_h, nByte));
  CHECK(cudaMallocHost((float **)&res_h, nByte));
  CHECK(cudaMallocHost((float **)&res_from_gpu_h, nByte));

  // 初始化数据
  memset(res_h, 0, nByte);
  memset(res_from_gpu_h, 0, nByte);
  initialData(a_h, nElem);
  initialData(b_h, nElem);

  // 申请设备显存
  float *a_d, *b_d, *res_d;
  CHECK(cudaMalloc((float **)&a_d, nByte));
  CHECK(cudaMalloc((float **)&b_d, nByte));
  CHECK(cudaMalloc((float **)&res_d, nByte));

  // 拷贝数据到设备显存
  CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
  CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));

  // 跑核函数
  dim3 block(1024);
  dim3 grid(nElem / block.x);
  sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d);
  printf("Execution configuration<<<%d,%d>>>\n", grid.x, block.x);

  // 结果拷贝回主机并检查两个设备计算结果是否一致
  CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));
  sumArrays(a_h, b_h, res_h, nElem);
  checkResult(res_h, res_from_gpu_h, nElem);

  // 释放内存
  // 注意这里的 cudaFreeHost 和 cudaFree 是不同的
  // 前者释放的是 host 固定内存，后者释放的是 device 显存
  cudaFree(a_d);
  cudaFree(b_d);
  cudaFree(res_d);
  cudaFreeHost(a_h);
  cudaFreeHost(b_h);
  cudaFreeHost(res_h);
  cudaFreeHost(res_from_gpu_h);

  return 0;
}

关键点

  1. 主存 普通内存 的分配和释放函数是 malloc 和 free
  2. 主存 固定内存 的分配和释放函数是 cudaMallocHost 和 cudaFreeHost
  3. 显存 的分配和释放函数是 cudaMalloc 和 cudaFree

9. 零拷贝内存 和 统一虚拟地址

  • Zero-Copy Memory 是 CUDA 中一种允许 GPU 直接访问主机内存的技术，避免了显式的数据拷贝操作（不需要用 cudaMemcpy 函数）
  • 实际上，Zero-Copy Memory 在很多时候并不快，因为 GPU 访问主机内存的速度远远低于访问显存的速度，因此，Zero-Copy Memory 只适用于一些特殊的场景，例如：
    • 主机内存中的数据只需要被 GPU 使用一次
    • 数据量太大，显存放不下
    • 调试用途
  • Zero-Copy Memory 的实现方式是将主机内存映射到 GPU 的地址空间中，GPU 通过访问这个地址空间来访问主机内存，实际上走的是 PCIe 总线
  • 由于不需要先完成所有数据的拷贝再开始执行核函数，因此 Zero-Copy Memory 使用异步拷贝的方式来实现，可将部分拷贝数据的时间和核函数执行的时间重叠，但并不多
  • Unified Virtual Addressing (UVA) 是 CUDA 中的一种内存管理机制，允许 CPU 和 GPU 共享同一虚拟地址空间

代码

1
2
3
4
5
6
7
8
9
10
11
12
13
float *a_host, *b_host, *res_d;
// 申请主机固定内存，添加特殊 flag cudaHostAllocMapped
CHECK(cudaMallocHost((float **)&a_host, nByte, cudaHostAllocMapped));
CHECK(cudaMallocHost((float **)&b_host, nByte, cudaHostAllocMapped));

// a_host 和 b_host 是可直接作为核函数的输入参数
// 也可以转成 Device 地址空间，如下：
float *a_dev, *b_dev;
// 映射主机内存到设备地址空间
CHECK(cudaHostGetDevicePointer((void **)&a_dev, (void *)a_host, 0));
CHECK(cudaHostGetDevicePointer((void **)&b_dev, (void *)b_host, 0));

// 用 a_dev 和 b_dev 作为核函数的输入参数计算
  • 和 Zero-Copy Memory 一样，UVA 甚至不需要将 a_host 转成 a_dev，直接用 a_host 就可以调用核函数