CUDA中的常量內存constant

2019-11-11 04:28:44

字體：大中小

來源：轉載

供稿：網友

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（）。運行效果：