cudaMallocPitch 和 cudaMemcpy2D

一、cudaMalloc()

cudaMalloc(void** devPtr, size_t cout);

devPtr: 在显存上分配数据的头指针

cout: 分配空间的大小，以字节为单位。

 

   在设备上分配count字节的线性内存，并返回分配内存的指针*devPtr。分配的内存适合任何类型的变量。如果分配失败，cudaMalloc()返回cudaErrorAlloction。

 

   例如，要分配行数为height，列数为width的float型矩阵空间

int height = 64,  width =64;

float *devPtr;

cudaMalloc((void**)&devPtr, 

sizeof(float) * width *height)

二、cudaMemcpy()

cudaMemcpy(void*dst,  

         const void*src,  

         size_t count, 

         enum cudaMemcpyKind kind);

dst: 目的矩阵内存头指针

src：源矩阵内存头指针

count: 拷贝数据的大小

kind：拷贝数据的方向

 

   从src指向的内存区域拷贝count字节的数据到dst指向的内存区域，kind表示拷贝方向：cudaMemcpyHostToHost,cudaMemcpyDeviceToHost, cudaMemcpyHostToDevice, 或者cudaMemcpyDeviceToDevice.

   例如：从主机内存拷贝一个行数为height，列数为width的float型矩阵到设备显存。

int height = 64,  width =64;

……..

float*src;//内存中的矩阵头指针，假设矩阵已经初始化

float *dst;//指向设备矩阵的头指针，假设设备内存空间已经开辟

…….

cudaMemcpy(dst, 

        src, 

        count, 

        cudaMemcpyHostToDevice);

 

三、cudaMallocPitch()

cudaMallocPitch(void**devPtr,  

              size_t*pitch,  

              size_twidthInBytes,  

              size_theight)

devPtr：开辟矩阵的数据的头指针

pitch：分配存储器的宽度，以字节为单位(cuda的返回值)

width：分配矩阵的列数

height：分配矩阵的行数

 

   在设备上分配widthInBytes * height字节的线性内存，并返回分配内存的指针*devPtr。函数将确保在任何给出的行中对应的指针是连续的。pitch返回的指针*pitch是分配的宽度。Pitch作为内存分配的一个分开的参数，用来计算2D数组中的地址。一个给定行和列的类型T的数组元素，地址等于：

    T*pElement = (T*)((char*)BaseAddress + Row * pitch) +Column;

   对于2D数组的分配，建议使用cudaMallocPitch()分配内存。由于pitch对列限制受限于硬件，特别是当应用程序从设备内存的不同区域执行一个2D的内存拷贝。

 

   例如：要在GPU上开辟一个行数为height，列数为width的float型矩阵空间。

Int width = 64, height = 64;

float* devPtr;

size_t pitch;

cudaMallocPitch((void**)&devPtr, 

&pitch, 

width *sizeof(float), 

height)

 

假设GPU中global memory被划分为128Byte的段（0-127，128-255,256-383,……），你要为内存分配float型矩阵数据。假设矩阵的行数为N，列数为M，在C和C++中数据是按行优先存储的。当N*4是128倍数的时候(float型数据占4个Byte)，那么用cudaMalloc分配出来的内存空间也是对齐的，也就是说你一行分配的字节数刚好是128Bytes的倍数的时候，cudaMalloc()也是对齐的；

再看另一种情况，N=33，此时33*4=132，不是128的倍数，当warp从第二行开头（从132开始读，第一行是0到131）读取globalmemory的时候，首地址并非globalmemory划分的对齐段的首地址，那么这样的访问就是非合并的，cudaMallocPitch()就是为了解决每行首地址是否是globalmemory对齐段的问题，如果用cudaMallocPitch()来分配N=33(即列数为33)的矩阵时，每一行大小会变成256个Bytes(0-131为我们需要使用的空间，132-255未使用)，而不是cudaMalloc中的132个Bytes，这样分配以后，每行的首地址将会是与globlamemory分段地址对齐的（都是128的整数倍），warp在访问的时候就可以对齐了！

四、cudaMemcpy2D()

cudaMemcpy2D(void* dst,

size_tdpitch,

const void*src,

size_tspitch,

size_twidth,

size_theight,

enum cudaMemcpyKindkind);

                  dst: 目的矩阵内存头指针

dpitch: dst指向的2D数组中的内存宽度，以字节为单位，是cuda为了读取方便，

对齐过的内存宽度，可能大于一行元素占据的实际内存。

src：源矩阵内存头指针

spitch: src指向的2D数组中的内存宽度，以字节为单位

width: src指向的2D数组中一行元素占据的实际宽度。以字节为单位，等于

width*sizeof(type)

     height: src指向的2D数组的行数。

kind：拷贝数据的方向

   从src指向的内存区域拷贝数据到dst指向的内存区域。kind表示拷贝方向：cudaMemcpyHostToHost,cudaMemcpyDeviceToHost, cudaMemcpyHostToDevice, 或者cudaMemcpyDeviceToDevice。