多种CUDA存储单元详解CUDA中的存储单元种类

1. 寄存器
2. 局部存储器
3. 共享存储器
4. 全局存储器
5. 主机端内存
6. 主机端页锁定内存
7. 常数存储器
8. 纹理存储器

共享存储器

示例：共享存储器的动态与静态分配与初始化

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);

}