一、cuda/gpu线程模型1、cuda线程模型核函数调用方法如下xxx_xxx_xxxgrid_size,block_size(函数参数);以上核函数调用会在gpu中创建grid_size * block_size个线程并行执行核函数内容一维时grid_size为一维时最大值block_size最大值1024三维时grid_size.x/grid_size.y/grid_size.z最大值/65535/65535block_size.x/block_size.y/block_size.z最大值1024/1024/64且block_size.x*block_size.y*block_size.z不能大于1024。2、GPU线程模型3、cuda线程模型与GPU线程模型对应关系二、cuda/gpu内存模型1、PageableMemoryHOST分页内存位于HOST DDR1、主机使用malloc()/free()2、主机使用GPU不能使用3、在stream中如果使用了异步数据传输不能使用该类型的Memory。2、PinnedMemory固定内存位于HOST DDR1、使用cudaMallocHost()/cudaHostAlloc()/cudaFreeHost()2、GPU可通过DMA访问的Host Memory3、在stream中如果使用了异步数据传输需要使用该类型的Memory。3、GlobalMemory全局内存位于GPU DDR1、全局内存cudaMallc()/cudaFree()2、静态全局内存变量核函数外__device__3、常量内存核函数外__constant__4、纹理内存/表面内存5、局部内存在核函数内定义的不加任何修定符的变量在寄存器内存装不下时。6、全局内存的合并传输4、SharedMemory共享内存位于GPU SM SRAM1、类似于寄存器内存区别是整个线程块可见2、静态共享内存使用__shared__修饰声明时需要指定大小3、动态共享内存使用extern __shared__修饰声明时不能指定大小在调用核函数时指定每个线程使用的共享内存大小xxx_xxx_xxxgrid_size, block_size,SharedMemorySize(函数参数);4、共享内存因为更靠近计算单元所以访问速度更快5、共享内存通常可以作为访问全局内存的缓存使用6、可以利用共享内存实现线程间的通信7、通常与__syncthreads同时出现这个函数是同步block内的所有线程全部执行到这一行才往下走8、使用方式通常是在线程id为0的时候从global memory取值然后syncthreads然后再使用。9、共享内存bank冲突问题5、RegisterMemory寄存器内存位于GPU SM1、在核函数内定义的不加任何修定符的变量。三、cuda runtime api以cuda为前缀CUDA Runtime API四、cuda driver api以cu为前缀CUDA Driver API五、原子函数atomicAdd、atomicSub、atomicExch、atomicCAS、atomicInc、atomicDec、atomicMax、atomicMin、atomicAnd、atomicOr、atomicXor六、线程束SM按照线程束包括32个线程进行运算一个线程块只能调度到一个SM如一个线程块有128个线程那就是4个线程束这时会被调度到一个SM上这个SM多度执行这4个线程束一个SM最多可以驻留32个线程束即最多32*321024个线程一个SM可以运行多个线程块如线程块有16个线程那么一个SM可以同时运行两个线程块。1、线程束内基本函数unsigned __ballot_sync(unsigned mask, int predicate);int __all_sync(unsigned mask, int predicate);int __any_sync(unsigned mask, int preducate);T __shfl_sync(unsigned mask, T v, int srcLane, int w warpSize);T __shfl_up_sync(unsigned mask, T v, unsigned d, int w warpSize);T __shfl_down_sync(unsigned mask, T v, unsigned d, int w warpSize);T __shfl_xor_sync(unsigned mask, T v, int laneMask, int w warpSize);__syncwarp();reduce_shfl();2、协作组-线程块片函数unsigned __ballot_sync(int predicate);int __all_sync(int predicate);int __any_sync(int predicate);T __shfl_sync(T v, int srcLane);T __shfl_up_sync(T v, unsigned d);T __shfl_down_sync(T v, unsigned d);T __shfl_xor_sync(T v, int laneMask);七、常用cuda库Thrust类似于C的标准模板库standard template librarycuBLAS基本线性代数子函数basic liner algebra subroutinescuFFT快速傅里叶变换fast Fourier transformscuSPARSE稀疏sparse矩阵cuRAND随机数生成器rendom number generatorcuSolver稠密dense矩阵和稀疏矩阵计算库cuDNN深度神经网络deep neural networks八、context、stream、event流是一种基于context之上的任务管道抽象未显式调用函数创建stream时使用默认streamdefault stream在一个stream里的执行步骤都是顺序的可以是阻塞也可以是非阻塞如下示例cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice); cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice); sumgrid_size, block_size(d_x, d_y, d_z, N); cudaMemcpy(h_z, d_z, cudaMemcpyDeviceToHost);stream并行实现计算与通信重叠通过创建stream可以同时运行多个CUDA stream多个stream可实现核函数和通信的并行执行stream的用法如下1、创建stream cudaError_t cudaStreamCreate(cudaStream_t *); // 创建stream cudaError_t cudaStreamDestroy(cudaStream_t); // 销毁stream 2、使用stream xxx_kernelgrid_size, block_size(函数参数): xxx_kernelgrid_size, block_size, SharedMemorySize(函数参数): xxx_kernelgrid_size, block_size, SharedMemorySize, stream_id(函数参数): 3、等待完成 cudaError_t cudaStreamSynchronize(cudaStream_t stream); // 阻塞等待完成 cudaError_t cudaStreamQuery(cudaStream_t stream); // 非阻塞检查完成状态 一个完成的示例如下 int main(){ int device_id 0; checkRuntime(cudaSetDevice(device_id)); cudaStream_t stream nullptr; checkRuntime(cudaStreamCreate(stream)); //创建一个流 // 在GPU上开辟空间 float* memory_device nullptr; checkRuntime(cudaMalloc(memory_device, 100 * sizeof(float))); // 在CPU上开辟空间并且放数据进去将数据复制到GPU float* memory_host new float[100]; memory_host[2] 520.25; // 异步复制操作主线程不需要等待复制结束才继续 checkRuntime(cudaMemcpyAsync(memory_device, memory_host, sizeof(float) * 100, cudaMemcpyHostToDevice, stream)); // 在CPU上开辟pin memory,并将GPU上的数据复制回来 float* memory_page_locked nullptr; checkRuntime(cudaMallocHost(memory_page_locked, 100 * sizeof(float))); checkRuntime(cudaMemcpyAsync(memory_page_locked, memory_device, sizeof(float) * 100, cudaMemcpyDeviceToHost, stream)); // 异步复制操作主线程不需要等待复制结束才继续 printf(%f\n, memory_page_locked[2]);//结果是0因为还没等待结果返回 checkRuntime(cudaStreamSynchronize(stream)); //统一等待流队列中的结果 printf(%f\n, memory_page_locked[2]); // 释放内存 checkRuntime(cudaFreeHost(memory_page_locked)); checkRuntime(cudaFree(memory_device)); checkRuntime(cudaStreamDestroy(stream)); delete [] memory_host; return 0; }event用以监控stream是否到达了某个检查点可实现多个stream同步举例说明如下在stream 1上面执行kernel 1 和 kernel 3在stream 2上面执行kernel 2但是必须等到stream 1上面的kernel 1执行结束之后才能开始// Create streams and event cudaStream_t stream1, stream2; cudaEvent_t event1; cudaEventCreate(event1); // Execute kernel1 in stream1 kernel1gridDim, blockDim, 0, stream1(d_data1, repeat); cudaEventRecord(event1, stream1); // Record event1 after kernel1 execution in stream1 // Execute kernel2 in stream2, waiting for event1 cudaStreamWaitEvent(stream2, event1, 0); kernel2gridDim, blockDim, 0, stream2(d_data1, repeat); // Execute kernel3 in stream1 on a different array kernel3gridDim, blockDim, 0, stream1(d_data2, repeat); // Synchronize streams cudaStreamSynchronize(stream1); cudaStreamSynchronize(stream2); // Free device memory and destroy streams and event cudaStreamDestroy(stream1); cudaStreamDestroy(stream2); cudaEventDestroy(event1);九、统一内存1、动态统一内存在HOST端使用cudaMallocManaged()函数进行申请使用cudaFree()函数释放HOST即Device端使用时无需进行Memory Copy而直接使用如下示例int main(void) { const int N 100000000; const int M sizeof(double) * N; double *x, *y, *z; CHECK(cudaMallocManaged((void **)x, M)); CHECK(cudaMallocManaged((void **)y, M)); CHECK(cudaMallocManaged((void **)z, M)); for (int n 0; n N; n) { x[n] a; y[n] b; } const int block_size 128; const int grid_size N / block_size; addgrid_size, block_size(x, y, z); CHECK(cudaDeviceSynchronize()); check(z, N); CHECK(cudaFree(x)); CHECK(cudaFree(x)); CHECK(cudaFree(x)); return 0 }2、静态统一内存使用__device__ __managed__进行修饰如下示例#include stdio.h __device__ __managed__ int ret[1000]; __globle__ void AplusB(int a, int b) { ret[threadIdx.x] a b threadIdx.x; } int main(void) { AplusB1, 1000(10, 100); CHECK(cudaDeviceSynchronize()); for (int i 0; i 1000; i) { printf(%d: AB %d\n, i, ret[i]); } return 0 }3、统一内存可以申请超出实际内存的容量因为实际使用时才会涉及到物理内存类似缺页异常申请的统一内存可以使用gpu_touch()核函数进行全部内存的初始化如果申请了超量的统一内存使用gpu_touch()核函数进行初始化时在超量的部分Memory会初始化失败4、统一内存如果涉及到HOST MemoryDevice Memory如果Device不访问这段统一内存那么HOST只能访问到HOST部分的内存及时让Device使用cudaMemAdvise()/cudaMemPreftchAsync()函数同步一下统一内存HOST才能访问到全部的统一内存参考代码如下int main(void) { int device_id 0; CHECK(cudaGetDevice(device_id)); const int N 100000000; const int M sizeof(double) * N; double *x, *y, *z; CHECK(cudaMallocManaged((void **)x, M)); CHECK(cudaMallocManaged((void **)y, M)); CHECK(cudaMallocManaged((void **)z, M)); for (int n 0; n N; n) { x[n] a; y[n] b; } const int block_size 128; const int grid_size N / block_size; CHECK(cudaMemPrefetchAsync(x, M, device_id, NULL)); CHECK(cudaMemPrefetchAsync(y, M, device_id, NULL)); CHECK(cudaMemPrefetchAsync(z, M, device_id, NULL)); addgrid_size, block_size(x, y, z); CHECK(cudaFree(x)); CHECK(cudaMemPrefetchAsync(z, M, cudaCpuDeviceId, NULL)); CHECK(cudaDeviceSynchronize()); check(z, N); CHECK(cudaFree(x)); CHECK(cudaFree(y)); CHECK(cudaFree(z)); return 0; }十、重点cuda多线程基础1、thread、block、grid和warpcuda内存模型2、内存模型3、统一内存可用于rdma吗4、多卡统一内存编程5、纹理内存与动态并行cuda同步6、多流编程event实现显示同步7、异步操作时怎么知道操作完成的8、事件同步与流同步9、节点间怎么划分流流之间怎么同步10、同步和原子操作11、对⽐__threadfence()、__threadfence_block()、__syncthreads()和cudaDeviceSynchronize()cuda卡间通信12、统一内存与cudaEnablePeerAccess()13、cuda程序调用ncclrdma场景编程使用nccl14、nvsharemem使用15、编程实现rdma跨卡传输显存数据16、cudaMemcpyAsync()与cudaMemcpyPeer()与多卡统一内存cuda其他17、张量tensor core编程18、rl详细训练推理流程多卡nccl训练和推理流程19、cuda编写深度学习训练代码20、onnx runtime。十一、疑问1、核函数怎么被加载到gpugpu怎么运行核函数2、runtime/driver api怎么与KMD交互3、卷积、激活、池化、归一化、标量、向量、矩阵、张量GPU软件抽象与硬件映射的理解Grid、Block、Warp、Thread与SM、SPGPU架构GPU架构与通信互联技术介绍CUDA运行APIRuntimeAPINVIDA CUDA-DirverAPI入门具体化讲解cuda runtime api是什么runtime driver是什么【YOLO系列】YOLOv5超详细解读源码详解入门实践改进一文读懂cuda stream与cuda event一文说清楚DeePSeek用的PTX与CUDA的区别自定义CUDA实现PyTorch算子的四种简单方法PhD 第四学期 年中随笔矩阵分解之: 特征值分解(EVD)、奇异值分解(SVD)、SVDcuda runtime/driver API解析【CUDA】Driver API 和 Runtime APIhttps://www.zhihu.com/people/li-wei-27-70-68/posts[CUDA]UVM机制的原理CUDA IPC 是什么CUDA Context IPC进程间通信