CUDA（三） 三种memory的活用

此博客为博主的自学笔记 ，欢迎大家共同交流，如果有错误的地方欢迎留言指正。

GPU编程时常用的高速访问内存有三种，分别为：register,   shared memory,   constant memory。

以下面一个程序段为例：

    设：  A_d[ 256 ] , B_d[ 256 * 256 ],  C_d[256]为global memory的内存空间。

     kernel函数中的程序语句：  

     for(int i = 0; i < 256; i++){

 A_d[ threadIdx.x ] = A_d [ threadIdx.x ]  + B_d[ threadIdx.x * 256 + i  ] * C_d[ i ];

}

   这个是将B,C矩阵相乘再将结果存到A中。

    REGISTER的活用

    很明显，上面的写法对于A距阵的每一个元素进行计算时，都要进行256次global memory内存空间的访问。而global memory的访问速度又非常的缓慢，大大降低了性能。

    但是，通过register的活用，我们可以轻松解决这个问题。因为上面的例子，线程之间没有交叉的内存访问，所以，我们可以将它的循环中的中间结果保存于register中的变量中。这样可以大大地加快访问速度，从而提高性能。  

   float temp = 0;

   for(int i = 0; i < 256; i++){

 temp = temp + B_d[ threadIdx.x * 256 + i  ] * C_d[ i ];

}

   A_d[ threadIdx.x ] = temp;

    SHARED MEMORY的活用

    不难看出，上面的程序语句中，C_d只有一维数组，所占内存空间较小，且被频繁访问， 这样我们可以将他复制到空间有限的shared memory当中，加快访问的速度，以提高性能。

    __shared__ float C[256];

    C[ threadIdx.x] = C_d[ threadIdx.x];

    for(int i = 0; i < 256; i++){

 A_d[ threadIdx.x ] = A_d [ threadIdx.x ]  + B_d[ threadIdx.x * 256 + i  ] * C[ i ];

}

    CONSTANT MEMORY的活用

     如果shared memory的空间不足，我们可以使用其它访问速度较快的内存来代替。因为此次的C_d只是被当做乘数来使用，无需要修改，所以我们可以将他保存在同样访问速度较快的只读constant memory中。只是他的声明和初始化同shared memory有些不同，不是在kernel函数中，而是主机侧。

    主机侧：

      __constant__   float C_c[ 256 ];

     初始化数组C之后，通过以下语句将C复制到C_c中

    cudaMemcpyToSymbol(C_c, C,  256*sizeof(float));

   GPU侧：

   for(int i = 0; i < 256; i++){

 A_d[ threadIdx.x ] = A_d [ threadIdx.x ]  + B_d[ threadIdx.x * 256 + i  ] * C_c[ i ];

}

      通过上述步骤，计算性能会有很大的提高。

      尽量充分利用这三种内存访问方式，以便发挥出GPU的最大性能。