From 5ff9d6184771970c14e682ecf55cbd9d27b2fada Mon Sep 17 00:00:00 2001 From: Aiden Gall Date: Tue, 16 Jan 2024 19:24:18 +0000 Subject: initial commit --- src/spirt.c | 357 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 357 insertions(+) create mode 100644 src/spirt.c (limited to 'src/spirt.c') diff --git a/src/spirt.c b/src/spirt.c new file mode 100644 index 0000000..d1e8399 --- /dev/null +++ b/src/spirt.c @@ -0,0 +1,357 @@ +/* Copyright (C) 2024 Aiden Gall + +This program is free software: you can redistribute it and/or modify +it under the terms of the GNU General Public License as published by +the Free Software Foundation, either version 3 of the License, or +(at your option) any later version. + +This program is distributed in the hope that it will be useful, +but WITHOUT ANY WARRANTY; without even the implied warranty of +MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the +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 . */ + +#include "farbfeld.h" +#include "util.h" + +#include +#include +#include +#include +#include +#include +#include +#include + +#ifndef SPIRV_START +# define SPIRV_START _binary_src_cl_spirt_spv_start +#endif +#ifndef SPIRV_END +# define SPIRV_END _binary_src_cl_spirt_spv_end +#endif +#define SPIRV_SIZE ((size_t)SPIRV_END - (size_t)SPIRV_START) + +struct ClRuntime { + cl_platform_id platform; + cl_device_id device; + cl_context context; + cl_command_queue command_queue; + cl_program program; + cl_kernel k_hello; +}; + +static cl_platform_id *get_platforms(cl_uint *num_platforms); +static cl_device_id *get_devices(cl_platform_id platform, cl_uint *num_devices); +static cl_program compile_spirv_program(cl_context context, cl_device_id device, + const void *spirv_start, + size_t spirv_size); + +static struct ClRuntime init(size_t platidx, size_t devidx); +static uint16_t *run_k_hello(struct ClRuntime *runtime, size_t *canvas_size_ret, + size_t image_width, size_t image_height); +static void clean(struct ClRuntime *runtime); + +static long argtonum(const char *arg, long minval, long maxval); + +extern const char *SPIRV_START[]; +extern const char *SPIRV_END[]; + +static cl_platform_id * +get_platforms(cl_uint *const num_platforms) +{ + cl_int err; + + cl_platform_id *platforms; + cl_uint len; + + err = clGetPlatformIDs(0, NULL, &len); + if (err != CL_SUCCESS) + die("clGetPlatformIDs: %s\n", cl_strerror(err)); + warn("number of platforms = %d\n", len); + if (len < 1) + die("clGetPlatformIDs: No OpenCL platforms\n"); + + platforms = calloc(len, sizeof(*platforms)); + if (!platforms) + die("calloc: Out of memory\n"); + err = clGetPlatformIDs(len, platforms, NULL); + if (err != CL_SUCCESS) + die("clGetPlatformIDs: %s\n", cl_strerror(err)); + + *num_platforms = len; + return platforms; +} + +static cl_device_id * +get_devices(const cl_platform_id platform, cl_uint *const num_devices) +{ + cl_int err; + + cl_device_id *devices; + cl_uint len; + + err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, NULL, &len); + if (err != CL_SUCCESS) + die("clGetDeviceIDs: %s\n", cl_strerror(err)); + warn("number of devices in platform = %d\n", len); + if (len < 1) + die("clGetDeviceIDs: No OpenCL devices in platform\n"); + + devices = calloc(len, sizeof(*devices)); + if (!devices) + die("calloc: Out of memory\n"); + err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, len, devices, NULL); + if (err != CL_SUCCESS) + die("clGetDeviceIDs: %s\n", cl_strerror(err)); + + *num_devices = len; + return devices; +} + +static cl_program +compile_spirv_program(const cl_context context, const cl_device_id device, + const void *const spirv_start, const size_t spirv_size) +{ + cl_int err, build_err; + cl_program program; + + program = clCreateProgramWithIL(context, spirv_start, spirv_size, &err); + if (err != CL_SUCCESS) + die("clCreateProgramWithIL: %s\n", cl_strerror(err)); + + build_err = clBuildProgram(program, 1, &device, NULL, NULL, NULL); + if (build_err != CL_SUCCESS) { + cl_char *log; + size_t buffer_size; + + err = clGetProgramBuildInfo(program, device, + CL_PROGRAM_BUILD_LOG, 0, NULL, + &buffer_size); + if (err != CL_SUCCESS) + die("clGetProgramBuildInfo: %s\n", cl_strerror(err)); + if (buffer_size < 1) + die("clGetProgramBuildInfo: " + "Build log buffer is empty\n"); + + log = calloc(buffer_size, sizeof(*log)); + if (!log) + die("calloc: Out of memory\n"); + err = clGetProgramBuildInfo( + program, device, CL_PROGRAM_BUILD_LOG, + sizeof(*log) * buffer_size, log, NULL); + if (err != CL_SUCCESS) + die("clGetProgramBuildInfo: %s\n", cl_strerror(err)); + + efwrite(log, sizeof(*log), buffer_size, stderr); + + free(log); + die("\nclBuildProgram: %s\n", cl_strerror(build_err)); + } + + return program; +} + +static struct ClRuntime +init(const size_t platidx, const size_t devidx) +{ + cl_int err; + + cl_platform_id *platforms; + cl_device_id *devices; + cl_uint num_platforms; + cl_uint num_devices; + + struct ClRuntime runtime; + + platforms = get_platforms(&num_platforms); + if (platidx >= num_platforms) + die("get_platforms: Index out of range\n"); + runtime.platform = platforms[platidx]; + free(platforms); + + devices = get_devices(runtime.platform, &num_devices); + if (devidx >= num_devices) + die("get_devices: Index out of range\n"); + runtime.device = devices[devidx]; + free(devices); + + runtime.context = + clCreateContext(NULL, 1, &runtime.device, NULL, NULL, &err); + if (err != CL_SUCCESS) + die("clCreateContext: %s\n", cl_strerror(err)); + + runtime.command_queue = clCreateCommandQueueWithProperties( + runtime.context, runtime.device, NULL, &err); + if (err != CL_SUCCESS) + die("clCreateCommandQueueWithProperties: %s\n", + cl_strerror(err)); + + runtime.program = compile_spirv_program(runtime.context, runtime.device, + SPIRV_START, SPIRV_SIZE); + + runtime.k_hello = clCreateKernel(runtime.program, "hello", &err); + if (err != CL_SUCCESS) + die("clCreateKernel: %s\n", cl_strerror(err)); + + return runtime; +} + +static uint16_t * +run_k_hello(struct ClRuntime *const runtime, size_t *const canvas_size_ret, + const size_t image_width, const size_t image_height) +{ + cl_int err; + + size_t global_item_size[2]; + + uint16_t *h_canvas; + cl_mem d_canvas; + size_t canvas_size; + + canvas_size = (sizeof("RGBA") - 1) * image_width * image_height; + + h_canvas = calloc(canvas_size, sizeof(*h_canvas)); + if (!h_canvas) + die("calloc: Out of memory\n"); + + d_canvas = clCreateBuffer(runtime->context, CL_MEM_WRITE_ONLY, + sizeof(*h_canvas) * canvas_size, NULL, &err); + if (err != CL_SUCCESS) + die("clCreateBuffer: %s\n", cl_strerror(err)); + + err = clSetKernelArg(runtime->k_hello, 0, sizeof(d_canvas), &d_canvas); + 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); + if (err != CL_SUCCESS) + die("clEnqueueNDRangeKernel: %s\n", cl_strerror(err)); + + err = clEnqueueReadBuffer(runtime->command_queue, d_canvas, CL_TRUE, 0, + sizeof(*h_canvas) * canvas_size, h_canvas, 0, + NULL, NULL); + if (err != CL_SUCCESS) + die("clEnqueueReadBuffer: %s\n", cl_strerror(err)); + + clReleaseMemObject(d_canvas); + + *canvas_size_ret = canvas_size; + return h_canvas; +} + +static void +clean(struct ClRuntime *const runtime) +{ + clFlush(runtime->command_queue); + clFinish(runtime->command_queue); + + clReleaseKernel(runtime->k_hello); + clReleaseProgram(runtime->program); + clReleaseCommandQueue(runtime->command_queue); + clReleaseContext(runtime->context); +} + +static long +argtonum(const char *const arg, const long minval, const long maxval) +{ + long ret; + char *endptr; + + if (minval > maxval) + die("argtonum: Invalid range\n"); + if (!arg) + die("argtonum: Missing argument\n"); + + errno = 0; + ret = strtol(arg, &endptr, 10); + if (errno) + die("strtol: %s\n", strerror(errno)); + if (arg == endptr || *endptr) + die("strtol: Invalid integer\n"); + + if (ret < minval || ret > maxval) + die("argtonum: Argument out of range\n"); + + return ret; +} + +int +main(int argc, char *argv[]) +{ + long opt_platidx, opt_devidx, opt_width, opt_height; + FILE *fd; + + struct ClRuntime runtime; + + uint16_t *h_canvas; + size_t canvas_size; + + opt_platidx = 0; + opt_devidx = 0; + opt_width = 256; + opt_height = 256; + + if (argv[0]) + argv++, argc--; + + for (; argv[0] && argv[0][0] == '-' && argv[0][1]; argv++, argc--) { + char flag; + + if (argv[0][2]) + die("arg_parse: Invalid flag '%s'\n", argv[0]); + + flag = argv[0][1]; + argv++, argc--; + + switch (flag) { + case 'p': + opt_platidx = argtonum(argv[0], 0, LONG_MAX); + break; + case 'd': + opt_devidx = argtonum(argv[0], 0, LONG_MAX); + break; + case 'w': + opt_width = argtonum(argv[0], 1, LONG_MAX); + break; + case 'h': + opt_height = argtonum(argv[0], 1, LONG_MAX); + break; + case '-': + goto end_flags; + default: + die("arg_parse: Unknown flag '-%c'\n", flag); + } + } + +end_flags: + if (argc > 1) + die("arg_parse: Unexpected argument\n"); + + fd = stdout; + if (argv[0]) { + errno = 0; + fd = fopen(argv[0], "w"); + if (errno) + die("fopen: %s\n", strerror(errno)); + if (!fd) + die("fopen: Returned NULL\n"); + } + + runtime = init(opt_platidx, opt_devidx); + + h_canvas = run_k_hello(&runtime, &canvas_size, opt_width, opt_height); + + farbfeld_write(fd, h_canvas, canvas_size, opt_width, opt_height); + fclose(fd); + + free(h_canvas); + clean(&runtime); + + return 0; +} -- cgit v1.2.3