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)

 

  

 

posted @ 2019-09-11 14:38 茶飘香~ 阅读(...) 评论(...) 编辑 收藏