欢迎您访问程序员文章站本站旨在为大家提供分享程序员计算机编程知识!
您现在的位置是: 首页  >  IT编程

CUDA -- 内存分配

程序员文章站 2022-06-16 22:01:58
CUDA可以认为是一个由软件和硬件构成的并行计算系统,其依赖于GPU的并行计算单元,CUDA有类C的API,方便程序编写。其依赖于CPU和GPU的异构体系,通过在CPU上串行执行环境初始化、内存分配、数据传输,然后在GPU上执行并行计算。 内存分配 1、一维 参数1:显存中开辟的空间的指针(术语:G ......

 

  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)