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@2
|
9 float radius;
|
nuclear@4
|
10 float4 kd, ks;
|
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@4
|
30 float4 shade(struct Ray ray, struct SurfPoint sp);
|
nuclear@4
|
31 bool intersect(struct Ray ray, global const struct Sphere *sph, struct SurfPoint *sp);
|
nuclear@2
|
32
|
nuclear@4
|
33
|
nuclear@4
|
34 kernel void render(global float4 *fb,
|
nuclear@4
|
35 global const struct RendInfo *rinf,
|
nuclear@4
|
36 global const struct Sphere *sphlist,
|
nuclear@4
|
37 global const struct Light *lights,
|
nuclear@4
|
38 global const struct Ray *primrays)
|
nuclear@2
|
39 {
|
nuclear@2
|
40 int idx = get_global_id(0);
|
nuclear@2
|
41
|
nuclear@2
|
42 struct Ray ray = primrays[idx];
|
nuclear@4
|
43 struct SurfPoint sp, sp0;
|
nuclear@2
|
44
|
nuclear@4
|
45 sp0.t = FLT_MAX;
|
nuclear@4
|
46
|
nuclear@4
|
47 for(int i=0; i<rinf->num_sph; i++) {
|
nuclear@4
|
48 if(intersect(ray, sphlist, &sp) && sp.t < sp0.t) {
|
nuclear@4
|
49 sp0 = sp;
|
nuclear@4
|
50 }
|
nuclear@2
|
51 }
|
nuclear@3
|
52
|
nuclear@4
|
53 fb[idx] = shade(ray, sp0);
|
nuclear@4
|
54 }
|
nuclear@4
|
55
|
nuclear@4
|
56 float4 shade(struct Ray ray, struct SurfPoint sp)
|
nuclear@4
|
57 {
|
nuclear@4
|
58 return sp.obj->kd;
|
nuclear@2
|
59 }
|
nuclear@2
|
60
|
nuclear@2
|
61 bool intersect(struct Ray ray,
|
nuclear@4
|
62 global const struct Sphere *sph,
|
nuclear@2
|
63 struct SurfPoint *sp)
|
nuclear@2
|
64 {
|
nuclear@2
|
65 float3 dir = ray.dir.xyz;
|
nuclear@2
|
66 float3 orig = ray.origin.xyz;
|
nuclear@2
|
67 float3 spos = sph->pos.xyz;
|
nuclear@2
|
68
|
nuclear@2
|
69 float a = dot(dir, dir);
|
nuclear@2
|
70 float b = 2.0 * dir.x * (orig.x - spos.x) +
|
nuclear@2
|
71 2.0 * dir.y * (orig.y - spos.y) +
|
nuclear@2
|
72 2.0 * dir.z * (orig.z - spos.z);
|
nuclear@2
|
73 float c = dot(spos, spos) + dot(orig, orig) +
|
nuclear@2
|
74 2.0 * dot(-spos, orig) - sph->radius * sph->radius;
|
nuclear@2
|
75
|
nuclear@2
|
76 float d = b * b - 4.0 * a * c;
|
nuclear@2
|
77 if(d < 0.0) return false;
|
nuclear@2
|
78
|
nuclear@2
|
79 float sqrt_d = sqrt(d);
|
nuclear@2
|
80 float t1 = (-b + sqrt_d) / (2.0 * a);
|
nuclear@2
|
81 float t2 = (-b - sqrt_d) / (2.0 * a);
|
nuclear@2
|
82
|
nuclear@2
|
83 if(t1 < EPSILON) t1 = t2;
|
nuclear@2
|
84 if(t2 < EPSILON) t2 = t1;
|
nuclear@2
|
85 float t = t1 < t2 ? t1 : t2;
|
nuclear@2
|
86
|
nuclear@2
|
87 if(t < EPSILON || t > 1.0) {
|
nuclear@2
|
88 return false;
|
nuclear@2
|
89 }
|
nuclear@2
|
90
|
nuclear@2
|
91 sp->t = t;
|
nuclear@2
|
92 sp->pos = orig + dir * sp->t;
|
nuclear@2
|
93 sp->norm = (sp->pos - spos) / sph->radius;
|
nuclear@2
|
94 return true;
|
nuclear@2
|
95 }
|