summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--src/cl/spirt.cl46
-rw-r--r--src/spirt.c105
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;
generated by cgit v1.2.3 (git 2.46.0) at 2025-10-30 02:53:09 +0000