欢迎来到从零开始的DFT工程师第七期。

本来第三周应该对工具的使用进行熟悉的，但是由于滚arch的时候出一点问题，导致vcs运行时出现linux kernel不支持的问题，这个我看了一下，似乎不是kernel真的不支持，好像是由于哪里的lib改动或者是环境改变导致的。这个问题报错和实际错误好像不符，比较难修，如果修不了就之后改到虚拟机上进行了，这一期就把特别篇提前放到今天了。

在这一期中我们将会完成CUDA工具+Nsight的安装，然后进行简单的CUDA编程和性能测试。

install CUDA

关于nvida其实很早之前就安装过了，详见此处，在这里面我当时安装了nvidia，nvidia-utils，如果是要使用cuda的话就需要多下一个cuda。这里要注意的是，nvidia-dkms是一个会在内核更新后根据内核重新编译个版本（Dynamic Kernel Module Support），他和nvidia是冲突的（起码在arch）；还有如果使用lts内核，应该下载nvidia-lts.

install Nsight tools

Nsight是一套调试分析的工具，我们一般下载nsys和ncu就可以了。在arch上我们可以

1
2
3
4
5
6
7

yay -S cuda-tools
yay -S nsight-systems
yay -S nsight-graphics

# 后两个可以在 https://developer.nvidia.com/nsight-systems/get-started中完成下载

cuda program

关于cuda编程的相关知识和语法我们可以找到很多tutorial，我这里给出几个我参考的：

在cuda编程中，一般把cpu及其相关的东西叫成host，gpu则是device，gpu是并行进行计算的。

一般我们需要完成：给host分配内存初始化数据->给devices分配内存->把输入的数据搬运到device->完成运行->结果数据搬回host.

  • device 内存管理
1
2
3
4
5
6
7
8
9
10
11
12

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

// 分配空间的原因是gpu不能直接访问分配给cpu的内存
// cudaMalloc在device上分配出count大小的内存，然后把devPtr指向device分配的内存区域
// cudaFree就是把devPtr上内存释放


cudaMallocManaged();
//这是分配一个可以被cpu和gpu访问的空间，释放空间同样使用cudaFree(),他属于虚拟内存，配置不好的话可能出现大量的缺页中断

  • 数据传输
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

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

//把count大小的数据从src复制到dst，cudaMemcpyKind代表传输方向是host->device还是反过来

unsigned int cudaGetDeviceFlags(&device)// device 是设备编号，这个函数的参数是unsigned int *.



cudaError_t cudaMemPrefetchAsync(const void* devPtr, size_t count, int dstDevice, cudaStream_t stream = 0);
// 如果提前就知道gpu需要什么数据可以进行预取来加快运行
// devPtr : 要迁移的内存（cudaMallocManaged分配的）
// count: bytes.
// deDevice: 目标设备ID，0 -> gpu0, 1 -> gpu1, cudaCpuDeviceId -> cpu
// cudaStream_t : cuda stream or not?

cudaError_t cudaMemPrefetchAsync(const void* devPtr, size_t count, cudaMemLocation location, cudaStream_t stream = 0);
// 比较新的语法是这样的

/*

typedef struct cudaMemLocation {
    enum cudaMemLocationType type; // 设备类型 (cudaMemLocationTypeDevice / Host / etc.)
    int id;                        // 设备号
} cudaMemLocation;
*/
//要获取location应该是：
/*
cudaMemLocation loc;
Loc.type = cudaMemLocationTypeDevice;
loc.id = 0;
*/
  • kernel 的调用
1
2
3
4
5
6
7
8
9
10

kernelname<<<gridDim, blockDim>>>(arg.);
// threadIdx.x : index of the current thread;
// blockDim.x : the number of threads in the block;
// blockIdx.x : the current thread block in the currnent grid;
// gridDim.x : the number of grids.

cudaDeviceSynchronize();// let CPU wait until kernel is done


  • 性能分析

可以使用nvidia自带的工具进行性能分析：

1
2
3
4
5
6

nsys profile ./executable # 全局分析
nsys profile -t cuda --stats=true ./add_cuda #

ncu ./executable # kernel级分析