CUDA中的常量内存__constant__

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实现一个简单场景的光线跟踪，光线跟踪是从三维场景生成二维图像的一种方式。主要思想为：在场景中选择一个位置放上一台假想的相机，该相机包含一个光传感器来生成图像，需要判断那些光将接触到这个传感器。图像中每个像素与命中传感器的光线有相同的颜色和强度。传感器中命中的光线可能来自场景中的任意位置，想象从该像素发出一道射线进入场景中，跟踪该光线穿过场景，直到光线命中某个物体。代码实现：

[cpp]  view plain  copy

1. #include "cuda_runtime.h"  
2. #include <highgui/highgui.hpp>  
3. #include <time.h>  
4.   
5. using namespace cv;  
6.   
7. #define INF 2e10f    
8. #define rnd(x) (x*rand()/RAND_MAX)    
9. #define SPHERES 100 //球体数量  
10. #define DIM 1024    //图像尺寸  
11.   
12. struct Sphere  
13. {  
14.     float r, g, b;  
15.     float radius;  
16.     float x, y, z;  
17.   
18.     __device__ float hit(float ox, float oy, float *n)  
19.     {  
20.         float dx = ox - x;  
21.         float dy = oy - y;  
22.   
23.         if (dx*dx + dy*dy < radius*radius)  
24.         {  
25.             float dz = sqrt(radius*radius - dx*dx - dy*dy);  
26.             *n = dz / sqrt(radius*radius);  
27.             return dz + z;  
28.         }  
29.   
30.         return -INF;  
31.     }  
32. };  
33.   
34. // Sphere *s;  
35. __constant__ Sphere s[SPHERES];  
36.   
37. /************************************************************************/  
38. //__global__ void rayTracing(unsigned char* ptr, Sphere* s)    
39. __global__ void rayTracing(unsigned char* ptr)  
40. {  
41.     int x = threadIdx.x + blockIdx.x * blockDim.x;  
42.     int y = threadIdx.y + blockIdx.y * blockDim.y;  
43.     int offset = x + y  * blockDim.x * gridDim.x;  
44.     float ox = (x - DIM / 2);  
45.     float oy = (y - DIM / 2);  
46.   
47.     float r = 0, g = 0, b = 0;  
48.     float maxz = -INF;  
49.     for (int i = 0; i < SPHERES; i++)  
50.     {  
51.         float n;  
52.         float t = s[i].hit(ox, oy, &n);  
53.         if (t > maxz)  
54.         {  
55.             float fscale = n;  
56.             r = s[i].r * fscale;  
57.             g = s[i].g * fscale;  
58.             b = s[i].b * fscale;  
59.             maxz = t;  
60.         }  
61.     }  
62.   
63.     ptr[offset * 3 + 2] = (int)(r * 255);  
64.     ptr[offset * 3 + 1] = (int)(g * 255);  
65.     ptr[offset * 3 + 0] = (int)(b * 255);  
66. }  
67. /************************************************************************/  
68.   
69.   
70. int main(int argc, char* argv[])  
71. {  
72.     cudaEvent_t start, stop;  
73.     cudaEventCreate(&start);  
74.     cudaEventCreate(&stop);  
75.     cudaEventRecord(start, 0);  
76.   
77.     Mat bitmap = Mat(Size(DIM, DIM), CV_8UC3, Scalar::all(0));  
78.     unsigned char *devBitmap;  
79.     (cudaMalloc((void**)&devBitmap, 3 * bitmap.rows*bitmap.cols));  
80.     //  cudaMalloc((void**)&s, sizeof(Sphere)*SPHERES);    
81.   
82.     Sphere *temps = (Sphere*)malloc(sizeof(Sphere)*SPHERES);  
83.   
84.     srand(time(0));  //随机数种子  
85.   
86.     for (int i = 0; i < SPHERES; i++)  
87.     {  
88.         temps[i].r = rnd(1.0f);  
89.         temps[i].g = rnd(1.0f);  
90.         temps[i].b = rnd(1.0f);  
91.         temps[i].x = rnd(1000.0f) - 500;  
92.         temps[i].y = rnd(1000.0f) - 500;  
93.         temps[i].z = rnd(1000.0f) - 500;  
94.         temps[i].radius = rnd(100.0f) + 20;  
95.     }  
96.   
97.     //  cudaMemcpy(s, temps, sizeof(Sphere)*SPHERES, cudaMemcpyHostToDevice);    
98.     cudaMemcpyToSymbol(s, temps, sizeof(Sphere)*SPHERES);  
99.     free(temps);  
100.   
101.     dim3 grids(DIM / 16, DIM / 16);  
102.     dim3 threads(16, 16);  
103.     //  rayTracing<<<grids, threads>>>(devBitmap, s);    
104.     rayTracing << <grids, threads >> > (devBitmap);  
105.   
106.     cudaMemcpy(bitmap.data, devBitmap, 3 * bitmap.rows*bitmap.cols, cudaMemcpyDeviceToHost);  
107.   
108.     cudaEventRecord(stop, 0);  
109.     cudaEventSynchronize(stop);  
110.   
111.     float elapsedTime;  
112.     cudaEventElapsedTime(&elapsedTime, start, stop);  
113.   
114.     printf("Processing time: %3.1f ms\n", elapsedTime);  
115.   
116.     imshow("CUDA常量内存使用示例", bitmap);  
117.     waitKey();  
118.     cudaFree(devBitmap);  
119.     //  cudaFree(s);    
120.     return 0;  
121. }  

程序里生成球体的大小和位置是随机的，为了产生随机数，加入了随机数种子srand（）。运行效果：