CUDA线性内存分配

雨点打透心脏的1/2处 2022-08-08 03:29 169阅读 0赞

**概述：线性存储器可以通过cudaMalloc()、cudaMallocPitch()和cudaMalloc3D()分配**

**1、1D线性内存分配**

1 cudaMalloc(void**,int)    //在设备端分配内存
    2 cudaMemcpy(void* dest,void* source,int size,enum direction)    //数据拷贝
    3 cudaMemcpyToSymbol       //将数据复制到__constant__变量中，或者__device__变量中
    4 cudaMemcpyFromSynbol   //同上相反
    5 cudaFree()               //内存释放
    6 cudaMemset()           //内存初始化

注意：主机和设备间的数据交换会自动同步，而设备与设备却不会，需要使用cudaThreadSynchronize()

**2、2D线性内存分配**

**2.1 分配**

1 cudaMallocPitch( void** devPtr，size_t* pitch，size_t widthInBytes，size_t height ) //在线性内存中分配二维数组，width的单位是字节，而height单位是数据类型

c语言申请2维内存时，一般是连续存放的。a\[y\]\[x\]存放在第y\*widthofx\*sizeof(元素)+x\*sizeof(元素)个字节。

但**在cuda的global memory访问中，从256字节对齐的地址(addr=0, 256, 512, ...)开始的连续访问是最有效率的**。这样，为了提高内存访问的效率，有了cudaMallocPitch函数。cudaMallocPitch函数分配的内存中，数组的每一行的第一个元素的开始地址都保证是对齐的。因为每行有多少个数据是不确定的，widthofx\*sizeof(元素)不一定是256的倍数。故此，为保证数组的每一行的第一个元素的开始地址对齐，**cudaMallocPitch在分配内存时，每行会多分配一些字节，以保证widthofx\*sizeof(元素)+多分配的字节是256的倍数(对齐)**。这样，上面的y\*widthofx\*sizeof(元素)+x\*sizeof(元素)来计算a\[y\]\[x\]的地址就不正确了。而应该是y\*\[widthofx\*sizeof(元素)+多分配的字节\]+x\*sizeof(元素)。而函数中返回的pitch的值就是widthofx\*sizeof(元素)+多分配的字节。说明：**widthInBytes作为输入参数，应该是widthofx\*sizeof(元素)**；这样的话，复制内容时也要作相应的修改。

**2.2 访问**

1 T* pElement = (T*)((char*)BaseAddress + Row * pitch) + Column;           //元素访问方式

cudaMallocPitch()以\*pitch的形式返回间距，即所分配存储器的宽度，以字节为单位。间距用作存储器分配的一个独立参数，用于在2D数组内计算地址。

**2.3 拷贝**

1 cudaMemcpy2D( void* dst，size_t dpitch，const void* src，size_t spitch，size_t width，size_t height，enum cudaMemcpyKind kind )

这里需要特别注意width与pitch的区别，width是实际需要拷贝的数据宽度而pitch是2D线性存储空间分配时对齐的行宽，而当数据传递发生在设备与主机之间时，主机端pitch==width.

综上我们可以看到，CUDA下对二维线性空间的访问是不提供多下标支持的，访问时依然是通过计算偏移量得到，不同的地方在于使用pitch对齐后非常利于实现coalesce访问

例：下面的代码分配了一个尺寸为width\*height的二维浮点数组，同时演示了怎样在设备代码中遍历数组元素

[![复制代码][copycode.gif]][copycode.gif 1]

1 // Host code
     2   int width = 64, height = 64;
     3   float* devPtr;
     4   int pitch;
     5   cudaMallocPitch((void**)&devPtr, &pitch, width * sizeof(float), height); 
     6   MyKernel<<<100, 512>>>(devPtr, pitch, width, height);
     7 // Device code
     8   __global__ void MyKernel(float* devPtr, int pitch, int width, int height){
     9    for (int r = 0; r < height; ++r) {
    10       float* row = (float*)((char*)devPtr + r * pitch);
    11       for (int c = 0; c < width; ++c) {
    12          float element = row[c];
    13       }
    14    }
    15 }

[![复制代码][copycode.gif]][copycode.gif 1]

**3、3D线性内存**

1 cudaError_t cudaMalloc3D(    
    2     struct cudaPitchedPtr *     pitchedDevPtr,
    3     struct cudaExtent             extent     
    4 )

例：下面的代码分配了一个尺寸为width\*height\*depth的三维浮点数组，同时演示了怎样在设备代码中遍历数组元素

[![复制代码][copycode.gif]][copycode.gif 1]

1 // Host code
     2 cudaPitchedPtr devPitchedPtr;
     3 cudaExtent extent = make_cudaExtent(64, 64, 64);
     4 cudaMalloc3D(&devPitchedPtr, extent);  
     5 MyKernel<<<100, 512>>>(devPitchedPtr, extent);
     6 // Device code
     7 __global__ void MyKernel(cudaPitchedPtr devPitchedPtr, cudaExtent extent) {
     8    char* devPtr = devPitchedPtr.ptr;
     9    size_t pitch = devPitchedPtr.pitch;
    10    size_t slicePitch = pitch * extent.height;
    11    for (int z = 0; z < extent.depth; ++z) {
    12      char* slice = devPtr + z * slicePitch;
    13      for (int y = 0; y < extent.height; ++y) {
    14         float* row = (float*)(slice + y * pitch); 
    15         for (int x = 0; x < extent.width; ++x) { float element = row[x];
    16      }
    17    }
    18 }

[![复制代码][copycode.gif]][copycode.gif 1]

[copycode.gif]: http://common.cnblogs.com/images/copycode.gif
[copycode.gif 1]: /images/20220805/6aa7a366ce724579bf21381394b44b6a.png