wine/vkd3d
0
0
Fork 0
mirror of https://gitlab.winehq.org/wine/vkd3d.git synced 2025-12-15 08:03:30 -08:00
Code Issues Packages Projects Releases Wiki Activity

vkd3d-shader/msl: Implement VKD3DSPR_LOCALTHREADINDEX.

Browse Source
This commit is contained in:
Giovanni Mascellani
2025-10-01 21:58:46 +02:00
committed by Henri Verbeet
parent a815e625d6
commit 3293d72f03
Notes: Henri Verbeet 2025-10-03 00:56:03 +02:00 
Approved-by: Henri Verbeet (@hverbeet)
Merge-Request: https://gitlab.winehq.org/wine/vkd3d/-/merge_requests/1765
2 changed files with 23 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

22
libs/vkd3d-shader/msl.c
View File

@@ -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_LOCALTHREADINDEX:
vkd3d_string_buffer_printf(buffer, "v_local_thread_index");
return MSL_DATA_UNION;
default:
msl_compiler_error(gen, VKD3D_SHADER_ERROR_MSL_INTERNAL,
"Internal compiler error: Unhandled register type %#x.", reg->type);
@@ -2022,6 +2026,12 @@ static void msl_generate_entrypoint_prologue(struct msl_generator *gen)
msl_print_indent(gen->buffer, 1);
vkd3d_string_buffer_printf(buffer, "v_thread_id.u = uint4(thread_id, 0u);\n");
}
if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADINDEX))
{
msl_print_indent(gen->buffer, 1);
vkd3d_string_buffer_printf(buffer, "v_local_thread_index.u = uint4(local_thread_index, 0u, 0u, 0u);\n");
}
}
static void msl_generate_entrypoint_epilogue(struct msl_generator *gen)
@@ -2114,6 +2124,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_LOCALTHREADINDEX))
{
msl_print_indent(gen->buffer, 2);
vkd3d_string_buffer_printf(gen->buffer, "uint local_thread_index [[thread_index_in_threadgroup]],\n");
}
msl_print_indent(gen->buffer, 2);
vkd3d_string_buffer_printf(gen->buffer, "vkd3d_%s_in input [[stage_in]])\n{\n", gen->prefix);
@@ -2125,6 +2141,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_LOCALTHREADINDEX))
vkd3d_string_buffer_printf(gen->buffer, " vkd3d_vec4 v_local_thread_index;\n");
vkd3d_string_buffer_printf(gen->buffer, "\n");
msl_generate_entrypoint_prologue(gen);
@@ -2138,6 +2156,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_LOCALTHREADINDEX))
vkd3d_string_buffer_printf(gen->buffer, ", v_local_thread_index");
if (gen->program->descriptors.descriptor_count)
vkd3d_string_buffer_printf(gen->buffer, ", descriptors");
vkd3d_string_buffer_printf(gen->buffer, ");\n\n");
@@ -2213,6 +2233,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_LOCALTHREADINDEX))
vkd3d_string_buffer_printf(gen->buffer, ", thread vkd3d_vec4 &v_local_thread_index");
if (gen->program->descriptors.descriptor_count)
vkd3d_string_buffer_printf(gen->buffer, ", constant descriptor *descriptors");
vkd3d_string_buffer_printf(gen->buffer, ")\n{\n");

2
tests/hlsl/compute.shader_test
View File

@@ -56,7 +56,7 @@ void main(uint3 thread_id : SV_DispatchThreadId, uint group_index : SV_GroupInde
}
[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, 38) u32(10, 0, 0, 1)
probe uav 0 (49, 49) u32(5, 0, 0, 1)