Commit: 684ec5b99db42cf252db5032e618b0c79e4739e1
Parent: b42585b2fd21afda9f14a25e875ec29cb83e1dac
Author: Randy Palamar
Date: Tue, 8 Oct 2024 10:52:49 -0600
reorganize code allowing for platform specific main()
Also the beamformer code is no longer allowed direct access to the
platform code. This prevents bugs such as allocating memory from
separate memory spaces and then trying to use them interchangeably.
Diffstat:
| M | beamformer.c | | | 12 | +++++++----- |
| M | beamformer.h | | | 88 | ++++++++++++++++++++++++++++++++++++++++++------------------------------------- |
| M | build.sh | | | 7 | +++---- |
| D | main.c | | | 334 | ------------------------------------------------------------------------------- |
| A | main_generic.c | | | 56 | ++++++++++++++++++++++++++++++++++++++++++++++++++++++++ |
| M | os_unix.c | | | 55 | +++++++++++++++++++++++-------------------------------- |
| M | os_win32.c | | | 103 | ++++++++++++++++++++++++++++++++++--------------------------------------------- |
| A | static.c | | | 317 | +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ |
| M | util.h | | | 33 | +++++++++++++++++++-------------- |
9 files changed, 516 insertions(+), 489 deletions(-)
diff --git a/beamformer.c b/beamformer.c
@@ -107,7 +107,7 @@ alloc_shader_storage(BeamformerCtx *ctx, Arena a)
full_rf_buf_size, map_flags);
break;
case GL_VENDOR_NVIDIA:
- cs->raw_data_arena = os_alloc_arena(cs->raw_data_arena, full_rf_buf_size);
+ cs->raw_data_arena = ctx->platform.alloc_arena(cs->raw_data_arena, full_rf_buf_size);
ctx->cuda_lib.register_cuda_buffers(cs->rf_data_ssbos, ARRAY_COUNT(cs->rf_data_ssbos),
cs->raw_data_ssbo);
ctx->cuda_lib.init_cuda_configuration(bp->rf_raw_dim.E, bp->dec_data_dim.E,
@@ -379,7 +379,7 @@ do_beamformer(BeamformerCtx *ctx, Arena arena)
BeamformerParameters *bp = &ctx->params->raw;
/* NOTE: Check for and Load RF Data into GPU */
- if (os_poll_pipe(ctx->data_pipe)) {
+ if (ctx->platform.poll_pipe(ctx->data_pipe)) {
ComputeShaderCtx *cs = &ctx->csctx;
if (!uv4_equal(cs->dec_data_dim, bp->dec_data_dim))
alloc_shader_storage(ctx, arena);
@@ -404,7 +404,7 @@ do_beamformer(BeamformerCtx *ctx, Arena arena)
size rf_raw_size = rf_raw_dim.x * rf_raw_dim.y * sizeof(i16);
void *rf_data_buf = cs->raw_data_arena.beg + raw_index * rf_raw_size;
- size rlen = os_read_pipe_data(ctx->data_pipe, rf_data_buf, rf_raw_size);
+ size rlen = ctx->platform.read_pipe(ctx->data_pipe, rf_data_buf, rf_raw_size);
if (rlen != rf_raw_size) {
ctx->partial_transfer_count++;
} else {
@@ -467,11 +467,11 @@ do_beamformer(BeamformerCtx *ctx, Arena arena)
ExportCtx *e = &ctx->export_ctx;
uv4 dim = e->volume_dim;
size volume_out_size = dim.x * dim.y * dim.z * sizeof(f32);
- e->volume_buf = os_alloc_arena(e->volume_buf, volume_out_size);
+ e->volume_buf = ctx->platform.alloc_arena(e->volume_buf, volume_out_size);
glGetTextureImage(e->volume_texture, 0, GL_RED, GL_FLOAT, volume_out_size,
e->volume_buf.beg);
s8 raw = {.len = volume_out_size, .data = e->volume_buf.beg};
- if (!os_write_file("raw_volume.bin", raw))
+ if (!ctx->platform.write_new_file("raw_volume.bin", raw))
TraceLog(LOG_WARNING, "failed to write output volume\n");
}
}
@@ -501,4 +501,6 @@ do_beamformer(BeamformerCtx *ctx, Arena arena)
if (IsKeyPressed(KEY_R))
ctx->flags |= RELOAD_SHADERS;
+ if (WindowShouldClose())
+ ctx->flags |= SHOULD_EXIT;
}
diff --git a/beamformer.h b/beamformer.h
@@ -2,6 +2,12 @@
#ifndef _BEAMFORMER_H_
#define _BEAMFORMER_H_
+#include <glad.h>
+
+#define GRAPHICS_API_OPENGL_43
+#include <raylib.h>
+#include <rlgl.h>
+
#include "util.h"
#define BG_COLOUR (v4){.r = 0.15, .g = 0.12, .b = 0.13, .a = 1.0}
@@ -21,7 +27,8 @@
#define RECT_BTN_BORDER_WIDTH 6.0f
enum program_flags {
- RELOAD_SHADERS = 1 << 0,
+ SHOULD_EXIT = 1 << 0,
+ RELOAD_SHADERS = 1 << 1,
GEN_MIPMAPS = 1 << 29,
DO_COMPUTE = 1 << 30,
};
@@ -56,55 +63,26 @@ typedef struct {
f32 cursor_blink_target;
} InputState;
-#if defined(__unix__)
- #include "os_unix.c"
-
- #ifdef _DEBUG
- #define DEBUG_EXPORT
- #define OS_DEBUG_LIB_NAME "./beamformer.so"
- #define OS_DEBUG_LIB_TEMP_NAME "./beamformer_temp.so"
- #else
- #define DEBUG_EXPORT static
- #endif
-
- #define OS_CUDA_LIB_NAME "./external/cuda_toolkit.so"
- #define OS_CUDA_LIB_TEMP_NAME "./external/cuda_toolkit_temp.so"
-
- #define OS_PIPE_NAME "/tmp/beamformer_data_fifo"
- #define OS_SMEM_NAME "/ogl_beamformer_parameters"
-#elif defined(_WIN32)
- #include "os_win32.c"
-
- #ifdef _DEBUG
- #define DEBUG_EXPORT __declspec(dllexport)
- #define OS_DEBUG_LIB_NAME "beamformer.dll"
- #define OS_DEBUG_LIB_TEMP_NAME "beamformer_temp.dll"
- #else
- #define DEBUG_EXPORT static
- #endif
-
- #define OS_CUDA_LIB_NAME "external\\cuda_toolkit.dll"
- #define OS_CUDA_LIB_TEMP_NAME "external\\cuda_toolkit_temp.dll"
-
- #define OS_PIPE_NAME "\\\\.\\pipe\\beamformer_data_fifo"
- #define OS_SMEM_NAME "Local\\ogl_beamformer_parameters"
-#else
- #error Unsupported Platform!
-#endif
-
#define MAX_FRAMES_IN_FLIGHT 3
#define INIT_CUDA_CONFIGURATION_FN(name) void name(u32 *input_dims, u32 *decoded_dims, u16 *channel_mapping, b32 rx_cols)
typedef INIT_CUDA_CONFIGURATION_FN(init_cuda_configuration_fn);
+INIT_CUDA_CONFIGURATION_FN(init_cuda_configuration_stub) {}
+
#define REGISTER_CUDA_BUFFERS_FN(name) void name(u32 *rf_data_ssbos, u32 rf_buffer_count, u32 raw_data_ssbo)
typedef REGISTER_CUDA_BUFFERS_FN(register_cuda_buffers_fn);
+REGISTER_CUDA_BUFFERS_FN(register_cuda_buffers_stub) {}
+
#define CUDA_DECODE_FN(name) void name(size_t input_offset, u32 output_buffer_idx)
typedef CUDA_DECODE_FN(cuda_decode_fn);
+CUDA_DECODE_FN(cuda_decode_stub) {}
+
#define CUDA_HILBERT_FN(name) void name(u32 input_buffer_idx, u32 output_buffer_idx)
typedef CUDA_HILBERT_FN(cuda_hilbert_fn);
+CUDA_HILBERT_FN(cuda_hilbert_stub) {}
typedef struct {
- os_library_handle lib;
+ void *lib;
u64 timestamp;
init_cuda_configuration_fn *init_cuda_configuration;
register_cuda_buffers_fn *register_cuda_buffers;
@@ -112,6 +90,33 @@ typedef struct {
cuda_hilbert_fn *cuda_hilbert;
} CudaLib;
+#define PLATFORM_ALLOC_ARENA_FN(name) Arena name(Arena old, size capacity)
+typedef PLATFORM_ALLOC_ARENA_FN(platform_alloc_arena_fn);
+
+#define PLATFORM_POLL_PIPE_FN(name) b32 name(Pipe p)
+typedef PLATFORM_POLL_PIPE_FN(platform_poll_pipe_fn);
+
+#define PLATFORM_READ_PIPE_FN(name) size name(Pipe p, void *buf, size len)
+typedef PLATFORM_READ_PIPE_FN(platform_read_pipe_fn);
+
+#define PLATFORM_WRITE_NEW_FILE_FN(name) b32 name(char *fname, s8 raw)
+typedef PLATFORM_WRITE_NEW_FILE_FN(platform_write_new_file_fn);
+
+typedef struct {
+ platform_alloc_arena_fn *alloc_arena;
+ platform_poll_pipe_fn *poll_pipe;
+ platform_read_pipe_fn *read_pipe;
+ platform_write_new_file_fn *write_new_file;
+} Platform;
+
+#include "beamformer_parameters.h"
+typedef struct {
+ BeamformerParameters raw;
+ enum compute_shaders compute_stages[16];
+ u32 compute_stages_count;
+ b32 upload;
+} BeamformerParametersFull;
+
typedef struct {
u32 programs[CS_LAST];
@@ -214,10 +219,11 @@ typedef struct {
FragmentShaderCtx fsctx;
ExportCtx export_ctx;
- os_pipe data_pipe;
- u32 partial_transfer_count;
+ Pipe data_pipe;
+ u32 partial_transfer_count;
- CudaLib cuda_lib;
+ CudaLib cuda_lib;
+ Platform platform;
BeamformerParametersFull *params;
} BeamformerCtx;
diff --git a/build.sh b/build.sh
@@ -9,10 +9,10 @@ debug=${DEBUG}
cc=${CC:-cc}
system_raylib=${USE_SYSTEM_RAYLIB}
+main=main_generic.c
-case $(uname -s) in
+case $(uname -sm) in
MINGW64*)
- os="win32"
ldflags="$ldflags -lgdi32 -lwinmm"
[! ${NO_MATLAB} ] && [ -d "C:/Program Files/MATLAB/R2022a/extern/lib/win64/microsoft" ] &&
libcflags="$libcflags -DMATLAB_CONSOLE"
@@ -22,7 +22,6 @@ MINGW64*)
-llibmat -llibmex
;;
Linux*)
- os="unix"
cflags="$cflags -D_DEFAULT_SOURCE"
libname="beamformer.so"
${cc} $libcflags helpers/ogl_beamformer_lib.c -o helpers/ogl_beamformer_lib.so
@@ -68,4 +67,4 @@ else
[ ! "$system_raylib" ] && ldflags="./external/lib/libraylib.a $ldflags"
fi
-${cc} $cflags -o ogl main.c $ldflags
+${cc} $cflags -o ogl $main $ldflags
diff --git a/main.c b/main.c
@@ -1,334 +0,0 @@
-/* See LICENSE for license details. */
-#include "beamformer.h"
-
-static s8 compute_shader_paths[CS_LAST] = {
- [CS_HADAMARD] = s8("shaders/hadamard.glsl"),
- [CS_HERCULES] = s8("shaders/hercules.glsl"),
- [CS_DEMOD] = s8("shaders/demod.glsl"),
- [CS_MIN_MAX] = s8("shaders/min_max.glsl"),
- [CS_SUM] = s8("shaders/sum.glsl"),
- [CS_UFORCES] = s8("shaders/uforces.glsl"),
-};
-
-#ifndef _DEBUG
-
-#include "beamformer.c"
-static void do_debug(void) { }
-
-#else
-static os_library_handle libhandle;
-
-typedef void do_beamformer_fn(BeamformerCtx *, Arena);
-static do_beamformer_fn *do_beamformer;
-
-static void
-do_debug(void)
-{
- static f32 updated_time;
- FileStats test_stats = os_get_file_stats(OS_DEBUG_LIB_NAME);
- if (test_stats.filesize > 32 && test_stats.timestamp > updated_time) {
- os_unload_library(libhandle);
- libhandle = os_load_library(OS_DEBUG_LIB_NAME, OS_DEBUG_LIB_TEMP_NAME);
- do_beamformer = os_lookup_dynamic_symbol(libhandle, "do_beamformer");
- updated_time = test_stats.timestamp;
- }
-}
-
-#endif /* _DEBUG */
-
-/* NOTE: cuda lib stubs */
-INIT_CUDA_CONFIGURATION_FN(init_cuda_configuration_stub) {}
-REGISTER_CUDA_BUFFERS_FN(register_cuda_buffers_stub) {}
-CUDA_DECODE_FN(cuda_decode_stub) {}
-CUDA_HILBERT_FN(cuda_hilbert_stub) {}
-
-static void
-gl_debug_logger(u32 src, u32 type, u32 id, u32 lvl, i32 len, const char *msg, const void *userctx)
-{
- (void)src; (void)type; (void)id; (void)userctx;
-
- u8 buf[128];
- Stream s = {.data = buf, .cap = ARRAY_COUNT(buf)};
- stream_append_s8(&s, s8("[GL DEBUG "));
- switch (lvl) {
- case GL_DEBUG_SEVERITY_HIGH: stream_append_s8(&s, s8("HIGH]: ")); break;
- case GL_DEBUG_SEVERITY_MEDIUM: stream_append_s8(&s, s8("MEDIUM]: ")); break;
- case GL_DEBUG_SEVERITY_LOW: stream_append_s8(&s, s8("LOW]: ")); break;
- case GL_DEBUG_SEVERITY_NOTIFICATION: stream_append_s8(&s, s8("NOTIFICATION]: ")); break;
- default: stream_append_s8(&s, s8("INVALID]: ")); break;
- }
- os_write_err_msg(stream_to_s8(s));
- os_write_err_msg((s8){.len = len, .data = (u8 *)msg});
- os_write_err_msg(s8("\n"));
-}
-
-static void
-get_gl_params(GLParams *gl)
-{
- char *vendor = (char *)glGetString(GL_VENDOR);
- if (!vendor) {
- os_write_err_msg(s8("Failed to determine GL Vendor\n"));
- os_fail();
- }
- switch (vendor[0]) {
- case 'A': gl->vendor_id = GL_VENDOR_AMD; break;
- case 'I': gl->vendor_id = GL_VENDOR_INTEL; break;
- case 'N': gl->vendor_id = GL_VENDOR_NVIDIA; break;
- default: {
- os_write_err_msg(s8("Unknown GL Vendor: "));
- os_write_err_msg(cstr_to_s8(vendor));
- os_write_err_msg(s8("\n"));
- os_fail();
- } break;
- }
-
- glGetIntegerv(GL_MAJOR_VERSION, &gl->version_major);
- glGetIntegerv(GL_MINOR_VERSION, &gl->version_minor);
- glGetIntegerv(GL_MAX_TEXTURE_SIZE, &gl->max_2d_texture_dim);
- glGetIntegerv(GL_MAX_3D_TEXTURE_SIZE, &gl->max_3d_texture_dim);
- glGetIntegerv(GL_MAX_SHADER_STORAGE_BLOCK_SIZE, &gl->max_ssbo_size);
- glGetIntegerv(GL_MAX_UNIFORM_BLOCK_SIZE, &gl->max_ubo_size);
-}
-
-static void
-validate_gl_requirements(GLParams *gl)
-{
- ASSERT(gl->max_ubo_size >= sizeof(BeamformerParameters));
- if (gl->version_major < 4 || (gl->version_major == 4 && gl->version_minor < 5)) {
- os_write_err_msg(s8("Only OpenGL Versions 4.5 or newer are supported!\n"));
- os_fail();
- }
-}
-
-static void
-dump_gl_params(GLParams *gl, Arena a)
-{
- (void)gl; (void)a;
-#ifdef _DEBUG
- Stream s = stream_alloc(&a, 1 * MEGABYTE);
- stream_append_s8(&s, s8("---- GL Parameters ----\n"));
- switch (gl->vendor_id) {
- case GL_VENDOR_AMD: stream_append_s8(&s, s8("Vendor: AMD\n")); break;
- case GL_VENDOR_INTEL: stream_append_s8(&s, s8("Vendor: Intel\n")); break;
- case GL_VENDOR_NVIDIA: stream_append_s8(&s, s8("Vendor: nVidia\n")); break;
- }
- stream_append_s8(&s, s8("Version: "));
- stream_append_i64(&s, gl->version_major);
- stream_append_s8(&s, s8("."));
- stream_append_i64(&s, gl->version_minor);
- stream_append_s8(&s, s8("\nMax 1D/2D Texture Dimension: "));
- stream_append_i64(&s, gl->max_2d_texture_dim);
- stream_append_s8(&s, s8("\nMax 3D Texture Dimension: "));
- stream_append_i64(&s, gl->max_3d_texture_dim);
- stream_append_s8(&s, s8("\nMax SSBO Size: "));
- stream_append_i64(&s, gl->max_ssbo_size);
- stream_append_s8(&s, s8("\nMax UBO Size: "));
- stream_append_i64(&s, gl->max_ubo_size);
- stream_append_s8(&s, s8("\n-----------------------\n"));
- if (!s.errors)
- os_write_err_msg(stream_to_s8(s));
-#endif
-}
-
-static u32
-compile_shader(Arena a, u32 type, s8 shader)
-{
- u32 sid = glCreateShader(type);
- glShaderSource(sid, 1, (const char **)&shader.data, (int *)&shader.len);
- glCompileShader(sid);
-
- i32 res = 0;
- glGetShaderiv(sid, GL_COMPILE_STATUS, &res);
-
- char *stype;
- switch (type) {
- case GL_COMPUTE_SHADER: stype = "Compute"; break;
- case GL_FRAGMENT_SHADER: stype = "Fragment"; break;
- }
-
- if (res == GL_FALSE) {
- TraceLog(LOG_WARNING, "SHADER: [ID %u] %s shader failed to compile", sid, stype);
- i32 len = 0;
- glGetShaderiv(sid, GL_INFO_LOG_LENGTH, &len);
- s8 err = s8alloc(&a, len);
- glGetShaderInfoLog(sid, len, (int *)&err.len, (char *)err.data);
- TraceLog(LOG_WARNING, "SHADER: [ID %u] Compile error: %s", sid, (char *)err.data);
- glDeleteShader(sid);
- } else {
- TraceLog(LOG_INFO, "SHADER: [ID %u] %s shader compiled successfully", sid, stype);
- }
-
- return sid;
-}
-
-static void
-init_fragment_shader_ctx(FragmentShaderCtx *ctx, uv4 out_data_dim)
-{
- ctx->output = LoadRenderTexture(out_data_dim.x, out_data_dim.y);
- ctx->db = -50.0f;
-}
-
-static void
-reload_shaders(BeamformerCtx *ctx, Arena a)
-{
- ComputeShaderCtx *csctx = &ctx->csctx;
- for (u32 i = 0; i < ARRAY_COUNT(csctx->programs); i++) {
- if (!compute_shader_paths[i].len)
- continue;
-
- Arena tmp = a;
- FileStats fs = os_get_file_stats((char *)compute_shader_paths[i].data);
- s8 shader_text = os_read_file(&tmp, (char *)compute_shader_paths[i].data, fs.filesize);
- if (shader_text.len == -1) {
- os_write_err_msg(s8("failed to read shader: "));
- os_write_err_msg(compute_shader_paths[i]);
- os_write_err_msg(s8("\n"));
- os_fail();
- }
- u32 shader_id = compile_shader(tmp, GL_COMPUTE_SHADER, shader_text);
-
- if (shader_id) {
- glDeleteProgram(csctx->programs[i]);
- csctx->programs[i] = rlLoadComputeShaderProgram(shader_id);
- ctx->flags |= DO_COMPUTE;
- }
-
- glDeleteShader(shader_id);
- }
-
- csctx->volume_export_pass_id = glGetUniformLocation(csctx->programs[CS_HERCULES],
- "u_volume_export_pass");
- csctx->volume_export_dim_offset_id = glGetUniformLocation(csctx->programs[CS_HERCULES],
- "u_volume_export_dim_offset");
- csctx->xdc_transform_id = glGetUniformLocation(csctx->programs[CS_UFORCES],
- "u_xdc_transform");
- csctx->xdc_index_id = glGetUniformLocation(csctx->programs[CS_UFORCES],
- "u_xdc_index");
-
- csctx->mips_level_id = glGetUniformLocation(csctx->programs[CS_MIN_MAX], "u_mip_map");
-
- csctx->sum_prescale_id = glGetUniformLocation(csctx->programs[CS_SUM], "u_prescale");
-
- Shader updated_fs = LoadShader(NULL, "shaders/render.glsl");
- if (updated_fs.id != rlGetShaderIdDefault()) {
- UnloadShader(ctx->fsctx.shader);
- ctx->fsctx.shader = updated_fs;
- ctx->fsctx.db_cutoff_id = GetShaderLocation(updated_fs, "u_db_cutoff");
- }
-}
-
-static void
-validate_cuda_lib(CudaLib *cl)
-{
- if (!cl->init_cuda_configuration) cl->init_cuda_configuration = init_cuda_configuration_stub;
- if (!cl->register_cuda_buffers) cl->register_cuda_buffers = register_cuda_buffers_stub;
- if (!cl->cuda_decode) cl->cuda_decode = cuda_decode_stub;
- if (!cl->cuda_hilbert) cl->cuda_hilbert = cuda_hilbert_stub;
-}
-
-static void
-check_and_load_cuda_lib(CudaLib *cl)
-{
- FileStats current = os_get_file_stats(OS_CUDA_LIB_NAME);
- if (cl->timestamp == current.timestamp || current.filesize < 32)
- return;
-
- TraceLog(LOG_INFO, "Loading CUDA lib: %s", OS_CUDA_LIB_NAME);
-
- cl->timestamp = current.timestamp;
- os_unload_library(cl->lib);
- cl->lib = os_load_library(OS_CUDA_LIB_NAME, OS_CUDA_LIB_TEMP_NAME);
-
- cl->init_cuda_configuration = os_lookup_dynamic_symbol(cl->lib, "init_cuda_configuration");
- cl->register_cuda_buffers = os_lookup_dynamic_symbol(cl->lib, "register_cuda_buffers");
- cl->cuda_decode = os_lookup_dynamic_symbol(cl->lib, "cuda_decode");
- cl->cuda_hilbert = os_lookup_dynamic_symbol(cl->lib, "cuda_hilbert");
-
- validate_cuda_lib(cl);
-}
-
-int
-main(void)
-{
- BeamformerCtx ctx = {0};
-
- Arena temp_memory = os_alloc_arena((Arena){0}, 8 * MEGABYTE);
-
- ctx.window_size = (uv2){.w = 1280, .h = 840};
-
- ctx.out_data_dim = (uv4){.x = 1, .y = 1, .z = 1};
- ctx.export_ctx.volume_dim = (uv4){.x = 1, .y = 1, .z = 1};
-
- SetConfigFlags(FLAG_VSYNC_HINT);
- InitWindow(ctx.window_size.w, ctx.window_size.h, "OGL Beamformer");
- /* NOTE: do this after initing so that the window starts out floating in tiling wm */
- SetWindowState(FLAG_WINDOW_RESIZABLE);
- SetWindowMinSize(INFO_COLUMN_WIDTH * 2, ctx.window_size.h);
-
- /* NOTE: Gather information about the GPU */
- get_gl_params(&ctx.gl);
- dump_gl_params(&ctx.gl, temp_memory);
- validate_gl_requirements(&ctx.gl);
-
- /* TODO: build these into the binary */
- ctx.font = LoadFontEx("assets/IBMPlexSans-Bold.ttf", 28, 0, 0);
- ctx.small_font = LoadFontEx("assets/IBMPlexSans-Bold.ttf", 22, 0, 0);
-
- ctx.is.cursor_blink_t = 1;
-
- init_fragment_shader_ctx(&ctx.fsctx, ctx.out_data_dim);
-
- ctx.data_pipe = os_open_named_pipe(OS_PIPE_NAME);
- ctx.params = os_open_shared_memory_area(OS_SMEM_NAME);
- /* TODO: properly handle this? */
- ASSERT(ctx.data_pipe.file != OS_INVALID_FILE);
- ASSERT(ctx.params);
-
- ctx.params->raw.output_points = ctx.out_data_dim;
- /* NOTE: default compute shader pipeline */
- ctx.params->compute_stages[0] = CS_HADAMARD;
- ctx.params->compute_stages[1] = CS_DEMOD;
- ctx.params->compute_stages[2] = CS_UFORCES;
- ctx.params->compute_stages[3] = CS_MIN_MAX;
- ctx.params->compute_stages_count = 4;
-
- /* NOTE: make sure function pointers are valid even if we are not using the cuda lib */
- validate_cuda_lib(&ctx.cuda_lib);
-
- /* NOTE: set up OpenGL debug logging */
- glDebugMessageCallback(gl_debug_logger, NULL);
-#ifdef _DEBUG
- glEnable(GL_DEBUG_OUTPUT);
-#endif
-
- /* NOTE: allocate space for Uniform Buffer but don't send anything yet */
- glCreateBuffers(1, &ctx.csctx.shared_ubo);
- glNamedBufferStorage(ctx.csctx.shared_ubo, sizeof(BeamformerParameters), 0, GL_DYNAMIC_STORAGE_BIT);
-
- glGenQueries(ARRAY_COUNT(ctx.csctx.timer_fences) * CS_LAST, (u32 *)ctx.csctx.timer_ids);
- glGenQueries(ARRAY_COUNT(ctx.export_ctx.timer_ids), ctx.export_ctx.timer_ids);
-
- /* NOTE: do not DO_COMPUTE on first frame */
- reload_shaders(&ctx, temp_memory);
- ctx.flags &= ~DO_COMPUTE;
-
- while(!WindowShouldClose()) {
- do_debug();
- if (ctx.gl.vendor_id == GL_VENDOR_NVIDIA)
- check_and_load_cuda_lib(&ctx.cuda_lib);
-
- if (ctx.flags & RELOAD_SHADERS) {
- ctx.flags &= ~RELOAD_SHADERS;
- reload_shaders(&ctx, temp_memory);
- }
-
- do_beamformer(&ctx, temp_memory);
- }
-
- /* NOTE: make sure this will get cleaned up after external
- * programs release their references */
- os_remove_shared_memory(OS_SMEM_NAME);
-
- /* NOTE: garbage code needed for Linux */
- os_close_named_pipe(ctx.data_pipe);
-}
diff --git a/main_generic.c b/main_generic.c
@@ -0,0 +1,56 @@
+/* See LICENSE for license details. */
+#include "beamformer.h"
+
+#if defined(__unix__)
+ #include "os_unix.c"
+
+ #define OS_DEBUG_LIB_NAME "./beamformer.so"
+ #define OS_DEBUG_LIB_TEMP_NAME "./beamformer_temp.so"
+
+ #define OS_CUDA_LIB_NAME "./external/cuda_toolkit.so"
+ #define OS_CUDA_LIB_TEMP_NAME "./external/cuda_toolkit_temp.so"
+
+ #define OS_PIPE_NAME "/tmp/beamformer_data_fifo"
+ #define OS_SMEM_NAME "/ogl_beamformer_parameters"
+
+#elif defined(_WIN32)
+ #include "os_win32.c"
+
+ #define OS_DEBUG_LIB_NAME "beamformer.dll"
+ #define OS_DEBUG_LIB_TEMP_NAME "beamformer_temp.dll"
+
+ #define OS_CUDA_LIB_NAME "external\\cuda_toolkit.dll"
+ #define OS_CUDA_LIB_TEMP_NAME "external\\cuda_toolkit_temp.dll"
+
+ #define OS_PIPE_NAME "\\\\.\\pipe\\beamformer_data_fifo"
+ #define OS_SMEM_NAME "Local\\ogl_beamformer_parameters"
+#else
+ #error Unsupported Platform!
+#endif
+
+#include "static.c"
+
+int
+main(void)
+{
+ BeamformerCtx ctx = {0};
+ Arena temp_memory = os_alloc_arena((Arena){0}, 8 * MEGABYTE);
+
+ ctx.platform.alloc_arena = os_alloc_arena;
+ ctx.platform.poll_pipe = os_poll_pipe;
+ ctx.platform.read_pipe = os_read_pipe;
+ ctx.platform.write_new_file = os_write_new_file;
+
+ setup_beamformer(&ctx, temp_memory);
+
+ while(!(ctx.flags & SHOULD_EXIT)) {
+ do_program_step(&ctx, temp_memory);
+ }
+
+ /* NOTE: make sure this will get cleaned up after external
+ * programs release their references */
+ os_remove_shared_memory(OS_SMEM_NAME);
+
+ /* NOTE: garbage code needed for Linux */
+ os_close_named_pipe(ctx.data_pipe);
+}
diff --git a/os_unix.c b/os_unix.c
@@ -1,3 +1,6 @@
+/* See LICENSE for license details. */
+#include "util.h"
+
#include <dlfcn.h>
#include <fcntl.h>
#include <poll.h>
@@ -5,15 +8,6 @@
#include <sys/stat.h>
#include <unistd.h>
-#define OS_INVALID_FILE (-1)
-typedef i32 os_file;
-typedef struct {
- os_file file;
- char *name;
-} os_pipe;
-
-typedef void *os_library_handle;
-
static void
os_write_err_msg(s8 msg)
{
@@ -27,27 +21,27 @@ os_fail(void)
unreachable();
}
-static Arena
-os_alloc_arena(Arena a, size capacity)
+static PLATFORM_ALLOC_ARENA_FN(os_alloc_arena)
{
+ Arena result;
size pagesize = sysconf(_SC_PAGESIZE);
if (capacity % pagesize != 0)
capacity += pagesize - capacity % pagesize;
- size oldsize = a.end - a.beg;
+ size oldsize = old.end - old.beg;
if (oldsize > capacity)
- return a;
+ return old;
- if (a.beg)
- munmap(a.beg, oldsize);
+ if (old.beg)
+ munmap(old.beg, oldsize);
- a.beg = mmap(0, capacity, PROT_READ|PROT_WRITE, MAP_ANONYMOUS|MAP_PRIVATE, -1, 0);
- if (a.beg == MAP_FAILED) {
+ result.beg = mmap(0, capacity, PROT_READ|PROT_WRITE, MAP_ANONYMOUS|MAP_PRIVATE, -1, 0);
+ if (result.beg == MAP_FAILED) {
os_write_err_msg(s8("os_alloc_arena: couldn't allocate memory\n"));
os_fail();
}
- a.end = a.beg + capacity;
- return a;
+ result.end = result.beg + capacity;
+ return result;
}
static s8
@@ -70,8 +64,7 @@ os_read_file(Arena *a, char *fname, size fsize)
return ret;
}
-static b32
-os_write_file(char *fname, s8 raw)
+static PLATFORM_WRITE_NEW_FILE_FN(os_write_new_file)
{
i32 fd = open(fname, O_WRONLY|O_TRUNC|O_CREAT, 0600);
if (fd < 0)
@@ -96,30 +89,28 @@ os_get_file_stats(char *fname)
};
}
-static os_pipe
+static Pipe
os_open_named_pipe(char *name)
{
mkfifo(name, 0660);
- return (os_pipe){.file = open(name, O_RDONLY|O_NONBLOCK), .name = name};
+ return (Pipe){.file = open(name, O_RDONLY|O_NONBLOCK), .name = name};
}
static void
-os_close_named_pipe(os_pipe p)
+os_close_named_pipe(Pipe p)
{
close(p.file);
unlink(p.name);
}
-static b32
-os_poll_pipe(os_pipe p)
+static PLATFORM_POLL_PIPE_FN(os_poll_pipe)
{
struct pollfd pfd = {.fd = p.file, .events = POLLIN};
poll(&pfd, 1, 0);
return !!(pfd.revents & POLLIN);
}
-static size
-os_read_pipe_data(os_pipe p, void *buf, size len)
+static PLATFORM_READ_PIPE_FN(os_read_pipe)
{
size r = 0, total_read = 0;
do {
@@ -188,14 +179,14 @@ ret:
return result;
}
-static os_library_handle
+static void *
os_load_library(char *name, char *temp_name)
{
if (temp_name) {
if (os_copy_file(name, temp_name))
name = temp_name;
}
- os_library_handle res = dlopen(name, RTLD_NOW|RTLD_LOCAL);
+ void *res = dlopen(name, RTLD_NOW|RTLD_LOCAL);
if (!res)
TraceLog(LOG_WARNING, "os_load_library(%s): %s\n", name, dlerror());
@@ -206,7 +197,7 @@ os_load_library(char *name, char *temp_name)
}
static void *
-os_lookup_dynamic_symbol(os_library_handle h, char *name)
+os_lookup_dynamic_symbol(void *h, char *name)
{
if (!h)
return 0;
@@ -217,7 +208,7 @@ os_lookup_dynamic_symbol(os_library_handle h, char *name)
}
static void
-os_unload_library(os_library_handle h)
+os_unload_library(void *h)
{
/* NOTE: glibc is buggy gnuware so we need to check this */
if (h)
diff --git a/os_win32.c b/os_win32.c
@@ -16,8 +16,6 @@
#define FILE_MAP_ALL_ACCESS 0x000F001F
-#define INVALID_HANDLE_VALUE (void *)-1
-
#define CREATE_ALWAYS 2
#define OPEN_EXISTING 3
@@ -49,37 +47,28 @@ typedef struct {
} w32_file_info;
#define W32(r) __declspec(dllimport) r __stdcall
-W32(b32) CloseHandle(void *);
+W32(b32) CloseHandle(iptr);
W32(b32) CopyFileA(c8 *, c8 *, b32);
-W32(void *) CreateFileA(c8 *, u32, u32, void *, u32, u32, void *);
-W32(void *) CreateFileMappingA(void *, void *, u32, u32, u32, c8 *);
-W32(void *) CreateNamedPipeA(c8 *, u32, u32, u32, u32, u32, u32, void *);
+W32(iptr) CreateFileA(c8 *, u32, u32, void *, u32, u32, void *);
+W32(iptr) CreateFileMappingA(iptr, void *, u32, u32, u32, c8 *);
+W32(iptr) CreateNamedPipeA(c8 *, u32, u32, u32, u32, u32, u32, void *);
W32(b32) DeleteFileA(c8 *);
W32(void) ExitProcess(i32);
W32(b32) FreeLibrary(void *);
-W32(b32) GetFileInformationByHandle(void *, void *);
+W32(b32) GetFileInformationByHandle(iptr, w32_file_info *);
W32(i32) GetLastError(void);
W32(void *) GetProcAddress(void *, c8 *);
-W32(void *) GetStdHandle(i32);
+W32(iptr) GetStdHandle(i32);
W32(void) GetSystemInfo(void *);
W32(void *) LoadLibraryA(c8 *);
-W32(void *) MapViewOfFile(void *, u32, u32, u32, u64);
-W32(b32) PeekNamedPipe(void *, u8 *, i32, i32 *, i32 *, i32 *);
-W32(b32) ReadFile(void *, u8 *, i32, i32 *, void *);
-W32(b32) WriteFile(void *, u8 *, i32, i32 *, void *);
+W32(void *) MapViewOfFile(iptr, u32, u32, u32, u64);
+W32(b32) PeekNamedPipe(iptr, u8 *, i32, i32 *, i32 *, i32 *);
+W32(b32) ReadFile(iptr, u8 *, i32, i32 *, void *);
+W32(b32) WriteFile(iptr, u8 *, i32, i32 *, void *);
W32(void *) VirtualAlloc(u8 *, size, u32, u32);
W32(b32) VirtualFree(u8 *, size, u32);
-#define OS_INVALID_FILE (INVALID_HANDLE_VALUE)
-typedef void *os_file;
-typedef struct {
- os_file file;
- char *name;
-} os_pipe;
-
-typedef void *os_library_handle;
-
-static void *win32_stderr_handle;
+static iptr win32_stderr_handle;
static void __attribute__((noreturn))
os_fail(void)
@@ -97,29 +86,29 @@ os_write_err_msg(s8 msg)
WriteFile(win32_stderr_handle, msg.data, msg.len, &wlen, 0);
}
-static Arena
-os_alloc_arena(Arena a, size capacity)
+static PLATFORM_ALLOC_ARENA_FN(os_alloc_arena)
{
+ Arena result;
w32_sys_info Info;
GetSystemInfo(&Info);
if (capacity % Info.dwPageSize != 0)
capacity += (Info.dwPageSize - capacity % Info.dwPageSize);
- size oldsize = a.end - a.beg;
+ size oldsize = old.end - old.beg;
if (oldsize > capacity)
- return a;
+ return old;
- if (a.beg)
- VirtualFree(a.beg, oldsize, MEM_RELEASE);
+ if (old.beg)
+ VirtualFree(old.beg, oldsize, MEM_RELEASE);
- a.beg = VirtualAlloc(0, capacity, MEM_RESERVE|MEM_COMMIT, PAGE_READWRITE);
- if (a.beg == NULL) {
+ result.beg = VirtualAlloc(0, capacity, MEM_RESERVE|MEM_COMMIT, PAGE_READWRITE);
+ if (result.beg == NULL) {
os_write_err_msg(s8("os_alloc_arena: couldn't allocate memory\n"));
os_fail();
}
- a.end = a.beg + capacity;
- return a;
+ result.end = result.beg + capacity;
+ return result;
}
static s8
@@ -134,8 +123,8 @@ os_read_file(Arena *a, char *fname, size fsize)
return (s8){.len = -1};
}
- void *h = CreateFileA(fname, GENERIC_READ, 0, 0, OPEN_EXISTING, 0, 0);
- if (h == INVALID_HANDLE_VALUE)
+ iptr h = CreateFileA(fname, GENERIC_READ, 0, 0, OPEN_EXISTING, 0, 0);
+ if (h == INVALID_FILE)
return (s8){.len = -1};
s8 ret = s8alloc(a, fsize);
@@ -149,16 +138,15 @@ os_read_file(Arena *a, char *fname, size fsize)
return ret;
}
-static b32
-os_write_file(char *fname, s8 raw)
+static PLATFORM_WRITE_NEW_FILE_FN(os_write_new_file)
{
if (raw.len > (size)U32_MAX) {
os_write_err_msg(s8("os_write_file: writing files > 4GB is not yet support on win32\n"));
return 0;
}
- void *h = CreateFileA(fname, GENERIC_WRITE, 0, 0, CREATE_ALWAYS, 0, 0);
- if (h == INVALID_HANDLE_VALUE)
+ iptr h = CreateFileA(fname, GENERIC_WRITE, 0, 0, CREATE_ALWAYS, 0, 0);
+ if (h == INVALID_FILE)
return 0;
i32 wlen;
@@ -170,10 +158,9 @@ os_write_file(char *fname, s8 raw)
static FileStats
os_get_file_stats(char *fname)
{
- void *h = CreateFileA(fname, 0, 0, 0, OPEN_EXISTING, 0, 0);
- if (h == INVALID_HANDLE_VALUE) {
+ iptr h = CreateFileA(fname, 0, 0, 0, OPEN_EXISTING, 0, 0);
+ if (h == INVALID_FILE)
return ERROR_FILE_STATS;
- }
w32_file_info fileinfo;
if (!GetFileInformationByHandle(h, &fileinfo)) {
@@ -189,29 +176,27 @@ os_get_file_stats(char *fname)
return (FileStats){.filesize = filesize, .timestamp = fileinfo.ftLastWriteTime};
}
-/* NOTE: win32 doesn't pollute the filesystem so no need to waste the user's time */
-static void
-os_close_named_pipe(os_pipe p)
+static Pipe
+os_open_named_pipe(char *name)
{
+ iptr h = CreateNamedPipeA(name, PIPE_ACCESS_INBOUND, PIPE_TYPE_BYTE, 1,
+ 0, 1 * MEGABYTE, 0, 0);
+ return (Pipe){.file = h, .name = name};
}
-static os_pipe
-os_open_named_pipe(char *name)
+/* NOTE: win32 doesn't pollute the filesystem so no need to waste the user's time */
+static void
+os_close_named_pipe(Pipe p)
{
- void *h = CreateNamedPipeA(name, PIPE_ACCESS_INBOUND, PIPE_TYPE_BYTE, 1,
- 0, 1 * MEGABYTE, 0, 0);
- return (os_pipe){.file = h, .name = name};
}
-static b32
-os_poll_pipe(os_pipe p)
+static PLATFORM_POLL_PIPE_FN(os_poll_pipe)
{
i32 bytes_available = 0;
return PeekNamedPipe(p.file, 0, 1 * MEGABYTE, 0, &bytes_available, 0) && bytes_available;
}
-static size
-os_read_pipe_data(os_pipe p, void *buf, size len)
+static PLATFORM_READ_PIPE_FN(os_read_pipe)
{
i32 total_read = 0;
ReadFile(p.file, buf, len, &total_read, 0);
@@ -221,9 +206,9 @@ os_read_pipe_data(os_pipe p, void *buf, size len)
static BeamformerParametersFull *
os_open_shared_memory_area(char *name)
{
- void *h = CreateFileMappingA(INVALID_HANDLE_VALUE, 0, PAGE_READWRITE, 0,
+ iptr h = CreateFileMappingA(-1, 0, PAGE_READWRITE, 0,
sizeof(BeamformerParametersFull), name);
- if (h == INVALID_HANDLE_VALUE)
+ if (h == INVALID_FILE)
return NULL;
BeamformerParametersFull *new;
@@ -238,7 +223,7 @@ os_remove_shared_memory(char *name)
{
}
-static os_library_handle
+static void *
os_load_library(char *name, char *temp_name)
{
if (temp_name) {
@@ -246,7 +231,7 @@ os_load_library(char *name, char *temp_name)
name = temp_name;
}
- os_library_handle res = LoadLibraryA(name);
+ void *res = LoadLibraryA(name);
if (!res)
TraceLog(LOG_WARNING, "os_load_library(%s): %d\n", name, GetLastError());
@@ -257,7 +242,7 @@ os_load_library(char *name, char *temp_name)
}
static void *
-os_lookup_dynamic_symbol(os_library_handle h, char *name)
+os_lookup_dynamic_symbol(void *h, char *name)
{
if (!h)
return 0;
@@ -268,7 +253,7 @@ os_lookup_dynamic_symbol(os_library_handle h, char *name)
}
static void
-os_unload_library(os_library_handle h)
+os_unload_library(void *h)
{
FreeLibrary(h);
}
diff --git a/static.c b/static.c
@@ -0,0 +1,317 @@
+/* See LICENSE for license details. */
+#include "beamformer.h"
+
+static s8 compute_shader_paths[CS_LAST] = {
+ [CS_HADAMARD] = s8("shaders/hadamard.glsl"),
+ [CS_HERCULES] = s8("shaders/hercules.glsl"),
+ [CS_DEMOD] = s8("shaders/demod.glsl"),
+ [CS_MIN_MAX] = s8("shaders/min_max.glsl"),
+ [CS_SUM] = s8("shaders/sum.glsl"),
+ [CS_UFORCES] = s8("shaders/uforces.glsl"),
+};
+
+#ifndef _DEBUG
+
+#include "beamformer.c"
+#define do_debug(...)
+
+#else
+static void *debug_lib;
+
+typedef void do_beamformer_fn(BeamformerCtx *, Arena);
+static do_beamformer_fn *do_beamformer;
+
+static void
+do_debug(void)
+{
+ static f32 updated_time;
+ FileStats test_stats = os_get_file_stats(OS_DEBUG_LIB_NAME);
+ if (test_stats.filesize > 32 && test_stats.timestamp > updated_time) {
+ os_unload_library(debug_lib);
+ debug_lib = os_load_library(OS_DEBUG_LIB_NAME, OS_DEBUG_LIB_TEMP_NAME);
+ do_beamformer = os_lookup_dynamic_symbol(debug_lib, "do_beamformer");
+ updated_time = test_stats.timestamp;
+ }
+}
+
+#endif /* _DEBUG */
+
+static void
+gl_debug_logger(u32 src, u32 type, u32 id, u32 lvl, i32 len, const char *msg, const void *userctx)
+{
+ (void)src; (void)type; (void)id; (void)userctx;
+
+ u8 buf[128];
+ Stream s = {.data = buf, .cap = ARRAY_COUNT(buf)};
+ stream_append_s8(&s, s8("[GL DEBUG "));
+ switch (lvl) {
+ case GL_DEBUG_SEVERITY_HIGH: stream_append_s8(&s, s8("HIGH]: ")); break;
+ case GL_DEBUG_SEVERITY_MEDIUM: stream_append_s8(&s, s8("MEDIUM]: ")); break;
+ case GL_DEBUG_SEVERITY_LOW: stream_append_s8(&s, s8("LOW]: ")); break;
+ case GL_DEBUG_SEVERITY_NOTIFICATION: stream_append_s8(&s, s8("NOTIFICATION]: ")); break;
+ default: stream_append_s8(&s, s8("INVALID]: ")); break;
+ }
+ os_write_err_msg(stream_to_s8(s));
+ os_write_err_msg((s8){.len = len, .data = (u8 *)msg});
+ os_write_err_msg(s8("\n"));
+}
+
+static void
+get_gl_params(GLParams *gl)
+{
+ char *vendor = (char *)glGetString(GL_VENDOR);
+ if (!vendor) {
+ os_write_err_msg(s8("Failed to determine GL Vendor\n"));
+ os_fail();
+ }
+ switch (vendor[0]) {
+ case 'A': gl->vendor_id = GL_VENDOR_AMD; break;
+ case 'I': gl->vendor_id = GL_VENDOR_INTEL; break;
+ case 'N': gl->vendor_id = GL_VENDOR_NVIDIA; break;
+ default: {
+ os_write_err_msg(s8("Unknown GL Vendor: "));
+ os_write_err_msg(cstr_to_s8(vendor));
+ os_write_err_msg(s8("\n"));
+ os_fail();
+ } break;
+ }
+
+ glGetIntegerv(GL_MAJOR_VERSION, &gl->version_major);
+ glGetIntegerv(GL_MINOR_VERSION, &gl->version_minor);
+ glGetIntegerv(GL_MAX_TEXTURE_SIZE, &gl->max_2d_texture_dim);
+ glGetIntegerv(GL_MAX_3D_TEXTURE_SIZE, &gl->max_3d_texture_dim);
+ glGetIntegerv(GL_MAX_SHADER_STORAGE_BLOCK_SIZE, &gl->max_ssbo_size);
+ glGetIntegerv(GL_MAX_UNIFORM_BLOCK_SIZE, &gl->max_ubo_size);
+}
+
+static void
+validate_gl_requirements(GLParams *gl)
+{
+ ASSERT(gl->max_ubo_size >= sizeof(BeamformerParameters));
+ if (gl->version_major < 4 || (gl->version_major == 4 && gl->version_minor < 5)) {
+ os_write_err_msg(s8("Only OpenGL Versions 4.5 or newer are supported!\n"));
+ os_fail();
+ }
+}
+
+static void
+dump_gl_params(GLParams *gl, Arena a)
+{
+ (void)gl; (void)a;
+#ifdef _DEBUG
+ Stream s = stream_alloc(&a, 1 * MEGABYTE);
+ stream_append_s8(&s, s8("---- GL Parameters ----\n"));
+ switch (gl->vendor_id) {
+ case GL_VENDOR_AMD: stream_append_s8(&s, s8("Vendor: AMD\n")); break;
+ case GL_VENDOR_INTEL: stream_append_s8(&s, s8("Vendor: Intel\n")); break;
+ case GL_VENDOR_NVIDIA: stream_append_s8(&s, s8("Vendor: nVidia\n")); break;
+ }
+ stream_append_s8(&s, s8("Version: "));
+ stream_append_i64(&s, gl->version_major);
+ stream_append_s8(&s, s8("."));
+ stream_append_i64(&s, gl->version_minor);
+ stream_append_s8(&s, s8("\nMax 1D/2D Texture Dimension: "));
+ stream_append_i64(&s, gl->max_2d_texture_dim);
+ stream_append_s8(&s, s8("\nMax 3D Texture Dimension: "));
+ stream_append_i64(&s, gl->max_3d_texture_dim);
+ stream_append_s8(&s, s8("\nMax SSBO Size: "));
+ stream_append_i64(&s, gl->max_ssbo_size);
+ stream_append_s8(&s, s8("\nMax UBO Size: "));
+ stream_append_i64(&s, gl->max_ubo_size);
+ stream_append_s8(&s, s8("\n-----------------------\n"));
+ if (!s.errors)
+ os_write_err_msg(stream_to_s8(s));
+#endif
+}
+
+static u32
+compile_shader(Arena a, u32 type, s8 shader)
+{
+ u32 sid = glCreateShader(type);
+ glShaderSource(sid, 1, (const char **)&shader.data, (int *)&shader.len);
+ glCompileShader(sid);
+
+ i32 res = 0;
+ glGetShaderiv(sid, GL_COMPILE_STATUS, &res);
+
+ char *stype;
+ switch (type) {
+ case GL_COMPUTE_SHADER: stype = "Compute"; break;
+ case GL_FRAGMENT_SHADER: stype = "Fragment"; break;
+ }
+
+ if (res == GL_FALSE) {
+ TraceLog(LOG_WARNING, "SHADER: [ID %u] %s shader failed to compile", sid, stype);
+ i32 len = 0;
+ glGetShaderiv(sid, GL_INFO_LOG_LENGTH, &len);
+ s8 err = s8alloc(&a, len);
+ glGetShaderInfoLog(sid, len, (int *)&err.len, (char *)err.data);
+ TraceLog(LOG_WARNING, "SHADER: [ID %u] Compile error: %s", sid, (char *)err.data);
+ glDeleteShader(sid);
+ } else {
+ TraceLog(LOG_INFO, "SHADER: [ID %u] %s shader compiled successfully", sid, stype);
+ }
+
+ return sid;
+}
+
+static void
+init_fragment_shader_ctx(FragmentShaderCtx *ctx, uv4 out_data_dim)
+{
+ ctx->output = LoadRenderTexture(out_data_dim.x, out_data_dim.y);
+ ctx->db = -50.0f;
+}
+
+static void
+reload_shaders(BeamformerCtx *ctx, Arena a)
+{
+ ComputeShaderCtx *csctx = &ctx->csctx;
+ for (u32 i = 0; i < ARRAY_COUNT(csctx->programs); i++) {
+ if (!compute_shader_paths[i].len)
+ continue;
+
+ Arena tmp = a;
+ FileStats fs = os_get_file_stats((char *)compute_shader_paths[i].data);
+ s8 shader_text = os_read_file(&tmp, (char *)compute_shader_paths[i].data, fs.filesize);
+ if (shader_text.len == -1) {
+ os_write_err_msg(s8("failed to read shader: "));
+ os_write_err_msg(compute_shader_paths[i]);
+ os_write_err_msg(s8("\n"));
+ os_fail();
+ }
+ u32 shader_id = compile_shader(tmp, GL_COMPUTE_SHADER, shader_text);
+
+ if (shader_id) {
+ glDeleteProgram(csctx->programs[i]);
+ csctx->programs[i] = rlLoadComputeShaderProgram(shader_id);
+ ctx->flags |= DO_COMPUTE;
+ }
+
+ glDeleteShader(shader_id);
+ }
+
+ csctx->volume_export_pass_id = glGetUniformLocation(csctx->programs[CS_HERCULES],
+ "u_volume_export_pass");
+ csctx->volume_export_dim_offset_id = glGetUniformLocation(csctx->programs[CS_HERCULES],
+ "u_volume_export_dim_offset");
+ csctx->xdc_transform_id = glGetUniformLocation(csctx->programs[CS_UFORCES],
+ "u_xdc_transform");
+ csctx->xdc_index_id = glGetUniformLocation(csctx->programs[CS_UFORCES],
+ "u_xdc_index");
+
+ csctx->mips_level_id = glGetUniformLocation(csctx->programs[CS_MIN_MAX], "u_mip_map");
+
+ csctx->sum_prescale_id = glGetUniformLocation(csctx->programs[CS_SUM], "u_prescale");
+
+ Shader updated_fs = LoadShader(NULL, "shaders/render.glsl");
+ if (updated_fs.id != rlGetShaderIdDefault()) {
+ UnloadShader(ctx->fsctx.shader);
+ ctx->fsctx.shader = updated_fs;
+ ctx->fsctx.db_cutoff_id = GetShaderLocation(updated_fs, "u_db_cutoff");
+ }
+}
+
+static void
+validate_cuda_lib(CudaLib *cl)
+{
+ if (!cl->init_cuda_configuration) cl->init_cuda_configuration = init_cuda_configuration_stub;
+ if (!cl->register_cuda_buffers) cl->register_cuda_buffers = register_cuda_buffers_stub;
+ if (!cl->cuda_decode) cl->cuda_decode = cuda_decode_stub;
+ if (!cl->cuda_hilbert) cl->cuda_hilbert = cuda_hilbert_stub;
+}
+
+static void
+check_and_load_cuda_lib(CudaLib *cl)
+{
+ FileStats current = os_get_file_stats(OS_CUDA_LIB_NAME);
+ if (cl->timestamp == current.timestamp || current.filesize < 32)
+ return;
+
+ TraceLog(LOG_INFO, "Loading CUDA lib: %s", OS_CUDA_LIB_NAME);
+
+ cl->timestamp = current.timestamp;
+ os_unload_library(cl->lib);
+ cl->lib = os_load_library(OS_CUDA_LIB_NAME, OS_CUDA_LIB_TEMP_NAME);
+
+ cl->init_cuda_configuration = os_lookup_dynamic_symbol(cl->lib, "init_cuda_configuration");
+ cl->register_cuda_buffers = os_lookup_dynamic_symbol(cl->lib, "register_cuda_buffers");
+ cl->cuda_decode = os_lookup_dynamic_symbol(cl->lib, "cuda_decode");
+ cl->cuda_hilbert = os_lookup_dynamic_symbol(cl->lib, "cuda_hilbert");
+
+ validate_cuda_lib(cl);
+}
+
+static void
+setup_beamformer(BeamformerCtx *ctx, Arena temp_memory)
+{
+ ctx->window_size = (uv2){.w = 1280, .h = 840};
+
+ ctx->out_data_dim = (uv4){.x = 1, .y = 1, .z = 1};
+ ctx->export_ctx.volume_dim = (uv4){.x = 1, .y = 1, .z = 1};
+
+ SetConfigFlags(FLAG_VSYNC_HINT);
+ InitWindow(ctx->window_size.w, ctx->window_size.h, "OGL Beamformer");
+ /* NOTE: do this after initing so that the window starts out floating in tiling wm */
+ SetWindowState(FLAG_WINDOW_RESIZABLE);
+ SetWindowMinSize(INFO_COLUMN_WIDTH * 2, ctx->window_size.h);
+
+ /* NOTE: Gather information about the GPU */
+ get_gl_params(&ctx->gl);
+ dump_gl_params(&ctx->gl, temp_memory);
+ validate_gl_requirements(&ctx->gl);
+
+ /* TODO: build these into the binary */
+ ctx->font = LoadFontEx("assets/IBMPlexSans-Bold.ttf", 28, 0, 0);
+ ctx->small_font = LoadFontEx("assets/IBMPlexSans-Bold.ttf", 22, 0, 0);
+
+ init_fragment_shader_ctx(&ctx->fsctx, ctx->out_data_dim);
+
+ ctx->data_pipe = os_open_named_pipe(OS_PIPE_NAME);
+ ctx->params = os_open_shared_memory_area(OS_SMEM_NAME);
+ /* TODO: properly handle this? */
+ ASSERT(ctx->data_pipe.file != INVALID_FILE);
+ ASSERT(ctx->params);
+
+ ctx->params->raw.output_points = ctx->out_data_dim;
+ /* NOTE: default compute shader pipeline */
+ ctx->params->compute_stages[0] = CS_HADAMARD;
+ ctx->params->compute_stages[1] = CS_DEMOD;
+ ctx->params->compute_stages[2] = CS_UFORCES;
+ ctx->params->compute_stages[3] = CS_MIN_MAX;
+ ctx->params->compute_stages_count = 4;
+
+ /* NOTE: make sure function pointers are valid even if we are not using the cuda lib */
+ validate_cuda_lib(&ctx->cuda_lib);
+
+ /* NOTE: set up OpenGL debug logging */
+ glDebugMessageCallback(gl_debug_logger, NULL);
+#ifdef _DEBUG
+ glEnable(GL_DEBUG_OUTPUT);
+#endif
+
+ /* NOTE: allocate space for Uniform Buffer but don't send anything yet */
+ glCreateBuffers(1, &ctx->csctx.shared_ubo);
+ glNamedBufferStorage(ctx->csctx.shared_ubo, sizeof(BeamformerParameters), 0, GL_DYNAMIC_STORAGE_BIT);
+
+ glGenQueries(ARRAY_COUNT(ctx->csctx.timer_fences) * CS_LAST, (u32 *)ctx->csctx.timer_ids);
+ glGenQueries(ARRAY_COUNT(ctx->export_ctx.timer_ids), ctx->export_ctx.timer_ids);
+
+ /* NOTE: do not DO_COMPUTE on first frame */
+ reload_shaders(ctx, temp_memory);
+ ctx->flags &= ~DO_COMPUTE;
+}
+
+static void
+do_program_step(BeamformerCtx *ctx, Arena temp_memory)
+{
+ do_debug();
+ if (ctx->gl.vendor_id == GL_VENDOR_NVIDIA)
+ check_and_load_cuda_lib(&ctx->cuda_lib);
+
+ if (ctx->flags & RELOAD_SHADERS) {
+ ctx->flags &= ~RELOAD_SHADERS;
+ reload_shaders(ctx, temp_memory);
+ }
+
+ do_beamformer(ctx, temp_memory);
+}
diff --git a/util.h b/util.h
@@ -7,12 +7,6 @@
#include <immintrin.h>
-#include <glad.h>
-
-#define GRAPHICS_API_OPENGL_43
-#include <raylib.h>
-#include <rlgl.h>
-
#ifndef asm
#define asm __asm__
#endif
@@ -26,8 +20,14 @@
#endif
#ifdef _DEBUG
+ #ifdef _WIN32
+ #define DEBUG_EXPORT __declspec(dllexport)
+ #else
+ #define DEBUG_EXPORT
+ #endif
#define ASSERT(c) do { if (!(c)) asm("int3; nop"); } while (0);
#else
+ #define DEBUG_EXPORT static
#define ASSERT(c)
#endif
@@ -59,12 +59,19 @@ typedef uint32_t b32;
typedef float f32;
typedef double f64;
typedef ptrdiff_t size;
+typedef ptrdiff_t iptr;
typedef struct { u8 *beg, *end; } Arena;
typedef struct { size len; u8 *data; } s8;
#define s8(s) (s8){.len = ARRAY_COUNT(s) - 1, .data = (u8 *)s}
+/* NOTE: raylib stubs */
+#ifndef RAYLIB_H
+typedef struct { f32 x, y; } Vector2;
+typedef struct { f32 x, y, w, h; } Rectangle;
+#endif
+
typedef union {
struct { i32 x, y; };
struct { i32 w, h; };
@@ -116,6 +123,12 @@ typedef union {
} Rect;
typedef struct {
+ iptr file;
+ char *name;
+} Pipe;
+#define INVALID_FILE (-1)
+
+typedef struct {
size filesize;
u64 timestamp;
} FileStats;
@@ -128,14 +141,6 @@ typedef struct {
b32 errors;
} Stream;
-#include "beamformer_parameters.h"
-typedef struct {
- BeamformerParameters raw;
- enum compute_shaders compute_stages[16];
- u32 compute_stages_count;
- b32 upload;
-} BeamformerParametersFull;
-
#include "util.c"
#endif /* _UTIL_H_ */