Avoid FXC's error X3694: race condition writing to shared resource detected

2026-04-22 03:02:01 -04:00 · 2023-10-05 23:46:02 +02:00
parent ca2e810492
commit d8ceeae3b2
6 changed files with 170 additions and 239 deletions
--- a/tests/in/const-exprs.wgsl
+++ b/tests/in/const-exprs.wgsl
@@ -1,32 +1,29 @@
-@group(0) @binding(0) var<storage, read_write> out: vec4<i32>;
-@group(0) @binding(1) var<storage, read_write> out2: i32;
-
 const TWO: u32 = 2u;
 const THREE: i32 = 3i;

@compute @workgroup_size(TWO, THREE, TWO - 1u)
 fn main() {
-   swizzle_of_compose();
-   index_of_compose();
-   compose_three_deep();
-   non_constant_initializers();
-   splat_of_constant();
-   compose_of_constant();
+    swizzle_of_compose();
+    index_of_compose();
+    compose_three_deep();
+    non_constant_initializers();
+    splat_of_constant();
+    compose_of_constant();
 }

 // Swizzle the value of nested Compose expressions.
 fn swizzle_of_compose() {
-   out = vec4(vec2(1, 2), vec2(3, 4)).wzyx; // should assign vec4(4, 3, 2, 1);
+    var out = vec4(vec2(1, 2), vec2(3, 4)).wzyx; // should assign vec4(4, 3, 2, 1);
 }

 // Index the value of nested Compose expressions.
 fn index_of_compose() {
-   out2 += vec4(vec2(1, 2), vec2(3, 4))[1]; // should assign 2
+    var out = vec4(vec2(1, 2), vec2(3, 4))[1]; // should assign 2
 }

 // Index the value of Compose expressions nested three deep
 fn compose_three_deep() {
-   out2 += vec4(vec3(vec2(6, 7), 8), 9)[0]; // should assign 6
+    var out = vec4(vec3(vec2(6, 7), 8), 9)[0]; // should assign 6
 }

 // While WGSL allows local variables to be declared anywhere in the function,
@@ -40,12 +37,12 @@ fn compose_three_deep() {
 // initializers that are constants as Naga locals with initializers. This test
 // checks that Naga local variable initializers are only used when safe.
 fn non_constant_initializers() {
-   var w = 10 + 20;
-   var x = w;
-   var y = x;
-   var z = 30 + 40;
+    var w = 10 + 20;
+    var x = w;
+    var y = x;
+    var z = 30 + 40;

-   out += vec4(w, x, y, z);
+    var out = vec4(w, x, y, z);
 }

 // Constant evaluation should be able to see through constants to
@@ -58,11 +55,11 @@ const TEST_CONSTANT_ADDITION: i32 = FOUR + FOUR;
 const TEST_CONSTANT_ALIAS_ADDITION: i32 = FOUR_ALIAS + FOUR_ALIAS;

 fn splat_of_constant() {
-    out = -vec4(FOUR);
+    var out = -vec4(FOUR);
 }

 fn compose_of_constant() {
-    out = -vec4(FOUR, FOUR, FOUR, FOUR);
+    var out = -vec4(FOUR, FOUR, FOUR, FOUR);
 }

 const PI: f32 = 3.141;
--- a/tests/out/glsl/const-exprs.main.Compute.glsl
+++ b/tests/out/glsl/const-exprs.main.Compute.glsl
@@ -18,26 +18,17 @@ const int TEXTURE_KIND_REGULAR = 0;
 const int TEXTURE_KIND_WARP = 1;
 const int TEXTURE_KIND_SKY = 2;

-layout(std430) buffer type_block_0Compute { ivec4 _group_0_binding_0_cs; };
-
-layout(std430) buffer type_1_block_1Compute { int _group_0_binding_1_cs; };
-

 void swizzle_of_compose() {
-    _group_0_binding_0_cs = ivec4(4, 3, 2, 1);
-    return;
+    ivec4 out_ = ivec4(4, 3, 2, 1);
 }

 void index_of_compose() {
-    int _e2 = _group_0_binding_1_cs;
-    _group_0_binding_1_cs = (_e2 + 2);
-    return;
+    int out_1 = 2;
 }

 void compose_three_deep() {
-    int _e2 = _group_0_binding_1_cs;
-    _group_0_binding_1_cs = (_e2 + 6);
-    return;
+    int out_2 = 6;
 }

 void non_constant_initializers() {
@@ -45,27 +36,25 @@ void non_constant_initializers() {
    int x = 0;
    int y = 0;
    int z = 70;
+    ivec4 out_3 = ivec4(0);
    int _e2 = w;
    x = _e2;
    int _e4 = x;
    y = _e4;
-    int _e9 = w;
-    int _e10 = x;
-    int _e11 = y;
-    int _e12 = z;
-    ivec4 _e14 = _group_0_binding_0_cs;
-    _group_0_binding_0_cs = (_e14 + ivec4(_e9, _e10, _e11, _e12));
+    int _e8 = w;
+    int _e9 = x;
+    int _e10 = y;
+    int _e11 = z;
+    out_3 = ivec4(_e8, _e9, _e10, _e11);
    return;
 }

 void splat_of_constant() {
-    _group_0_binding_0_cs = ivec4(-4, -4, -4, -4);
-    return;
+    ivec4 out_4 = ivec4(-4, -4, -4, -4);
 }

 void compose_of_constant() {
-    _group_0_binding_0_cs = ivec4(-4, -4, -4, -4);
-    return;
+    ivec4 out_5 = ivec4(-4, -4, -4, -4);
 }

 uint map_texture_kind(int texture_kind) {
--- a/tests/out/hlsl/const-exprs.hlsl
+++ b/tests/out/hlsl/const-exprs.hlsl
@@ -11,27 +11,22 @@ static const int TEXTURE_KIND_REGULAR = 0;
 static const int TEXTURE_KIND_WARP = 1;
 static const int TEXTURE_KIND_SKY = 2;

-RWByteAddressBuffer out_ : register(u0);
-RWByteAddressBuffer out2_ : register(u1);
-
 void swizzle_of_compose()
 {
-    out_.Store4(0, asuint(int4(4, 3, 2, 1)));
-    return;
+    int4 out_ = int4(4, 3, 2, 1);
+
 }

 void index_of_compose()
 {
-    int _expr2 = asint(out2_.Load(0));
-    out2_.Store(0, asuint((_expr2 + 2)));
-    return;
+    int out_1 = 2;
+
 }

 void compose_three_deep()
 {
-    int _expr2 = asint(out2_.Load(0));
-    out2_.Store(0, asuint((_expr2 + 6)));
-    return;
+    int out_2 = 6;
+
 }

 void non_constant_initializers()
@@ -40,30 +35,30 @@ void non_constant_initializers()
    int x = (int)0;
    int y = (int)0;
    int z = 70;
+    int4 out_3 = (int4)0;

    int _expr2 = w;
    x = _expr2;
    int _expr4 = x;
    y = _expr4;
-    int _expr9 = w;
-    int _expr10 = x;
-    int _expr11 = y;
-    int _expr12 = z;
-    int4 _expr14 = asint(out_.Load4(0));
-    out_.Store4(0, asuint((_expr14 + int4(_expr9, _expr10, _expr11, _expr12))));
+    int _expr8 = w;
+    int _expr9 = x;
+    int _expr10 = y;
+    int _expr11 = z;
+    out_3 = int4(_expr8, _expr9, _expr10, _expr11);
    return;
 }

 void splat_of_constant()
 {
-    out_.Store4(0, asuint(int4(-4, -4, -4, -4)));
-    return;
+    int4 out_4 = int4(-4, -4, -4, -4);
+
 }

 void compose_of_constant()
 {
-    out_.Store4(0, asuint(int4(-4, -4, -4, -4)));
-    return;
+    int4 out_5 = int4(-4, -4, -4, -4);
+
 }

 uint map_texture_kind(int texture_kind)
--- a/tests/out/msl/const-exprs.msl
+++ b/tests/out/msl/const-exprs.msl
@@ -18,60 +18,47 @@ constant int TEXTURE_KIND_WARP = 1;
 constant int TEXTURE_KIND_SKY = 2;

 void swizzle_of_compose(
-    device metal::int4& out
 ) {
-    out = metal::int4(4, 3, 2, 1);
-    return;
+    metal::int4 out = metal::int4(4, 3, 2, 1);
 }

 void index_of_compose(
-    device int& out2_
 ) {
-    int _e2 = out2_;
-    out2_ = _e2 + 2;
-    return;
+    int out_1 = 2;
 }

 void compose_three_deep(
-    device int& out2_
 ) {
-    int _e2 = out2_;
-    out2_ = _e2 + 6;
-    return;
+    int out_2 = 6;
 }

 void non_constant_initializers(
-    device metal::int4& out
 ) {
    int w = 30;
    int x = {};
    int y = {};
    int z = 70;
+    metal::int4 out_3 = {};
    int _e2 = w;
    x = _e2;
    int _e4 = x;
    y = _e4;
-    int _e9 = w;
-    int _e10 = x;
-    int _e11 = y;
-    int _e12 = z;
-    metal::int4 _e14 = out;
-    out = _e14 + metal::int4(_e9, _e10, _e11, _e12);
+    int _e8 = w;
+    int _e9 = x;
+    int _e10 = y;
+    int _e11 = z;
+    out_3 = metal::int4(_e8, _e9, _e10, _e11);
    return;
 }

 void splat_of_constant(
-    device metal::int4& out
 ) {
-    out = metal::int4(-4, -4, -4, -4);
-    return;
+    metal::int4 out_4 = metal::int4(-4, -4, -4, -4);
 }

 void compose_of_constant(
-    device metal::int4& out
 ) {
-    out = metal::int4(-4, -4, -4, -4);
-    return;
+    metal::int4 out_5 = metal::int4(-4, -4, -4, -4);
 }

 uint map_texture_kind(
@@ -94,14 +81,12 @@ uint map_texture_kind(
 }

 kernel void main_(
-  device metal::int4& out [[user(fake0)]]
-, device int& out2_ [[user(fake0)]]
 ) {
-    swizzle_of_compose(out);
-    index_of_compose(out2_);
-    compose_three_deep(out2_);
-    non_constant_initializers(out);
-    splat_of_constant(out);
-    compose_of_constant(out);
+    swizzle_of_compose();
+    index_of_compose();
+    compose_three_deep();
+    non_constant_initializers();
+    splat_of_constant();
+    compose_of_constant();
    return;
 }
--- a/tests/out/spv/const-exprs.spvasm
+++ b/tests/out/spv/const-exprs.spvasm
@@ -1,28 +1,19 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 105
+; Bound: 91
 OpCapability Shader
-OpExtension "SPV_KHR_storage_buffer_storage_class"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint GLCompute %95 "main"
-OpExecutionMode %95 LocalSize 2 3 1
-OpDecorate %20 DescriptorSet 0
-OpDecorate %20 Binding 0
-OpDecorate %21 Block
-OpMemberDecorate %21 0 Offset 0
-OpDecorate %23 DescriptorSet 0
-OpDecorate %23 Binding 1
-OpDecorate %24 Block
-OpMemberDecorate %24 0 Offset 0
+OpEntryPoint GLCompute %83 "main"
+OpExecutionMode %83 LocalSize 2 3 1
 %2 = OpTypeVoid
+%3 = OpTypeInt 32 0
 %4 = OpTypeInt 32 1
-%3 = OpTypeVector %4 4
-%5 = OpTypeInt 32 0
+%5 = OpTypeVector %4 4
 %6 = OpTypeFloat 32
 %7 = OpTypeVector %6 4
-%8 = OpConstant  %5  2
+%8 = OpConstant  %3  2
 %9 = OpConstant  %4  3
 %10 = OpConstant  %4  4
 %11 = OpConstant  %4  8
@@ -34,126 +25,107 @@ OpMemberDecorate %24 0 Offset 0
 %17 = OpConstant  %4  0
 %18 = OpConstant  %4  1
 %19 = OpConstant  %4  2
-%21 = OpTypeStruct %3
-%22 = OpTypePointer StorageBuffer %21
-%20 = OpVariable  %22  StorageBuffer
-%24 = OpTypeStruct %4
-%25 = OpTypePointer StorageBuffer %24
-%23 = OpVariable  %25  StorageBuffer
-%28 = OpTypeFunction %2
-%29 = OpTypePointer StorageBuffer %3
-%30 = OpConstant  %5  0
-%32 = OpConstantComposite  %3  %10 %9 %19 %18
-%36 = OpTypePointer StorageBuffer %4
-%44 = OpConstant  %4  6
-%51 = OpConstant  %4  30
-%52 = OpConstant  %4  70
-%54 = OpTypePointer Function %4
-%56 = OpConstantNull  %4
-%58 = OpConstantNull  %4
-%73 = OpConstant  %4  -4
-%74 = OpConstantComposite  %3  %73 %73 %73 %73
-%83 = OpTypeFunction %5 %4
-%84 = OpConstant  %5  10
-%85 = OpConstant  %5  20
-%86 = OpConstant  %5  30
-%93 = OpConstantNull  %5
-%27 = OpFunction  %2  None %28
+%22 = OpTypeFunction %2
+%23 = OpConstantComposite  %5  %10 %9 %19 %18
+%25 = OpTypePointer Function %5
+%30 = OpTypePointer Function %4
+%34 = OpConstant  %4  6
+%39 = OpConstant  %4  30
+%40 = OpConstant  %4  70
+%43 = OpConstantNull  %4
+%45 = OpConstantNull  %4
+%48 = OpConstantNull  %5
+%59 = OpConstant  %4  -4
+%60 = OpConstantComposite  %5  %59 %59 %59 %59
+%70 = OpTypeFunction %3 %4
+%71 = OpConstant  %3  10
+%72 = OpConstant  %3  20
+%73 = OpConstant  %3  30
+%74 = OpConstant  %3  0
+%81 = OpConstantNull  %3
+%21 = OpFunction  %2  None %22
+%20 = OpLabel
+%24 = OpVariable  %25  Function %23
+OpBranch %26
 %26 = OpLabel
-%31 = OpAccessChain  %29  %20 %30
-OpBranch %33
-%33 = OpLabel
-OpStore %31 %32
 OpReturn
 OpFunctionEnd
-%35 = OpFunction  %2  None %28
-%34 = OpLabel
-%37 = OpAccessChain  %36  %23 %30
-OpBranch %38
-%38 = OpLabel
-%39 = OpLoad  %4  %37
-%40 = OpIAdd  %4  %39 %19
-OpStore %37 %40
+%28 = OpFunction  %2  None %22
+%27 = OpLabel
+%29 = OpVariable  %30  Function %19
+OpBranch %31
+%31 = OpLabel
 OpReturn
 OpFunctionEnd
-%42 = OpFunction  %2  None %28
-%41 = OpLabel
-%43 = OpAccessChain  %36  %23 %30
-OpBranch %45
-%45 = OpLabel
-%46 = OpLoad  %4  %43
-%47 = OpIAdd  %4  %46 %44
-OpStore %43 %47
+%33 = OpFunction  %2  None %22
+%32 = OpLabel
+%35 = OpVariable  %30  Function %34
+OpBranch %36
+%36 = OpLabel
 OpReturn
 OpFunctionEnd
-%49 = OpFunction  %2  None %28
-%48 = OpLabel
-%55 = OpVariable  %54  Function %56
-%59 = OpVariable  %54  Function %52
-%53 = OpVariable  %54  Function %51
-%57 = OpVariable  %54  Function %58
-%50 = OpAccessChain  %29  %20 %30
-OpBranch %60
-%60 = OpLabel
-%61 = OpLoad  %4  %53
-OpStore %55 %61
-%62 = OpLoad  %4  %55
-OpStore %57 %62
-%63 = OpLoad  %4  %53
-%64 = OpLoad  %4  %55
-%65 = OpLoad  %4  %57
-%66 = OpLoad  %4  %59
-%67 = OpCompositeConstruct  %3  %63 %64 %65 %66
-%68 = OpLoad  %3  %50
-%69 = OpIAdd  %3  %68 %67
-OpStore %50 %69
+%38 = OpFunction  %2  None %22
+%37 = OpLabel
+%47 = OpVariable  %25  Function %48
+%42 = OpVariable  %30  Function %43
+%46 = OpVariable  %30  Function %40
+%41 = OpVariable  %30  Function %39
+%44 = OpVariable  %30  Function %45
+OpBranch %49
+%49 = OpLabel
+%50 = OpLoad  %4  %41
+OpStore %42 %50
+%51 = OpLoad  %4  %42
+OpStore %44 %51
+%52 = OpLoad  %4  %41
+%53 = OpLoad  %4  %42
+%54 = OpLoad  %4  %44
+%55 = OpLoad  %4  %46
+%56 = OpCompositeConstruct  %5  %52 %53 %54 %55
+OpStore %47 %56
 OpReturn
 OpFunctionEnd
-%71 = OpFunction  %2  None %28
-%70 = OpLabel
-%72 = OpAccessChain  %29  %20 %30
+%58 = OpFunction  %2  None %22
+%57 = OpLabel
+%61 = OpVariable  %25  Function %60
+OpBranch %62
+%62 = OpLabel
+OpReturn
+OpFunctionEnd
+%64 = OpFunction  %2  None %22
+%63 = OpLabel
+%65 = OpVariable  %25  Function %60
+OpBranch %66
+%66 = OpLabel
+OpReturn
+OpFunctionEnd
+%69 = OpFunction  %3  None %70
+%68 = OpFunctionParameter  %4
+%67 = OpLabel
 OpBranch %75
 %75 = OpLabel
-OpStore %72 %74
-OpReturn
-OpFunctionEnd
-%77 = OpFunction  %2  None %28
-%76 = OpLabel
-%78 = OpAccessChain  %29  %20 %30
-OpBranch %79
+OpSelectionMerge %76 None
+OpSwitch %68 %80 0 %77 1 %78 2 %79
+%77 = OpLabel
+OpReturnValue %71
+%78 = OpLabel
+OpReturnValue %72
 %79 = OpLabel
-OpStore %78 %74
-OpReturn
-OpFunctionEnd
-%82 = OpFunction  %5  None %83
-%81 = OpFunctionParameter  %4
+OpReturnValue %73
 %80 = OpLabel
-OpBranch %87
-%87 = OpLabel
-OpSelectionMerge %88 None
-OpSwitch %81 %92 0 %89 1 %90 2 %91
-%89 = OpLabel
-OpReturnValue %84
-%90 = OpLabel
-OpReturnValue %85
-%91 = OpLabel
-OpReturnValue %86
-%92 = OpLabel
-OpReturnValue %30
-%88 = OpLabel
-OpReturnValue %93
+OpReturnValue %74
+%76 = OpLabel
+OpReturnValue %81
 OpFunctionEnd
-%95 = OpFunction  %2  None %28
-%94 = OpLabel
-%96 = OpAccessChain  %29  %20 %30
-%97 = OpAccessChain  %36  %23 %30
-OpBranch %98
-%98 = OpLabel
-%99 = OpFunctionCall  %2  %27
-%100 = OpFunctionCall  %2  %35
-%101 = OpFunctionCall  %2  %42
-%102 = OpFunctionCall  %2  %49
-%103 = OpFunctionCall  %2  %71
-%104 = OpFunctionCall  %2  %77
+%83 = OpFunction  %2  None %22
+%82 = OpLabel
+OpBranch %84
+%84 = OpLabel
+%85 = OpFunctionCall  %2  %21
+%86 = OpFunctionCall  %2  %28
+%87 = OpFunctionCall  %2  %33
+%88 = OpFunctionCall  %2  %38
+%89 = OpFunctionCall  %2  %58
+%90 = OpFunctionCall  %2  %64
 OpReturn
 OpFunctionEnd
--- a/tests/out/wgsl/const-exprs.wgsl
+++ b/tests/out/wgsl/const-exprs.wgsl
@@ -11,26 +11,19 @@ const TEXTURE_KIND_REGULAR: i32 = 0;
 const TEXTURE_KIND_WARP: i32 = 1;
 const TEXTURE_KIND_SKY: i32 = 2;

-@group(0) @binding(0) 
-var<storage, read_write> out: vec4<i32>;
-@group(0) @binding(1) 
-var<storage, read_write> out2_: i32;
-
 fn swizzle_of_compose() {
-    out = vec4<i32>(4, 3, 2, 1);
-    return;
+    var out: vec4<i32> = vec4<i32>(4, 3, 2, 1);
+
 }

 fn index_of_compose() {
-    let _e2 = out2_;
-    out2_ = (_e2 + 2);
-    return;
+    var out_1: i32 = 2;
+
 }

 fn compose_three_deep() {
-    let _e2 = out2_;
-    out2_ = (_e2 + 6);
-    return;
+    var out_2: i32 = 6;
+
 }

 fn non_constant_initializers() {
@@ -38,28 +31,28 @@ fn non_constant_initializers() {
    var x: i32;
    var y: i32;
    var z: i32 = 70;
+    var out_3: vec4<i32>;

    let _e2 = w;
    x = _e2;
    let _e4 = x;
    y = _e4;
-    let _e9 = w;
-    let _e10 = x;
-    let _e11 = y;
-    let _e12 = z;
-    let _e14 = out;
-    out = (_e14 + vec4<i32>(_e9, _e10, _e11, _e12));
+    let _e8 = w;
+    let _e9 = x;
+    let _e10 = y;
+    let _e11 = z;
+    out_3 = vec4<i32>(_e8, _e9, _e10, _e11);
    return;
 }

 fn splat_of_constant() {
-    out = vec4<i32>(-4, -4, -4, -4);
-    return;
+    var out_4: vec4<i32> = vec4<i32>(-4, -4, -4, -4);
+
 }

 fn compose_of_constant() {
-    out = vec4<i32>(-4, -4, -4, -4);
-    return;
+    var out_5: vec4<i32> = vec4<i32>(-4, -4, -4, -4);
+
 }

 fn map_texture_kind(texture_kind: i32) -> u32 {