GPU包含數百個數學計算單元，具有強大的處理運算能力，可以強大到計算速率高于輸入數據的速率，即充分利用帶寬，滿負荷向GPU傳輸數據還不夠它計算的。CUDA C除全局內存和共享內存外，還支持常量內存，常量內存用于保存在核函數執行期間不會發生變化的數據，使用常量內存在一些情況下，能有效減少內存帶寬，降低GPU運算單元的空閑等待。
使用常量內存可以提升運算性能的原因如下：
對常量內存的單次讀操作可以廣播到其他的“鄰近(nearby)”線程，這將節約15次讀取操作；高速緩存。常量內存的數據將緩存起來，因此對于相同地址的連續操作將不會產生額外的內存通信量；在CUDA架構中，線程束是指一個包含32個線程的集合，這個線程集合被“編織在一起”并且以“步調一致(Lockstep)”的形式執行。當處理常量內存時，NVIDIA硬件將把單次內存讀取操作廣播到每個半線程束(Half-Warp)。在半線程束中包含16個線程，即線程束中線程數量的一半。如果在半線程束中的每個線程從常量內存的相同地址上讀取數據，那么GPU只會產生一次讀取請求并在隨后將數據廣播到每個線程。如果從常量內存中讀取大量數據，那么這種方式產生的內存流量只是使用全局內存時的1/16。#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（）。運行效果： 
新聞熱點
疑難解答
