光线跟踪的 GPU 程序解读

#include <stdio.h>

#include "common/cpu_bitmap.h"

//是否使用__constan__常量内存的开关

#define CONSTANT

#define INF 2e10f

#define rnd(x) (x*rand()/RAND_MAX)

#define SPHERES 200

#define DIM 800

//求结构体

struct

Sphere {

float

r, g, b;

//球的颜色

float

radius;

//球的半径

float

x, y, z;

//球心坐标

//判断从像素点(ox,oy)射出的射线是否与该球相交，并返回交点的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 = sqrtf(radius*radius - dx*dx - dy*dy);

*n = dz/sqrtf(radius*radius);

return

dz + z;

}

return

-INF;

}

};

#ifdef CONSTANT

//__constant__ 常量内存

__constant__ Sphere s[SPHERES];

#else

Sphere *s;

#endif

#ifdef CONSTANT

//使用constant常量内存时，不能将其当作参数传到global函数

__global__

void

kernel(unsigned

char

* ptr)

#else

//普通全局变量必须用传参的形式传递到global函数

__global__

void

kernel(unsigned

char

* ptr, Sphere *s)

#endif

{

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(

void

)

{

//使用cuda事件来测试性能

cudaEvent_t start, stop;

cudaEventCreate(&start);

cudaEventCreate(&stop);

cudaEventRecord(start, 0);

CPUBitmap bitmap(DIM, DIM);

unsigned

char

* dev_bitmap;

cudaMalloc((

void

**) &dev_bitmap, bitmap.image_size());

#ifdef CONSTANT

// __constant__常量内存不需要动态分配内存

#else

// 在GPU设备上分配内存给球数组

cudaMalloc((

void

**) &s,

sizeof

(Sphere) * SPHERES);

#endif

// 在CPU上生成求数据数据

Sphere *temp_s = (Sphere *)

malloc

(

sizeof

(Sphere) * SPHERES);

for

(

int

i=0; i<SPHERES; i++)

{

temp_s[i].r = rnd(1.0f);

temp_s[i].g = rnd(1.0f);

temp_s[i].b = rnd(1.0f);

temp_s[i].x = rnd(1000.f) - 500;

temp_s[i].y = rnd(1000.f) - 500;

temp_s[i].z = rnd(1000.f) - 500;

temp_s[i].radius = rnd(100.f) + 20;

}

#ifdef CONSTANT

//从CPU拷贝到__constant__常量内存

cudaMemcpyToSymbol(s, temp_s,

sizeof

(Sphere) * SPHERES);

#else

//从CPU拷贝到GPU

cudaMemcpy(s, temp_s,

sizeof

(Sphere) * SPHERES, cudaMemcpyHostToDevice);

#endif

free

(temp_s);

dim3 grids(DIM/16, DIM/16);

dim3 threads(16, 16);

#ifdef CONSTANT

kernel<<<grids, threads>>>(dev_bitmap);

#else

kernel<<<grids, threads>>>(dev_bitmap, s);

#endif

cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(), cudaMemcpyDeviceToHost);

cudaEventRecord(stop, 0);

//事件同步

cudaEventSynchronize(stop);

float

elapseTime;

cudaEventElapsedTime(&elapseTime, start, stop);

printf

(

"Time to generate: %3.1f ms\n"

, elapseTime);

cudaEventDestroy(start);

cudaEventDestroy(stop);

bitmap.display_and_exit();

cudaFree(dev_bitmap);

#ifdef CONSTANT

// __constant__ 常量内存不需要free

#else

cudaFree(s);

#endif

return

1;

}