rev |
line source |
nuclear@2
|
1 struct RendInfo {
|
nuclear@2
|
2 int xsz, ysz;
|
nuclear@3
|
3 int num_sph, num_lights;
|
nuclear@2
|
4 int max_iter;
|
nuclear@2
|
5 };
|
nuclear@2
|
6
|
nuclear@2
|
7 struct Sphere {
|
nuclear@2
|
8 float4 pos;
|
nuclear@5
|
9 float4 kd, ks;
|
nuclear@2
|
10 float radius;
|
nuclear@4
|
11 float spow, kr, kt;
|
nuclear@2
|
12 };
|
nuclear@2
|
13
|
nuclear@3
|
14 struct Light {
|
nuclear@3
|
15 float4 pos, color;
|
nuclear@3
|
16 };
|
nuclear@3
|
17
|
nuclear@2
|
18 struct Ray {
|
nuclear@2
|
19 float4 origin, dir;
|
nuclear@2
|
20 };
|
nuclear@2
|
21
|
nuclear@2
|
22 struct SurfPoint {
|
nuclear@2
|
23 float t;
|
nuclear@2
|
24 float3 pos, norm;
|
nuclear@4
|
25 global const struct Sphere *obj;
|
nuclear@2
|
26 };
|
nuclear@2
|
27
|
nuclear@2
|
28 #define EPSILON 1e-6
|
nuclear@2
|
29
|
nuclear@5
|
30 float4 shade(struct Ray ray, struct SurfPoint sp,
|
nuclear@5
|
31 global const struct Light *lights, int num_lights);
|
nuclear@4
|
32 bool intersect(struct Ray ray, global const struct Sphere *sph, struct SurfPoint *sp);
|
nuclear@5
|
33 float3 reflect(float3 v, float3 n);
|
nuclear@2
|
34
|
nuclear@4
|
35
|
nuclear@4
|
36 kernel void render(global float4 *fb,
|
nuclear@4
|
37 global const struct RendInfo *rinf,
|
nuclear@4
|
38 global const struct Sphere *sphlist,
|
nuclear@4
|
39 global const struct Light *lights,
|
nuclear@4
|
40 global const struct Ray *primrays)
|
nuclear@2
|
41 {
|
nuclear@2
|
42 int idx = get_global_id(0);
|
nuclear@2
|
43
|
nuclear@2
|
44 struct Ray ray = primrays[idx];
|
nuclear@4
|
45 struct SurfPoint sp, sp0;
|
nuclear@2
|
46
|
nuclear@4
|
47 sp0.t = FLT_MAX;
|
nuclear@5
|
48 sp0.obj = 0;
|
nuclear@4
|
49
|
nuclear@4
|
50 for(int i=0; i<rinf->num_sph; i++) {
|
nuclear@5
|
51 if(intersect(ray, sphlist + i, &sp) && sp.t < sp0.t) {
|
nuclear@4
|
52 sp0 = sp;
|
nuclear@4
|
53 }
|
nuclear@2
|
54 }
|
nuclear@3
|
55
|
nuclear@5
|
56 if(sp0.obj) {
|
nuclear@5
|
57 fb[idx] = shade(ray, sp0, lights, rinf->num_lights);
|
nuclear@5
|
58 } else {
|
nuclear@5
|
59 fb[idx] = (float4)(0, 0, 0, 0);
|
nuclear@5
|
60 }
|
nuclear@4
|
61 }
|
nuclear@4
|
62
|
nuclear@5
|
63 float4 shade(struct Ray ray, struct SurfPoint sp,
|
nuclear@5
|
64 global const struct Light *lights, int num_lights)
|
nuclear@4
|
65 {
|
nuclear@5
|
66 float3 dcol = (float3)(0, 0, 0);
|
nuclear@5
|
67 float3 scol = (float3)(0, 0, 0);
|
nuclear@5
|
68
|
nuclear@5
|
69 for(int i=0; i<num_lights; i++) {
|
nuclear@5
|
70 float3 ldir = normalize(lights[i].pos.xyz - sp.pos);
|
nuclear@5
|
71 float3 vdir = -normalize(ray.dir.xyz);
|
nuclear@5
|
72 float3 vref = reflect(vdir, sp.norm);
|
nuclear@5
|
73
|
nuclear@5
|
74 float diff = fmax(dot(ldir, sp.norm), 0.0f);
|
nuclear@5
|
75 float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.obj->spow);
|
nuclear@5
|
76
|
nuclear@5
|
77 dcol += sp.obj->kd.xyz * diff * lights[i].color.xyz;
|
nuclear@5
|
78 scol += sp.obj->ks.xyz * spec * lights[i].color.xyz;
|
nuclear@5
|
79 }
|
nuclear@5
|
80
|
nuclear@5
|
81 return (float4)(dcol + scol, 1.0f);
|
nuclear@2
|
82 }
|
nuclear@2
|
83
|
nuclear@2
|
84 bool intersect(struct Ray ray,
|
nuclear@4
|
85 global const struct Sphere *sph,
|
nuclear@2
|
86 struct SurfPoint *sp)
|
nuclear@2
|
87 {
|
nuclear@2
|
88 float3 dir = ray.dir.xyz;
|
nuclear@2
|
89 float3 orig = ray.origin.xyz;
|
nuclear@2
|
90 float3 spos = sph->pos.xyz;
|
nuclear@2
|
91
|
nuclear@2
|
92 float a = dot(dir, dir);
|
nuclear@2
|
93 float b = 2.0 * dir.x * (orig.x - spos.x) +
|
nuclear@2
|
94 2.0 * dir.y * (orig.y - spos.y) +
|
nuclear@2
|
95 2.0 * dir.z * (orig.z - spos.z);
|
nuclear@2
|
96 float c = dot(spos, spos) + dot(orig, orig) +
|
nuclear@2
|
97 2.0 * dot(-spos, orig) - sph->radius * sph->radius;
|
nuclear@2
|
98
|
nuclear@2
|
99 float d = b * b - 4.0 * a * c;
|
nuclear@2
|
100 if(d < 0.0) return false;
|
nuclear@2
|
101
|
nuclear@2
|
102 float sqrt_d = sqrt(d);
|
nuclear@2
|
103 float t1 = (-b + sqrt_d) / (2.0 * a);
|
nuclear@2
|
104 float t2 = (-b - sqrt_d) / (2.0 * a);
|
nuclear@2
|
105
|
nuclear@2
|
106 if(t1 < EPSILON) t1 = t2;
|
nuclear@2
|
107 if(t2 < EPSILON) t2 = t1;
|
nuclear@2
|
108 float t = t1 < t2 ? t1 : t2;
|
nuclear@2
|
109
|
nuclear@2
|
110 if(t < EPSILON || t > 1.0) {
|
nuclear@2
|
111 return false;
|
nuclear@2
|
112 }
|
nuclear@2
|
113
|
nuclear@2
|
114 sp->t = t;
|
nuclear@2
|
115 sp->pos = orig + dir * sp->t;
|
nuclear@2
|
116 sp->norm = (sp->pos - spos) / sph->radius;
|
nuclear@5
|
117 sp->obj = sph;
|
nuclear@2
|
118 return true;
|
nuclear@2
|
119 }
|
nuclear@5
|
120
|
nuclear@5
|
121 float3 reflect(float3 v, float3 n)
|
nuclear@5
|
122 {
|
nuclear@5
|
123 return 2.0f * dot(v, n) * n - v;
|
nuclear@5
|
124 }
|