ogl_beamforming

Ultrasound Beamforming Implemented with OpenGL
git clone anongit@rnpnr.xyz:ogl_beamforming.git
Log | Files | Refs | Feed | Submodules | LICENSE

Commit: d6e1c420bb9429c530ba9d57f53c54e3374696f1
Parent: 0f0c7a6bff0d66d753f43639fb023cc1e02d313d
Author: Randy Palamar
Date:   Wed, 26 Jun 2024 12:57:22 -0600

compute min max values and plot on a grayscale

The output rf data is now stored in an image to make calculating
the min and max on the gpu easy. A binary search can be performed
in 3D with the results being stored to higher mip map levels.

Diffstat:
Mbeamformer.c | 63++++++++++++++++++++++++++++++++++++++++++++++++++++++---------
Mmain.c | 125+++++++++++++++++++++++++++++++++++++++----------------------------------------
Mshaders/hadamard.glsl | 3++-
Ashaders/min_max.glsl | 27+++++++++++++++++++++++++++
Mshaders/render.glsl | 26++++++++++++++++----------
Mshaders/uforces.glsl | 28+++++++++++++---------------
Mutil.h | 17+++++++++++++----
7 files changed, 187 insertions(+), 102 deletions(-)

diff --git a/beamformer.c b/beamformer.c @@ -13,7 +13,9 @@ do_compute_shader(BeamformerCtx *ctx, u32 rf_ssbo_idx, enum compute_shaders shad glUseProgram(csctx->programs[shader]); glUniform3uiv(csctx->rf_data_dim_id, 1, csctx->rf_data_dim.E); - glUniform3uiv(csctx->out_data_dim_id, 1, ctx->out_data_dim.E); + glBindImageTexture(ctx->out_texture_unit, ctx->out_texture, 0, GL_FALSE, 0, + GL_WRITE_ONLY, GL_RG32F); + glUniform1i(csctx->out_data_tex_id, ctx->out_texture_unit); u32 decoded_ssbo_idx = 2; switch (shader) { @@ -23,9 +25,27 @@ do_compute_shader(BeamformerCtx *ctx, u32 rf_ssbo_idx, enum compute_shaders shad glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, csctx->hadamard_ssbo); glDispatchCompute(csctx->rf_data_dim.x / 32, csctx->rf_data_dim.y / 32, csctx->rf_data_dim.z); break; + case CS_MIN_MAX: + for (u32 i = 1; i < ctx->out_texture_mips; i++) { + u32 otu = ctx->out_texture_unit; + glBindImageTexture(otu + 1, ctx->out_texture, i - 1, + GL_FALSE, 0, GL_READ_ONLY, GL_RG32F); + glBindImageTexture(otu + 2, ctx->out_texture, i, + GL_FALSE, 0, GL_WRITE_ONLY, GL_RG32F); + glUniform1i(csctx->out_data_tex_id, otu + 1); + glUniform1i(csctx->mip_view_tex_id, otu + 2); + glUniform1i(csctx->mips_level_id, i); + + #define ORONE(x) ((x)? (x) : 1) + u32 width = ctx->out_data_dim.w >> i; + u32 height = ctx->out_data_dim.h >> i; + u32 depth = ctx->out_data_dim.d >> i; + glDispatchCompute(ORONE(width), ORONE(height), ORONE(depth)); + glMemoryBarrier(GL_SHADER_IMAGE_ACCESS_BARRIER_BIT); + } + break; case CS_UFORCES: glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, csctx->rf_data_ssbos[decoded_ssbo_idx]); - glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2, ctx->out_data_ssbo); glDispatchCompute(ctx->out_data_dim.x / 32, ctx->out_data_dim.y / 32, ctx->out_data_dim.z); break; default: ASSERT(0); @@ -39,17 +59,30 @@ draw_debug_overlay(BeamformerCtx *ctx, Arena arena) u32 fontsize = 32; u32 fontspace = 1; + s8 db_txt = s8alloc(&arena, 64); s8 compute_txt = s8alloc(&arena, 64); + snprintf((char *)db_txt.data, db_txt.len, "Dynamic Range: %0.01f [db]", ctx->fsctx.db); snprintf((char *)compute_txt.data, compute_txt.len, "Compute: %d", !!(ctx->flags & DO_COMPUTE)); + v2 db_fs = {.rl = MeasureTextEx(ctx->font, (char *)db_txt.data, fontsize, fontspace)}; v2 compute_fs = {.rl = MeasureTextEx(ctx->font, (char *)compute_txt.data, fontsize, fontspace)}; v2 scale = {.x = 90, .y = 20 }; + /* NOTE: Dynamic Range */ + { + v2 dpos = {.x = 20, .y = ctx->window_size.y - db_fs.y - compute_fs.y - 20}; + v2 dposa = {.x = dpos.x + db_fs.x / scale.x, .y = dpos.y + db_fs.y / scale.y }; + DrawTextEx(ctx->font, (char *)db_txt.data, dposa.rl, fontsize, fontspace, Fade(BLACK, 0.8)); + DrawTextEx(ctx->font, (char *)db_txt.data, dpos.rl, fontsize, fontspace, RED); + } - v2 dpos = {.x = 20, .y = ctx->window_size.y - compute_fs.y - 20}; - v2 dposa = {.x = dpos.x + compute_fs.x / scale.x, .y = dpos.y + compute_fs.y / scale.y }; - DrawTextEx(ctx->font, (char *)compute_txt.data, dposa.rl, fontsize, fontspace, Fade(BLACK, 0.8)); - DrawTextEx(ctx->font, (char *)compute_txt.data, dpos.rl, fontsize, fontspace, RED); + /* NOTE: Compute Status */ + { + v2 dpos = {.x = 20, .y = ctx->window_size.y - compute_fs.y - 20}; + v2 dposa = {.x = dpos.x + compute_fs.x / scale.x, .y = dpos.y + compute_fs.y / scale.y }; + DrawTextEx(ctx->font, (char *)compute_txt.data, dposa.rl, fontsize, fontspace, Fade(BLACK, 0.8)); + DrawTextEx(ctx->font, (char *)compute_txt.data, dpos.rl, fontsize, fontspace, RED); + } } @@ -100,14 +133,26 @@ do_beamformer(BeamformerCtx *ctx, Arena arena, s8 rf_data) if (ctx->flags & DO_COMPUTE) { do_compute_shader(ctx, rf_ssbo_idx, CS_HADAMARD); do_compute_shader(ctx, rf_ssbo_idx, CS_UFORCES); + do_compute_shader(ctx, rf_ssbo_idx, CS_MIN_MAX); } + /* NOTE: check mouse wheel for adjusting dynamic range of image */ + ctx->fsctx.db += GetMouseWheelMove(); + CLAMP(ctx->fsctx.db, -120, 0); + + /* NOTE: draw output image texture using render fragment shader */ BeginTextureMode(ctx->fsctx.output); ClearBackground(ctx->bg); BeginShaderMode(ctx->fsctx.shader); glUseProgram(ctx->fsctx.shader.id); - glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, ctx->out_data_ssbo); - glUniform3uiv(ctx->fsctx.out_data_dim_id, 1, ctx->out_data_dim.E); + u32 otu = ctx->out_texture_unit; + glBindImageTexture(otu + 1, ctx->out_texture, ctx->out_texture_mips - 1, + GL_FALSE, 0, GL_READ_ONLY, GL_RG32F); + glBindImageTexture(otu, ctx->out_texture, 0, + GL_FALSE, 0, GL_READ_ONLY, GL_RG32F); + glUniform1i(ctx->fsctx.out_data_tex_id, otu); + glUniform1i(ctx->fsctx.mip_view_tex_id, otu + 1); + glUniform1f(ctx->fsctx.db_cutoff_id, ctx->fsctx.db); DrawTexture(ctx->fsctx.output.texture, 0, 0, WHITE); EndShaderMode(); EndTextureMode(); @@ -119,7 +164,7 @@ do_beamformer(BeamformerCtx *ctx, Arena arena, s8 rf_data) Rectangle rect = { 0.0f, 0.0f, (f32)rtext->width, -(f32)rtext->height }; DrawTextureRec(*rtext, rect, (Vector2){0}, WHITE); - DrawTextEx(ctx->font, txt[txt_idx], pos.rl, fontsize, fontspace, BLACK); + DrawTextEx(ctx->font, txt[txt_idx], pos.rl, fontsize, fontspace, RED); draw_debug_overlay(ctx, arena); EndDrawing(); diff --git a/main.c b/main.c @@ -1,12 +1,13 @@ /* See LICENSE for license details. */ +#include <immintrin.h> #include <raylib.h> #include <rlgl.h> #include "util.h" static char *compute_shader_paths[CS_LAST] = { - //[CS_MIN_MAX] = "shaders/min_max.glsl", + [CS_MIN_MAX] = "shaders/min_max.glsl", [CS_HADAMARD] = "shaders/hadamard.glsl", [CS_UFORCES] = "shaders/uforces.glsl", }; @@ -72,8 +73,7 @@ do_debug(void) static void fill_hadamard(i32 *m, u32 dim) { - /* bit hack to check if dim is power of 2 */ - ASSERT(!(dim & (dim - 1)) && dim); + ASSERT(dim && ISPOWEROF2(dim)); #define IND(i, j) ((i) * dim + (j)) m[0] = 1; @@ -90,37 +90,52 @@ fill_hadamard(i32 *m, u32 dim) #undef IND } -#if 0 -static void -update_output_image_dimensions(BeamformerCtx *ctx, uv2 new_size) +static u32 +compile_shader(Arena a, u32 type, s8 shader) { - UnloadTexture(ctx->fsctx.output); - rlUnloadShaderBuffer(ctx->csctx.out_img_ssbo); + u32 sid = glCreateShader(type); + glShaderSource(sid, 1, (const char **)&shader.data, (int *)&shader.len); + glCompileShader(sid); - size out_img_size = new_size.w * new_size.h * sizeof(f32); - ctx->csctx.out_img_ssbo = rlLoadShaderBuffer(out_img_size, NULL, GL_DYNAMIC_COPY); + i32 res = 0; + glGetShaderiv(sid, GL_COMPILE_STATUS, &res); - Texture2D t = ctx->fsctx.output; - t.width = new_size.w; - t.height = new_size.h; - t.id = rlLoadTexture(0, t.width, t.height, t.format, t.mipmaps); - ctx->fsctx.output = t; + 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; } -#endif static void init_compute_shader_ctx(ComputeShaderCtx *ctx, Arena a, uv3 rf_data_dim) { for (u32 i = 0; i < ARRAY_COUNT(ctx->programs); i++) { - char *shader_text = LoadFileText(compute_shader_paths[i]); - u32 shader_id = rlCompileShader(shader_text, RL_COMPUTE_SHADER); - ctx->programs[i] = rlLoadComputeShaderProgram(shader_id); + Arena tmp = a; + os_file_stats fs = os_get_file_stats(compute_shader_paths[i]); + s8 shader_text = os_read_file(&tmp, compute_shader_paths[i], fs.filesize); + u32 shader_id = compile_shader(tmp, GL_COMPUTE_SHADER, shader_text); + ctx->programs[i] = rlLoadComputeShaderProgram(shader_id); glDeleteShader(shader_id); - UnloadFileText(shader_text); } ctx->rf_data_dim_id = glGetUniformLocation(ctx->programs[CS_UFORCES], "u_rf_data_dim"); - ctx->out_data_dim_id = glGetUniformLocation(ctx->programs[CS_UFORCES], "u_out_data_dim"); + ctx->out_data_tex_id = glGetUniformLocation(ctx->programs[CS_UFORCES], "u_out_data_tex"); + ctx->mip_view_tex_id = glGetUniformLocation(ctx->programs[CS_MIN_MAX], "u_mip_view_tex"); ctx->rf_data_dim = rf_data_dim; size rf_data_size = rf_data_dim.w * rf_data_dim.h * rf_data_dim.d * sizeof(i32); @@ -146,44 +161,12 @@ init_compute_shader_ctx(ComputeShaderCtx *ctx, Arena a, uv3 rf_data_dim) static void init_fragment_shader_ctx(FragmentShaderCtx *ctx, uv3 out_data_dim) { - ctx->shader = LoadShader(NULL, "shaders/render.glsl"); - - ctx->out_data_dim_id = glGetUniformLocation(ctx->shader.id, "u_out_data_dim"); - /* TODO: add min max uniform */ - - /* output texture for image blitting */ - ctx->output = LoadRenderTexture(out_data_dim.w, out_data_dim.h); -} - -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; + ctx->shader = LoadShader(NULL, "shaders/render.glsl"); + ctx->output = LoadRenderTexture(out_data_dim.w, out_data_dim.h); + ctx->out_data_tex_id = glGetUniformLocation(ctx->shader.id, "u_out_data_tex"); + ctx->mip_view_tex_id = glGetUniformLocation(ctx->shader.id, "u_mip_view_tex"); + ctx->db_cutoff_id = glGetUniformLocation(ctx->shader.id, "u_db_cutoff"); + ctx->db = -50.0f; } static void @@ -205,13 +188,17 @@ reload_shaders(BeamformerCtx *ctx, Arena a) } csctx->rf_data_dim_id = glGetUniformLocation(csctx->programs[CS_UFORCES], "u_rf_data_dim"); - csctx->out_data_dim_id = glGetUniformLocation(csctx->programs[CS_UFORCES], "u_out_data_dim"); + csctx->out_data_tex_id = glGetUniformLocation(csctx->programs[CS_UFORCES], "u_out_data_tex"); + csctx->mip_view_tex_id = glGetUniformLocation(csctx->programs[CS_MIN_MAX], "u_mip_view_tex"); + csctx->mips_level_id = glGetUniformLocation(csctx->programs[CS_MIN_MAX], "u_mip_map"); Shader updated_fs = LoadShader(NULL, "shaders/render.glsl"); if (updated_fs.id != rlGetShaderIdDefault()) { UnloadShader(ctx->fsctx.shader); ctx->fsctx.shader = updated_fs; - ctx->fsctx.out_data_dim_id = GetShaderLocation(updated_fs, "u_out_data_dim"); + ctx->fsctx.out_data_tex_id = GetShaderLocation(updated_fs, "u_out_data_tex"); + ctx->fsctx.mip_view_tex_id = GetShaderLocation(updated_fs, "u_mip_view_tex"); + ctx->fsctx.db_cutoff_id = GetShaderLocation(updated_fs, "u_db_cutoff"); } } @@ -237,8 +224,20 @@ main(void) ctx.font = GetFontDefault(); - size out_data_size = ctx.out_data_dim.w * ctx.out_data_dim.h * ctx.out_data_dim.d * sizeof(f32); - ctx.out_data_ssbo = rlLoadShaderBuffer(out_data_size, NULL, GL_DYNAMIC_COPY); + /* NOTE: allocate storage for beamformed output data; + * this is shared between compute and fragment shaders */ + { + uv3 odim = ctx.out_data_dim; + u32 max_dim = MAX(odim.x, MAX(odim.y, odim.z)); + /* TODO: does this actually matter or is 0 fine? */ + ctx.out_texture_unit = 0; + ctx.out_texture_mips = _tzcnt_u32(max_dim) + 1; + glActiveTexture(GL_TEXTURE0 + ctx.out_texture_unit); + glGenTextures(1, &ctx.out_texture); + glBindTexture(GL_TEXTURE_3D, ctx.out_texture); + glTexStorage3D(GL_TEXTURE_3D, ctx.out_texture_mips, GL_RG32F, odim.x, odim.y, odim.z); + } + init_compute_shader_ctx(&ctx.csctx, temp_memory, (uv3){.w = 3456, .h = 128, .d = 8}); init_fragment_shader_ctx(&ctx.fsctx, ctx.out_data_dim); diff --git a/shaders/hadamard.glsl b/shaders/hadamard.glsl @@ -1,3 +1,4 @@ +/* See LICENSE for license details. */ #version 460 core layout(local_size_x = 32, local_size_y = 32, local_size_z = 1) in; @@ -13,7 +14,7 @@ layout(std430, binding = 3) readonly restrict buffer buffer_3 { int hadamard[]; }; -layout(location = 3) uniform uvec3 u_rf_data_dim; +layout(location = 2) uniform uvec3 u_rf_data_dim; void main() { diff --git a/shaders/min_max.glsl b/shaders/min_max.glsl @@ -0,0 +1,27 @@ +/* See LICENSE for license details. */ + +/* NOTE: Does a binary search in 3D for smallest and largest output values */ + +#version 460 core +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(rg32f, location = 1) uniform image3D u_out_data_tex; +layout(rg32f, location = 2) uniform image3D u_mip_view_tex; +layout(location = 3) uniform int u_mip_map = 0; + +void main() +{ + ivec3 out_coord = ivec3(gl_GlobalInvocationID.xyz); + + ivec3 idx = out_coord * 2; + vec2 min_max = vec2(1000000000, 0); + for (int i = 0; i < 2; i++) { + for (int j = 0; j < 2; j++) { + vec2 a = imageLoad(u_out_data_tex, idx + ivec3(i, j, 0)).xy; + vec2 b = imageLoad(u_out_data_tex, idx + ivec3(i, j, 1)).xy; + min_max.x = min(min_max.x, min(a.x, b.x)); + min_max.y = max(min_max.y, max(a.y, b.y)); + } + } + imageStore(u_mip_view_tex, out_coord, vec4(min_max, 0, 1)); +} diff --git a/shaders/render.glsl b/shaders/render.glsl @@ -1,14 +1,12 @@ +/* See LICENSE for license details. */ #version 430 in vec2 fragTexCoord; out vec4 v_out_colour; -layout(std430, binding = 1) readonly buffer beamformed_data -{ - float out_data[]; -}; - -layout(location = 1) uniform uvec3 u_out_data_dim; +layout(rg32f, location = 1) uniform image3D u_out_data_tex; +layout(rg32f, location = 2) uniform image3D u_mip_view_tex; +layout(location = 3) uniform float u_db_cutoff = -60; /* input: h [0,360] | s,v [0, 1] * * output: rgb [0,1] */ @@ -21,9 +19,17 @@ vec3 hsv2rgb(vec3 hsv) void main() { - ivec2 coord = ivec2(fragTexCoord * u_out_data_dim.xy); - float smp = out_data[coord.y * u_out_data_dim.x + coord.x]; - smp = 20 * log(abs(smp)) + 50; + ivec3 out_data_dim = imageSize(u_out_data_tex); + ivec2 coord = ivec2(fragTexCoord * out_data_dim.xy); + vec2 min_max = imageLoad(u_mip_view_tex, ivec3(0, 0, 0)).xy; + + float smp = imageLoad(u_out_data_tex, ivec3(coord.x, coord.y, 0)).x; + float absmax = max(abs(min_max.y), abs(min_max.x)); + + smp = 20 * log(abs(smp) / absmax); + smp = clamp(smp, u_db_cutoff, 0) / u_db_cutoff; + smp = 1 - smp; - v_out_colour = vec4(hsv2rgb(vec3(smp, 0.8, 0.95)), 1); + //v_out_colour = vec4(hsv2rgb(vec3(360 * smp + 120, 0.8, 0.95)), 1); + v_out_colour = vec4(smp, smp, smp, 1); } diff --git a/shaders/uforces.glsl b/shaders/uforces.glsl @@ -1,3 +1,4 @@ +/* See LICENSE for license details. */ #version 460 core layout(local_size_x = 32, local_size_y = 32, local_size_z = 1) in; @@ -5,18 +6,14 @@ layout(std430, binding = 1) readonly restrict buffer buffer_1 { float rf_data[]; }; -layout(std430, binding = 2) writeonly restrict buffer buffer_2 { - float out_data[]; -}; - #define C_SPLINE 0.5 -layout(location = 3) uniform uvec3 u_rf_data_dim; -layout(location = 4) uniform uvec3 u_out_data_dim; -layout(location = 5) uniform float u_sound_speed = 1452; -layout(location = 6) uniform float u_sampling_frequency = 2.0833e7; -layout(location = 7) uniform float u_focal_depth = 0.07; -//layout(location = 9) uniform sampler2D u_element_positions; +layout(rg32f, location = 1) uniform image3D u_out_data_tex; +layout(location = 2) uniform uvec3 u_rf_data_dim; +layout(location = 3) uniform float u_sound_speed = 1452; +layout(location = 4) uniform float u_sampling_frequency = 2.0833e7; +layout(location = 5) uniform float u_focal_depth = 0.07; +//layout(location = 6) uniform sampler2D u_element_positions; /* NOTE: See: https://en.wikipedia.org/wiki/Cubic_Hermite_spline */ float cubic(uint ridx, float x) @@ -47,7 +44,9 @@ float cubic(uint ridx, float x) void main() { vec2 pixel = vec2(gl_GlobalInvocationID.xy); - ivec2 out_coord = ivec2(gl_GlobalInvocationID.xy); + ivec3 out_coord = ivec3(gl_GlobalInvocationID.xyz); + + ivec3 out_data_dim = imageSize(u_out_data_tex); /* NOTE: Convert pixel to physical coordinates */ /* TODO: Send these in like the 3D program */ @@ -60,9 +59,9 @@ void main() /* TODO: image extent can be different than xdc_size */ /* TODO: for now assume y-dimension is along transducer center */ vec3 image_point = vec3( - xdc_upper_left.x + pixel.x * xdc_size.x / u_out_data_dim.x, + xdc_upper_left.x + pixel.x * xdc_size.x / out_data_dim.x, 0, - pixel.y * 60e-3 / u_out_data_dim.y + 10e-3 + pixel.y * 60e-3 / out_data_dim.y + 10e-3 ); /* TODO: Send this into the GPU */ @@ -91,6 +90,5 @@ void main() } ridx += u_rf_data_dim.y * u_rf_data_dim.x; } - uint oidx = u_out_data_dim.x * out_coord.y + out_coord.x; - out_data[oidx] = sum; + imageStore(u_out_data_tex, out_coord, vec4(sum, sum, 0, 0)); } diff --git a/util.h b/util.h @@ -56,7 +56,7 @@ enum compute_shaders { // CS_FORCES, CS_HADAMARD, // CS_HERCULES, -// CS_MIN_MAX, + CS_MIN_MAX, CS_UFORCES, CS_LAST }; @@ -82,13 +82,18 @@ typedef struct { uv3 rf_data_dim; i32 rf_data_dim_id; - i32 out_data_dim_id; + i32 out_data_tex_id; + i32 mip_view_tex_id; + i32 mips_level_id; } ComputeShaderCtx; typedef struct { Shader shader; RenderTexture2D output; - i32 out_data_dim_id; + i32 out_data_tex_id; + i32 mip_view_tex_id; + i32 db_cutoff_id; + f32 db; } FragmentShaderCtx; typedef struct { @@ -99,8 +104,10 @@ typedef struct { Color bg, fg; - u32 out_data_ssbo; uv3 out_data_dim; + u32 out_texture; + u32 out_texture_unit; + u32 out_texture_mips; ComputeShaderCtx csctx; FragmentShaderCtx fsctx; @@ -113,7 +120,9 @@ typedef struct { #define ARRAY_COUNT(a) (sizeof(a) / sizeof(*a)) #define ABS(x) ((x) < 0 ? (-x) : (x)) +#define MAX(a, b) ((a) > (b) ? (a) : (b)) #define CLAMP(x, a, b) ((x) = (x) < (a) ? (a) : (x) > (b) ? (b) : (x)) +#define ISPOWEROF2(a) (((a) & ((a) - 1)) == 0) #include "util.c"