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

Browse Source
This commit is contained in:
Giovanni Mascellani
2025-10-01 22:06:26 +02:00
committed by Henri Verbeet
parent f15e5ac228
commit e524c36916
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_THREADGROUPID:
vkd3d_string_buffer_printf(buffer, "v_thread_group_id");
return MSL_DATA_UNION;
case VKD3DSPR_LOCALTHREADID:
vkd3d_string_buffer_printf(buffer, "v_local_thread_id");
return MSL_DATA_UNION;
@@ -2031,6 +2035,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_THREADGROUPID))
{
msl_print_indent(gen->buffer, 1);
vkd3d_string_buffer_printf(buffer, "v_thread_group_id.u = uint4(thread_group_id, 0u);\n");
}
if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADID))
{
msl_print_indent(gen->buffer, 1);
@@ -2134,6 +2144,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_THREADGROUPID))
{
msl_print_indent(gen->buffer, 2);
vkd3d_string_buffer_printf(gen->buffer, "uint3 thread_group_id [[threadgroup_position_in_grid]],\n");
}
if (bitmap_is_set(gen->program->io_dcls, VKD3DSPR_LOCALTHREADID))
{
msl_print_indent(gen->buffer, 2);
@@ -2157,6 +2173,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_THREADGROUPID))
vkd3d_string_buffer_printf(gen->buffer, " vkd3d_vec4 v_thread_group_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))
@@ -2174,6 +2192,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_THREADGROUPID))
vkd3d_string_buffer_printf(gen->buffer, ", v_thread_group_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))
@@ -2253,6 +2273,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_THREADGROUPID))
vkd3d_string_buffer_printf(gen->buffer, ", thread vkd3d_vec4 &v_thread_group_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))

2
tests/hlsl/compute.shader_test
View File

@@ -72,7 +72,7 @@ void main(uint3 thread_id : SV_DispatchThreadId, uint3 group_id : SV_GroupID)
}
[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(3, 9, 0, 1)
probe uav 0 (49, 49) u32(12, 12, 0, 1)