2. 内存结构 & 核函数基本说明

内存结构

对于 CPU + GPU 的异构模型，两者内存是隔离的 （当然了，CPU有自己的内存，GPU有自己的显存），因此我们会有对应的 CUDA 函数对设备进行内存分配与复制：

  • cudaMalloc: 分配内存
  • cudaMemcpy: 复制内存
  • cudaMemset: 设置内存
  • cudaFree: 释放内存

cuda 函数统一返回 cudaError_t 方便主机端进行错误处理，为了更好显示，runtime api 也提供了一些必要的转换方法，比如，具体参考官方文档

1
2
3
4
// Returns the string representation of an error code enum name. 
__host__ __device__ const char* cudaGetErrorName (  cudaError_t error )
// Returns the description string for an error code. 
__host__ __device__ const char* cudaGetErrorString (  cudaError_t error )

从软件层面理解 GPU 内存结构，是这样的：

GPU软件架构.png

一个核函数对应一个 Grid，而一个 Grid 中可以用多个 Block 和 Thread (这里 Block 对应硬件层面的 SM，而多个 Thread 会捆绑为一个 Warp 对应硬件层面的 SP 进行执行)，同时注意上图，在同一个 Block 内多个 Thread 是共享内存的，而不同 Block 的线程不能相互影响，这可以和硬件层面的 SM 架构做直接的对应。

当然，从前面 Hello World 的示例程序我们也知道，不同 Thread 之间通过

  • blockIdx: 线程所在的 Block 对应 Grid 内部的索引号（最多可以有 3 维，xyz）
  • threadIdx: 线程对应 Block 内部的索引号（最多可以有 3 维，xyz）有索引

核函数基本说明

简单举个代码例子，注意在计算 Block 数量的时候应当取上确界。

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

__global__ void kernel() {    
		printf("Block ID: (%d, %d, %d), Thread ID: (%d, %d, %d)\n", 
			blockIdx.x, blockIdx.y, blockIdx.z, 
			threadIdx.x, threadIdx.y, threadIdx.z);
}

int main() {    
		fmt::println("CUDA Block 和 Thread 数量说明");    
		int n_elements = 1024; // 总共的元素数量    
		dim3 block_size(256); // 每个 Block 的线程数量    
		dim3 grid_size((n_elements + block_size.x - 1) / block_size.x); // Grid 的 Block 数量，注意这里是取上确界    
		fmt::println("总元素数量: {}", n_elements);    
		fmt::println("每个 Block 的线程数量: {}", block_size.x);    
		fmt::println("需要的 Block 数量: {}", grid_size.x);
    kernel<<<grid_size, block_size>>>();    
    cudaDeviceSynchronize(); // 等待 GPU 完成
}

输出结果为：

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
CUDA Block 和 Thread 数量说明
总元素数量: 32
每个 Block 的线程数量: 16
需要的 Block 数量: 2
Block ID: (0, 0, 0), Thread ID: (0, 0, 0)
Block ID: (0, 0, 0), Thread ID: (1, 0, 0)
Block ID: (0, 0, 0), Thread ID: (2, 0, 0)
Block ID: (0, 0, 0), Thread ID: (3, 0, 0)
Block ID: (0, 0, 0), Thread ID: (4, 0, 0)
Block ID: (0, 0, 0), Thread ID: (5, 0, 0)
Block ID: (0, 0, 0), Thread ID: (6, 0, 0)
Block ID: (0, 0, 0), Thread ID: (7, 0, 0)
Block ID: (0, 0, 0), Thread ID: (8, 0, 0)
Block ID: (0, 0, 0), Thread ID: (9, 0, 0)
Block ID: (0, 0, 0), Thread ID: (10, 0, 0)
Block ID: (0, 0, 0), Thread ID: (11, 0, 0)
Block ID: (0, 0, 0), Thread ID: (12, 0, 0)
Block ID: (0, 0, 0), Thread ID: (13, 0, 0)
Block ID: (0, 0, 0), Thread ID: (14, 0, 0)
Block ID: (0, 0, 0), Thread ID: (15, 0, 0)
Block ID: (1, 0, 0), Thread ID: (0, 0, 0)
Block ID: (1, 0, 0), Thread ID: (1, 0, 0)
Block ID: (1, 0, 0), Thread ID: (2, 0, 0)
Block ID: (1, 0, 0), Thread ID: (3, 0, 0)
Block ID: (1, 0, 0), Thread ID: (4, 0, 0)
Block ID: (1, 0, 0), Thread ID: (5, 0, 0)
Block ID: (1, 0, 0), Thread ID: (6, 0, 0)
Block ID: (1, 0, 0), Thread ID: (7, 0, 0)
Block ID: (1, 0, 0), Thread ID: (8, 0, 0)
Block ID: (1, 0, 0), Thread ID: (9, 0, 0)
Block ID: (1, 0, 0), Thread ID: (10, 0, 0)
Block ID: (1, 0, 0), Thread ID: (11, 0, 0)
Block ID: (1, 0, 0), Thread ID: (12, 0, 0)
Block ID: (1, 0, 0), Thread ID: (13, 0, 0)
Block ID: (1, 0, 0), Thread ID: (14, 0, 0)
Block ID: (1, 0, 0), Thread ID: (15, 0, 0)

从上面代码可以看到，我们采用 __global__ 来声明核函数，这是 CUDA 规定的，具体而言：

说明符 执行 调用
global 设备执行 主机/设备 都可调用
device 设备执行 设备调用
host 主机执行 主机调用

而且这里有个特殊的情况就是有些函数可以同时定义为 device 和 host ，这种函数可以同时被设备和主机端的代码调用，主机端代码调用函数很正常，设备端调用函数与C语言一致，但是要声明成设备端代码，告诉nvcc编译成设备机器码，同时声明主机端设备端函数，那么就要告诉编译器，生成两份不同设备的机器码。

调用核函数时就是简单的 kernel<<<GridDim, BlockDim>>>() 非常符合直觉，唯一值得注意的是：

错误处理

我们前面提到了 C++ 常见的通过返回值进行错误处理，我们可以简单定义一个宏，帮助快速定位错误地方：

1
2
3
4
5
6
7
8
9
10
#define CHECK(call)\
{\
  const cudaError_t error=call;\
  if(cudaError_t::cudaSuccess != error)\
  {\
      printf("ERROR: %s:%d,",__FILE__,__LINE__);\
      printf("code:%d,reason:%s\n",error,cudaGetErrorString(error));\
      exit(1);\
  }\
}

Profiling

2026了，当然得采用现代的 Nsight 而不是 nvprof 来做性能分析，直接在终端使用：

1
nsys profile --stats=true ./memory

--status=true 会在终端直接打印 API 的耗时和核函数执行时间统计汇总，比如像下面文本：

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
[4/8] Executing 'osrt_sum' stats report

 Time (%)  Total Time (ns)  Num Calls   Avg (ns)    Med (ns)   Min (ns)  Max (ns)   StdDev (ns)       Name     
 --------  ---------------  ---------  ----------  ----------  --------  ---------  -----------  --------------
     78.9        348889231         24  14537051.3  10061604.0      1300  147873503   29645623.8  poll          
     20.5         90671745        432    209888.3      6045.0      1002    5784257     520051.3  ioctl         
      0.3          1357574         10    135757.4    122988.5     47363     354565      91872.8  sem_timedwait 
      0.1           588535         26     22636.0      3518.0      1265     362989      70685.6  mmap64        
      0.0            97718         16      6107.4      3282.0      1006      26967       6770.8  mmap          
      0.0            81426         33      2467.5      1794.0      1001       8270       1790.3  open64        
      0.0            68668          3     22889.3     24316.0     16023      28329       6275.8  pthread_create
      0.0            61140         16      3821.3      2503.0      1221      13949       3287.1  fopen         
      0.0            60226          1     60226.0     60226.0     60226      60226          0.0  fgets         
      0.0            27569          2     13784.5     13784.5      1866      25703      16855.3  socket        
      0.0            15904          6      2650.7      1660.0      1008       6620       2188.8  read          
      0.0            14178          4      3544.5      3417.0      3373       3971        285.3  open          
      0.0            12998          4      3249.5      3315.5      1919       4448       1208.9  munmap        
      0.0            11931          2      5965.5      5965.5      4483       7448       2096.6  fread         
      0.0             8470          3      2823.3      2801.0      2440       3229        395.0  close         
      0.0             8078          3      2692.7      2921.0      1223       3934       1369.8  pipe2         
      0.0             8004          6      1334.0      1261.5      1019       1911        316.7  write         
      0.0             6579          4      1644.8      1542.0      1290       2205        435.7  fclose        
      0.0             5928          2      2964.0      2964.0      1636       4292       1878.1  fwrite        
      0.0             5854          1      5854.0      5854.0      5854       5854          0.0  connect       
      0.0             1152          1      1152.0      1152.0      1152       1152          0.0  bind          

[5/8] Executing 'cuda_api_sum' stats report

 Time (%)  Total Time (ns)  Num Calls   Avg (ns)     Med (ns)    Min (ns)   Max (ns)   StdDev (ns)           Name         
 --------  ---------------  ---------  -----------  -----------  ---------  ---------  -----------  ----------------------
     97.5        102239930          1  102239930.0  102239930.0  102239930  102239930          0.0  cudaLaunchKernel      
      1.9          1940297          1    1940297.0    1940297.0    1940297    1940297          0.0  cuLibraryLoadData     
      0.4           434286          1     434286.0     434286.0     434286     434286          0.0  cudaDeviceSynchronize 
      0.2           218750          1     218750.0     218750.0     218750     218750          0.0  cuKernelGetName       
      0.0              535          1        535.0        535.0        535        535          0.0  cuModuleGetLoadingMode
      0.0              223          1        223.0        223.0        223        223          0.0  cuLibraryGetKernel    

[6/8] Executing 'cuda_gpu_kern_sum' stats report

 Time (%)  Total Time (ns)  Instances  Avg (ns)  Med (ns)  Min (ns)  Max (ns)  StdDev (ns)    Name  
 --------  ---------------  ---------  --------  --------  --------  --------  -----------  --------
    100.0            27840          1   27840.0   27840.0     27840     27840          0.0  kernel()

同时还会生成对应的以 .nsys-rep 结尾的文件，可以使用 nsight-compute 打开：

1
nsys-ui ./report1.nsys-rep

profile.png

说些什么吧！

beaudar