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

…tected

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;

tests/out/glsl/const-exprs.main.Compute.glsl

@@ -18,54 +18,43 @@ 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() {
     int w = 30;
     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) {

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)

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;
 }