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

Browse Source
This commit is contained in:
Henri Verbeet
2025-10-12 12:39:58 +02:00
parent 3660a5a79c
commit 51cce61770
Notes: Henri Verbeet 2025-11-04 15:55:42 +01:00 
Approved-by: Giovanni Mascellani (@giomasce)
Approved-by: Henri Verbeet (@hverbeet)
Merge-Request: https://gitlab.winehq.org/wine/vkd3d/-/merge_requests/1808
2 changed files with 44 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -1362,6 +1362,46 @@ static void msl_dcl_indexable_temp(struct msl_generator *gen, const struct vkd3d
ins->declaration.indexable_temp.register_size);
}
static void msl_barrier(struct msl_generator *gen, const struct vkd3d_shader_instruction *ins)
{
uint32_t flags = ins->flags;
if (flags & (VKD3DSSF_GLOBAL_UAV | VKD3DSSF_THREAD_GROUP_UAV))
{
const char *scope = flags & VKD3DSSF_GLOBAL_UAV ? "thread_scope_device" : "thread_scope_threadgroup";
const char *mem_flags = "mem_flags::mem_device | mem_flags::mem_texture";
if (flags & VKD3DSSF_GROUP_SHARED_MEMORY)
{
mem_flags = "mem_flags::mem_device | mem_flags::mem_texture | mem_flags::mem_threadgroup";
flags &= ~VKD3DSSF_GROUP_SHARED_MEMORY;
}
msl_print_indent(gen->buffer, gen->indent);
vkd3d_string_buffer_printf(gen->buffer,
"atomic_thread_fence(%s, memory_order_seq_cst, %s);\n", mem_flags, scope);
flags &= ~(VKD3DSSF_GLOBAL_UAV | VKD3DSSF_THREAD_GROUP_UAV);
}
else if (flags & VKD3DSSF_GROUP_SHARED_MEMORY)
{
msl_print_indent(gen->buffer, gen->indent);
vkd3d_string_buffer_printf(gen->buffer,
"atomic_thread_fence(mem_flags::mem_threadgroup, memory_order_seq_cst, thread_scope_threadgroup);\n");
flags &= ~VKD3DSSF_GROUP_SHARED_MEMORY;
}
if (flags & VKD3DSSF_THREAD_GROUP)
{
msl_print_indent(gen->buffer, gen->indent);
vkd3d_string_buffer_printf(gen->buffer, "threadgroup_barrier(mem_flags::mem_none);\n");
flags &= ~VKD3DSSF_THREAD_GROUP;
}
if (flags)
msl_compiler_error(gen, VKD3D_SHADER_ERROR_MSL_INTERNAL,
"Internal compiler error: Unhandled synchronisation flags %#x.", flags);
}
static void msl_handle_instruction(struct msl_generator *gen, const struct vkd3d_shader_instruction *ins)
{
gen->location = ins->location;
@@ -1596,6 +1636,9 @@ static void msl_handle_instruction(struct msl_generator *gen, const struct vkd3d
case VSIR_OP_SWITCH:
msl_switch(gen, ins);
break;
case VSIR_OP_SYNC:
msl_barrier(gen, ins);
break;
case VSIR_OP_TAN:
msl_intrinsic(gen, ins, "tan");
break;

2
tests/hlsl/barriers.shader_test
View File

@@ -88,7 +88,7 @@ void main(uint local_idx : SV_GroupIndex, uint group_id : SV_GroupID)
}
[test]
todo(glsl | msl) dispatch 1 1 1
todo(glsl) dispatch 1 1 1
probe uav 2 (0) rgba(42.0, 42.0, 42.0, 42.0)
[require]