在GPU上实现光线跟踪

#include "cuda.h"
#include "book.h"
#include "cpu_bitmap.h"

#define DIM 1024   //生成图像的大小，DIM*DIM
#define SPHERES 20   //生成的图像中球体的个数

#define rnd( x ) (x * rand() / RAND_MAX)    //产生随机数
#define INF     2e10f    //定义的无穷大数
//球体的定义
struct Sphere {
    float   r,b,g;   //颜色分量的定义
    float   radius;
    float   x,y,z;   //位置信息，球心三维坐标
    __device__ float hit( float ox, float oy, float *n ) {
        float dx = ox - x;
        float dy = oy - y;
  //x*x+sy*sy+z*z=r*r,等价于z*z=r*r-x*xs-y*y
        if (dx*dx + dy*dy < radius*radius) {
            float dz = sqrtf( radius*radius - dx*dx - dy*dy );
            *n = dz /radius;   //n代表计算的Z分量与半径的比值，作用比值决定灰度的大小，比值大，灰度大
            return dz + z;
        }
        return -INF;
    }
};

__global__ void kernel( Sphere *s, 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;   //深度距投影面距离为无穷远，不用显示该点的颜色
 //绘制SPHERES指定个数的球体
    for(int i=0; i<SPHERES; i++) {
        float   n;
        float   t = s[i].hit( ox, oy, &n );
  //如果t不是无穷远的话，就根据n的值计算该像素点的颜色
        if (t > maxz) {
            float fscale = n;
            r = s[i].r * fscale;
            g = s[i].g * fscale;
            b = s[i].b * fscale;
            maxz = t;
        }
    }
//将颜色变换到0~255之间，并写回输出图像中
    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;
}

// 图像块的定义
struct DataBlock {
    unsigned char   *dev_bitmap;
    Sphere          *s;
};

int main( void ) {
    DataBlock   data;   //图像数据块变量的定义
    //定义记录起始时间的事件对象
    cudaEvent_t     start, stop;
    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );
    HANDLE_ERROR( cudaEventRecord( start, 0 ) );

    CPUBitmap bitmap( DIM, DIM, &data );
    unsigned char   *dev_bitmap;
    Sphere   *s;

    // 在GPU上分配内存，以输出计算所得的位图
    HANDLE_ERROR( cudaMalloc( (void**)&dev_bitmap,bitmap.image_size() ) );
    // 为球体数据集分配内存，SPHERES个球体的内存空间
    HANDLE_ERROR( cudaMalloc( (void**)&s,sizeof(Sphere) * SPHERES ) );

    // 分配临时主机内存，对初始化球体的初始数据并赋值到GPU上的内存，然后释放临时内存
    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.0f ) - 500;
        temp_s[i].y = rnd( 1000.0f ) - 500;
        temp_s[i].z = rnd( 1000.0f ) - 500;
        temp_s[i].radius = rnd( 100.0f ) + 20;
    }
    HANDLE_ERROR( cudaMemcpy( s, temp_s,
                        sizeof(Sphere) * SPHERES,
                              cudaMemcpyHostToDevice ) );
    free( temp_s );

    // 建立与球体对应的的位图显存空间所需要的线程组织结构
    dim3    grids(DIM/16,DIM/16);
    dim3    threads(16,16);
    kernel<<<grids,threads>>>( s, dev_bitmap );

    //从GPU将数据拷回到主机内存以显示
    HANDLE_ERROR( cudaMemcpy( bitmap.get_ptr(), dev_bitmap,
                              bitmap.image_size(),
                              cudaMemcpyDeviceToHost ) );

    // 获取终止时间，以计算图像生成运行的时间
    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
    HANDLE_ERROR( cudaEventSynchronize( stop ) );
    float   elapsedTime;
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,
                                        start, stop ) );
    printf( "Time to generate:  %3.1f ms\n", elapsedTime );
     HANDLE_ERROR( cudaEventDestroy( start ) );
    HANDLE_ERROR( cudaEventDestroy( stop ) );
//释放显存空间
    HANDLE_ERROR( cudaFree( dev_bitmap ) );
    HANDLE_ERROR( cudaFree( s ) );

    // 显示图像
    bitmap.display_and_exit();
}