1 struct RendInfo {

2 int xsz, ysz;

3 int num_sph, num_lights;

4 int max_iter;

5 };

6

7 struct Sphere {

8 float4 pos;

9 float radius;

10 float4 color;

11 };

12

13 struct Light {

14 float4 pos, color;

15 };

16

17 struct Ray {

18 float4 origin, dir;

19 };

20

21 struct SurfPoint {

22 float t;

23 float3 pos, norm;

24 };

25

26 #define EPSILON 1e-6

27

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

29

30 __kernel void render(__global float4 *fb,

31 __global const struct RendInfo *rinf,

32 __global const struct Sphere *sphlist,

33 __global const struct Light *lights,

34 __global const struct Ray *primrays)

35 {

36 int idx = get_global_id(0);

37

38 struct Ray ray = primrays[idx];

39 struct SurfPoint sp;

40

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

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

43 } else {

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

45 }

46

47 fb[idx] = primrays[idx].dir * 0.5 + 0.5;

48 }

49

50 bool intersect(struct Ray ray,

51 __global const struct Sphere *sph,

52 struct SurfPoint *sp)

53 {

54 float3 dir = ray.dir.xyz;

55 float3 orig = ray.origin.xyz;

56 float3 spos = sph->pos.xyz;

57

58 float a = dot(dir, dir);

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

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

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

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

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

64

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

66 if(d < 0.0) return false;

67

68 float sqrt_d = sqrt(d);

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

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

71

72 if(t1 < EPSILON) t1 = t2;

73 if(t2 < EPSILON) t2 = t1;

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

75

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

77 return false;

78 }

79

80 sp->t = t;

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

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

83 return true;

84 }