From f15e5ac22866ccb0b95e72e0acf2d86c415b1188 Mon Sep 17 00:00:00 2001 From: Giovanni Mascellani Date: Wed, 1 Oct 2025 22:03:30 +0200 Subject: [PATCH] vkd3d-shader/msl: Implement VKD3DSPR_LOCALTHREADID. --- libs/vkd3d-shader/msl.c | 22 ++++++++++++++++++++++ tests/hlsl/compute.shader_test | 2 +- 2 files changed, 23 insertions(+), 1 deletion(-) diff --git a/libs/vkd3d-shader/msl.c b/libs/vkd3d-shader/msl.c index 8d724c30c..322822d1a 100644 --- a/libs/vkd3d-shader/msl.c +++ b/libs/vkd3d-shader/msl.c @@ -490,6 +490,10 @@ static enum msl_data_type msl_print_register_name(struct vkd3d_string_buffer *bu vkd3d_string_buffer_printf(buffer, "v_thread_id"); return MSL_DATA_UNION; + case VKD3DSPR_LOCALTHREADID: + vkd3d_string_buffer_printf(buffer, "v_local_thread_id"); + return MSL_DATA_UNION; + case VKD3DSPR_LOCALTHREADINDEX: vkd3d_string_buffer_printf(buffer, "v_local_thread_index"); return MSL_DATA_UNION; @@ -2027,6 +2031,12 @@ static void msl_generate_entrypoint_prologue(struct msl_generator *gen) vkd3d_string_buffer_printf(buffer, "v_thread_id.u = uint4(thread_id, 0u);\n"); } + if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADID)) + { + msl_print_indent(gen->buffer, 1); + vkd3d_string_buffer_printf(buffer, "v_local_thread_id.u = uint4(local_thread_id, 0u);\n"); + } + if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADINDEX)) { msl_print_indent(gen->buffer, 1); @@ -2124,6 +2134,12 @@ static void msl_generate_entrypoint(struct msl_generator *gen) vkd3d_string_buffer_printf(gen->buffer, "uint3 thread_id [[thread_position_in_grid]],\n"); } + if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADID)) + { + msl_print_indent(gen->buffer, 2); + vkd3d_string_buffer_printf(gen->buffer, "uint3 local_thread_id [[thread_position_in_threadgroup]],\n"); + } + if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADINDEX)) { msl_print_indent(gen->buffer, 2); @@ -2141,6 +2157,8 @@ static void msl_generate_entrypoint(struct msl_generator *gen) vkd3d_string_buffer_printf(gen->buffer, " vkd3d_scalar o_mask;\n"); if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_THREADID)) vkd3d_string_buffer_printf(gen->buffer, " vkd3d_vec4 v_thread_id;\n"); + if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADID)) + vkd3d_string_buffer_printf(gen->buffer, " vkd3d_vec4 v_local_thread_id;\n"); if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADINDEX)) vkd3d_string_buffer_printf(gen->buffer, " vkd3d_vec4 v_local_thread_index;\n"); vkd3d_string_buffer_printf(gen->buffer, "\n"); @@ -2156,6 +2174,8 @@ static void msl_generate_entrypoint(struct msl_generator *gen) vkd3d_string_buffer_printf(gen->buffer, ", o_mask"); if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_THREADID)) vkd3d_string_buffer_printf(gen->buffer, ", v_thread_id"); + if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADID)) + vkd3d_string_buffer_printf(gen->buffer, ", v_local_thread_id"); if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADINDEX)) vkd3d_string_buffer_printf(gen->buffer, ", v_local_thread_index"); if (gen->program->descriptors.descriptor_count) @@ -2233,6 +2253,8 @@ static int msl_generator_generate(struct msl_generator *gen, struct vkd3d_shader vkd3d_string_buffer_printf(gen->buffer, ", thread vkd3d_scalar &o_mask"); if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_THREADID)) vkd3d_string_buffer_printf(gen->buffer, ", thread vkd3d_vec4 &v_thread_id"); + if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADID)) + vkd3d_string_buffer_printf(gen->buffer, ", thread vkd3d_vec4 &v_local_thread_id"); if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADINDEX)) vkd3d_string_buffer_printf(gen->buffer, ", thread vkd3d_vec4 &v_local_thread_index"); if (gen->program->descriptors.descriptor_count) diff --git a/tests/hlsl/compute.shader_test b/tests/hlsl/compute.shader_test index 438785cea..53161bf60 100644 --- a/tests/hlsl/compute.shader_test +++ b/tests/hlsl/compute.shader_test @@ -88,7 +88,7 @@ void main(uint3 thread_id : SV_DispatchThreadId, uint3 group_thread_id : SV_Grou } [test] -todo(glsl | msl) dispatch 13 13 1 +todo(glsl) dispatch 13 13 1 probe uav 0 (0, 0) u32(0, 0, 0, 1) probe uav 0 (14, 39) u32(2, 3, 0, 1) probe uav 0 (49, 49) u32(1, 1, 0, 1)