raytraced background gradient

2 files changed, 127 insertions, 24 deletions

diff --git a/src/cl/spirt.cl b/src/cl/spirt.cl
index 12022bb..26bdf00 100644
--- a/src/cl/spirt.cl
+++ b/src/cl/spirt.cl
@@ -13,21 +13,51 @@ GNU General Public License for more details.
 You should have received a copy of the GNU General Public License
 along with this program.  If not, see <http://www.gnu.org/licenses/>. */
 
+#define RAY_ORIG(RAYS, INDEX) (RAYS[INDEX])
+#define RAY_DIR(RAYS, INDEX) (RAYS[INDEX + 1])
+
+__kernel void
+gen_rays(__global float3 *const rays, float3 camera_centre,
+         float3 pixel_delta_u, float3 pixel_delta_v, float3 corner00)
+{
+	size_t idx, w, i, j;
+
+	w = get_global_size(0);
+
+	i = get_global_id(0);
+	j = get_global_id(1);
+
+	idx = (w * j + i) * 2;
+
+	RAY_ORIG(rays, idx) = camera_centre;
+	RAY_DIR(rays, idx) = corner00 +
+	                     ((i * pixel_delta_u) + (j * pixel_delta_v)) -
+	                     camera_centre;
+}
+
 __kernel void
-hello(__global unsigned char *const canvas)
+ray_colour(__global uchar *const canvas, __global float3 *const rays)
 {
-	size_t idx, w, h, i, j;
+	size_t canvas_idx, ray_idx, w, i, j;
+	float3 unit_direction, colour;
+	float a;
 
 	w = get_global_size(0);
-	h = get_global_size(1);
 
 	i = get_global_id(0);
 	j = get_global_id(1);
 
-	idx = (w * j + i) * 4;
+	canvas_idx = (w * j + i) * 4;
+	ray_idx = (w * j + i) * 2;
+
+	unit_direction = normalize(RAY_DIR(rays, ray_idx));
+	a = 0.5f * (unit_direction.y + 1.0f);
+
+	colour = (1.0f - a) * (float3)(1.0f) + a * (float3)(0.5f, 0.7f, 1.0f);
+	colour = 256.0f * clamp(sqrt(colour), 0.0f, 0.999f);
 
-	canvas[idx] = (unsigned char)((255 * i) / (w - 1));
-	canvas[idx + 1] = (unsigned char)((255 * j) / (h - 1));
-	canvas[idx + 2] = 63;
-	canvas[idx + 3] = 255;
+	canvas[canvas_idx] = colour.x;
+	canvas[canvas_idx + 1] = colour.y;
+	canvas[canvas_idx + 2] = colour.z;
+	canvas[canvas_idx + 3] = 255;
 }
diff --git a/src/spirt.c b/src/spirt.c
index fd6457b..527981f 100644
--- a/src/spirt.c
+++ b/src/spirt.c
@@ -33,7 +33,8 @@ struct kernel_context {
 	cl_context context;
 	cl_command_queue command_queue;
 	cl_program program;
-	cl_kernel k_hello;
+	cl_kernel k_gen_rays;
+	cl_kernel k_ray_colour;
 };
 
 struct opts {
@@ -43,8 +44,16 @@ struct opts {
 	long devidx;
 };
 
-static unsigned char *kernel_hello(struct kernel_context *runtime,
-                                   size_t image_width, size_t image_height);
+struct camera {
+	cl_float3 centre;
+	cl_float3 pixel_delta_u;
+	cl_float3 pixel_delta_v;
+	cl_float3 corner00;
+};
+
+static unsigned char *kernel_render(struct kernel_context *runtime,
+                                    size_t image_width, size_t image_height,
+                                    struct camera cam);
 
 static struct kernel_context kernel_context_init(size_t platidx, size_t devidx);
 static void kernel_context_clean(struct kernel_context *runtime);
@@ -55,18 +64,24 @@ extern const char *SPIRV_START[];
 extern const char *SPIRV_END[];
 
 static unsigned char *
-kernel_hello(struct kernel_context *const runtime, const size_t image_width,
-             const size_t image_height)
+kernel_render(struct kernel_context *const runtime, const size_t image_width,
+              const size_t image_height, const struct camera cam)
 {
 	cl_int err;
 
 	size_t global_item_size[2];
 
 	unsigned char *h_canvas;
-	cl_mem d_canvas;
-	size_t canvas_size;
+	cl_mem d_canvas, d_rays;
+	size_t rays_size, canvas_size;
 
 	canvas_size = (sizeof("RGBA") - 1) * image_width * image_height;
+	rays_size = 2 * image_width * image_height;
+
+	d_rays = clCreateBuffer(runtime->context, CL_MEM_READ_WRITE,
+	                        sizeof(cl_float3) * rays_size, NULL, &err);
+	if (err != CL_SUCCESS)
+		die("clCreateBuffer: %s\n", cl_strerror(err));
 
 	h_canvas = calloc(canvas_size, sizeof(*h_canvas));
 	if (!h_canvas)
@@ -77,15 +92,38 @@ kernel_hello(struct kernel_context *const runtime, const size_t image_width,
 	if (err != CL_SUCCESS)
 		die("clCreateBuffer: %s\n", cl_strerror(err));
 
-	err = clSetKernelArg(runtime->k_hello, 0, sizeof(d_canvas), &d_canvas);
+	err = clSetKernelArg(runtime->k_gen_rays, 0, sizeof(d_rays), &d_rays);
+	err |= clSetKernelArg(runtime->k_gen_rays, 1, sizeof(cl_float3),
+	                      &cam.centre);
+	err |= clSetKernelArg(runtime->k_gen_rays, 2, sizeof(cl_float3),
+	                      &cam.pixel_delta_u);
+	err |= clSetKernelArg(runtime->k_gen_rays, 3, sizeof(cl_float3),
+	                      &cam.pixel_delta_v);
+	err |= clSetKernelArg(runtime->k_gen_rays, 4, sizeof(cl_float3),
+	                      &cam.corner00);
 	if (err != CL_SUCCESS)
 		die("clSetKernelArg: %s\n", cl_strerror(err));
 
 	global_item_size[0] = image_width;
 	global_item_size[1] = image_height;
-	err = clEnqueueNDRangeKernel(runtime->command_queue, runtime->k_hello,
-	                             2, NULL, global_item_size, NULL, 0, NULL,
-	                             NULL);
+	err = clEnqueueNDRangeKernel(runtime->command_queue,
+	                             runtime->k_gen_rays, 2, NULL,
+	                             global_item_size, NULL, 0, NULL, NULL);
+	if (err != CL_SUCCESS)
+		die("clEnqueueNDRangeKernel: %s\n", cl_strerror(err));
+
+	err = clSetKernelArg(runtime->k_ray_colour, 0, sizeof(d_canvas),
+	                     &d_canvas);
+	err |= clSetKernelArg(runtime->k_ray_colour, 1, sizeof(d_rays),
+	                      &d_rays);
+	if (err != CL_SUCCESS)
+		die("clSetKernelArg: %s\n", cl_strerror(err));
+
+	global_item_size[0] = image_width;
+	global_item_size[1] = image_height;
+	err = clEnqueueNDRangeKernel(runtime->command_queue,
+	                             runtime->k_ray_colour, 2, NULL,
+	                             global_item_size, NULL, 0, NULL, NULL);
 	if (err != CL_SUCCESS)
 		die("clEnqueueNDRangeKernel: %s\n", cl_strerror(err));
 
@@ -95,6 +133,7 @@ kernel_hello(struct kernel_context *const runtime, const size_t image_width,
 	if (err != CL_SUCCESS)
 		die("clEnqueueReadBuffer: %s\n", cl_strerror(err));
 
+	clReleaseMemObject(d_rays);
 	clReleaseMemObject(d_canvas);
 
 	return h_canvas;
@@ -138,7 +177,11 @@ kernel_context_init(const size_t platidx, const size_t devidx)
 	runtime.program = compile_spirv_program(runtime.context, runtime.device,
 	                                        SPIRV_START, SPIRV_SIZE);
 
-	runtime.k_hello = clCreateKernel(runtime.program, "hello", &err);
+	runtime.k_gen_rays = clCreateKernel(runtime.program, "gen_rays", &err);
+	if (err != CL_SUCCESS)
+		die("clCreateKernel: %s\n", cl_strerror(err));
+	runtime.k_ray_colour =
+	        clCreateKernel(runtime.program, "ray_colour", &err);
 	if (err != CL_SUCCESS)
 		die("clCreateKernel: %s\n", cl_strerror(err));
 
@@ -151,7 +194,8 @@ kernel_context_clean(struct kernel_context *const runtime)
 	clFlush(runtime->command_queue);
 	clFinish(runtime->command_queue);
 
-	clReleaseKernel(runtime->k_hello);
+	clReleaseKernel(runtime->k_gen_rays);
+	clReleaseKernel(runtime->k_ray_colour);
 	clReleaseProgram(runtime->program);
 	clReleaseCommandQueue(runtime->command_queue);
 	clReleaseContext(runtime->context);
@@ -164,8 +208,8 @@ get_args(int argc, char *argv[])
 
 	opts.platidx = 0;
 	opts.devidx = 0;
-	opts.width = 256;
-	opts.height = 256;
+	opts.width = 1024;
+	opts.height = 576;
 
 	if (argv[0])
 		argv++, argc--;
@@ -213,6 +257,8 @@ main(int argc, char *argv[])
 
 	struct kernel_context runtime;
 	unsigned char *h_canvas;
+	struct camera cam;
+	float focal_length, viewport_width, viewport_height;
 
 	Image img;
 	Texture2D texture;
@@ -221,7 +267,34 @@ main(int argc, char *argv[])
 
 	runtime = kernel_context_init(opts.platidx, opts.devidx);
 
-	h_canvas = kernel_hello(&runtime, opts.width, opts.height);
+	focal_length = 1.0f;
+	viewport_height = 2.0f;
+	viewport_width =
+	        viewport_height * ((float)opts.width / (float)opts.height);
+
+	cam.centre.s[0] = 0.0f;
+	cam.centre.s[1] = 0.0f;
+	cam.centre.s[2] = 0.0f;
+
+	cam.pixel_delta_u.s[0] = viewport_width / (float)opts.width;
+	cam.pixel_delta_u.s[1] = 0.0f;
+	cam.pixel_delta_u.s[2] = 0.0f;
+
+	cam.pixel_delta_v.s[0] = 0.0f;
+	cam.pixel_delta_v.s[1] = -viewport_height / (float)opts.height;
+	cam.pixel_delta_v.s[2] = 0.0f;
+
+	cam.corner00.s[0] =
+	        cam.centre.s[0] - viewport_width / 2.0f +
+	        0.5f * (cam.pixel_delta_u.s[0] + cam.pixel_delta_v.s[0]);
+	cam.corner00.s[1] =
+	        cam.centre.s[1] + viewport_height / 2.0f +
+	        0.5f * (cam.pixel_delta_u.s[1] + cam.pixel_delta_v.s[1]);
+	cam.corner00.s[2] =
+	        cam.centre.s[2] - focal_length +
+	        0.5f * (cam.pixel_delta_u.s[2] + cam.pixel_delta_v.s[2]);
+
+	h_canvas = kernel_render(&runtime, opts.width, opts.height, cam);
 
 	img.data = h_canvas;
 	img.width = opts.width;