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){
testKernel<<<1, 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);
myKernel<<<100, 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);
myKernel<<<100, 512>>>(devPitchedPtr, extent);

//设备端代码
__global__ void myKernel(cudaPitchedPtr devPitchedPtr, cudaExtent extent){
char* devPtr = devPitchedPtr.ptr;
size_t pitch = devPitchedPtr.pitch;
size_t slicePitch = pitch * extent.height;
for (int z = 0; z < extent.depth; ++z){
char* slice = devPtr + z * slicePitch;
for (int y = 0; y < extent.height; ++y){
float* row = (float*)(slice + y * pitch);
for (int x = 0; x < extent.width; ++x){
float element = row[x];
}
}
}
}

例：二维数组和CUDA数组间的数据拷贝。

cudaMemcpy2DToArray(cuArray, 0, 0, devPtr, pitch, width * sizeof(float), height, cudaMemcpyDeviceToDevice);

主机端页锁定内存
通过 cudaHostAlloc() 和 cudaFreeHost() 来分配和释放 pinned memory。

常数存储器
定义常数存储器时，需要将其定义在所有函数之外，作用范围为整个文件，并且对主机端和设备端函数都可见。下面两段代码说明了两种常数存储器的使用方法。

第一种方法是直接在定义时直接初始化常熟存储器，然后再Kernel里面直接使用就可以了。
__constant__ int t_HelloCUDA[11] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; __constant__ int num = 11
第二种方法是定义一个costant 数组，然后使用函数进行赋值。
__constant__ char p_HelloCUDA[11];
CUDA_SAFE_CALL(cudaMemcpyToSymbol(p_HelloCUDA, helloCUDA, sizeof(char) * 11));
纹理存储器

在Kernel 中访问纹理存储器的操作称为纹理拾取（texture fetching）。纹理拾取使用的坐标与数据在显存中的位置可以不同，我们通过纹理参照系（texture reference）约定二者的映射方式。将显存中的数据与纹理参照系关联的操作，称为将数据与纹理绑定（texture binding）。显存中可以绑定到纹理的数据有两种，分别是普通的线性存储器（Linear Memory）和 CUDA 数组（CUDA Array）。
纹理存储器的使用

声明CUDA数组，分配空间

声明纹理参照系

设置运行时纹理参照系属性

纹理绑定

纹理拾取

例：简单的纹理使用。

/*
声明纹理参照系：texture<Type, Dim, ReadMode> texRef;
*/
//2D float texture
texture<float, 2, cudaReadModeElementType> texRef;
//设备端代码，一个简单的转换kernel
__global__ void transformKernel(float* output, int width, int height, float theta){
//计算归一化的纹理坐标
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
float u = x / (float)width;
float v = y / (float)height;
//坐标转换
u -= 0.5f;
v -= 0.5f;
float tu = u * cosf(theta) - v * sinf(theta) + 0.5f;
float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;
//读纹理并向全局存储器写会
output[y * width + x] = tex2D(tex, tu, tv);
}

//主机端代码
int main()
{
//在先存上为CUDA array 分配空间
cudaChannelFormatDesc channelDesc = cudaCreatChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
cudaArray* cuArray;
cudaMallocArray(&cuArray, &channelDesc, width, height);
//内存中h_data地址处的数据向显存进行拷贝
cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);
//设置纹理参数
texRef.addressMode[0] = cudaAddressModeWrap;
texRef.addressMode[1] = cudaAddressModeWrap;
texRef.filterMode = cudaFilterModeLinear;
texRef.normalized = true;
//数组绑定到纹理
cudaBindTextureToArray(texRef, cuArray, &channelDesc);
//转换结果分配显存空间
float* output;
cudaMalloc((void**)&output, width * height * sizeof(float));
//启动 kernel
dim3 dimBlock(16, 16);
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y - 1) / dimBlock.y);
transformKernel<<<dimGrid, dimBlock>>>(output, width, height, angle);
//释放显存空间
cudaFreeArray(cuArray);
cudaFree(output);
}