使用全局内存的光线跟踪实验

  1 /*
  2 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
  3 *
  4 * NVIDIA Corporation and its licensors retain all intellectual property and
  5 * proprietary rights in and to this software and related documentation.
  6 * Any use, reproduction, disclosure, or distribution of this software
  7 * and related documentation without an express license agreement from
  8 * NVIDIA Corporation is strictly prohibited.
  9 *
 10 * Please refer to the applicable NVIDIA end user license agreement (EULA)
 11 * associated with this source code for terms and conditions that govern
 12 * your use of this NVIDIA software.
 13 *
 14 */
 15 #include <GL\glut.h>
 16 #include "cuda.h"
 17 #include "cuda_runtime.h"
 18 #include "device_launch_parameters.h"
 19 #include "../common/book.h"
 20 #include "../common/cpu_bitmap.h"
 21 #include "device_functions.h"
 22 
 23 #define DIM 1024
 24 
 25 #define rnd( x ) (x * rand() / RAND_MAX)
 26 #define INF     2e10f
 27 
 28 struct Sphere {
 29     float   r, b, g;
 30     float   radius;
 31     float   x, y, z;
 32     __device__ float hit(float ox, float oy, float *n) {
 33         //将中心坐标移动到图像中间，dx和dy就是相对于新的中心坐标ox和oy的新坐标
 34         float dx = ox - x;
 35         float dy = oy - y;
 36         //只处理点在圆内的Sphere对象
 37         if (dx*dx + dy*dy < radius*radius) {
 38             /*
 39             计算的dz为离圆心轴的距离
 40             */
 41             float dz = sqrtf(radius*radius - dx*dx - dy*dy);
 42             /*
 43             和半径相除，由于dz是float类型，所以结果也为float类型
 44             也就是说结果为0.xxx这样的数字，
 45             n为一个指针，*n为解析指针n，存放的也就是值0.xxx这样的值
 46             呈现在最后的结果就是图像颜色的渐变效果
 47             */
 48             *n = dz / sqrtf(radius * radius);
 49             /*
 50             在三维空间中，已经将xoy投影到为图上，z轴垂直于位图
 51             距离圆心的距离dz再加上原来的Z轴坐标就是当前坐标对应于xoy面的Z轴方向距离
 52             */
 53             return dz + z;
 54         }
 55         return -INF;
 56     }
 57 };
 58 #define SPHERES 20
 59 
 60 
 61 __global__ void kernel(Sphere *s, unsigned char *ptr) {
 62     // 映射到图像像素上的位置
 63     int x = threadIdx.x + blockIdx.x * blockDim.x;
 64     int y = threadIdx.y + blockIdx.y * blockDim.y;
 65     int offset = x + y * blockDim.x * gridDim.x;//步长
 66     //移动使得Z轴在图像中心
 67     float   ox = (x - DIM / 2);
 68     float   oy = (y - DIM / 2);
 69 
 70     float   r = 0, g = 0, b = 0;
 71     float   maxz = -INF;
 72     //每个像素递归判断SPHERES个对象在这个像素点上的值
 73     for (int i = 0; i<SPHERES; i++) {
 74         float   n;
 75         float   t = s[i].hit(ox, oy, &n);
 76         if (t > maxz) {
 77             //这里取n的地址，hit函数将结果存放在&n地址所指的空间，不同的n对应不同的颜色及深度
 78             float fscale = n;
 79             r = s[i].r * fscale;
 80             g = s[i].g * fscale;
 81             b = s[i].b * fscale;
 82             maxz = t;
 83         }
 84     }
 85 
 86     ptr[offset * 4 + 0] = (int)(r * 255);
 87     ptr[offset * 4 + 1] = (int)(g * 255);
 88     ptr[offset * 4 + 2] = (int)(b * 255);
 89     ptr[offset * 4 + 3] = 255;
 90 }
 91 
 92 
 93 // globals needed by the update routine
 94 struct DataBlock {
 95     unsigned char   *dev_bitmap;
 96     Sphere          *s;
 97 };
 98 
 99 int main(void) {
100     DataBlock   data;
101     //获取时间
102     cudaEvent_t     start, stop;
103     HANDLE_ERROR(cudaEventCreate(&start));
104     HANDLE_ERROR(cudaEventCreate(&stop));
105     HANDLE_ERROR(cudaEventRecord(start, 0));
106 
107     CPUBitmap bitmap(DIM, DIM, &data);
108     unsigned char   *dev_bitmap;
109     Sphere          *s;
110 
111 
112     // allocate memory on the GPU for the output bitmap
113     HANDLE_ERROR(cudaMalloc((void**)&dev_bitmap,
114         bitmap.image_size()));
115     //在全局内存中分配s
116     HANDLE_ERROR(cudaMalloc((void**)&s,
117         sizeof(Sphere)* SPHERES));
118 
119     //主机上申请存储空间
120     Sphere *temp_s = (Sphere*)malloc(sizeof(Sphere)* SPHERES);
121     for (int i = 0; i<SPHERES; i++) {
122         temp_s[i].r = rnd(1.0f);
123         temp_s[i].g = rnd(1.0f);
124         temp_s[i].b = rnd(1.0f);
125         temp_s[i].x = rnd(1000.0f) - 500;
126         temp_s[i].y = rnd(1000.0f) - 500;
127         temp_s[i].z = rnd(1000.0f) - 500;
128         temp_s[i].radius = rnd(100.0f) + 20;
129     }
130     HANDLE_ERROR(cudaMemcpy(s, temp_s,
131         sizeof(Sphere)* SPHERES,
132         cudaMemcpyHostToDevice));
133     free(temp_s);
134 
135     // generate a bitmap from our sphere data
136     dim3    grids(DIM / 16, DIM / 16);
137     dim3    threads(16, 16);
138     kernel <<<grids, threads >>>(s, dev_bitmap);
139 
140     // copy our bitmap back from the GPU for display
141     HANDLE_ERROR(cudaMemcpy(bitmap.get_ptr(), dev_bitmap,
142         bitmap.image_size(),
143         cudaMemcpyDeviceToHost));
144 
145     // get stop time, and display the timing results
146     HANDLE_ERROR(cudaEventRecord(stop, 0));
147     HANDLE_ERROR(cudaEventSynchronize(stop));
148     float   elapsedTime;
149     HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,
150         start, stop));
151     printf("Time to generate:  %3.1f ms\n", elapsedTime);
152 
153     HANDLE_ERROR(cudaEventDestroy(start));
154     HANDLE_ERROR(cudaEventDestroy(stop));
155 
156     HANDLE_ERROR(cudaFree(dev_bitmap));
157     HANDLE_ERROR(cudaFree(s));
158 
159     // display
160     bitmap.display_and_exit();
161 }