标签:球面 false 图形 3.1 style display isp udaf host
本章介绍了常量内存的使用,并给光线追踪的一个例子。
章节代码
1 #include <stdio.h> 2 #include "cuda_runtime.h" 3 #include "device_launch_parameters.h" 4 #include "D:\Code\CUDA\book\common\book.h" 5 #include "D:\Code\CUDA\book\common\cpu_bitmap.h" 6 7 #define DIM 1024 8 #define rnd( x ) (x * rand() / RAND_MAX) 9 #define INF 2e10f 10 #define SPHERES 40 11 #define USE_CONSTANT_MEMORY false 12 13 struct Sphere 14 { 15 float r, b, g; 16 float radius; 17 float x, y, z; 18 __device__ float hit(float ox, float oy, float *n)//计算球体上一点处的法向量n并返回该点到画面平面的距离 19 { 20 float dx = ox - x; 21 float dy = oy - y; 22 if (dx*dx + dy*dy < radius*radius) 23 { 24 float dz = sqrtf(radius*radius - dx*dx - dy*dy); 25 *n = dz / sqrtf(radius * radius);//球上该点法向量与画面法向量夹角的余弦值 26 return dz + z; 27 } 28 return -INF; 29 } 30 }; 31 32 #if USE_CONSTANT_MEMORY 33 __constant__ Sphere s[SPHERES]; 34 struct DataBlock 35 { 36 unsigned char *dev_bitmap; 37 }; 38 #else 39 struct DataBlock 40 { 41 unsigned char *dev_bitmap; 42 Sphere *s; 43 }; 44 #endif 45 46 #if USE_CONSTANT_MEMORY 47 __global__ void kernel(unsigned char *ptr) 48 #else 49 __global__ void kernel(Sphere *s, unsigned char *ptr) 50 #endif 51 { 52 int x = threadIdx.x + blockIdx.x * blockDim.x; 53 int y = threadIdx.y + blockIdx.y * blockDim.y; 54 int offset = x + y * blockDim.x * gridDim.x; 55 float ox = (x - DIM / 2); 56 float oy = (y - DIM / 2); 57 58 float r = 0, g = 0, b = 0;//计算该像素应该显示什么颜色 59 float maxz = -INF; 60 for (int i = 0; i<SPHERES; i++) 61 { 62 float n; 63 float t = s[i].hit(ox, oy, &n);//o理解成"observation",即当前线程代表的坐标 64 if (t > maxz)//距离画面最近,更新该像素的显示 65 { 66 float fscale = n; 67 r = s[i].r * fscale; 68 g = s[i].g * fscale; 69 b = s[i].b * fscale; 70 maxz = t; 71 } 72 } 73 74 ptr[offset * 4 + 0] = (int)(r * 255); 75 ptr[offset * 4 + 1] = (int)(g * 255); 76 ptr[offset * 4 + 2] = (int)(b * 255); 77 ptr[offset * 4 + 3] = 255; 78 79 return; 80 } 81 82 int main(void) 83 { 84 DataBlock data; 85 cudaEvent_t start, stop;// 计时器 86 cudaEventCreate(&start); 87 cudaEventCreate(&stop); 88 cudaEventRecord(start, 0); 89 90 CPUBitmap bitmap(DIM, DIM, &data); 91 unsigned char *dev_bitmap; 92 #if !USE_CONSTANT_MEMORY 93 Sphere *s; 94 #endif 95 cudaMalloc((void**)&dev_bitmap,bitmap.image_size()); 96 cudaMalloc((void**)&s,sizeof(Sphere) * SPHERES); 97 Sphere *temp_s = (Sphere*)malloc(sizeof(Sphere) * SPHERES); 98 99 for (int i = 0; i<SPHERES; i++) 100 { 101 temp_s[i].r = rnd(1.0f); 102 temp_s[i].g = rnd(1.0f); 103 temp_s[i].b = rnd(1.0f); 104 temp_s[i].x = rnd(1000.0f) - 500; 105 temp_s[i].y = rnd(1000.0f) - 500; 106 temp_s[i].z = rnd(1000.0f) - 500; 107 temp_s[i].radius = rnd(100.0f) + 20; 108 } 109 110 #if USE_CONSTANT_MEMORY 111 cudaMemcpyToSymbol(s, temp_s, sizeof(Sphere) * SPHERES); 112 kernel << < dim3(DIM / 16, DIM / 16), dim3(16, 16) >> > (dev_bitmap); 113 #else 114 cudaMemcpy(s, temp_s, sizeof(Sphere) * SPHERES, cudaMemcpyHostToDevice); 115 kernel << < dim3(DIM / 16, DIM / 16), dim3(16, 16) >> > (s, dev_bitmap); 116 #endif 117 118 cudaMemcpy(bitmap.get_ptr(), dev_bitmap,bitmap.image_size(),cudaMemcpyDeviceToHost); 119 120 cudaEventRecord(stop, 0);//测量计算耗时 121 cudaEventSynchronize(stop); 122 float elapsedTime; 123 cudaEventElapsedTime(&elapsedTime,start, stop); 124 printf("Time to generate: %3.1f ms\n", elapsedTime); 125 cudaEventDestroy(start); 126 cudaEventDestroy(stop); 127 128 free(temp_s); 129 cudaFree(dev_bitmap); 130 cudaFree(s); 131 132 bitmap.display_and_exit(); 133 134 getchar(); 135 return; 136 }
? 算法大意:有SPHERES个球分布在原点附近,其球心坐标在每个坐标轴方向上分量绝对值不大于500,其半径介于20到120;观察者(画面平面)位于z正半轴充分远处(z>500),现将所有的球体平行投影到画面平面上,考虑遮挡关系,并考虑球面与画面平面的夹角给球体绘制阴影。
? 使用常量内存时球数组定义在所有函数外部,核函数只需图形参数就够了;不使用常量内存时球数组定义在结构DataBlock内部,核函数需要球数组和图形参数
? 使用了结构cudaEvent_t用于计时,并介绍了与此相关的时间控制函数,按顺序使用如下。
1 cudaEvent_t start, stop; 2 cudaEventCreate(&start); 3 cudaEventRecord(start, 0); 4 5 //Do something 6 7 cudaEventCreate(&stop); 8 cudaEventRecord(stop, 0); 9 cudaEventSynchronize(stop); 10 11 float elapsedTime; 12 cudaEventElapsedTime(&elapsedTime, start, stop); 13 cudaEventDestroy(start); 14 cudaEventDestroy(stop);
? 使用cudaMemcpyToSymbol()函数复制内存到到常量内存中(可以反向从显存复制到内存中,但由于只读一般没有情况会这样做)。其定义于cuda_runtime.h中
1 template<class T> 2 static __inline__ __host__ cudaError_t cudaMemcpyToSymbol( 3 const T &symbol, 4 const void *src, 5 size_t count, 6 size_t offset = 0, 7 enum cudaMemcpyKind kind = cudaMemcpyHostToDevice) 8 { 9 return ::cudaMemcpyToSymbol((const void*)&symbol, src, count, offset, kind); 10 }
标签:球面 false 图形 3.1 style display isp udaf host
原文地址:http://www.cnblogs.com/cuancuancuanhao/p/7640613.html