使用全局内存来存储球面对象，项目打包下载
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 <GLglut.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
", 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 }
结果如下图所示：
