1 struct RendInfo {

2 int xsz, ysz;

3 int num_sph;

4 int max_iter;

5 };

6

7 struct Sphere {

8 float4 pos;

9 float radius;

10 float4 color;

11 };

12

13 struct Ray {

14 float4 origin, dir;

15 };

16

17 struct SurfPoint {

18 float t;

19 float3 pos, norm;

20 };

21

22 #define EPSILON 1e-6

23

24 bool intersect(struct Ray ray, __global const struct Sphere *sph, struct SurfPoint *sp);

25

26 __kernel void render(__global float4 *fb,

27 __global const struct RendInfo *rinf,

28 __global const struct Sphere *sphlist,

29 __global const struct Ray *primrays)

30 {

31 int idx = get_global_id(0);

32

33 struct Ray ray = primrays[idx];

34 struct SurfPoint sp;

35

36 if(intersect(ray, sphlist, &sp)) {

37 fb[idx] = (float4)(1, 0, 0, 1);

38 } else {

39 fb[idx] = (float4)(0, 0, 0, 1);

40 }

41 }

42

43 bool intersect(struct Ray ray,

44 __global const struct Sphere *sph,

45 struct SurfPoint *sp)

46 {

47 float3 dir = ray.dir.xyz;

48 float3 orig = ray.origin.xyz;

49 float3 spos = sph->pos.xyz;

50

51 float a = dot(dir, dir);

52 float b = 2.0 * dir.x * (orig.x - spos.x) +

53 2.0 * dir.y * (orig.y - spos.y) +

54 2.0 * dir.z * (orig.z - spos.z);

55 float c = dot(spos, spos) + dot(orig, orig) +

56 2.0 * dot(-spos, orig) - sph->radius * sph->radius;

57

58 float d = b * b - 4.0 * a * c;

59 if(d < 0.0) return false;

60

61 float sqrt_d = sqrt(d);

62 float t1 = (-b + sqrt_d) / (2.0 * a);

63 float t2 = (-b - sqrt_d) / (2.0 * a);

64

65 if(t1 < EPSILON) t1 = t2;

66 if(t2 < EPSILON) t2 = t1;

67 float t = t1 < t2 ? t1 : t2;

68

69 if(t < EPSILON || t > 1.0) {

70 return false;

71 }

72

73 sp->t = t;

74 sp->pos = orig + dir * sp->t;

75 sp->norm = (sp->pos - spos) / sph->radius;

76 return true;

77 }