CUDA 存储器种类分析及使用方法指南.docVIP

CUDA 存储器种类分析及使用方法指南 寄存器 局部存储器 共享存储器 全局存储器 主机端内存 主机端页锁定内存 常数存储器 纹理存储器 存储器 位置 拥有缓存 访问权限 变量生存周期 register GPU （芯）片内 N/A device 可读/写 与thread相同 local memory 板载显存 无 device 可读/写 与thread相同 shared memory GPU 片内 N/A device 可读/写 与block相同 constant memory 板载显存 有 device 可读，host可读/写 可在程序中保持 texture memory 板载显存 有 device 可读，host可读/写 可在程序中保持 global memory 板载显存 无 device 可读/写，host可读/写 可在程序中保持 host memory host 内存 无 host 可读/写 可在程序中保持 pinned memory host 内存 无 host 可读/写 可在程序中保持 共享存储器 示例：共享存储器的动态与静态分配与初始化 int main(int argc, char** argv) { testKernel1, 10, mem_size(d_idata, d_odata); CUT_EXIT(argc, argv); } __global__ void testKernel(float* g_idata, float* g_odata){ extern __shared__ float sdata_dynamic[]; //extern 声明，大小由主机端程序决定。动态声明 __shared__ int sdata_static[16]; //静态声明数组大小 sdata_static[tid] = 0; //shared memory 不能在定义时初始化 } 将共享存储器中的变量声明为外部数组时，数组的大小将在Kernel 启动时确定，通过其执行参数确定。通过这种方式定义的所有变量都开始于相同的地址，因此数组中的变量的布局必须通过偏移量显式管理。例：如果希望在动态分配的共享存储器内获得与以下代码对应的内容： short array0[128]; float array1[64]; int array2[256]; 应该按照下面的方式定义： extern __shared__ char array[]; __device__ void func() { short* array0 = (short*)array; float* array1 = (float*)array0[128]; int* array2 = (int*)array1[64]; } 全局存储器 显存中的全局存储器也称为线性内存。线性内存通常使用 cudaMalloc() 函数分配， cudaFree() 函数释放，并由 cudaMemcpy() 进行主机端与设备端的数据传输。通过CUDA API分配的空间未经过初始化，初始化全局存储器需要调用 cudaMemset 函数。 对于二维、三维数组，我们使用 cudaMallocPitch() 和 cudaMalloc3D() 分配线性存储空间。这些函数能够确保分配满足对齐要求。 例：分配一个尺寸为 width * height 的 float 型2D 数组，以及遍历数组元素。 //主机端代码 float* devPtr; int pitch; cudaMallocPitch((void**)devPtr, pitch, width * sizeof(float), height); myKernel100, 512(devPtr, pitch); //设备端代码 __global__ void myKernel(float* devPtr, int pitch) { for (int r = 0; r height; ++r){ float* row = (float*)((char*)devPtr + r * pitch); for (int c = 0; c width; ++c){ float element = row[c]; } } } 例：分配一个 width * height * depth 的 float 型3D 数组，以及遍历数组元素。 //主机端代码 cudaPitchedPtr devPitchedPtr; cudaExtent extent = make_cudaExtent(64, 64, 64); cudaMalloc3D(devPitchedPtr, extent); m