文章目录1 CUDA内存模型-多种可编程内存的类型1.1 多种内存的生命周期1.2 静态全局内存声明和赋值1.3 内存的管理和释放1.4 内存之间的带宽速度1.5 固定内存1.6 零拷贝内存1.6.1 使用零拷贝内存的方法1.7 统一内存寻址1.8 统一虚拟寻址1.9 内存访问模式, ♤♠1.9.1 对齐与合并内存加载1.9.2 非对齐与未合并内存加载1.9.3 全局内存加载和访问模式1.10 启用一级缓存加载模式最小单元是1281.10.1 未合并内存加载1.10.2 总线利用率很低的内存加载1.10.3 N个内存事务情况1.11 禁用一级缓存加载模式最小单元是321.11.1 内存对齐合并内存访问总线利用率还是100%用满1.11.2 未合并内存访问有可能提高总线利用率1.12 非对齐读取1.13 结构体数组和数组结构体1.13.1 AoS结构体数组1.13.2 SoA数组结构体1.14 循环展开1.15 尝试最佳的启动参数1.16 内存带宽1.16.1 内存带宽公式1.16.2 矩阵转置和内存带宽1.16.3 *对角坐标*1.16.3.1 分区冲突和bank冲突1.16.4 瘦块1.17 统一内存空间博主公司重组求推荐大模型部署ai infra base上海的工作。my phone 15601237103todo第三章还剩下最后一小节代码没有看。1 CUDA内存模型-多种可编程内存的类型寄存器共享内存本地内存常量内存纹理内存全局内存。线程有自己私有的本地内存。block有共享内存block中所有线程可见。所有的线程都可以访问全局内存纹理内存常量内存。pascal架构 中单个线程的最大寄存器个数是 2551.1 多种内存的生命周期存储器作用域生命周期存放内容其他片上片即SMRegisterThread和Kernel 函数生命周期相同无修饰的本地变量, 数组的常量索引编译器相关的限制字符 maxrregcount 64K个reg/SM1reg32bit是Local MemoryThreadKernel数组结构体本地内存和global memory在一块区域高延迟低带宽否Shared MemoryBlock线程块___shared__修饰高带宽低延迟可编程是Global MemoryGridApplication___device__或 cudamalloc延时最大内存事务32/64/128字节访问否ConstantGridApplication数学公式的系数__constant__修饰否纹理内存否L1 缓存 *这行是模型生成不确定正确todoblockblockblock的共享数据和局部缓存每个SM都有自己私有的L1缓存。-Xptxas -dlcmca 启用 L1 缓存或通过 -Xptxas -dlcmcg 禁用 L1是L2 缓存 *这行是模型生成不确定正确todoGlobalApp全局数据的缓存大小M级别否位于DRAM和和显存控制器之间设备无法直接访问主机内存反之设备也是。1.2 静态全局内存声明和赋值__device__floatdevData;//声明一个静态全局变量符号__global__voidcheckGlobalVariable(){devData2.0f;}intmain(void){floatvalue3.14f;// 从主机内存拷贝到设备变量符号CHECK(cudaMemcpyToSymbol(devData,value,sizeof(float)));checkGlobalVariable1,1();// 拷贝回主机CHECK(cudaMemcpyFromSymbol(value,devData,sizeof(float)));CHECK(cudaDeviceReset());returnEXIT_SUCCESS;}1.3 内存的管理和释放内存管理注意cudaMalloc分配设备上的全局内存cudaMallocHost分配主机的锁页内存页面锁定并且设备可访问释放必须用cudaFreeHostcudaHostAlloc分配主机的锁页内存 flag为cudaHostAllocDefault时这个和cudaMallocHost相同 flag为cudaHostAllocMapped时是零拷贝内存释放必须用cudaFreeHostcudaMallocManaged分配统一内存空间抽象的概念并不指定某个设备的内存主机和设备都可以直接访问无需显式调用cudaMemcpycudaMemsetcudaFreecudaMemcpycudaDeviceResetcudaDevicePropcudaSetDevice1.4 内存之间的带宽速度CPU和GPU之间通过PCIe连接之间的带宽速度只有 8GB/s。GPU芯片和GPU显存之间的带宽速度48 GB/s。1.5 固定内存从可分页主机内存传输数据到设备时候cuda驱动首先要分配临时的页面锁定内存。Page 256如果在主机端使用cudaMallocHost分配锁页内存那么从HtoD的拷贝16M的数据速度能够提升25%左右。1.6 零拷贝内存主机和设备都可以访问该内存可以减少H和D之间的显性传输。申请零拷贝内存用cudaHostAllocflag需要设置cudaHostAllocMapped。 零拷贝内存每次映射要经过PCIe总线比全局内存要慢。1.6.1 使用零拷贝内存的方法cudaHostAlloc申请内存指定flagcudaHostAllocMapped 获取主机pinned内存映射到设备可用的指针。CHECK(cudaHostAlloc((void**)h_A,nBytes,cudaHostAllocMapped));CHECK(cudaHostGetDevicePointer((void**)d_A,(void*)h_A,0));cuda用主机上的零拷贝内存会明显降低运算速度和申请主机分页内存HtoD做拷贝用设别内存运算拷贝回设备只有数据量很小10K时候影响小建议用。1.7 统一内存寻址1.8 统一虚拟寻址这两个小结和零拷贝内存有什么区别todo1.9 内存访问模式, ♤♠对全局内存的访问只使用二级缓存那么内存访问是由一个32字节的内存访问事务实现。对全局内存的访问使用一级和二级缓存那么内存访问是由一个128字节的内存访问事务实现。一级缓存一行是128个字节。二级缓存一行是 32个字节。1.9.1 对齐与合并内存加载下图一个线程warp只需要一个128字节的内存事务从设备内存读取数据因为warp中所有线程请求的地址都在128字节的缓存行内。总线利用率100%。1.9.2 非对齐与未合并内存加载下图一个线程warp需要3个128字节的内存事务从设备内存读取数据。1.9.3 全局内存加载和访问模式缓存加载启用一级缓存没有缓存的加载禁用一级缓存对齐与非对齐加载 第一个址是32字节的倍数合并与非合并加载线程束访问一个连续的数据块1.10 启用一级缓存加载模式最小单元是128启用一级缓存编译时候传递 -Xptxas -dlcmca1.10.1 未合并内存加载下图线程warp请求128字节的地址在落2个128字节段范围内当启用一级缓要求用两个128字节事务来执行总线利用率是50%加载的一半字节是未使用的。1.10.2 总线利用率很低的内存加载下图warp线程都在请求一个相同的地址只需要一个内存事务但是总线利用率很低。1.10.3 N个内存事务情况warp线程请求的地址分布在全局内存不连续地方完成一次内存加载需要N0N32个内存事务。1.11 禁用一级缓存加载模式最小单元是32加载的内存段更小32字节, 禁用一级缓存编译时候传递 -Xptxas -dlcmcg1.11.1 内存对齐合并内存访问总线利用率还是100%用满下图总线利用率100%128个字节请求占用4个内存段。1.11.2 未合并内存访问有可能提高总线利用率下图请求的128字节位于5个32字节段内总线利用率 4/580%。段更加细粒度加载的无效内存可以更少。1.12 非对齐读取4.3.2.3 这节提到内存索引加上一定的偏移地址会影响全局内存读取效率原因是全局内存没有对齐要多次读取但是我的显卡已经测试不到任何性能影响认为编译器已经优化掉了。1.13 结构体数组和数组结构体1.13.1 AoS结构体数组structinnerStruct{floatx;floaty;}structinnerStructmyAoS[N]1.13.2 SoA数组结构体structinnerArray{floatx[N];floaty[N];}数组结构体相同字段顺序存放全局带宽利用率更高。cuda编程倾向于SoA。对于AoS布局的内存加载请求和内存存储请求是重复的因为x和y在内存中是相邻存储的并且有相同大小被加载的一半也属于其他字段。因此加载和存储的50%带宽是未使用的。1.14 循环展开循环展开 可以增加每个线程更多的内存操作。gld_transaction和gst_transaction指标会更小说明读写的事务次数明显变少。__global__voidreadOffset(float*A,float*B,float*C,constintn,intoffset){unsignedintiblockIdx.x*blockDim.xthreadIdx.x;unsignedintkioffset;if(kn)C[i]A[k]B[k];}__global__voidreadOffsetUnroll4(float*A,float*B,float*C,constintn,intoffset){unsignedintiblockIdx.x*blockDim.x*4threadIdx.x;unsignedintkioffset;if(k3*blockDim.xn){C[i]A[k]B[k];C[iblockDim.x]A[kblockDim.x]B[kblockDim.x];C[i2*blockDim.x]A[k2*blockDim.x]B[k2*blockDim.x];C[i3*blockDim.x]A[k3*blockDim.x]B[k3*blockDim.x];}}1.15 尝试最佳的启动参数我的mx150显卡SM支持的最大并发线程数量是2048, warp中有32个线程那么SM支持的最大并发warp数量是64, 支持的最大并发block数量不明。 不同的kernel函数的配置会影响并发的warp数量通过achieved_occupancy可以看到该指标数值 所以配置不同的启动参数影响并行性。1.16 内存带宽1.16.1 内存带宽公式有效带宽(GB/s)(读字节数写字节数) / 运行时间秒 / 10^91.16.2 矩阵转置和内存带宽不同读写方式和速度指标CopyRow和CopyCol意味者内存带宽利用率的上下限分别是83% 34%因为iy * nx ix] 是行优先访问满足合并内存访问条件内存事务数量少宽带利用率高。[ix * ny iy] 是列优先访问不满足合并访问条件访问的内存是不连续的内存事务多宽带利用低。CopyRow elapsed0.000839secgrid(128,128)block(16,16)effective bandwidth39.993603GB/s CopyCol elapsed0.002030secgrid(128,128)block(16,16)effective bandwidth16.530125GB/s NaiveRow elapsed0.001969secgrid(128,128)block(16,16)effective bandwidth17.040501GB/s NaiveCol elapsed0.000830secgrid(128,128)block(16,16)effective bandwidth40.430187GB/s Unroll4Row elapsed0.001967secgrid(32,128)block(16,16)effective bandwidth17.057022GB/s Unroll4Col elapsed0.000841secgrid(32,128)block(16,16)effective bandwidth39.902889GB/s DiagonalRow elapsed0.001971secgrid(128,128)block(16,16)effective bandwidth17.024010GB/s DiagonalCol elapsed0.000886secgrid(128,128)block(16,16)effective bandwidth37.873383GB/s 这里的effective bandwidth计算方式是 有效的搬运字节数×2/kernel运行时间[有效字节数矩阵长×宽×sizeof(元素]transposeNaiveRow的 out[ix *ny iy] in[iy *nx ix] 是按照行主序读按照列主序写。__global__voidtransposeNaiveRow(float*out,float*in,constintnx,constintny){unsignedintixblockDim.x*blockIdx.xthreadIdx.x;unsignedintiyblockDim.y*blockIdx.ythreadIdx.y;if(ixnxiyny){out[ix*nyiy]in[iy*nxix];}}transposeNaiveCol分析 out[iy *nx ix] in[ix *ny iy] 这行是列主序读行主序写。读取的数据虽然这次没有被使用但是仍然存在缓存中下次触发了缓存命中。禁用L1缓存之后-Xptxas -dlcmcgeffective bandwidth从40GB/s掉到 33 GB/s 。__global__voidtransposeNaiveCol(float*out,float*in,constintnx,constintny){unsignedintixblockDim.x*blockIdx.xthreadIdx.x;unsignedintiyblockDim.y*blockIdx.ythreadIdx.y;if(ixnxiyny){out[iy*nxix]in[ix*nyiy];}}下面是打开L1缓存之后的一些指标gld_throughput,gst_throughput, gld_efficiency,gst_efficiency可以分析出列主序存储的效率列读写25.00%12.50%的效率明显不高也许是L1的缓存原因列读效率更高25.00%。transposeNaiveCol的gld_throughput 虽然很大76.575GB/s但是gld_efficiency却不高说明虽然读的速度很高但是有效读的数据很少。Kernel:transposeNaiveRow(float*,float*,int,int)gld_throughput Global Load Throughput7.9960GB/s7.9960GB/s7.9960GB/s gst_throughput Global Store Throughput63.968GB/s63.968GB/s63.968GB/s Kernel:transposeNaiveCol(float*,float*,int,int)gld_throughput Global Load Throughput76.575GB/s76.575GB/s76.575GB/s gst_throughput Global Store Throughput19.144GB/s19.144GB/s19.144GB/s Kernel:transposeNaiveRow(float*,float*,int,int)gld_efficiency Global Memory Load Efficiency100.00%100.00%100.00%gst_efficiency Global Memory Store Efficiency12.50%12.50%12.50%Kernel:transposeNaiveCol(float*,float*,int,int)gld_efficiency Global Memory Load Efficiency25.00%25.00%25.00%gst_efficiency Global Memory Store Efficiency100.00%100.00%100.00%1.16.3对角坐标从直角坐标系到对角坐标系有一层映射大概理解就是水平平移blockIdx.y个位置y轴方向换成x的坐标。原来的x坐标相当于在位置blk_x blockIdx.x % gridDim.x即blk_x blockIdx.x现在是在blk_x (blockIdx.x blockIdx.y) % gridDim.x__global__voidtransposeDiagonalRow(float*out,float*in,constintnx,constintny){unsignedintblk_yblockIdx.x;unsignedintblk_x(blockIdx.xblockIdx.y)%gridDim.x;unsignedintixblockDim.x*blk_xthreadIdx.x;unsignedintiyblockDim.y*blk_ythreadIdx.y;if(ixnxiyny){out[ix*nyiy]in[iy*nxix];}}这个实现我理解的不是很好对角坐标系看代码应该是给原来的线程块分配了不同位置的内存或者说在内存块和线程块之间做了位置映射。我在显卡mx150上没有测到DiagonalRow比NaiveRow有什么实质的性能提升但是据书本上说直角坐标系排布block请求的全局内存无法均匀分配到DRAM从分区中这就有可能发生分区冲突。对角坐标系映射后block请求的内存会更加均匀的分配到整个DRAM从分区中因此会有性能提升这里是不是和oneflow公司tranpose 做padding的思路是类似的。这里涉及到一个知识点1.16.3.1。1.16.3.1 分区冲突和bank冲突分区冲突DRAM partition conflict) , 全局内存按照一定字节划分到不同分区如果多个内存访问地址在同一个分区这些请求需要串行执行。bank冲突 Bank conflict 共享内存在直角坐标的x轴方向划分到不同bank不同请求地址落在同一个bank需串行执行另外需要多个内存事务。所以这是发生在不同硬件上的冲突。1.16.4 瘦块使用不同的kernel启动配置也许可以获得小幅度的性能提升比如使用瘦块block.xblock.y )可以提高存储效率。NaiveCol elapsed0.000826secgrid(256,64)block(8,32)effective bandwidth40.6GB/s NaiveCol elapsed0.000832secgrid(128,128)block(16,16)effective bandwidth40.3GB/s不同block形状的带宽性能数据尤其efficiency差异大最后还是要优先考虑耗时grid(256,64)block(8,32)Kernel:transposeNaiveCol(float*,float*,int,int)1gst_throughput Global Store Throughput38.573GB/s38.573GB/s38.573GB/s1gld_throughput Global Load Throughput38.573GB/s38.573GB/s38.573GB/s1gld_efficiency Global Memory Load Efficiency50.00%50.00%50.00%1gst_efficiency Global Memory Store Efficiency50.00%50.00%50.00%grid(128,128)block(16,16)Kernel:transposeNaiveCol(float*,float*,int,int)1gst_throughput Global Store Throughput19.135GB/s19.135GB/s19.135GB/s1gld_throughput Global Load Throughput76.541GB/s76.541GB/s76.541GB/s1gld_efficiency Global Memory Load Efficiency25.00%25.00%25.00%1gst_efficiency Global Memory Store Efficiency100.00%100.00%100.00%1.17 统一内存空间apicudaMallocManaged分配位置 抽象概念不具体哪块空间会根据访问需要切换。访问方式 主机和设备都可以直接访问。自动管理 CUDA运行时会自动管理内存的迁移无需显式调用cudaMemcpy。缺点看书上cudaMallocManaged要比手动管理内存的方式慢。