[菜鸟每天来段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).