cuda手记

1.cuda核心基本概念：

显卡的CUDA Core数量=SM(Streaming Multiprocessor) * 64,一个SM包含两个SMP(SM Processing Block),SMP为显卡调用的基本单元(warp)，通常包含32个核心，一个核心处理fp32的数据，处理fp64的数据用两个warp。NVIDIA提供了很多内置的核函数，计算能力（如4090的8.9）代表的就是这些核函数的端口号。

2.不同显卡之间的数据传输：

        显卡间的通信通常可以用无线带宽技术(IB)，处理速度大约10Gb/s，如果使用NVLink可以到达200Gb/s。可以用的网卡接口包括1、4、8、16。显卡之间数据通讯的方式通常使用三种，下面将介绍这三种方式

        DP：首先由主卡cuda:0将需要运算的数据平均分发到每个子卡上，分别完成运算之后再传回主卡，因此传输数据量为D所需时间公式为2(N - 1) * D / S，N为显卡数量，D为参数总量，S为传输速率，乘以2是因为一次传输包含收发两部。

        DDP：Allgather:1.将显卡环形连接，2.将数据均分成n份，3.循环N-1次，每次将当前层的一份数据发送给相邻显卡，这样做的传输耗时为2 * (N - 1) * D / N / S，优势是传输数据的时长最大为2*D/S

        

3.代码相关：

__global__修饰的函数由host启动，在device执行，返回值类型为void

__device__修饰的函数只能由device启动，也在device上执行

线程的数量为gridDim.x *gridDim.y *gridDim.z * blockDim.x * blockDim.y * blockDim.z

grideDim最大值范围（2^31, 2^15, 2^15）

blockDim最大范围(2^10, 2^10, 2^10)

gride、block的x、y、z分别对应图像的列数、行数、通道数，等同于pytorch对tensor的处理

最终启动的线程为：gridDim.z  gridDim.y  gridDim.x  blockDim.z blockDim.y blockDim.x

对单线程的索引为：blockIdx.z blockIdx.y blockIdx.x threadIdx.z threadIdx.y threadIdx.x

转化成position: ((blockIdx.z * gridDim.y + blockIdx.y) *gridDim.x + blockIdx.x) * blockDim.z + ...

如果用二维的则更方便理解，也能满足绝大部分使用情况：

最终启动的线程为：gridDim.y  gridDim.x  blockDim.y blockDim.x

对单线程的索引为：blockIdx.y blockIdx.x threadIdx.y threadIdx.x

转化成position: ((blockIdx.y * gridDim.x + blockIdx.x) *blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x

一维情况是最常用的：gridDim.x blockDim.x

对应的线程索引为    :  blockIdx.x threadIdx.x

转化成position：(blockIdx.x * blockDim.x) + threadIdx.x

规律还是比较明显的：从最高级索引开始，(当前索引号*次级Dim + 次级索引号)*再次级Dim + 再次级索引号)...

4.cuda性能测试工具：

首先从官网下载nsight:

Nsight Systems - Get Started | NVIDIA Developer | NVIDIA Developer

安装：sudo ./...run

添加环境变量： vim ~/.basgrc 

        export PATH="/path/to/Nsight_Systems_2021.4.1/bin:$PATH"

使用工具分析名为test的可执行文件：

./nsys profile -o .output_report /path/to/test

5.cuda流

流是任务级别的线程

流是一个任务队列

流是异步的

指定nullptr则执行默认流

可以使用cudaMemcpy完成设备同步，可以使用流同步完成任务的局部同步，流同步也是使用最多的