sphere without shading

src/editortest.cu

@@ -1,12 +0,0 @@
-
-__global__ void some_global_kernel() {
-
-   some_kernel<<<32, 32>>>(0,0);
-}
-
-
-__device__ void some_kernel(int a, int b)
-{
-
-  return;
-}

src/main.cu

@@ -48,6 +48,12 @@ union Vec3F32
     F32 y;
     F32 z;
   };
+  struct 
+  {
+    F32 r;
+    F32 g;
+    F32 b;
+  };
   F32 v[3];
 };
 
@@ -144,6 +150,12 @@ __host__ __device__ function Vec3F32 scale_V3F32(F32 s, Vec3F32 v)
   return out;
 }
 
+__host__ __device__ function F32 dot_V3F32(Vec3F32 a, Vec3F32 b)
+{
+  return a.x*b.x + a.y*b.y + a.z*b.z;
+}
+
+
 __device__ function Vec3F32 ray_point_F32(F32 t, RayF32 *ray)
 {
   Vec3F32 out = add_V3F32(ray->origin, scale_V3F32(t, ray->direction));
@@ -171,10 +183,7 @@ __device__ function Vec3F32 lerp_V3F32(F32 s, Vec3F32 a, Vec3F32 b)
 }
 
 
-
-
-
-function void write_buffer_to_ppm(Vec3F32 *buffer, U32 image_width, U32 image_height, U32 *idx_buffer) 
+__host__ function void write_buffer_to_ppm(Vec3F32 *buffer, U32 image_width, U32 image_height, U32 *idx_buffer) 
 {
 
   const char *filename = "output.ppm";
@@ -201,9 +210,9 @@ function void write_buffer_to_ppm(Vec3F32 *buffer, U32 image_width, U32 image_he
       if(idx_buffer[idx] != 0) {
         //LOG("idx %i, idxbuffer[idx] = %i \n", idx, idx_buffer[idx]);
       }
-      F32 r = buffer[idx].x;
-      F32 g = buffer[idx].y;
-      F32 b = buffer[idx].z;
+      F32 r = buffer[idx].r;
+      F32 g = buffer[idx].g;
+      F32 b = buffer[idx].b;
 
       U32 ir = int(255.999f * r);
       U32 ig = int(255.999f * g);
@@ -217,6 +226,31 @@ function void write_buffer_to_ppm(Vec3F32 *buffer, U32 image_width, U32 image_he
   fclose(file);
 }
 
+__device__ function U32 hit_sphere(Vec3F32 center, F32 radius, RayF32 r)
+{
+  // We take the quadratic formula -b +- sqrt(b*b-4ac) / 2a,
+  // and we calculate only the sqrt part. If there is a hit with the sphere we either
+  // have two solutions (positive sqrt), one solution (zero sqrt) or no solution (negative sqrt).
+  // If we have no solution we have no hit on the sphere centered at center, with the given radius.
+  // No hit -> return 0
+  // Otherwise we do have a hit, return 1
+
+  // Compare lines with RTIOW
+  // (C-Q)
+  Vec3F32 oc = sub_V3F32(center, r.origin);
+  // a = d.d
+  F32 a = dot_V3F32(r.direction, r.direction);
+  // b = -2d . (C-Q)
+  F32 b = dot_V3F32(scale_V3F32(-2.0f, r.direction), oc);
+  // c = (C-Q) . (C-Q) - r*r
+  F32 c = dot_V3F32(oc, oc) - radius*radius;
+
+  F32 discriminant = b*b - 4*a*c;
+  U32 out = discriminant >= 0.0f ? 1 : 0;
+
+  return out;
+}
+
 
 
 __global__ function void cuda_main(Vec3F32 *pixelbuffer, U32 *idxbuffer) 
@@ -232,6 +266,8 @@ __global__ function void cuda_main(Vec3F32 *pixelbuffer, U32 *idxbuffer)
     Vec3F32 px_u = scale_V3F32((F32)x, viewport.pixel_delta_u);
     Vec3F32 px_v = scale_V3F32((F32)y, viewport.pixel_delta_v);
     Vec3F32 pixel_center = add_V3F32(viewport.pixel_origin, add_V3F32(px_u, px_v));
+    
+    // TODO(anton): Maybe we dont need some ray structure here..
     Vec3F32 ray_direction = sub_V3F32(pixel_center, camera.center);
     RayF32 r = {0};
     r.origin = camera.center;
@@ -244,7 +280,20 @@ __global__ function void cuda_main(Vec3F32 *pixelbuffer, U32 *idxbuffer)
 
     // Lerp between white and light blue depending on y position
     F32 blend = 0.5f*(unit_dir.y + 1.0f);
-    pixelbuffer[idx] = lerp_V3F32(blend, white, light_blue);
+    Vec3F32 pixel_color = {0};
+
+    Vec3F32 sphere_center = vec3F32(0.0f, 0.0f, -1.0f);
+    F32 sphere_radius = 0.5f;
+    if(hit_sphere(sphere_center, sphere_radius, r)) 
+    {
+      pixel_color = vec3F32(1.0f, 0.0f, 0.0f);
+    } 
+    else
+    {
+      pixel_color = lerp_V3F32(blend, white, light_blue);
+    }
+
+    pixelbuffer[idx] = pixel_color;
 
     //pixelbuffer[idx].x = (F32)x/(F32)image.width;
     //pixelbuffer[idx].y = (F32)y/(F32)image.height;