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_LOCALTHREADID.

Browse Source
This commit is contained in:
Giovanni Mascellani
2025-10-01 22:03:30 +02:00
committed by Henri Verbeet
parent 3293d72f03
commit f15e5ac228
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_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)

2
tests/hlsl/compute.shader_test
View File

@@ -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)