diff --git a/src/editortest.cu b/src/editortest.cu deleted file mode 100644 index 64ea95c..0000000 --- a/src/editortest.cu +++ /dev/null @@ -1,12 +0,0 @@ - -__global__ void some_global_kernel() { - - some_kernel<<<32, 32>>>(0,0); -} - - -__device__ void some_kernel(int a, int b) -{ - - return; -} \ No newline at end of file diff --git a/src/main.cu b/src/main.cu index e956113..1373299 100644 --- a/src/main.cu +++ b/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; diff --git a/timeBuild.ctm b/timeBuild.ctm index c8a2e1e..957b33d 100644 Binary files a/timeBuild.ctm and b/timeBuild.ctm differ