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 bitmap113     HANDLE_ERROR(cudaMalloc((void**)&dev_bitmap,114         bitmap.image_size()));115     //在全局内存中分配s116     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 data136     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 display141     HANDLE_ERROR(cudaMemcpy(bitmap.get_ptr(), dev_bitmap,142         bitmap.image_size(),143         cudaMemcpyDeviceToHost));144 145     // get stop time, and display the timing results146     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     // display160     bitmap.display_and_exit();161 }