Commit: de24b8d286728efe9eb375d63b284f8d0aa4af0b
Parent: b98b2a7bc6a8de677ef12f992543bff73fa5ce48
Author: Randy Palamar
Date: Mon, 19 Oct 2026 06:35:00 -0700
core/lib: internal api for compute push parameters
there are a number of parameters that can be specified per
compute. Voxel transform definitely needs to be and the transmit
focus is useful for common use cases of FORCES/HERCULES. So while
this exact interface is not set in stone (not even visible
externally) this is a start.
Diffstat:
10 files changed, 175 insertions(+), 117 deletions(-)
diff --git a/beamformer.c b/beamformer.c
@@ -242,6 +242,10 @@ beamformer_init(BeamformerInput *input)
ctx->shared_memory->version = BEAMFORMER_SHARED_MEMORY_VERSION;
ctx->shared_memory->reserved_parameter_blocks = 1;
+ {
+ BeamformerParameterBlock *pb = beamformer_parameter_block(ctx->shared_memory, 0);
+ zero_struct(pb);
+ }
ctx->shared_memory->beamformed_frame_buffer_size = cs->backlog.buffer->size;
diff --git a/beamformer.meta b/beamformer.meta
@@ -450,12 +450,15 @@
@PushConstants
{
- [voxel_transform M4]
- [xdc_element_pitch V2]
- [rf_data U64]
- [output_frame U64]
- [channel_offset U32]
- [readi_group U32]
+ [voxel_transform M4]
+ [xdc_element_pitch V2]
+ [rf_data U64]
+ [output_frame U64]
+ [channel_offset U32]
+ [focus_depth F32]
+ [steering_angle F32]
+ [tx_rx_orientation U16]
+ [readi_group U16]
}
}
diff --git a/beamformer_core.c b/beamformer_core.c
@@ -946,11 +946,8 @@ plan_compute_pipeline(BeamformerComputePlan *cp, BeamformerParameterBlock *pb, A
db->OutputSizeZ = cp->output_points.z;
db->TransmitReceiveOrientation = pb->parameters.transmit_receive_orientation;
- cp->voxel_transform = m4_mul(cp->ui_voxel_transform, pb->parameters.das_voxel_transform);
cp->xdc_element_pitch = pb->parameters.xdc_element_pitch;
- memory_copy(cp->das_voxel_transform.E, cp->voxel_transform.E, sizeof(cp->voxel_transform));
-
u32 id = pb->parameters.acquisition_kind;
u32 acquisition_kind_flags = beamformer_acquisition_kind_flags[id];
@@ -1255,6 +1252,7 @@ beamformer_commit_parameter_block(BeamformerCtx *ctx, BeamformerComputePlan *cp,
switch (region) {
case BeamformerParameterDirtyFlag_NotifyUI:{
atomic_store_u32(&ctx->ui_dirty_parameter_blocks, 1u << block);
+ cp->ui_voxel_transform = m4_identity();
}break;
case BeamformerParameterDirtyFlag_Parameters:{
@@ -1333,9 +1331,12 @@ do_compute_shader(GPUCommandList cmd, BeamformerComputePlan *cp, BeamformerFrame
.rf_data = input_pointer,
.output_frame = frame->gpu_pointer,
.channel_offset = channel_offset,
+ .focus_depth = frame->focus_depth,
+ .steering_angle = frame->steering_angle,
+ .tx_rx_orientation = frame->transmit_receive_orientation,
.readi_group = cp->readi_group,
};
- memory_copy(pc.voxel_transform.E, cp->das_voxel_transform.E, sizeof(pc.voxel_transform));
+ memory_copy(pc.voxel_transform.E, frame->voxel_transform.E, sizeof(pc.voxel_transform));
gpu_command_push_constants(cmd, 0, sizeof(pc), &pc);
gpu_command_dispatch_compute(cmd, dispatch);
@@ -1490,11 +1491,11 @@ complete_queue(BeamformerCtx *ctx, BeamformWorkQueue *q, Arena *arena)
push_compute_timing_info(ctx->compute_timing_table,
(ComputeTimingInfo){.kind = ComputeTimingInfoKind_ComputeFrameBegin});
- BeamformerComputePlan *cp = beamformer_compute_plan_for_block(cs, work->compute_context.parameter_block, arena);
- if unlikely(beamformer_parameter_block_dirty(sm, work->compute_context.parameter_block)) {
- u32 block = work->compute_context.parameter_block;
+ BeamformerComputePlan *cp = beamformer_compute_plan_for_block(cs, work->compute_context.parameter_set, arena);
+ if unlikely(beamformer_parameter_block_dirty(sm, work->compute_context.parameter_set)) {
+ u32 set = work->compute_context.parameter_set;
Temp scratch = temp_begin(arena);
- beamformer_commit_parameter_block(ctx, cp, block, arena);
+ beamformer_commit_parameter_block(ctx, cp, set, arena);
temp_end(scratch);
}
@@ -1542,9 +1543,22 @@ complete_queue(BeamformerCtx *ctx, BeamformWorkQueue *q, Arena *arena)
frame->acquisition_kind = cp->acquisition_kind;
frame->contrast_mode = cp->contrast_mode;
frame->compound_count = cp->acquisition_count;
- frame->parameter_block = work->compute_context.parameter_block;
+ frame->parameter_block = work->compute_context.parameter_set;
frame->view_plane_tag = work->compute_context.view_plane;
- memory_copy(frame->voxel_transform.E, cp->voxel_transform.E, sizeof(cp->voxel_transform));
+ frame->focus_depth = work->compute_context.focus_depth;
+ frame->steering_angle = work->compute_context.steering_angle;
+ frame->transmit_receive_orientation = work->compute_context.tx_rx_orientation;
+
+ b32 transform_changed = !m4_equal(cp->last_voxel_transform, work->compute_context.voxel_transform);
+ memory_copy(&cp->last_voxel_transform, &work->compute_context.voxel_transform, sizeof(cp->last_voxel_transform));
+ memory_copy(&frame->voxel_transform, &cp->last_voxel_transform, sizeof(frame->voxel_transform));
+
+ if (transform_changed) {
+ atomic_store_u32(&beamformer_context->ui_dirty_parameter_blocks, 1u << work->compute_context.parameter_set);
+ cp->ui_voxel_transform = m4_identity();
+ } else if (u64_hash_from_str8(str8_struct(&cp->last_voxel_transform)) == cp->ui_last_voxel_transform_expected_hash) {
+ frame->voxel_transform = m4_mul(cp->ui_voxel_transform, frame->voxel_transform);
+ }
GPUCommandList cmd = gpu_command_list_begin(GPUTimeline_Compute);
gpu_command_timestamp(cmd);
@@ -1833,7 +1847,7 @@ DEBUG_EXPORT BEAMFORMER_RF_UPLOAD_FN(beamformer_rf_upload)
}
function void
-beamformer_queue_compute(BeamformerCtx *ctx, BeamformerFrame *frame, u32 parameter_block)
+beamformer_queue_compute(BeamformerCtx *ctx, BeamformerComputeWorkContext cw)
{
BeamformerSharedMemory *sm = ctx->shared_memory;
BeamformerSharedMemoryLockKind dispatch_lock = BeamformerSharedMemoryLockKind_DispatchCompute;
@@ -1842,8 +1856,7 @@ beamformer_queue_compute(BeamformerCtx *ctx, BeamformerFrame *frame, u32 paramet
BeamformWork *work = beamform_work_queue_push(ctx->beamform_work_queue);
if (work) {
work->kind = BeamformerWorkKind_Compute;
- work->compute_context.view_plane = frame ? frame->view_plane_tag : 0;
- work->compute_context.parameter_block = parameter_block;
+ memory_copy(&work->compute_context, &cw, sizeof(cw));
beamform_work_queue_push_commit(ctx->beamform_work_queue);
}
}
@@ -1885,8 +1898,18 @@ beamformer_process_input_events(BeamformerCtx *ctx, BeamformerInput *input,
}
// TODO(rnp): track latest parameter block
- if (ctx->latest_frame)
- beamformer_queue_compute(ctx, ctx->latest_frame, 0);
+ if (ctx->latest_frame) {
+ BeamformerFrame *f = ctx->latest_frame;
+ beamformer_queue_compute(ctx, (BeamformerComputeWorkContext){
+ .voxel_transform = f->voxel_transform,
+ .rf_offset = 0,
+ .focus_depth = f->focus_depth,
+ .steering_angle = f->steering_angle,
+ .tx_rx_orientation = f->transmit_receive_orientation,
+ .parameter_set = 0,
+ .view_plane = f->view_plane_tag,
+ });
+ }
}break;
case BeamformerFileReloadKind_RenderShader:{
diff --git a/beamformer_internal.h b/beamformer_internal.h
@@ -325,11 +325,10 @@ struct BeamformerComputePlan {
u32 rf_size;
b32 iq_pipeline;
+ u64 ui_last_voxel_transform_expected_hash;
m4 ui_voxel_transform;
- // NOTE(rnp): final voxel transform determining frame dimensions
- m4 voxel_transform;
- // NOTE(rnp): final voxel transform used for beamforming
- m4 das_voxel_transform;
+ // NOTE(rnp): voxel transform passed in from work item
+ m4 last_voxel_transform;
iv3 output_points;
i32 average_frames;
@@ -428,7 +427,11 @@ typedef struct {
u32 id;
u32 compound_count;
- u32 parameter_block;
+ u16 parameter_block;
+ u16 transmit_receive_orientation;
+ f32 focus_depth;
+ f32 steering_angle;
+
BeamformerDataKind data_kind;
BeamformerAcquisitionKind acquisition_kind;
BeamformerContrastMode contrast_mode;
diff --git a/beamformer_shared_memory.c b/beamformer_shared_memory.c
@@ -1,5 +1,5 @@
/* See LICENSE for license details. */
-#define BEAMFORMER_SHARED_MEMORY_VERSION (36UL)
+#define BEAMFORMER_SHARED_MEMORY_VERSION (37UL)
typedef enum {
BeamformerWorkKind_Compute,
@@ -38,9 +38,15 @@ typedef enum {BEAMFORMER_SHARED_MEMORY_LOCKS BeamformerSharedMemoryLockKind_Coun
#undef X
typedef struct {
+ m4 voxel_transform;
+ u64 rf_offset;
+ f32 focus_depth;
+ f32 steering_angle;
+ u16 tx_rx_orientation;
+ u16 parameter_set;
BeamformerViewPlaneTag view_plane;
- u32 parameter_block;
- u64 rf_offset;
+
+ static_assert(BeamformerMaxParameterSets <= U16_MAX, "");
} BeamformerComputeWorkContext;
/* NOTE: discriminated union based on type */
diff --git a/generated/beamformer.c b/generated/beamformer.c
@@ -276,7 +276,10 @@ typedef struct {
u64 rf_data;
u64 output_frame;
u32 channel_offset;
- u32 readi_group;
+ f32 focus_depth;
+ f32 steering_angle;
+ u16 tx_rx_orientation;
+ u16 readi_group;
} BeamformerDASPushConstants;
typedef struct {
@@ -739,12 +742,15 @@ read_only str8 beamformer_shader_global_header_strings[] = {
"\n"),
str8_comp(""
"layout(push_constant, std430) uniform PushConstants {\n"
- " f32mat4 voxel_transform;\n"
- " f32vec2 xdc_element_pitch;\n"
- " uint64_t rf_data;\n"
- " uint64_t output_frame;\n"
- " uint32_t channel_offset;\n"
- " uint32_t readi_group;\n"
+ " f32mat4 voxel_transform;\n"
+ " f32vec2 xdc_element_pitch;\n"
+ " uint64_t rf_data;\n"
+ " uint64_t output_frame;\n"
+ " uint32_t channel_offset;\n"
+ " float32_t focus_depth;\n"
+ " float32_t steering_angle;\n"
+ " uint16_t tx_rx_orientation;\n"
+ " uint16_t readi_group;\n"
"};\n"
"\n"),
str8_comp(""
diff --git a/lib/ogl_beamformer_lib.c b/lib/ogl_beamformer_lib.c
@@ -572,13 +572,20 @@ beamformer_start_compute(u64 offset, u32 image_plane_tag, u32 parameter_slot)
if (lib_error_check(parameter_slot < reserved_blocks, ParameterBlockUnallocated) &&
lib_error_check(image_plane_tag < BeamformerViewPlaneTag_Count, InvalidImagePlane))
{
+ // TODO(rnp): this stuff will be passed in rather than taken from the existing set
+ BeamformerParameters *bp = &beamformer_parameter_block(g_beamformer_library_context.bp, parameter_slot)->parameters;
+
BeamformWork *work = try_push_work_queue();
if (work) {
b64 new_rf = atomic_swap_u64(&g_beamformer_library_context.bp->new_rf_upload, 0ull);
work->kind = new_rf ? BeamformerWorkKind_WaitThenCompute : BeamformerWorkKind_Compute;
- work->compute_context.view_plane = image_plane_tag;
- work->compute_context.parameter_block = parameter_slot;
- work->compute_context.rf_offset = offset;
+ work->compute_context.view_plane = image_plane_tag;
+ work->compute_context.parameter_set = parameter_slot;
+ work->compute_context.rf_offset = offset;
+ work->compute_context.steering_angle = bp->focal_vector.x;
+ work->compute_context.focus_depth = bp->focal_vector.y;
+ work->compute_context.tx_rx_orientation = bp->transmit_receive_orientation;
+ memory_copy(&work->compute_context.voxel_transform, &bp->das_voxel_transform, sizeof(bp->das_voxel_transform));
beamform_work_queue_push_commit(&g_beamformer_library_context.bp->external_work_queue);
beamformer_flush_commands();
result = 1;
diff --git a/math.c b/math.c
@@ -427,8 +427,18 @@ function b32
m4_equal(m4 a, m4 b)
{
b32 result = 1;
+ #ifdef __AVX512F__
+ __m512 one = _mm512_set1_ps(1.f);
+ __m512 eps = _mm512_set1_ps(F32_EPSILON);
+ __m512 av = _mm512_loadu_ps(&a);
+ __m512 bv = _mm512_loadu_ps(&b);
+ __m512 delta = _mm512_abs_ps(_mm512_sub_ps(av, bv));
+ __m512 test = _mm512_mul_ps(eps, _mm512_max_ps(one, _mm512_max_ps(_mm512_abs_ps(av), _mm512_abs_ps(bv))));
+ result = _mm512_cmp_ps_mask(delta, test, _CMP_GT_OS) == 0;
+ #else
for EachElement(a.E, it)
result &= f32_equal(a.E[it], b.E[it]);
+ #endif
return result;
}
diff --git a/shaders/das.glsl b/shaders/das.glsl
@@ -304,14 +304,12 @@ RESULT_TYPE RCA(const vec3 world_point)
RESULT_TYPE HERCULES(const vec3 world_point)
{
// TODO(rnp): this whole thing needs have an outer loop over transmit tiles
- const mat4 xdc_transform = transducer_transform(0);
+ const mat4 xdc_transform = transducer_transform(0);
+ const bool rx_cols = RX_ORIENTATION(tx_rx_orientation) == RCAOrientation_Columns;
+ const vec2 focal_vector = f32vec2(steering_angle, focus_depth);
+ const vec3 xdc_world_point = (xdc_transform * vec4(world_point, 1)).xyz;
- const u8 tx_rx_orientation = tx_rx_orientation_for_acquisition(0);
- const bool rx_cols = RX_ORIENTATION(tx_rx_orientation) == RCAOrientation_Columns;
- const vec2 focal_vector = focal_vector_for_acquisition(0);
- const vec3 xdc_world_point = (xdc_transform * vec4(world_point, 1)).xyz;
-
- const f32 transmit_index = sample_index(rca_transmit_distance(world_point, focal_vector, tx_rx_orientation));
+ const f32 transmit_index = sample_index(rca_transmit_distance(world_point, focal_vector, u8(tx_rx_orientation)));
const f32 z_delta_squared = xdc_world_point.z * xdc_world_point.z;
const f32 f_number_over_z = abs(FNumber / xdc_world_point.z);
const f32 apodization_test = 0.25f / (f_number_over_z * f_number_over_z);
@@ -357,12 +355,11 @@ RESULT_TYPE FORCES(const vec3 world_point)
{
RESULT_TYPE result = RESULT_TYPE(0);
- const mat4 xdc_transform = transducer_transform(0);
- const u8 tx_rx_orientation = tx_rx_orientation_for_acquisition(0);
- const u32 rx_axis = u32(RX_ORIENTATION(tx_rx_orientation) == RCAOrientation_Rows);
- const f32 element_pitch = xdc_element_pitch[rx_axis];
+ const mat4 xdc_transform = transducer_transform(0);
+ const u32 rx_axis = u32(RX_ORIENTATION(tx_rx_orientation) == RCAOrientation_Rows);
+ const f32 element_pitch = xdc_element_pitch[rx_axis];
// TODO(rnp): the sign of the origin offset might be flipped
- const f32 origin_offset = FocusDepth * tan(radians(TransmitAngle));
+ const f32 origin_offset = focus_depth * tan(radians(steering_angle));
vec3 xdc_world_point = (xdc_transform * vec4(world_point, 1)).xyz;
f32 lateral_coord = xdc_world_point[rx_axis];
@@ -411,12 +408,11 @@ RESULT_TYPE READI_FORCES(const vec3 world_point)
{
RESULT_TYPE result = RESULT_TYPE(0);
- const mat4 xdc_transform = transducer_transform(0);
- const u8 tx_rx_orientation = tx_rx_orientation_for_acquisition(0);
- const u32 rx_axis = u32(RX_ORIENTATION(tx_rx_orientation) == RCAOrientation_Rows);
- const f32 element_pitch = xdc_element_pitch[rx_axis];
+ const mat4 xdc_transform = transducer_transform(0);
+ const u32 rx_axis = u32(RX_ORIENTATION(tx_rx_orientation) == RCAOrientation_Rows);
+ const f32 element_pitch = xdc_element_pitch[rx_axis];
// TODO(rnp): the sign of the origin offset might be flipped
- const f32 origin_offset = FocusDepth * tan(radians(TransmitAngle));
+ const f32 origin_offset = focus_depth * tan(radians(steering_angle));
vec3 xdc_world_point = (xdc_transform * vec4(world_point, 1)).xyz;
f32 lateral_coord = xdc_world_point[rx_axis];
diff --git a/ui.c b/ui.c
@@ -375,6 +375,7 @@ typedef struct {
BeamformerFrame latest_plane[BeamformerViewPlaneTag_Count];
b32 flush_parameters;
+ u32 flush_dirty_flags;
u32 selected_parameter_block;
// TODO(rnp): this should be per parameter block
@@ -2322,17 +2323,9 @@ function UI_CUSTOM_DRAW_FUNCTION(beamformer_custom_draw_frame_view)
}
function b32
-ui_rebuild_das_transform(u32 parameter_block, i32 dimension, v3 min, v3 max, f32 off_axis_position)
+ui_rebuild_das_transform(m4 das_transform, u32 parameter_block, i32 dimension, v3 min, v3 max, f32 off_axis_position)
{
- BeamformerUI *ui = ui_context;
-
- b32 result = 0;
m4 new_transform = m4_identity();
-
- BeamformerParameterBlock *pb = beamformer_parameter_block(beamformer_context->shared_memory, parameter_block);
-
- m4 das_transform = pb->parameters.das_voxel_transform;
-
switch (dimension) {
case 1:{new_transform = das_transform_1d(min, max);}break;
case 3:{new_transform = das_transform_3d(min, max);}break;
@@ -2348,7 +2341,7 @@ ui_rebuild_das_transform(u32 parameter_block, i32 dimension, v3 min, v3 max, f32
v3 rotation_axis = cross(v3_normalize(new_transform.c[0].xyz), N);
- m4 R = m4_rotation_about_axis(rotation_axis, ui->beamform_plane);
+ m4 R = m4_rotation_about_axis(rotation_axis, ui_context->beamform_plane);
m4 T = m4_translation(v3_scale(m4_mul_v3(R, N), off_axis_position));
new_transform = m4_mul(T, m4_mul(R, new_transform));
@@ -2358,16 +2351,12 @@ ui_rebuild_das_transform(u32 parameter_block, i32 dimension, v3 min, v3 max, f32
new_transform = m4_mul(new_transform, m4_inverse(das_transform));
BeamformerComputePlan *cp = beamformer_context->compute_context.compute_plans[parameter_block];
+ b32 result = 0;
if (cp) {
result |= !m4_equal(new_transform, cp->ui_voxel_transform);
- memory_copy(cp->ui_voxel_transform.E, new_transform.E, sizeof(new_transform));
- }
-
- if (result) {
- beamformed_shared_memory_dirty_flag(beamformer_context->shared_memory, parameter_block,
- BeamformerParameterDirtyFlag_Parameters);
+ memory_copy(&cp->ui_voxel_transform, &new_transform, sizeof(new_transform));
+ cp->ui_last_voxel_transform_expected_hash = u64_hash_from_str8(str8_struct(&das_transform));
}
-
return result;
}
@@ -3167,11 +3156,17 @@ ui_build_frame_view(UINode *container, BeamformerFrameView *view)
if (rebuild_transform) {
min_coordinate.E[0] = min_2d.E[0]; min_coordinate.E[1] = min_2d.E[1];
max_coordinate.E[0] = max_2d.E[0]; max_coordinate.E[1] = max_2d.E[1];
- if (ui_rebuild_das_transform(frame->parameter_block, iv3_dimension(frame->points),
- min_coordinate, max_coordinate,
+
+ BeamformerComputePlan *cp = beamformer_context->compute_context.compute_plans[frame->parameter_block];
+ m4 last_transform = m4_identity();
+ if (cp) last_transform = cp->last_voxel_transform;
+
+ if (ui_rebuild_das_transform(last_transform, frame->parameter_block,
+ iv3_dimension(frame->points), min_coordinate, max_coordinate,
plane_offset_from_transform(frame->voxel_transform)))
{
ui->flush_parameters = 1;
+ ui->selected_parameter_block = frame->parameter_block;
}
}
}
@@ -3498,14 +3493,20 @@ ui_build_parameters_listing(BeamformerUIPanel *panel)
UIFlags(UINodeFlag_Scroll|UINodeFlag_TextInputNumeric)
{
UISignal signal = ui_text_boxf("%0.2f###sound", bp->speed_of_sound);
- if (ui_tweak_f32_compute_variable(signal, &bp->speed_of_sound, 1.f, 10.f, (v2){{0, inf32()}}))
- ui->flush_parameters = 1;
+ if (ui_tweak_f32_compute_variable(signal, &bp->speed_of_sound, 1.f, 10.f, (v2){{0, inf32()}})) {
+ ui->selected_parameter_block = parameter_block;
+ ui->flush_parameters = 1;
+ ui->flush_dirty_flags |= 1 << BeamformerParameterDirtyFlag_Parameters;
+ }
}
b32 rebuild_transform = 0;
- m4 das_transform = pb->parameters.das_voxel_transform;
- if (cp) das_transform = m4_mul(cp->ui_voxel_transform, das_transform);
-
+ m4 das_transform = m4_identity();
+ m4 last_transform = m4_identity();
+ if (cp) {
+ last_transform = cp->last_voxel_transform;
+ das_transform = m4_mul(cp->ui_voxel_transform, last_transform);
+ }
f32 off_axis_position = plane_offset_from_transform(das_transform);
v3 coordinates[2] = {
@@ -3623,8 +3624,11 @@ ui_build_parameters_listing(BeamformerUIPanel *panel)
UIFlags(UINodeFlag_Scroll|UINodeFlag_TextInputNumeric)
{
UISignal signal = ui_text_boxf("%0.2f###f_number", bp->f_number);
- if (ui_tweak_f32_compute_variable(signal, &bp->f_number, 1.f, 0.05f, (v2){{0, inf32()}}))
- ui->flush_parameters = 1;
+ if (ui_tweak_f32_compute_variable(signal, &bp->f_number, 1.f, 0.05f, (v2){{0, inf32()}})) {
+ ui->selected_parameter_block = parameter_block;
+ ui->flush_parameters = 1;
+ ui->flush_dirty_flags |= 1 << BeamformerParameterDirtyFlag_Parameters;
+ }
}
UIParent(label_column) ui_label(str8("Interpolation"));
@@ -3638,7 +3642,9 @@ ui_build_parameters_listing(BeamformerUIPanel *panel)
i32 delta = signal.scroll.y + ui_pressed(signal);
bp->interpolation_mode = circular_add(bp->interpolation_mode, delta,
BeamformerInterpolationMode_Count);
- ui->flush_parameters = 1;
+ ui->selected_parameter_block = parameter_block;
+ ui->flush_parameters = 1;
+ ui->flush_dirty_flags |= 1 << BeamformerParameterDirtyFlag_Parameters;
}
}
@@ -3652,13 +3658,18 @@ ui_build_parameters_listing(BeamformerUIPanel *panel)
str8("False###coherency_weighting"));
if (signal.flags & (UISignalFlag_Pressed|UISignalFlag_Scrolled)) {
bp->coherency_weighting = !bp->coherency_weighting;
- ui->flush_parameters = 1;
+ ui->selected_parameter_block = parameter_block;
+ ui->flush_parameters = 1;
+ ui->flush_dirty_flags |= 1 << BeamformerParameterDirtyFlag_Parameters;
}
}
if (rebuild_transform) {
- if (ui_rebuild_das_transform(parameter_block, dimension, coordinates[0], coordinates[1], off_axis_position))
+ if (ui_rebuild_das_transform(last_transform, parameter_block, dimension, coordinates[0], coordinates[1], off_axis_position))
+ {
ui->flush_parameters = 1;
+ ui->selected_parameter_block = parameter_block;
+ }
}
}
}
@@ -5293,32 +5304,11 @@ beamformer_ui_frame(void)
u32 selected_block = ui->selected_parameter_block % BeamformerMaxParameterBlocks;
u32 selected_mask = 1 << selected_block;
- if (beamformer_context->ui_dirty_parameter_blocks & selected_mask) {
- BeamformerParameterBlock *pb = beamformer_parameter_block_lock(beamformer_context->shared_memory, selected_block, 0);
- if (pb) {
- ui->flush_parameters = 0;
-
- m4 das_transform;
- memory_copy(das_transform.E, pb->parameters.das_voxel_transform.E, sizeof(das_transform));
-
- atomic_and_u32(&beamformer_context->ui_dirty_parameter_blocks, ~selected_mask);
- beamformer_parameter_block_unlock(beamformer_context->shared_memory, selected_block);
-
- BeamformerComputePlan *cp = beamformer_context->compute_context.compute_plans[selected_block];
- m4 identity = m4_identity();
- b32 recompute = !m4_equal(identity, cp->ui_voxel_transform);
- memory_copy(cp->ui_voxel_transform.E, identity.E, sizeof(identity));
-
- if (recompute) {
- beamformed_shared_memory_dirty_flag(beamformer_context->shared_memory, selected_block,
- BeamformerParameterDirtyFlag_Parameters);
- beamformer_queue_compute(beamformer_context,
- beamformer_frame_from_index(beamformer_registers()->frame),
- selected_block);
- }
-
- ui->beamform_plane = 0;
- }
+ if (atomic_load_u32(&beamformer_context->ui_dirty_parameter_blocks) & selected_mask) {
+ atomic_and_u32(&beamformer_context->ui_dirty_parameter_blocks, ~selected_mask);
+ ui->flush_parameters = 0;
+ ui->flush_dirty_flags = 0;
+ ui->beamform_plane = 0;
}
/* NOTE: process interactions first because the user interacted with
@@ -5334,13 +5324,23 @@ beamformer_ui_frame(void)
if (ui->flush_parameters && beamformer_frame_valid(beamformer_registers()->frame)) {
BeamformerParameterBlock *pb = beamformer_parameter_block_lock(beamformer_context->shared_memory, selected_block, 0);
if (pb) {
- ui->flush_parameters = 0;
- beamformed_shared_memory_dirty_flag(beamformer_context->shared_memory, selected_block,
- BeamformerParameterDirtyFlag_Parameters);
+ atomic_or_u32(&pb->region_update_flags, ui->flush_dirty_flags);
+ ui->flush_parameters = 0;
+ ui->flush_dirty_flags = 0;
+
+ BeamformerComputePlan *cp = beamformer_context->compute_context.compute_plans[selected_block];
+ BeamformerFrame *f = beamformer_frame_from_index(beamformer_registers()->frame);
+ beamformer_queue_compute(beamformer_context, (BeamformerComputeWorkContext){
+ .voxel_transform = cp->last_voxel_transform,
+ .rf_offset = 0,
+ .focus_depth = f->focus_depth,
+ .steering_angle = f->steering_angle,
+ .tx_rx_orientation = f->transmit_receive_orientation,
+ .parameter_set = selected_block,
+ .view_plane = f->view_plane_tag,
+ });
+
beamformer_parameter_block_unlock(beamformer_context->shared_memory, selected_block);
- beamformer_queue_compute(beamformer_context,
- beamformer_frame_from_index(beamformer_registers()->frame),
- selected_block);
}
}