writer/hlsl: Use vector write helper for dynamic indices

This uses FXC compilation failure mitigation for _any_ vector index assignment that has a non-constant index. FXC can still fall over if the loop calls a function that performs the dynamic index.

Use some vector swizzle logic to avoid branches in the helper.

Fixed: tint:980
Change-Id: I2a759d88a7d884bc61b4631cf57feb4acc8178de
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57882
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>

test/bug/tint/534.wgsl.expected.hlsl

@@ -1,10 +1,5 @@
-void Set_uint4(inout uint4 vec, int idx, uint val) {
-  switch(idx) {
-    case 0: vec[0] = val; break;
-    case 1: vec[1] = val; break;
-    case 2: vec[2] = val; break;
-    case 3: vec[3] = val; break;
-  }
+void set_uint4(inout uint4 vec, int idx, uint val) {
+  vec = (idx.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : vec;
 }
 
 Texture2D<float4> src : register(t0, space0);
@@ -40,7 +35,7 @@ void main(tint_symbol_2 tint_symbol_1) {
   uint4 dstColorBits = uint4(dstColor);
   {
     for(uint i = 0u; (i < uniforms[0].w); i = (i + 1u)) {
-      Set_uint4(srcColorBits, i, ConvertToFp16FloatValue(srcColor[i]));
+      set_uint4(srcColorBits, i, ConvertToFp16FloatValue(srcColor[i]));
       bool tint_tmp_1 = success;
       if (tint_tmp_1) {
         tint_tmp_1 = (srcColorBits[i] == dstColorBits[i]);

test/bug/tint/980.wgsl

@@ -0,0 +1,13 @@
+// Fails with "D3D compile failed with value cannot be NaN, isnan() may not be necessary.  /Gis may force isnan() to be performed"
+fn Bad(index: u32, rd: vec3<f32>) -> vec3<f32> {
+  var normal: vec3<f32> = vec3<f32>(0.0);
+  normal[index] = -sign(rd[index]);
+  return normalize(normal);
+}
+
+[[block]] struct S { v : vec3<f32>; i : u32; };
+[[binding(0), group(0)]] var<storage, read_write> io : S;
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+    io.v = Bad(io.i, io.v);
+}

test/bug/tint/980.wgsl.expected.hlsl

@@ -0,0 +1,22 @@
+void set_float3(inout float3 vec, int idx, float val) {
+  vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec;
+}
+
+float3 Bad(uint index, float3 rd) {
+  float3 normal = float3((0.0f).xxx);
+  set_float3(normal, index, -(sign(rd[index])));
+  return normalize(normal);
+}
+
+RWByteAddressBuffer io : register(u0, space0);
+
+struct tint_symbol_1 {
+  uint idx : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint idx = tint_symbol.idx;
+  io.Store3(0u, asuint(Bad(io.Load(12u), asfloat(io.Load3(0u)))));
+  return;
+}

test/bug/tint/980.wgsl.expected.msl

@@ -0,0 +1,19 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  /* 0x0000 */ packed_float3 v;
+  /* 0x000c */ uint i;
+};
+
+float3 Bad(uint index, float3 rd) {
+  float3 normal = float3(0.0f);
+  normal[index] = -(sign(rd[index]));
+  return normalize(normal);
+}
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], device S& io [[buffer(0)]]) {
+  io.v = Bad(io.i, io.v);
+  return;
+}
+

test/bug/tint/980.wgsl.expected.spvasm

@@ -0,0 +1,72 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 41
+; Schema: 0
+               OpCapability Shader
+         %23 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "v"
+               OpMemberName %S 1 "i"
+               OpName %io "io"
+               OpName %tint_symbol "tint_symbol"
+               OpName %Bad "Bad"
+               OpName %index "index"
+               OpName %rd "rd"
+               OpName %normal "normal"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 12
+               OpDecorate %io Binding 0
+               OpDecorate %io DescriptorSet 0
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+       %uint = OpTypeInt 32 0
+          %S = OpTypeStruct %v3float %uint
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+         %io = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+          %9 = OpTypeFunction %v3float %uint %v3float
+    %float_0 = OpConstant %float 0
+         %15 = OpConstantComposite %v3float %float_0 %float_0 %float_0
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %18 = OpConstantNull %v3float
+%_ptr_Function_float = OpTypePointer Function %float
+       %void = OpTypeVoid
+         %27 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+     %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+        %Bad = OpFunction %v3float None %9
+      %index = OpFunctionParameter %uint
+         %rd = OpFunctionParameter %v3float
+         %13 = OpLabel
+     %normal = OpVariable %_ptr_Function_v3float Function %18
+               OpStore %normal %15
+         %20 = OpAccessChain %_ptr_Function_float %normal %index
+         %24 = OpVectorExtractDynamic %float %rd %index
+         %22 = OpExtInst %float %23 FSign %24
+         %21 = OpFNegate %float %22
+               OpStore %20 %21
+         %26 = OpLoad %v3float %normal
+         %25 = OpExtInst %v3float %23 Normalize %26
+               OpReturnValue %25
+               OpFunctionEnd
+       %main = OpFunction %void None %27
+         %30 = OpLabel
+         %33 = OpAccessChain %_ptr_StorageBuffer_v3float %io %uint_0
+         %37 = OpAccessChain %_ptr_StorageBuffer_uint %io %uint_1
+         %38 = OpLoad %uint %37
+         %39 = OpAccessChain %_ptr_StorageBuffer_v3float %io %uint_0
+         %40 = OpLoad %v3float %39
+         %34 = OpFunctionCall %v3float %Bad %38 %40
+               OpStore %33 %34
+               OpReturn
+               OpFunctionEnd

test/bug/tint/980.wgsl.expected.wgsl

@@ -0,0 +1,18 @@
+fn Bad(index : u32, rd : vec3<f32>) -> vec3<f32> {
+  var normal : vec3<f32> = vec3<f32>(0.0);
+  normal[index] = -(sign(rd[index]));
+  return normalize(normal);
+}
+
+[[block]]
+struct S {
+  v : vec3<f32>;
+  i : u32;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> io : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+  io.v = Bad(io.i, io.v);
+}