有关padding和pitch

最新推荐文章于 2024-06-11 16:43:37 发布

fb_help

最新推荐文章于 2024-06-11 16:43:37 发布

阅读量1.1k

点赞数 1

CC 4.0 BY-SA版权

分类专栏： CUDA 文章标签：笔记

本文链接：https://blog.youkuaiyun.com/fb_help/article/details/79796796

CUDA 专栏收录该内容

51 篇文章

订阅专栏

本文深入探讨了CUDA中Padding和Pitch的概念与作用，详细解释了Padding如何通过填充确保二维数组对齐，以及Pitch如何作为二维数组的领先维度来优化内存访问效率。文章还介绍了与Pitch相关的CUDA API，并通过实例说明了Padding的重要性。

摘要生成于 C知道，由 DeepSeek-R1 满血版支持，前往体验 >

#有关padding和pitch
首先解释Padding和Pitch
##Padding：
即填充，在cuda程序中，其用在两个地方

对二维数组的逻辑行进行填充
为防止shared memory的冲突而填充shared memory

本篇只介绍第一中用法。
##Pitch（间距）：
对于二维数组，c/c++使用行主导的方式，即行控制数组的访问
（如：baseAddress + N×row+col）
Pitch在这里就可以是：the leading dimension of an array A（called Ida）

##为什么要Padding
主要针对二维数组的对齐

为什么使用二维内存：
在CUDA中，二维内存在物理上也是以一维形式存储的，所以二维数组其实是逻辑上的二维数组，因为有些操作（如图片操作）与二维逻辑有关，如卷积。所以使用二维数组。

在逻辑中二维内存的访问中，为保证对齐对内存访问的对齐，我们必须保证我们的leading dimension是内存事物数（如32字节，L1：128字节）的整倍数。这样我们的内存操作就会是对齐的（这在处理跟图像行相关的问题时很有效），如31×31的二维数组，我们将其padding（填充）为32×31的二维数组（32就是pitch），定义一个add操作，第一行+1，第二行+2。。。第n行+n
如果padding后，每一行都对齐访问，内存效率是100%，但如果不处理，那么第一行是0～31，只处理0～30，而在处理第二行的时候，需要第一行剩下的那个元素，此时还需要多一个内存事物去取剩下的那一个元素，这会导致内存效率不是100%。

##pitch相关的CUDA API

申请二维内存：
cudaMallocPitch((void**)&devPtr, size_t* &pitch, size_t width*sizeof(type), size_t height);
在主机和设备之间copy二维内存：
cudaMemcpy2D(void * dst, size_t dpitch, const void * src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind)

##补充
内存事物的大小跟设备和起步启用一级缓存有关，而内存效率和内存事物大小和每个warp（线程束）合并的内存请求的大小和分布有关。内存效率 =
每个warp请求的内存量/实际的内存事物数×内存事物的大小

##实例
关于padding的实验及其相关内容，在下一篇实验