[cuda]常量内存

    我们在cuda编程中定义变量的时候，可以用__constant__来修饰定义，代表着定义了一个常量内存。向常量内存中拷贝数据的时候不再使用cudaMemcpy，而改为使用cudaMemcpyToSymbol。对于常量内存，不需要在使用完成后用cudafree释放空间。

    常量内存带来的性能提升

    __constant__把对变量的访问限制为只读，跟从全局内存读取数据相比，从常量内存读取相同的数据可以节约内存带宽，原因是对常量内存的单次读操作可以广播到同一个线程块内的其他线程(节约15次读操作)。且常量内存的数据将被缓存起来，对相同地址的连续读操作将不会产生额外的内存通信量。 

     为了方便理解，写了一个用常量内存给数组赋值的程序。

#include<cuda_runtime.h>
#include<windows.h>
#include<iostream>
using namespace std;
const int nMax = 50;

__constant__ float num1[40];
__global__ void exchangeKernel(float *aaa)
{
	int offset = threadIdx.x + blockDim.x * blockIdx.x;
	aaa[offset] = num1[offset];
}

int main(){
	float *devA,tmp[40],res[40];
	cudaMalloc((void**)&devA, 40*sizeof(float));
	for (int i = 0; i < 40; i++)tmp[i] = i*1.5f;
	cudaMemcpyToSymbol(num1, tmp, 40 * sizeof(float));
	exchangeKernel << <4, 10 >> >(devA);
	cudaMemcpy(res, devA, 40 * sizeof(float), cudaMemcpyDeviceToHost);
	for (int i = 0; i < 40; i++){
		cout << res[i] << " " << endl;
	}
	cin >> res[1];
	return 0;
}