[菜鳥每天來段CUDA_C]GPU上通過常量記憶體實作光線跟蹤

        光線跟蹤是從三維場景生成二維圖像的一種方式。主要思想為：在場景中選擇一個位置放上一台假想的相機，該相機包含一個光傳感器來生成圖像，需要判斷那些光将接觸到這個傳感器。圖像中每個像素與命中傳感器的光線有相同的顔色和強度。傳感器中命中的光線可能來自場景中的任意位置，想象從該像素發出一道射線進入場景中，跟蹤該光線穿過場景，直到光線命中某個物體。 

        本文實作一個簡單場景的光線跟蹤，場景中隻有一組不同半徑的球，沒有任何光源，假想相機固定在Z軸。從每個像素發射一道光線，跟蹤這些光線會命中哪些球面。當一束光線穿越多個球面時，最接近相機的球面才會被看到(判斷命中的位置與相機之間的距離是否比上一次命中的距離更加接近)。如果沒有命中任何球面，則改點的顔色值為初始值(背景為黑色)。

運作結果如下：

主要代碼如下：

/********************************************************************
*  rayTracing.cu
*********************************************************************/

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cutil_inline.h>
#include "CPUBitmap.h"

#define INF 2e10f
#define rnd(x) (x*rand()/RAND_MAX)
#define SPHERES 100
#define DIM 512

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;
	}
};


/************************************************************************/
/* Init CUDA                                                            */
/************************************************************************/
bool InitCUDA(void)
{
     ......
}

//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*4 + 0] = (int)(r*255);
	ptr[offset*4 + 1] = (int)(g*255);
	ptr[offset*4 + 2] = (int)(b*255);
	ptr[offset*4 + 3] = 255;
}
/************************************************************************/


int main(int argc, char* argv[])
{

	if(!InitCUDA()) {
		return 0;
	}

	cudaEvent_t start, stop;
	cutilSafeCall(cudaEventCreate(&start));
	cutilSafeCall(cudaEventCreate(&stop));
	cutilSafeCall(cudaEventRecord(start, 0));

	CPUBitmap bitmap(DIM, DIM);
	unsigned char *devBitmap;
	cutilSafeCall(cudaMalloc((void**)&devBitmap, bitmap.image_size()));
//	cutilSafeCall(cudaMalloc((void**)&s, sizeof(Sphere)*SPHERES));

	Sphere *temps = (Sphere*)malloc(sizeof(Sphere)*SPHERES);
	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;
	}

//	cutilSafeCall(cudaMemcpy(s, temps, sizeof(Sphere)*SPHERES, cudaMemcpyHostToDevice));
	cutilSafeCall(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);

	cutilSafeCall(cudaMemcpy(bitmap.get_ptr(), devBitmap, bitmap.image_size(), cudaMemcpyDeviceToHost));
	
	cutilSafeCall(cudaEventRecord(stop, 0));
	cutilSafeCall(cudaEventSynchronize(stop));

	float elapsedTime;
	cutilSafeCall(cudaEventElapsedTime(&elapsedTime, start, stop));

	printf("Processing time: %3.1f ms\n", elapsedTime);
	
	bitmap.display_and_exit();

	cudaFree(devBitmap);
//	cudaFree(s);
	return 0;
}
           

實驗中取50個球進行計算，采用常量記憶體時耗時3.0ms, 不采用時耗時8.2ms。 常量記憶體明顯提高了運算性能。

常量記憶體帶來的性能提升

        __constant__把對變量的通路限制為隻讀，跟從全局記憶體讀取資料相比，從常量記憶體讀取相同的資料可以節約記憶體帶寬，原因是對常量記憶體的單次讀操作可以廣播到同一個線程塊内的其他線程(節約15次讀操作)。且常量記憶體的資料将被緩存起來，對相同位址的連續讀操作将不會産生額外的記憶體通信量。

        上述光線跟蹤中每個線程都要讀取球面的相應資料進而計算它與光線的相交情況，是以将球面資料儲存在常量記憶體以後，硬體隻需要請求這個資料一次。緩存這個資料後，其他每個線程将不會産生記憶體流量。降低通信量進而提高了性能。

參考資源：

Jason Sanders, Edward Kandrot, CUDA By Example: An Introduction toGeneral-Purpose GPU Programming (2011).