CUDA -- 内存分配

 

　　cuda可以认为是一个由软件和硬件构成的并行计算系统，其依赖于gpu的并行计算单元，cuda有类c的api，方便程序编写。其依赖于cpu和gpu的异构体系，通过在cpu上串行执行环境初始化、内存分配、数据传输，然后在gpu上执行并行计算。

内存分配

　　1、一维

int *dev_ans = 0;
cudamalloc((void**)&dev_ans, d.y * sizeof(int));

　　参数1：显存中开辟的空间的指针（术语：gpu设备端数据指针）

　　参数2：空间大小，字节为单位

　　2、二维

int *dev_mat = 0;
int pitch;
cudamallocpitch((void**)&dev_mat, (size_t *)&pitch, d.x * sizeof(int), d.y);

　　参数1：gpu设备端数据指针

　　参数2：一行数据的真实空间大小（字节）【此参数是获取返回值】，gpu中从256字节对齐的地址（address=0,256,512……）连续访问最有效率，故每行实际分配的大小要大于需要分配的大小

　　参数3：每行需要分配的空间大小

　　参数4：矩阵行数

 

内存拷贝

　　1、一维

cudamemcpy(ans, dev_ans, d.y * sizeof(int), cudamemcpydevicetohost);

　　参数1：目标数据地址

　　参数2：源数据地址

　　参数3：数据大小

　　参数4：拷贝类型（主机至主机，主机至设备，设备至主机，设备至设备）

　　2、二维

cudamemcpy2d(dev_mat, pitch, mat, d.x*sizeof(int), d.x*sizeof(int), d.y, cudamemcpyhosttodevice);

　　参数1：目标数据地址

　　参数2：pitch，分配空间的行宽（字节单位）

　　参数3：源数据地址

　　参数4：pitch，分配空间的行宽（字节单位）

　　参数5：需要拷贝数据的真实行宽（字节单位）

　　参数6：数据的行数（非字节单位哦！）

　　参数7：数据拷贝类型

　　注：pitch是线性存储空间的行宽不是数据的行宽，在设备端 pitch大于等于数据行宽，在主机端pitch==数据行宽。

 

内存访问

　　主机中的内存访问就是c++的访存没什么好说的，现在看看显存中的访问方式（也就是在kernel中的访存）。

__global__ void addkernel(int *mat, int *ans, size_t pitch)
{
    int bid = blockidx.x;
    int tid = threadidx.x;
    __shared__ int data[8];
    int *row = (int*)((char*)mat + bid*pitch);
    data[tid] = row[tid];
    __syncthreads();
    for (int i = 4; i > 0; i /= 2) {
        if (tid < i)
            data[tid] = data[tid] + data[tid + i];
        __syncthreads();
    }
    if (tid == 0)
        ans[bid] = data[0];
}

　　一维：

　　　　ans[index]直接访问

　　二维：

　　　　先计算访问的行的初始地址　int *row = (int*)((char*)mat + bid*pitch)

　　　　然后访问此行的对应元素 row[index]

 

内存释放

cudafree(dev_mat)