GPU包含数百个数学计算单元，具有强大的处理运算能力，可以强大到计算速率高于输入数据的速率，即充分利用带宽，满负荷向GPU传输数据还不够它计算的。CUDA C除全局内存和共享内存外，还支持常量内存，常量内存用于保存在核函数执行期间不会发生变化的数据，使用常量内存在一些情况下，能有效减少内存带宽，降低GPU运算单元的空闲等待。
使用常量内存提升性能
使用常量内存可以提升运算性能的原因如下：
- 对常量内存的单次读操作可以广播到其他的“邻近(nearby)”线程，这将节约15次读取操作；
- 高速缓存。常量内存的数据将缓存起来，因此对于相同地址的连续操作将不会产生额外的内存通信量；
在CUDA架构中，线程束是指一个包含32个线程的集合，这个线程集合被“编织在一起”并且以“步调一致(Lockstep)”的形式执行。
当处理常量内存时，NVIDIA硬件将把单次内存读取操作广播到每个半线程束(Half-Warp)。在半线程束中包含16个线程，即线程束中线程数量的一半。如果在半线程束中的每个线程从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求并在随后将数据广播到每个线程。如果从常量内存中读取大量数据，那么这种方式产生的内存流量只是使用全局内存时的1/16。
常量内存的声明
为普通变量分配内存时是先声明一个指针，然后通过cudaMalloc()来为指针分配GPU内存。而当我们将其改为常量内存时，则要将这个声明修改为在常量内存中静态地分配空间。我们不再需要对变量指针调用cudaMalloc()或者cudaFree()，而是在编译时为这个变量（如一个数组）提交固定的大小。首先用“___constant_”声明一个常量内存变量，然后使用cudaMemcpyToSymbol（而不是cudaMemcpy）把数据从主机拷贝到设备GPU中。
常量内存使用示例
以下程序用CUDA+OpenCv实现一个简单场景的光线跟踪，光线跟踪是从三维场景生成二维图像的一种方式。主要思想为：在场景中选择一个位置放上一台假想的相机，该相机包含一个光传感器来生成图像，需要判断那些光将接触到这个传感器。图像中每个像素与命中传感器的光线有相同的颜色和强度。传感器中命中的光线可能来自场景中的任意位置，想象从该像素发出一道射线进入场景中，跟踪该光线穿过场景，直到光线命中某个物体。代码实现：
#include "cuda_runtime.h"
#include <highgui/highgui.hpp>
#include <time.h>
using namespace cv;
#define INF 2e10f
#define rnd(x) (x*rand()/RAND_MAX)
#define SPHERES 100 //球体数量
#define DIM 1024 //图像尺寸
struct Sphere
{
float r, g, b;
float radius;
float x, y, z;
__device__ float hit(float ox, float oy, float *n)
{
float dx = ox - x;
float dy = oy - y;
if (dx*dx + dy*dy < radius*radius)
{
float dz = sqrt(radius*radius - dx*dx - dy*dy);
*n = dz / sqrt(radius*radius);
return dz + z;
}
return -INF;
}
};
// Sphere *s;
__constant__ Sphere s[SPHERES];
/************************************************************************/
//__global__ void rayTracing(unsigned char* ptr, Sphere* s)
__global__ void rayTracing(unsigned char* ptr)
{
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int offset = x + y * blockDim.x * gridDim.x;
float ox = (x - DIM / 2);
float oy = (y - DIM / 2);
float r = 0, g = 0, b = 0;
float maxz = -INF;
for (int i = 0; i < SPHERES; i++)
{
float n;
float t = s[i].hit(ox, oy, &n);
if (t > maxz)
{
float fscale = n;
r = s[i].r * fscale;
g = s[i].g * fscale;
b = s[i].b * fscale;
maxz = t;
}
}
ptr[offset * 3 + 2] = (int)(r * 255);
ptr[offset * 3 + 1] = (int)(g * 255);
ptr[offset * 3 + 0] = (int)(b * 255);
}
/************************************************************************/
int main(int argc, char* argv[])
{
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
Mat bitmap = Mat(Size(DIM, DIM), CV_8UC3, Scalar::all(0));
unsigned char *devBitmap;
(cudaMalloc((void**)&devBitmap, 3 * bitmap.rows*bitmap.cols));
// cudaMalloc((void**)&s, sizeof(Sphere)*SPHERES);
Sphere *temps = (Sphere*)malloc(sizeof(Sphere)*SPHERES);
srand(time(0)); //随机数种子
for (int i = 0; i < SPHERES; i++)
{
temps[i].r = rnd(1.0f);
temps[i].g = rnd(1.0f);
temps[i].b = rnd(1.0f);
temps[i].x = rnd(1000.0f) - 500;
temps[i].y = rnd(1000.0f) - 500;
temps[i].z = rnd(1000.0f) - 500;
temps[i].radius = rnd(100.0f) + 20;
}
// cudaMemcpy(s, temps, sizeof(Sphere)*SPHERES, cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(s, temps, sizeof(Sphere)*SPHERES);
free(temps);
dim3 grids(DIM / 16, DIM / 16);
dim3 threads(16, 16);
// rayTracing<<<grids, threads>>>(devBitmap, s);
rayTracing << <grids, threads >> > (devBitmap);
cudaMemcpy(bitmap.data, devBitmap, 3 * bitmap.rows*bitmap.cols, cudaMemcpyDeviceToHost);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("Processing time: %3.1f ms\n", elapsedTime);
imshow("CUDA常量内存使用示例", bitmap);
waitKey();
cudaFree(devBitmap);
// cudaFree(s);
return 0;
}
程序里生成球体的大小和位置是随机的，为了产生随机数，加入了随机数种子srand（）。运行效果：
