alex/vello
1
0
Fork 0
mirror of https://github.com/italicsjenga/vello.git synced 2025-01-09 20:31:29 +11:00
Code Issues Packages Projects Releases Wiki Activity

Fix write-after-read in prefix test

Browse source 
Thanks to Jeff Bolz for spotting the write-after-read hazard on the
sh_flag accesses. This fixes observed failures on Nvidia Turing and
Ampere on DX12.
This commit is contained in:
Raph Levien 2021-11-13 22:31:26 -08:00
parent 76a6f1fec8
commit 69b6632085
10 changed files with 10 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

BIN
tests/shader/gen/prefix.dxil
View file

Binary file not shown.

2
tests/shader/gen/prefix.hlsl
View file

@ -109,6 +109,7 @@ void comp_main()
GroupMemoryBarrierWithGroupSync();
DeviceMemoryBarrier();
uint flag_1 = sh_flag;
GroupMemoryBarrierWithGroupSync();
if (flag_1 == 2u)
{
if (gl_LocalInvocationID.x == 511u)
@ -174,6 +175,7 @@ void comp_main()
}
GroupMemoryBarrierWithGroupSync();
flag_1 = sh_flag;
GroupMemoryBarrierWithGroupSync();
if (flag_1 == 2u)
{
break;

2
tests/shader/gen/prefix.msl
View file

@ -160,6 +160,7 @@ kernel void main0(const device InBuf& _67 [[buffer(0)]], device OutBuf& _372 [[b
threadgroup_barrier(mem_flags::mem_threadgroup);
threadgroup_barrier(mem_flags::mem_device);
uint flag_1 = sh_flag;
threadgroup_barrier(mem_flags::mem_threadgroup);
if (flag_1 == 2u)
{
if (gl_LocalInvocationID.x == 511u)
@ -219,6 +220,7 @@ kernel void main0(const device InBuf& _67 [[buffer(0)]], device OutBuf& _372 [[b
}
threadgroup_barrier(mem_flags::mem_threadgroup);
flag_1 = sh_flag;
threadgroup_barrier(mem_flags::mem_threadgroup);
if (flag_1 == 2u)
{
break;

BIN
tests/shader/gen/prefix.spv
View file

Binary file not shown.

BIN
tests/shader/gen/prefix_atomic.dxil
View file

Binary file not shown.

2
tests/shader/gen/prefix_atomic.hlsl
View file

@ -112,6 +112,7 @@ void comp_main()
GroupMemoryBarrierWithGroupSync();
DeviceMemoryBarrier();
uint flag_1 = sh_flag;
GroupMemoryBarrierWithGroupSync();
if (flag_1 == 2u)
{
if (gl_LocalInvocationID.x == 511u)
@ -177,6 +178,7 @@ void comp_main()
}
GroupMemoryBarrierWithGroupSync();
flag_1 = sh_flag;
GroupMemoryBarrierWithGroupSync();
if (flag_1 == 2u)
{
break;

2
tests/shader/gen/prefix_atomic.msl
View file

@ -161,6 +161,7 @@ kernel void main0(const device InBuf& _67 [[buffer(0)]], device OutBuf& _372 [[b
threadgroup_barrier(mem_flags::mem_threadgroup);
threadgroup_barrier(mem_flags::mem_device);
uint flag_1 = sh_flag;
threadgroup_barrier(mem_flags::mem_threadgroup);
if (flag_1 == 2u)
{
if (gl_LocalInvocationID.x == 511u)
@ -220,6 +221,7 @@ kernel void main0(const device InBuf& _67 [[buffer(0)]], device OutBuf& _372 [[b
}
threadgroup_barrier(mem_flags::mem_threadgroup);
flag_1 = sh_flag;
threadgroup_barrier(mem_flags::mem_threadgroup);
if (flag_1 == 2u)
{
break;

BIN
tests/shader/gen/prefix_atomic.spv
View file

Binary file not shown.

BIN
tests/shader/gen/prefix_vkmm.spv
View file

Binary file not shown.

2
tests/shader/prefix.comp
View file

@ -144,6 +144,7 @@ void main() {
memoryBarrierBuffer();
#endif
uint flag = sh_flag;
barrier();
if (flag == FLAG_PREFIX_READY) {
if (gl_LocalInvocationID.x == WG_SIZE - 1) {
@ -185,6 +186,7 @@ void main() {
}
barrier();
flag = sh_flag;
barrier();
if (flag == FLAG_PREFIX_READY) {
break;
}