CUDA笔记
CUDA笔记
CUDA基础
grid block thread划分
32个线程为一个Warp，每个指令都是以最小一个Warp运行的
一个Block最大1024个线程，即threadPerBlock=1024，也即是BlockSize=2014
一个Grid多少个Block根据具体来定，比如10000个长度，那么blockPerGrid=(10000 + threadPerBlock - 1) / threadPerBlock, blockPerGrid也即是GridSize
所以kernel调用时Kernel<<<blockPerGrid, threadPerBlock, DS, Stream>>>(param list)
其含义：
DS:可选参数，用于设置每个Block除了静态分配的shared memeory之外最多能动态分配的shared memory大小，单位：Byte，不需要动态分配时该值为0或者省略
Stream: cudaStream_t类型流索引，初始值为零，表示该核函数处在哪个流之中运行
cmake示例：
find_package(CUDA REQUIRED)
include_directories(${CUDA_INCLUDE_DIRS})
cuda_add_executable(test test.cu OPTIONS -arch=sm_75)
target_link_libraries(test ${CUDA_LIBRARIES})
CUDA进阶一
数据拷贝，CPU到GPU加速
锁页内存使用
在GPU上运行加快访存
CPU到GPU可以通过锁页内存加快传输速率，此时依然是拷贝到了GPU的global_memory
此时运行kernel，访存是从global_memory读数据计算，慢！
使用共享内存shared_memory，共享内存是GPU一个Block内的线程共享内存，不同Block不共享
使用方法有俩种，一种静态分配共享内存，一种是动态分配共享内存
由于共享内存的大小有限，大概只有几十K，T4是48K，所以只能分多次拷贝数据
申请共享内存关键字：__shared__
块内共享内存同步：__syncthreads()函数（用于块内不同线程之间同步）
静态分配共享内存
#include "cuda_runtime.h"
__global__ void staticFun(int* d, int n)
{
__shared__ int s[64]; //静态申请，需要指定申请内存的大小
int t = treadIdx.x；
s[t] = d[t]; //将全局内存数据拷贝到申请的共享内存中，之后利用共享内存中的数据参与运算将会比调
//用全局内存数据参与运算快(由于共享内存有限，不能全部拷贝到共享内存，这其中就涉及到分批拷贝问题了）
__syncthreads();//需要等所有线程块都拷贝完成后再进行计算
}
staticFun <<<1, n>>> (d, n);
动态分配共享内存
#include "cuda_runtime.h"
__global__ void dynamicFun(int *d, int n)
{
extern __shared__ int s[]; //动态申请，不需要指定大小，需要加上extern关键字
int t = threadIdx.x;
s[t] = d[t];
__syncthreads();
}
dynamicFun <<<1, n, n*sizeof(int) >>> (d, n); //动态申请需要在外部指定共享内存大小
CUDA进阶二
利用stream加速大批量文件IO读写耗时
CUDA的stream流，类似我们经常使用CPU时开多线程。
当我们使用GPU进行计算时，如果我们没有主动开启stream流，GPU会自动创建默认流来执行核函数，默认流和CPU端的计算是同步的。（也即在CPU执行任务过程中，必须等GPU执行完核函数后，才能继续往下执行）
当我们使用GPU进行计算时，我们可以主动开启多个stream流，类似CPU开启多线程。我们可以将大批量文件读写分给多个流去执行，或者用不同的流分别计算不同的核函数。开启的多个流之间是异步的，流与CPU端的计算也是异步的。所以我们需要注意加上同步操作。
值得注意的是，受PCIe总线带宽的限制，当一个流在进行读写操作时，另外一个流不能同时进行读写操作，但是其他流可以进行数值计算任务。
CUDA流的API函数
#include "cuda_runtime.h"
// 创建一个stream
cudaStream_t stream;
cudaStreamCreate(&stream);
// 将host数据拷贝到device
cudaMemcpyAsync(dst, src, size, type, stream);
// kernel在流中执行
kernel_name<<<grid, block, stream>>>(praments);
// 流同步
cudaError_t cudaStreamSynchronize(cudaStream_t stream);
// 流查询
cudaError_t cudaStreamQuery(cudaStream_t stream);
// 销毁流
cudaError_t cudaStreamDestroy(cudaStream_t stream);
CUDA进阶三
cuBLAS库API进行矩阵计算
cuBLAS只是CUDA kernel级别的GPU API函数，所以话需要自己手动分配显存、传输数据等等工作
cuBLAS使用模板范例
int main(int argc, char **argv)
{
......
cublasStatus_t status;
cublasHandle_t handle;
cublasCreate(&handle);
float a = 1, b = 0;
cublasSgemm(
handle,
CUBLAS_OP_T, //矩阵A的属性参数，转置，按行优先
CUBLAS_OP_T, //矩阵B的属性参数，转置，按行优先
M, //矩阵A、C的行数
N, //矩阵B、C的列数
K, //A的列数，B的行数，此处也可为B_ROW,一样的
&a, //alpha的值
d_A, //左矩阵，为A
K, //A的leading dimension，此时选择转置，按行优先，则leading dimension为A的列数
d_B, //右矩阵，为B
N, //B的leading dimension，此时选择转置，按行优先，则leading dimension为B的列数
&b, //beta的值
d_C, //结果矩阵C
M //C的leading dimension，C矩阵一定按列优先，则leading dimension为C的行数
);
cudaMemcpy(deviceRef, d_C, Cxy * sizeof(float), cudaMemcpyDeviceToHost);
cudaDeviceSynchronize();
......
}
cuBLAS矩阵乘法实例
//计算矩阵乘法
#include <stdio.h>
#include "cublas_v2.h"
//主函数
int main(int argc, char *argv[]) {
int M = 3;
int N = 3;
int K = 2;
int n = 6;
size_t size_A = sizeof(float)*M*K;
size_t size_B = sizeof(float)*K*N;
size_t size_C = sizeof(float)*M*N;
// 在CPU上分配内存
float *h_A = (float *)malloc(size_A);
float *h_B = (float *)malloc(size_B);
float *h_C = (float *)malloc(size_C);
//初始化a b的值（将需要计算的向量放到分配好的内存中）
for (int i = 0; i < n; i++) {
h_A[i] = i+1;
h_B[i] = i+1;
}
printf(">>>>>>> h_A:\n");
for (int i = 0; i < n; i++)
{
printf("%.2f ", h_A[i]);
if ((i+1)%2 == 0) printf("\n");
}
printf("\n-----------------------\n\n");
printf(">>>>>>> h_B:\n");
for (int i = 0; i < n; i++)
{
printf("%.2f ", h_B[i]);
if ((i+1)%3 == 0) printf("\n");
}
printf("\n-----------------------\n\n");
//在GPU上分配显存（格式按照 参考下面代码，size为需要分配的显存大小）
float *d_A = NULL;
float *d_B = NULL;
float *d_C = NULL;
cudaMalloc((void **)&d_A, size_A);
cudaMalloc((void **)&d_B, size_B);
cudaMalloc((void **)&d_C, size_C);
//将CPU上初始化的a b值拷贝到GPU上
cudaMemcpy(d_A, h_A, size_A, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size_B, cudaMemcpyHostToDevice);
// Param for Sgemm
// cublasStatus_t status;
cublasHandle_t handle;
cublasCreate(&handle);
float alpha = 1.0f;
float beta = 0.0f;
// 按行优先
cublasSgemm(
handle,
CUBLAS_OP_T, //矩阵A的属性参数，转置，按行优先
CUBLAS_OP_T, //矩阵B的属性参数，转置，按行优先
M, //矩阵A、C的行数
N, //矩阵B、C的列数
K, //A的列数，B的行数，此处也可为B_ROW,一样的
&alpha, //alpha的值
d_A, //左矩阵，为A
K, //A的leading dimension，此时选择转置，按行优先，则leading dimension为A的列数
d_B, //右矩阵，为B
N, //B的leading dimension，此时选择转置，按行优先，则leading dimension为B的列数
&beta, //beta的值
d_C, //结果矩阵C
M //C的leading dimension，C矩阵一定按列优先，则leading dimension为C的行数
);
// 按列优先
cublasSgemm(
handle,
CUBLAS_OP_N, //矩阵A的属性参数，不转置，按列优先
CUBLAS_OP_N, //矩阵B的属性参数，不转置，按列优先
N, //矩阵B^T、C^T的行数
M, //矩阵A^T、C^T的列数
K, //B^T的列数，A^T的行数，此处也可为A_COL,一样的
&alpha, //alpha的值
d_B, //左矩阵，为B^T
N, //B^T的leading dimension，按列优先，则leading dimension为B^T的行数(B的列数)
d_A, //右矩阵，为A^T
K, //A^T的leading dimension，按列优先，则leading dimension为A^T的行数(A的列数)
&beta, //beta的值
d_C, //结果矩阵C
N //C^T的leading dimension，C^T矩阵一定按列优先，则leading dimension为C^T的行数(C的列数)
);
//将GPU上的计算结果拷贝回CPU
cudaMemcpy(h_C, d_C, size_C, cudaMemcpyDeviceToHost);
for(int i = 0; i < M*N; i++)
{
printf(">>>>>>> h_C[%d]: %.2f\n", i, h_C[i]);
}
//释放GPU显存资源
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
//释放CPU内存资源
free(h_A);
free(h_B);
free(h_C);
return 0;
}
使用cuBLAS的时候引用头文件#include <cublas_v2.h>
cmake链接cuBLAS示例：
find_package(CUDA REQUIRED)
include_directories(${CUDA_INCLUDE_DIRS})
cuda_add_executable(test test.cu OPTIONS -arch=sm_75)
target_link_libraries(test ${CUDA_LIBRARIES} cublas)
进阶四 CPU内存与GPU内存相互传输数据
经典的基础传输过程
cudaMalloc((void **)&devInData, memSize);
cudaMemcpy(devInData, hInData, memSize, cudaMemcpyHostToDevice)
// or
cudaMemcpyAsync(devInData, hInData, memSize, cudaMemcpyHostToDevice)
>. hostToDeviceTransfer bandwith: 4.739540 GB/s
>. deviceToHostTransfer bandwith: 4.991877 GB/s
锁页内存
锁页内存本质是在CPU内存上，分配锁页内存可以使用cudaMallocHost,cudaHostAlloc俩个API不同之处：
函数声明：
__host__ cudaError_t cudaMallocHost ( void** ptr, size_t size )
__host__ cudaError_t cudaHostAlloc ( void** pHost, size_t size, unsigned int flags )
cudaMallocHost((void **)&hOutData, memSize)
cudaHostAlloc((void **)&hOutData, memSize, flags)
flags可选择：
#define cudaHostAllocDefault 0x00
Default page-locked allocation flag
#define cudaHostAllocPortable 0x01
Pinned memory accessible by all CUDA contexts
锁页内存空间可以被多个GPU设备使用
#define cudaHostAllocMapped 0x02
Map allocation into device space
把锁页内存地址映射到设备地址空间，这块内存空间会有了2个内存地址，
一个是CPU内存地址，一个是映射到GPU显存的地址，GPU显存映射地址指针
通过cudaHostGetDevicePointer()获取，注意：这种情况不是没有数据拷贝，而是GPU核函数在执行的时候隐式传输数据。
cudaHostAlloc() 返回的地址指针一个的例外情况是，主机和设备使用统一地址空间UVA Unified Virtual Address Space）
#define cudaHostAllocWriteCombined 0x04
Write-combined memory
锁页主机存储是可缓存的，其被分配为写结合的（Write-Combining Memory）。写结合存储不使用L1 和L2 cache，所以程序的其它部分就有更多的缓存可用。此外，写结合内存通过PCI-E传输数据时不会被监视（snoop），这能够获得高达40%的传输加速。 从主机读取写结合存储非常慢（因为没有使用L1、L2cache），所以写结合存储应当只用于那些主机只写的存储，就是只适用分配储存GPU上和函数运算结果的，也就是用来GPU到CPU的。
将普通CPU内存变成CPU锁页内存
__host__ cudaError_t cudaHostRegister ( void* ptr, size_t size, unsigned int flags )
flags可选值如下，含义也和上述的分配锁页内存的flags含义一致，只不过这个是从普通内存注册成锁页内存的时使用的flags
#define cudaHostRegisterDefault 0x00
Default host memory registration flag
#define cudaHostRegisterPortable 0x01
Pinned memory accessible by all CUDA contexts
#define cudaHostRegisterMapped 0x02
Map registered memory into device space
#define cudaHostRegisterIoMemory 0x04
Memory-mapped I/O space
使用锁页内存的吞吐量
>. hostToDeviceTransferWithPinned bandwith: 12.305291 GB/s
>. deviceToHostTransferWithPinned bandwith: 13.156867 GB/s
使用零拷贝(Zero copy)
Zero copy(零拷贝) 是GPU计算单元直接从系统内存读取数据，不需要将数据从系统内存转移到GPU的显存。通常系统内存的数据需要经过GPU内存才能进入计算的缓存中，但通过zero copy 能够实现直接的读写。
float *a_h, *a_map; // 定义两个指针：a_h 内存原指针，a_map映射指针
...
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0); // 获取GPU的特性，看是否支持地址映射
if (!prop.canMapHostMemory)
exit(0);
cudaSetDeviceFlags(cudaDeviceMapHost); // 设置设备属性，打开地址映射
cudaHostAlloc(&a_h, nBytes, cudaHostAllocMapped); // 开辟cudaHostAllocMapped模式下的pinned memory
cudaHostGetDevicePointer(&a_map, a_h, 0); // 地址映射 a_h -> a_map
kernel<<<gridSize, blockSize>>>(a_map);
zero copy 需要借助pinned memory
zero copy 适用于只需要一次读取或者写入的数据操作，要频繁读写的数据，不建议用zero copy。
这种方法的吞吐量是
>. Data tranfer via zero copy. VectorAdd throughput: 833.333313 GB/s
使用统一虚拟内存(UVA)
避免缺页中断
// 方法一
__global__ void kernel_func(int *array) {
...
}
int main() {
char * array = nullptr;
cudaMallocManaged(&array, N) //分配内存
fill_data(array);
cudaMemPerfechAsync(array, N, GPU_DEVICE ,NULL); //告诉GPU预取的数据，使其可以一次性DMA读取,避免缺页中断
kernel_func<<<...>>>(array); //GPU process
cudaMemPerfechAsync(array, N, cudaCpuDeviceId ,NULL); //告诉CPU预取的数据，使其可以一次性DMA读取,避免缺页中断
cudaDeviceSynchronize();
use_data(array); //CPU process
cudaFree(array);
return 0;
}
// 方法二
__device__ __managed__ char array[1000];
__global__ void kernel_func(int a, int b) {
...
}
int main() {
fill_data(array);
cudaMemPerfechAsync(array, N, GPU_DEVICE ,NULL); //告诉GPU预取的数据，使其可以一次性DMA读取
kernel_func<<<...>>>(array); //GPU process
cudaMemPerfechAsync(array, N, cudaCpuDeviceId ,NULL); //告诉CPU预取的数据，使其可以一次性DMA读取
cudaDeviceSynchronize();
use_data(array); //CPU process
cudaFree(array);
return 0;
}
指定内存读写特性
cudaMemAdvise告知分配内存的读写特性
避免同时读写同一片内存，一般cudaMemAdvise配合cudaMemPrefetchAsync使用
有2种flag cudaMemAdviseSetReadMostly和cudaMemAdviseUnSetReadMostly
#define GPU_DEVICE 0
{
char * array = nullptr;
cudaMallocManaged(&array, N) //分配内存
fill_data(array);
cudaMemAdvise(array, N, cudaMemAdviseSetReadMostly, GPU_DEVICE); //提示GPU端几乎仅用于读取这片数据
cudaMemPrefetchAsync(array, N, GPU_DEVICE, NULL); // GPU prefetch
qsort<<<...>>>(array); //GPU 无缺页中断，产生read-only副本
//cudaDeviceSynchronize();
use_data(array); //CPU process 没有page-fault.
cudaFree(array);
}
用cudaMemAdvisePreferredLocation来指定数据存储位置，数据只在指定设备上一个副本
#define GPU_DEVICE 0
{
char * array = nullptr;
cudaMallocManaged(&array, N) //分配内存
//fill_data(array); //也可以uncommit
cudaMemAdvise(array, N, cudaMemAdvisePreferredLocation, GPU_DEVICE); //从此，这片空间仅可以存在CPU上。
qsort<<<...>>>(array); //GPU发生缺页中断，将数据populate到CPU，建立一个访问CPU内存的映射表
cudaDeviceSynchronize();
use_data(array);
cudaFree(array);
}
程序在CPU端调用fill_data(array)但实际上CPU没有为array实际分配内存空间，仅仅是有保留的页表存在，所以必然会产生缺页中断，缺页中断会促使GPU内存的内容通过PCIE总线migrate到CPU内存当中，待CPU处理完缺页，fill_data才函数会继续处理。
当GPU的页传输到CPU当中，为保证数据一致性，GPU的页就标记为失效。之后程序调用cudaKernel: kernel_func<<<....,s>>> 此时，轮到GPU发生缺页中断和数据migrate了，待kernel_func处理完毕，cudaDeviceSyncornize等待同步，CPU再执行use_data首先会发生缺页中断，数据再传回CPU。内存页的数据量小，页的数量很多。导致缺页中断次数多，CPU在用户态和内核态之间来回摇摆，而且内存页的传输不一定用到了DMA。
为了提升传输性能，这时就要引入cudaMemPrefetchAsync，调用DMA来异步传输，再通过cudaStream同步。这样指定的device端由于知道预取多大的数据，就不会频繁的发生缺页中断了。
cudaMemPerfechAsync是防止频繁产生分页中断,参考链接https://blog.csdn.net/weixin_41172895/article/details/115403922
对于使用cudaMallocManaged开辟的内存，在GPU运行时，CPU端的内存页百年城失效状态，在CPU运行时，GPU端的内存页变成失效状态，但有些应用场景不要求这么严格的数据一致性，比如CPU和GPU都对同一片地址空间进行读取操作而没有写入操作就不存在数据竞争，CPU和GPU原本能同时进行操作，然而却被UVM子系统却杜绝了这样的并行操作。
cudaMemAdvise能提前告知一片地址空间的特性
设备与设备之间数据相互传输
普通的经典方法PCIe
// 速度测试的主要代码段：
unsigned char *d_idata;
checkCudaErrors(cudaMalloc((void **)&d_idata, memSize));
unsigned char *d_odata;
checkCudaErrors(cudaMalloc((void **)&d_odata, memSize));
// initialize memory
checkCudaErrors(
cudaMemcpy(d_idata, h_idata, memSize, cudaMemcpyHostToDevice));
// run the memcopy
sdkStartTimer(&timer);
checkCudaErrors(cudaEventRecord(start, 0));
for (unsigned int i = 0; i < MEMCOPY_ITERATIONS; i++) {
checkCudaErrors(
cudaMemcpy(d_odata, d_idata, memSize, cudaMemcpyDeviceToDevice)); // 选择不同传输形式
}
传输速率，A100大概约25GB/s，因为cudaMemcpyHostToDevice是走的PCIe通道
NVLink方法
使用NVLink进行不同GPU之间的数据copy需要使用cudaMemcpyPeer函数
// NVlink COPY 代码差异
cudaSetDevice(0);
unsigned char *d_idata;
checkCudaErrors(cudaDeviceEnablePeerAccess(1, 0));
checkCudaErrors(cudaMalloc((void **)&d_idata, memSize));
cudaSetDevice(1);
unsigned char *d_odata;
checkCudaErrors(cudaDeviceEnablePeerAccess(0, 0));
checkCudaErrors(cudaMalloc((void **)&d_odata, memSize));
cudaSetDevice(0);
// initialize memory
checkCudaErrors(
cudaMemcpy(d_idata, h_idata, memSize, cudaMemcpyHostToDevice));
// run the memcopy
sdkStartTimer(&timer);
checkCudaErrors(cudaEventRecord(start, 0));
for (unsigned int i = 0; i < MEMCOPY_ITERATIONS; i++) {
checkCudaErrors(cudaMemcpyPeer(d_odata, 1, d_idata, 0, memSize));// 这一句是会调用NVLINK通道操作
}
测试第三代的NVLink通信，其传输速度为241GB/s，为所测的PCIe速度的近10倍
NVLink 的卡间通信还可以结合NVSwitch达到节点内GPU之间全量通信。NVSwitch可以理解为一个NVLink连接之间的交换机，目的是提高同一个时刻多组GPU之间的通信带宽。举个例子，在8卡H100 GPU的服务器内部，同一时刻通过NVLink进行数据交换的GPU卡不能做到全量的带宽吞吐，比如有GPU0<->GPU1, GPU2<->GPU3同时交换数据时，所有GPU的带宽小于最大带宽900GB/s。增加NVSwitch之后的GPU之间能够以最大带宽进行数据交换，同一个时刻4对GPU以全量通信，总的带宽吞吐可达到3.6T/s。
