GPU通信——内存

                                                                  GPU内存分类
全局内存：通俗意义上的设备内存。
共享内存
  1. 位置：设备内存。
  2. 形式：关键字shared添加到变量声明中。如shared float cache[10]。
  3. 目的：对于GPU上启动的每个线程块，CUDA C编译器都将创建该共享变量的一个副本。线程块中的每个线程都共享这块内存，但线程却无法看到也不能修改其他线程块的变量副本。这样使得一个线程块中的多个线程能够在计算上通信和协作。
常量内存
  1. 位置：设备内存
  2. 形式：关键字constant添加到变量声明中。如constant float s[10];。
  3. 目的：为了提升性能。常量内存采取了不同于标准全局内存的处理方式。在某些情况下，用常量内存替换全局内存能有效地减少内存带宽。
  4. 特点：常量内存用于保存在核函数执行期间不会发生变化的数据。变量的访问限制为只读。NVIDIA硬件提供了64KB的常量内存。不再需要cudaMalloc()或者cudaFree(),而是在编译时，静态地分配空间。
  5. 要求：当我们需要拷贝数据到常量内存中应该使用cudaMemcpyToSymbol()，而cudaMemcpy()会复制到全局内存。
  6. 性能提升的原因： 
      ○ 对常量内存的单次读操作可以广播到其他的“邻近”线程。这将节约15次读取操作。（为什么是15，因为“邻近”指半个线程束，一个线程束包含32个线程的集合。）
      ○ 常量内存的数据将缓存起来，因此对相同地址的连续读操作将不会产生额外的内存通信量。
纹理内存
  1. 位置：设备内存
  2. 目的：能够减少对内存的请求并提供高效的内存带宽。是专门为那些在内存访问模式中存在大量空间局部性的图形应用程序设计，意味着一个线程读取的位置可能与邻近线程读取的位置“非常接近”。如下图：
  3. 纹理变量（引用）必须声明为文件作用域内的全局变量。
  4. 形式：分为一维纹理内存 和 二维纹理内存。 
      ○ 一维纹理内存 
          ■ 用texture<类型>类型声明，如texture<float> texIn。
          ■ 通过cudaBindTexture()绑定到纹理内存中。
          ■ 通过tex1Dfetch()来读取纹理内存中的数据。
          ■ 通过cudaUnbindTexture()取消绑定纹理内存。
      ○ 二维纹理内存 
          ■ 用texture<类型,数字>类型声明，如texture<float，2> texIn。
          ■ 通过cudaBindTexture2D()绑定到纹理内存中。
          ■ 通过tex2D()来读取纹理内存中的数据。
          ■ 通过cudaUnbindTexture()取消绑定纹理内存。                             

                                   一：共享内存

对于GPU的编程来说，最重要的一个方面就是解决 并行执行的各个部分如何相互协作的问题，共享内存是其中的一种解决方案，

有必要来研究一下！共享内存中的变量（核函数中用__shared__声明），在GPU上启动的每个线程块编译器都创建该变量的副本，若启动N个线程块，则有N个该变量副本，为每个线程块私有，也就是说线程块中的每个线程共享这块内存；同步则是使线程块中所有的线程能够在执行完某些语句后，才执行后续语句。

(1)以下程序实现了点积运算，计算公式为 f(n) = 1+2*2+ 3*3+ … (n-1)*(n-1)，使用共享变量计算各个程序块内所有线程的求和运算结果。

#include <cuda_runtime.h>  

#include <iostream>  
//main1.cu
#include "book.h"

using namespace std;

#define N 33*1024          //数组长度

const int threadsPerBlock = 64;   //每个线程块的线程数量
const int blocksPerGrid = 64;     //第一维线程格内线程块数量

__global__ void add(float *a, float *b, float *c)
{       //在设备中开辟共享内存
	__shared__ float cache[threadsPerBlock];       //__shared__声明共享变量，每个线程块均有自己的副本，被其所
	                                             //有线程共享,这里用于存放每个线程块内各个线程所计算得的点积和                
	int index =threadIdx.x + blockIdx.x *blockDim.x;    //将线程块、线程索引转换为数组的索引
	int cacheIdx = threadIdx.x;

	float temp = 0;
	while (index < N){
		temp += a[index] * b[index];
		index += gridDim.x * blockDim.x;
	}

	cache[cacheIdx] = temp;      //存放每个线程块内各个线程所计算得的点积和

	__syncthreads();               //cuda内置函数，使所有线程均执行完该命令前代码，才执行后面语句，也即保持同步
	                               //目的为获得各个cache副本，此时共有64个cache副本
	
	//规约运算，将每个cache副本求和，结果保存于cache[0]
	int i = blockDim.x / 2;
	while (i != 0){
		if (cacheIdx < i){
			cache[cacheIdx] += cache[i + cacheIdx];
		}
		__syncthreads();    //所有线程完成一次规约运算，方可进行下一次
		i /= 2;
	}

	if (cacheIdx == 2)     //一个操作只需一个线程完成即可
		c[blockIdx.x] = cache[0]; //所有副本的cache[0] 存放于数组c

}

int main()
{
	float a[N], b[N];
	float *c = new float[blocksPerGrid];
	float *dev_a, *dev_b, *dev_c;

	//gpu上分配内存
	HANDLE_ERROR(cudaMalloc((void**)&dev_a, N*sizeof(float)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b, N*sizeof(float)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c, N*sizeof(float)));

	//为数组a,b初始化
	for (int i = 0; i < N; ++i){
		a[i] = i;
		b[i] = i;
	}

	//讲数组a,b数据复制至gpu
	(cudaMemcpy(dev_a, a, N*sizeof(float), cudaMemcpyHostToDevice));
	(cudaMemcpy(dev_b, b, N*sizeof(float), cudaMemcpyHostToDevice));
	
	add <<< blocksPerGrid, threadsPerBlock >> >(dev_a, dev_b, dev_c);

	//将数组dev_c复制至cpu
	HANDLE_ERROR(cudaMemcpy(c, dev_c, blocksPerGrid*sizeof(float), cudaMemcpyDeviceToHost));
	
	//进一步求和
	double sums = 0.0;
	for (int i = 0; i < blocksPerGrid; ++i){
		sums += c[i];
	}

	//显示结果
	cout << "gpu dot compute result:" << sums << "\n";

	sums = 0.0;
	for (int i = 0; i < N; ++i){
		sums += i*i;
	}
	cout << "cpu dot compute result:" << sums << "\n";

	//释放在gpu分配的内存
	cudaFree( dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);

	delete c;
	return 0;
}

(2)基于共享内存的位图：以下程序使用二维程序块共享变量计算图像数据，生成图像

//main2.cu
#include <cuda_runtime.h>  

#include <iostream>  

#include "book.h"

#include <opencv2/opencv.hpp>

using namespace cv;
using namespace std;

#define PI 3.1415926

#define DIM 1024       //灰度图像的长与宽

__global__ void kernel(uchar * _ptr )
{
	int x = threadIdx.x + blockIdx.x * blockDim.x;
	int y = threadIdx.y + blockIdx.y * blockDim.y;
	int idx = x + y *gridDim.x *blockDim.x;

	__shared__ float shared [16][16] ;   //每个线程块中每个线程的共享内存缓冲区
	const float period = 128.0f;
	shared[threadIdx.x][threadIdx.y] = 255 * (sinf(x*2.0f*PI / period) + 1.0f)*(sinf(y*2.0f*PI / period) + 1.0f) / 4.0f;

	__syncthreads();   //使所有shared副本均被计算完成

	_ptr[idx] = shared[15 - threadIdx.x][15 - threadIdx.y];

}

int main()
{
	Mat src(DIM,DIM , CV_8UC1 , Scalar::all(0));

	uchar *ptr_dev;
	HANDLE_ERROR(cudaMalloc((void**)&ptr_dev, DIM * DIM*sizeof(uchar)));

	dim3 blocks(DIM / 16, DIM / 16);
	dim3 threads(16 ,16);
	kernel << < blocks, threads >> >( ptr_dev );

	HANDLE_ERROR(cudaMemcpy(src.data, ptr_dev, DIM * DIM*sizeof(uchar), cudaMemcpyDeviceToHost));

	cudaFree(ptr_dev);

	namedWindow("Demo",  0);
	imshow("Demo" , src);
	waitKey(0);
	return 0;

}