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)
本文来自网络收集，不代表计算机技术网立场，如涉及侵权请联系管理员删除。
ctvol管理联系方式QQ:251552304
本文章地址：https://www.ctvol.com/c-cdevelopment/601954.html