存储器位置拥有缓存访问权限变量生存周期registerGPU （芯）片内N/Adevice 可读/写与thread相同local memory板载显存无device 可读/写与thread相同shared memoryGPU 片内N/Adevice 可读/写与block相同constant memory板载显存有device 可读，host可读/写可在程序中保持texture memory板载显存有device 可读，host可读/写可在程序中保持global memory板载显存无device 可读/写，host可读/写可在程序中保持host memoryhost 内存无host 可读/写可在程序中保持pinned memoryhost 内存无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 texturetexture<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));//启动 kerneldim3 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);}
