From 69b6632085518aaa8778258b42de58d510123578 Mon Sep 17 00:00:00 2001
From: Raph Levien <raph.levien@gmail.com>
Date: Sat, 13 Nov 2021 22:31:26 -0800
Subject: [PATCH] Fix write-after-read in prefix test

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.
---
 tests/shader/gen/prefix.dxil        | Bin 4860 -> 4876 bytes
 tests/shader/gen/prefix.hlsl        |   2 ++
 tests/shader/gen/prefix.msl         |   2 ++
 tests/shader/gen/prefix.spv         | Bin 9796 -> 9828 bytes
 tests/shader/gen/prefix_atomic.dxil | Bin 4864 -> 4884 bytes
 tests/shader/gen/prefix_atomic.hlsl |   2 ++
 tests/shader/gen/prefix_atomic.msl  |   2 ++
 tests/shader/gen/prefix_atomic.spv  | Bin 9820 -> 9852 bytes
 tests/shader/gen/prefix_vkmm.spv    | Bin 9984 -> 10016 bytes
 tests/shader/prefix.comp            |   2 ++
 10 files changed, 10 insertions(+)

diff --git a/tests/shader/gen/prefix.dxil b/tests/shader/gen/prefix.dxil
index 34f3d6affd966a3b2998becd3a99eb17fe69ce2b..73f1ba1fa807b2ce21f5bb6500358e7e1bbb8e3a 100644
GIT binary patch
delta 1257
zcmY*Ze@qi+82;{hJ+AGwJ&HhE3~8Nhx(eG-m;>FgrC1SBt8N6QWLKmNr`9@{!Rf~J
z7nIN$@W7#{xovJ^NSF({6r#zsWuOQIHt<J=t|$?+xX3cMWSNOO<KnXKpYMI&=Xu{J
z@Ao9%V_itEE-Wjs7Cm84HXLv`kf9~@MNn1`1OUL3+SF<^o&dn(gtk=^x@i;u3JAdk
z9xN<dUjq1>;CUDMdV!U@t;GT8?aze%f<%X@J7iqD8A6yf;)=C$o@7ezT%t-WX$^_=
z>xL04oDFm&Y{#8jkbs@Wgi1TD`C{vzaG;c!miQ@Y2(+GtH2DH^B|rS25&#c#Q*33F
z2%%Sx%w4@#?FFx;DGJK5N(!T_lO;s>NGx@H41s7p*1FO{f5&V>@}R3s5UFh_9GZ{L
z4o{BW(_`lEmyPc|;&ORG+43J|@ADOIt-lX{*z#*;C?u*--M+r`R@V+y%V5r{g99R%
z-ph@0_yoWwHPUH-U!y3&m(tm|G_H}&Q!-+e`4>*sbMD)l7Y5F>bqq{x6nTwR^)DT8
zzE`_zpKBZZnMifIrNW`?syIH|&URZH7feklwo?7|N4C=8l{)ii>)gQQj`2$~S3d3-
z-ze%U=9}vdx}6r6YhRtCj-OGx4epbvPH#A#c4|lZ^!5{Jw`SWZe30)8U78t~>bSzU
zJplbb3zOlzceH{X=wRORl9W69OI(}(cxn+Qd1hG@{=a35;{Plg=+DHS2$OSpH<OTE
zgaiSAM**^=#_V7$vLHk=qq0w+Ii^yc2@6*<ZSrM&>AalGK6FkzV(g6?&qp#Mrb$RE
zJ{m+!l@!oms4U}A<i#-)xsjZ;qdfrqQQFOkEvGz8Dq)8xriL&y^cEOHN+rO&p1YE>
zc1fX(Q)ZEXKm-e^WtQZ|G<ZAkm7~q#YL%Ys7CVGrv&qyDfXoDLa8rSN?64z83yf~+
z7{D9TWNx0^Xp~f|dUvl_mvYAYCY>lifQ4koEGagvrfJF#x@m8@cQ*?Y8LxDs#s(1E
z>g(xwi!O)vurS$|wTK$QUA17}^W>-SFxRE4uU6$e#EWBrvu2<8+MEq101}~3G`(FP
z(4^avU@9DGNgz92txY8MF2+XLbSkV}6q23J8y8tLo1SYlg?Zc=AEa*3n&uCbyx=!u
z+|mZ>7?);%@QM3X7ACNVMclbx)#j96MvKAU@zAD}XW>kI@pB&^+Hjng<y4}2cS5m<
z`uE|iir*GxF=-^9j%37|g#8FbPC~)xyhQQ)xXJHPKClQs0x@)`7!%kplJc$4>s-(x
zpf$=ey|zI<PqYu&6}l~o1#*K16(kdD{GJc#NYDfEj9lY>xi;%6iB=KQFW>OJK`6+|
z$*0bi7Ej0W5JQ&2;87SpmGvW4yk_psco1sIc>zU%mx`bF=f`7iU+&35Worr8IjB|o
z%h7CuueQ66ODKy{Nh=cVU|UFxb?Qi;M{G6-0(iA(f@0yt#O#_0>X2_WS7?SvqlOK>
zPm3$>*l_jJ?W3Zj;ikC0u`Ar%bvYq>4~)j7mLh(4f%>F{hOz)NqR=G!)+5ey&iZDS
J|BjxYe*hhBwyOXD

delta 1313
zcmYjRZ%h+s7=Q2X?e!=}IS-M80xL+n@^2ok0c5zelrVv^l*wE%COgn}QyFNgiB2-^
z%0F2;h1`ZLu%roU8f1j6%LHZS3eGV^2TG<NItJ5p)$xJ3Wyuyzb|{*?U*7lkJiq68
z@;uKkuUG$AmsIK~v6f{g9)6OsaTW$Ai%;bqRUrU?$8=9p5*|?k0EwQWo4ODV00$J{
zqU<bnl-q!OQ!?+WY%j4&i)kbPBY|AxAEbpx`U=To(>vEF&yLjn8$h+|fKQ9XAioZP
zHC0LQ&xswy8z=zRUF30t77*)zdE_v^0No1oS#e`(ACtJj*6oebzkUr`Q4`RyC~GI3
z)v(>p6AH%8YjxFGx-Iti6qawz#7`0d9wdTaO(4VdH1Tl;0FOA2?bPr7e))RX<pISS
zx}=t<=LpU^XQK*FPbJ)#g$S=F=JGB4m{<X~A-Bag=x&LFzVvi`LO8!;SAKy7<ZRmC
z=sx0(gG=c>;j#kR!loZ_9($*uscxSwGP8Om99g_S<5)O1K7XT%;v2ACxa$zk@%Vlp
z&W+-j=4(MSfh1oQl4p?=?NR=2Fr}t$Z{VQ2d~>#Yu%Cu+O|R}*nYkanwHP`NOXFO|
z;)Kg={$=BMrlq=*{L!h-^$y!c9(+9PeZPYLByGywlih#ff~d*v@2B^CDhEd%EzNjB
zp~c8#M9vXD5BH=Pi_aJfipGom7q=E%8qW+eN94G~6Bpcr*`gd4LO>vygs2!&Qblhj
z-?Q{gN}Z8`%$kbp5AUprU{&=@6ch43{x6EVvi-Rz0;u>3@^YYxxTA<S%8QPJ`>AMo
z$1(-Ljb#EZZd0E`4vCZMGe}U}9h-}`7l~c5=}5OY9Lplj;?J?kn^uFFvqrXx4cYM_
zfZvYsODfCfK{1xH%aI&qXz`b1c^L)}9^qc}kUhFoEisKm146nIGqU=eeUdNhZ2VdD
zZGHm=fra%+Lj1ql)%N)|;FpA2lw}$Vr8V+UPM>ph<copYZq+WNjFiO~<c-Q&rkzhN
zi$i^=$`v&tsLZ1az9c^&D`SR+4j`_mQK3a8ja`#8e&9VwP6*FI!@VS|sdwo!=iG$J
z5R)*1#$-7<HF;)0LF_>4hs*JvV@wsnFXLNMCGL_XbyDk3<5jI6@XH-ig2eUK{pm@X
z1s$51CE4FA{O+m4JQ)w`fGcXkh`IS%w8;nA!a-3{)w1a;ftHpBZqDyZ(-@GIR#Q>!
z+|y@G>(puWE%}+pu9bZ^V=<Uw(D8X8Y6W9vjm-|W65v=0Z?Ij_=KYdmWljcDYnGI3
zgL$}(1>Z&grHx;}yv)|M(K_izX0@f9)K$~E!|DL6XYI%zR_3mfX;~lsApJK*2gto<
z-11~{`wm4KvKfy79*#R#x_!XP)0+j_AkYulO7mJx6Ml*r-C?TdfR@1DsJxH8Mi!i$
zDsMx6fb4mV$I|CyTaAUC70PB_N9KP?BMQmS%Waz&KWH=Br@7uG63+rJ<+Z*VCch~d
z!)$OP7(~&J97?}d-+0xBm?<)6K-+bKXsjRW>L2TBW~cNV8=~5v+`BKB2pI|!@6g(K
RYdLg|IuCXE<ZtZ7`41m%s51Zn

diff --git a/tests/shader/gen/prefix.hlsl b/tests/shader/gen/prefix.hlsl
index 3af5a96..72cfa90 100644
--- a/tests/shader/gen/prefix.hlsl
+++ b/tests/shader/gen/prefix.hlsl
@@ -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;
diff --git a/tests/shader/gen/prefix.msl b/tests/shader/gen/prefix.msl
index 8e402a9..24bee60 100644
--- a/tests/shader/gen/prefix.msl
+++ b/tests/shader/gen/prefix.msl
@@ -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;
diff --git a/tests/shader/gen/prefix.spv b/tests/shader/gen/prefix.spv
index d2c1aadd6ee5592e1f023414f72843340b903e73..8e7db4a4638e3f8ad517015323935706f323bdc1 100644
GIT binary patch
delta 25
hcmX@&^TcO^gT&+%UWv^f5+dxA4df*@r^r8G0|0!*2&n)7

delta 17
YcmaFjbHrzZgT&?l2^sdyHS#am07I_^!vFvP

diff --git a/tests/shader/gen/prefix_atomic.dxil b/tests/shader/gen/prefix_atomic.dxil
index 68f47e55574ed40ff0301c6808f4953e3ed345a2..45a7dd8759ebdd7b2ce08e53aeaa163ef913d810 100644
GIT binary patch
delta 1205
zcmY*ZZA@EL7(VCr_R^LMy%sudDIbOLS)eN)p@1y8rMSRW7F-+wGxfqk-6k+1S)%(f
zZd<@jyRFNej5%WN#A!vX1%5!cska}S9TX{S>V7N)gfONMP>E6J4=*m6@%%Z@dEe)K
z-#qV=946~xma4k3%5-NuJ22dH>D{fj2Mqf^-8}>W0Pz=8YVleW0Fn#kGnoq*05lMT
z2_aS8SZfBtO$e-z5U(=v^JzE$UyK#NKOtn;Y>^}V+JP;kw?QU0eo0Bb=Oy5g3jA0M
zGAu$30B@22*x5yyi6v1Y>&0&tr&cUBP^nf_HAz}%32lm)iL&#7gy}E}nJh`9)W}S*
z`_XcEmpzSkbyo+M{SR(OCRef~)UEXC_H!K_Hqe+c_h40s@Xz(Q@r`%Z3mk1WK$}-5
zH@{k0onM}8ym!qzJ6%UGZK4Lu(uvUwc9g?t4@UEtSv}!X^Two*FDYV7W~|8%jLOo|
z8>N>@MvDyBbB*=A_{=T0xkI-;-M*JPR+1Lz8i_tUmd9RE#GeRl9$Ja`LzBzz&4(uM
ze}vYkiYgPlqu%if@5rb(*LUy?KGUm{yYoZS>t}V2?H(O_`R{}txm`o-l>OwsuAwVe
z6b@rTQRStP@e0HBBCk-EyuGNgd>lpCwJ4$@ba!7<@<XTJeaOGH@$36(+a{pVPfYP4
z|2Jh#_Ma&e2&;qHV=qa55GxIWX>sULA{qripa3-f2-VEi#-0{Dn~Uv-z3frB4z95S
z@^lE<8*;sX*W^hFk*JULQ+$c9)1-1Cq}jkBdFTo36q`zz5q~b*PULiZWS{6;XrPf`
z=XQNoRX~)3!$BTU1U~o*Ej~zgNZKH7QDtZcS_gP;*slE0*1$X?LV;1w@mA$un$6zX
z9^m4vO^BZCDC4)u!}$ZP9{1;oO-|X%a1TxZ)W;M>HIcnc(w;bkvq>#dw^c@YB3z=Y
z8{p9~!yj0nMe4E2_=(%R+65f1;WwW?+LjVU$NLv9w?=LjyTUu%HeB2Imb#Dv1S&d~
za%xd&dxNZ#FpGV<6rR3rOyM=IGzPJs6e3&#Pxm+fW=os|ed^GHkngwtY-<+TGJa|X
z7U@l`==E-<2|3Q8`m!OrIA+JFRlzR~emhI2sdM0(y||(=@-M8n2hPXN39wMVJ@(a%
z1z&PbCR6xiaj~_*z_YjkeGITMJatD{!T6{sjVx|gP1rC+U@e{~<sT4N6M1Pk^^3o<
zM!W@fVllu*)0ffiN~ft_VXD1YT{{@ax03h^)ZYa5JH&Qkn7~+Y1$XxZ^8=*QSpPLG
zKVYokPx}x?L+T5jmp`QBa4#3E2`ik!KyVgjV7fU6Ndk$@x}<pQlm}v#7)3Kx3>}a9
zXu_fjb9m(Hy~Pw-8mRMBQ&6wf)YJVuj%mSFJicfkv&?Rl0dtrux39-+5OGk<yBH?@
zKGZtTs0)v5OV7*RgQOyZWG9l{S(O@|kdlIM4HSeQ>BJxw1neC}y3$mOwt8AS&v3%$
H^7Q-z!3C=j

delta 1172
zcmY*Ze@qj16o22l>rvX;o(1H($}G)V2tzo2iB+@phoT^EY%U@hY}Y~&7%NC&!?=Xf
z!wRlC*wqeUF{j(IP7_ll*^s&9x^>L)!;#D-M$o{vpps?jGL17!ww*IB+56}F-sgQ@
z-uu4KH{X!iXVz`qYj>=j2oJ2R8%f*mTH5yf4_)UFNH(O?%MwWdsv*mFrk_aw7G!Wy
zN^RX+UJ24kNu(JGw>!j}nJR$j*n0E;p*}tf4;oD`<jC@?nm($ct}6T3RY4d$i44!+
zR+m%&2WWsk{)WQA(*$K6`DST)xzP_>^jdq3+{MaS$6O^f{?1qFvoi$pndMZe%{Re6
zM_AB2-%ogZQV79lG^a0b2<+KVYJ=Q0%^q*3H-%U<4lS3JN(f`@ba%hn+U9IXAq@5Q
z>cIVy?f1>qlc)1$k9s_QXwpXu8LkNM;vIMia3z3mXu?ao@s0hsWf<QS!8Zf9S>$qB
zTi<v+uerhP>D=uX6uOgd*PU3Xz1}i9AM!QX1Gil1la1BpQd9QI;L^pJQ-_9*e_Qg!
z(DC!Tk_QVLy7M~R-nOoedbfAC|EyMbQd#$8ZO!25yuOaEyFNO<5b~XM1a7^!T$c4-
z(!KD*!A&%mC)L%rb$L479p2s8ma#<{jamMDQ1s!>{daC8H_!j2oWAq*LjitKJQCz{
z?EivXmGH-cTo@~m|1C=^khDr6nkgj#AZ39_E$b!uZR&p1&R<c7k&Jh1^3Vz&(ioAA
z|4hT61b<tzHeF2mn3)!fW9<%I1fwfU1?*{}i|pa`R3(Px@Xb_qkD&O3aj{^KlgIKS
z_Vttn>c=yqGG^F47-7Rsj5~8|&!D0fm8mEoVm>Ts&uvajFH6A$KiQZpvg1pTB#|Q*
z*l=v()0ReCPPdFMumVNkQ)?TXsaOZSOYT#chfU*T4#ezSi-0k?#gQW;vSE`*_T<1%
zPvC75cCH{H5grB6qpzth)%f>?ZNFa;3SS;q6h_z#PaKeyQbi8@aj07Uqq8^@0{V$=
z;O>wIfD;0sP_3YV-v0`4HGx@WOEyk>4BSB}BF`xUtK{!S4x?dH;*9DtYD~V=s}&c6
z>NUkZdoPm}aM-(PJ9ns6AflQpw^`f2u633A*-bknH!gzI_Whymq&I~SOtw3s8kU((
zIAghpR&~N@<~M)*tud9;Kpw2>T_;A-dF_{wnm@<ARHb2Da}QS=E*EhNVJZl1Ef2*l
zGur-`RTkGc*>!Bs<K<UebL^NY3?dzIPkG;v)<*>CT#)<G_u6~y>bsOcR0QL>NGCO%
zqphw*<7|H~@)_ujd8afFm<2|iy_zge5r;i;(ix7arXZst<5v`B7KoihDtv%pjma~t
zfKgM{Os>XXL`Oi3t44~hZM>OK7p$&=phBMQ&k(Q^<EuFmY)rJn^;9BL96(v^Sk~QG
kaMUL~BG5_~_mTT3)~e!fXpL!jIpL1Fn+62w@A%LD0Yw0cMgRZ+

diff --git a/tests/shader/gen/prefix_atomic.hlsl b/tests/shader/gen/prefix_atomic.hlsl
index 10f7081..a75448f 100644
--- a/tests/shader/gen/prefix_atomic.hlsl
+++ b/tests/shader/gen/prefix_atomic.hlsl
@@ -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;
diff --git a/tests/shader/gen/prefix_atomic.msl b/tests/shader/gen/prefix_atomic.msl
index 6d7d155..910e842 100644
--- a/tests/shader/gen/prefix_atomic.msl
+++ b/tests/shader/gen/prefix_atomic.msl
@@ -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;
diff --git a/tests/shader/gen/prefix_atomic.spv b/tests/shader/gen/prefix_atomic.spv
index acca5452de5de6728d604a997c43cfad7b177987..d7dac5bcca704dbeee8902b63dc1a54d731f3c7d 100644
GIT binary patch
delta 25
hcmccP^T%gHfW+h!UWv^S5+dxA9poi8m&m_h0|1022_XOg

delta 17
Ycmez4bH`^xfW+nm2^sdyJ@Oye07y><D*ylh

diff --git a/tests/shader/gen/prefix_vkmm.spv b/tests/shader/gen/prefix_vkmm.spv
index 0b8d47505ddfe036afd31f31fbe9f3010124a2ea..cef3965f647503b83d6f6a63ff09657add1e08b0 100644
GIT binary patch
delta 36
ocmZqhTj00hfW+h}N)nq-NSLtmJz!vAFk)Z;V`GNRzvS820p7z2)&Kwi

delta 17
YcmZ4B*WkC|fW+nt5;p9cWfWN00Yird9RL6T

diff --git a/tests/shader/prefix.comp b/tests/shader/prefix.comp
index a6a0d57..b41d1b1 100644
--- a/tests/shader/prefix.comp
+++ b/tests/shader/prefix.comp
@@ -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;
             }