writer/hlsl: Fix continuing block emission

Inline the `continuing` block in the places where `continue` is called.

Simplifies the emission, and fixes emission of `let` statements in the loop.

Also fix random indenting of intrinsic functions.

Fixed: tint:744
Fixed: tint:818
Change-Id: I06994dbc724bc646e0435a1035b00760eaf5f5ab
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51784
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>

src/writer/hlsl/generator_impl.cc

@@ -34,6 +34,7 @@
 #include "src/sem/variable.h"
 #include "src/transform/calculate_array_length.h"
 #include "src/transform/decompose_storage_access.h"
+#include "src/utils/scoped_assignment.h"
 #include "src/writer/append_vector.h"
 #include "src/writer/float_to_string.h"
 
@@ -422,20 +423,14 @@ bool GeneratorImpl::EmitBinary(std::ostream& pre,
 
 bool GeneratorImpl::EmitBlock(std::ostream& out,
                               const ast::BlockStatement* stmt) {
-  out << "{" << std::endl;
-  increment_indent();
-
-  for (auto* s : *stmt) {
-    if (!EmitStatement(out, s)) {
-      return false;
+  return EmitBlockBraces(out, [&] {
+    for (auto* s : *stmt) {
+      if (!EmitStatement(out, s)) {
+        return false;
+      }
     }
-  }
-
-  decrement_indent();
-  make_indent(out);
-  out << "}";
-
-  return true;
+    return true;
+  });
 }
 
 bool GeneratorImpl::EmitBlockAndNewline(std::ostream& out,
@@ -614,7 +609,6 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
       return false;
     }
 
-    make_indent(out);
     out << name << "(";
 
     bool first = true;
@@ -1338,6 +1332,9 @@ bool GeneratorImpl::EmitTypeConstructor(std::ostream& pre,
 }
 
 bool GeneratorImpl::EmitContinue(std::ostream& out, ast::ContinueStatement*) {
+  if (!emit_continuing_(out)) {
+    return false;
+  }
   make_indent(out);
   out << "continue;" << std::endl;
   return true;
@@ -2192,100 +2189,30 @@ bool GeneratorImpl::EmitZeroValue(std::ostream& out, const sem::Type* type) {
 }
 
 bool GeneratorImpl::EmitLoop(std::ostream& out, ast::LoopStatement* stmt) {
-  loop_emission_counter_++;
+  make_indent(out);
 
-  std::string guard = generate_name("tint_hlsl_is_first_" +
-                                    std::to_string(loop_emission_counter_));
+  auto emit_continuing = [this, stmt](std::ostream& o) {
+    if (stmt->has_continuing()) {
+      make_indent(o);
+      if (!EmitBlock(o, stmt->continuing())) {
+        return false;
+      }
+      o << std::endl;
+    }
+    return true;
+  };
 
-  if (stmt->has_continuing()) {
-    make_indent(out);
-
-    // Continuing variables get their own scope.
-    out << "{" << std::endl;
-    increment_indent();
-
-    make_indent(out);
-    out << "bool " << guard << " = true;" << std::endl;
-
-    // A continuing block may use variables declared in the method body. As a
-    // first pass, if we have a continuing, we pull all declarations outside
-    // the for loop into the continuing scope. Then, the variable declarations
-    // will be turned into assignments.
-    for (auto* s : *stmt->body()) {
-      if (auto* v = s->As<ast::VariableDeclStatement>()) {
-        if (!EmitVariable(out, v->variable(), true)) {
-          return false;
-        }
+  TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
+  bool ok = EmitBlockBraces(out, "while (true)", [&] {
+    for (auto* s : stmt->body()->statements()) {
+      if (!EmitStatement(out, s)) {
+        return false;
       }
     }
-  }
-
-  make_indent(out);
-  out << "for(;;) {" << std::endl;
-  increment_indent();
-
-  if (stmt->has_continuing()) {
-    make_indent(out);
-    out << "if (!" << guard << ") ";
-
-    if (!EmitBlockAndNewline(out, stmt->continuing())) {
-      return false;
-    }
-
-    make_indent(out);
-    out << guard << " = false;" << std::endl;
-    out << std::endl;
-  }
-
-  for (auto* s : *(stmt->body())) {
-    // If we have a continuing block we've already emitted the variable
-    // declaration before the loop, so treat it as an assignment.
-    if (auto* decl = s->As<ast::VariableDeclStatement>()) {
-      if (stmt->has_continuing()) {
-        make_indent(out);
-
-        auto* var = decl->variable();
-
-        std::ostringstream pre;
-        std::ostringstream constructor_out;
-        if (var->constructor() != nullptr) {
-          if (!EmitExpression(pre, constructor_out, var->constructor())) {
-            return false;
-          }
-        }
-        out << pre.str();
-
-        out << builder_.Symbols().NameFor(var->symbol()) << " = ";
-        if (var->constructor() != nullptr) {
-          out << constructor_out.str();
-        } else {
-          auto* type = builder_.Sem().Get(var)->Type()->UnwrapRef();
-          if (!EmitZeroValue(out, type)) {
-            return false;
-          }
-        }
-        out << ";" << std::endl;
-        continue;
-      }
-    }
-
-    if (!EmitStatement(out, s)) {
-      return false;
-    }
-  }
-
-  decrement_indent();
-  make_indent(out);
-  out << "}" << std::endl;
-
-  // Close the scope for any continuing variables.
-  if (stmt->has_continuing()) {
-    decrement_indent();
-    make_indent(out);
-    out << "}" << std::endl;
-  }
-
-  return true;
+    return emit_continuing(out);
+  });
+  out << std::endl;
+  return ok;
 }
 
 bool GeneratorImpl::EmitMemberAccessor(std::ostream& pre,
@@ -2379,7 +2306,7 @@ bool GeneratorImpl::EmitStatement(std::ostream& out, ast::Statement* stmt) {
     return EmitSwitch(out, s);
   }
   if (auto* v = stmt->As<ast::VariableDeclStatement>()) {
-    return EmitVariable(out, v->variable(), false);
+    return EmitVariable(out, v->variable());
   }
 
   diagnostics_.add_error("unknown statement type: " + builder_.str(stmt));
@@ -2662,9 +2589,7 @@ bool GeneratorImpl::EmitUnaryOp(std::ostream& pre,
   return true;
 }
 
-bool GeneratorImpl::EmitVariable(std::ostream& out,
-                                 ast::Variable* var,
-                                 bool skip_constructor) {
+bool GeneratorImpl::EmitVariable(std::ostream& out, ast::Variable* var) {
   make_indent(out);
 
   auto* sem = builder_.Sem().Get(var);
@@ -2677,19 +2602,17 @@ bool GeneratorImpl::EmitVariable(std::ostream& out,
   }
 
   std::ostringstream constructor_out;
-  if (!skip_constructor) {
-    constructor_out << " = ";
+  constructor_out << " = ";
 
-    if (var->constructor()) {
-      std::ostringstream pre;
-      if (!EmitExpression(pre, constructor_out, var->constructor())) {
-        return false;
-      }
-      out << pre.str();
-    } else {
-      if (!EmitZeroValue(constructor_out, type)) {
-        return false;
-      }
+  if (var->constructor()) {
+    std::ostringstream pre;
+    if (!EmitExpression(pre, constructor_out, var->constructor())) {
+      return false;
+    }
+    out << pre.str();
+  } else {
+    if (!EmitZeroValue(constructor_out, type)) {
+      return false;
     }
   }
 
@@ -2789,6 +2712,23 @@ std::string GeneratorImpl::get_buffer_name(ast::Expression* expr) {
   return "";
 }
 
+template <typename F>
+bool GeneratorImpl::EmitBlockBraces(std::ostream& out,
+                                    const std::string& prefix,
+                                    F&& cb) {
+  out << prefix << (prefix.empty() ? "{" : " {") << std::endl;
+  increment_indent();
+
+  if (!cb()) {
+    return false;
+  }
+
+  decrement_indent();
+  make_indent(out);
+  out << "}";
+  return true;
+}
+
 }  // namespace hlsl
 }  // namespace writer
 }  // namespace tint

src/writer/hlsl/generator_impl.h

@@ -18,6 +18,7 @@
 #include <string>
 #include <unordered_map>
 #include <unordered_set>
+#include <utility>
 
 #include "src/ast/assignment_statement.h"
 #include "src/ast/bitcast_expression.h"
@@ -319,11 +320,8 @@ class GeneratorImpl : public TextGenerator {
   /// Handles generating a variable
   /// @param out the output stream
   /// @param var the variable to generate
-  /// @param skip_constructor set true if the constructor should be skipped
   /// @returns true if the variable was emitted
-  bool EmitVariable(std::ostream& out,
-                    ast::Variable* var,
-                    bool skip_constructor);
+  bool EmitVariable(std::ostream& out, ast::Variable* var);
   /// Handles generating a program scope constant variable
   /// @param out the output stream
   /// @param var the variable to emit
@@ -387,10 +385,31 @@ class GeneratorImpl : public TextGenerator {
     return builder_.TypeOf(type);
   }
 
+  /// Emits `prefix`, followed by an opening brace `{`, then calls `cb` to emit
+  /// the block body, then finally emits the closing brace `}`.
+  /// @param out the output stream
+  /// @param prefix the string to emit before the opening brace
+  /// @param cb a function or function-like object with the signature `bool()`
+  /// that emits the block body.
+  /// @returns the return value of `cb`.
+  template <typename F>
+  bool EmitBlockBraces(std::ostream& out, const std::string& prefix, F&& cb);
+
+  /// Emits an opening brace `{`, then calls `cb` to emit the block body, then
+  /// finally emits the closing brace `}`.
+  /// @param out the output stream
+  /// @param cb a function or function-like object with the signature `bool()`
+  /// that emits the block body.
+  /// @returns the return value of `cb`.
+  template <typename F>
+  bool EmitBlockBraces(std::ostream& out, F&& cb) {
+    return EmitBlockBraces(out, "", std::forward<F>(cb));
+  }
+
   ProgramBuilder builder_;
   Symbol current_ep_sym_;
   bool generating_entry_point_ = false;
-  uint32_t loop_emission_counter_ = 0;
+  std::function<bool(std::ostream& out)> emit_continuing_;
   ScopeStack<const sem::Variable*> global_variables_;
   std::unordered_map<Symbol, EntryPointData> ep_sym_to_in_data_;
   std::unordered_map<Symbol, EntryPointData> ep_sym_to_out_data_;

src/writer/hlsl/generator_impl_continue_test.cc

@@ -22,15 +22,18 @@ namespace {
 using HlslGeneratorImplTest_Continue = TestHelper;
 
 TEST_F(HlslGeneratorImplTest_Continue, Emit_Continue) {
-  auto* c = create<ast::ContinueStatement>();
-  WrapInFunction(Loop(Block(c)));
+  auto* loop = Loop(Block(create<ast::ContinueStatement>()));
+  WrapInFunction(loop);
 
   GeneratorImpl& gen = Build();
 
   gen.increment_indent();
 
-  ASSERT_TRUE(gen.EmitStatement(out, c)) << gen.error();
-  EXPECT_EQ(result(), "  continue;\n");
+  ASSERT_TRUE(gen.EmitStatement(out, loop)) << gen.error();
+  EXPECT_EQ(result(), R"(  while (true) {
+    continue;
+  }
+)");
 }
 
 }  // namespace

src/writer/hlsl/generator_impl_intrinsic_test.cc

@@ -268,7 +268,7 @@ TEST_F(HlslGeneratorImplTest_Intrinsic, Intrinsic_Call) {
 
   gen.increment_indent();
   ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
-  EXPECT_EQ(result(), "  dot(param1, param2)");
+  EXPECT_EQ(result(), "dot(param1, param2)");
 }
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Pack4x8Snorm) {

src/writer/hlsl/generator_impl_loop_test.cc

@@ -34,7 +34,7 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_Loop) {
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(out, l)) << gen.error();
-  EXPECT_EQ(result(), R"(  for(;;) {
+  EXPECT_EQ(result(), R"(  while (true) {
     discard;
   }
 )");
@@ -52,15 +52,10 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_LoopWithContinuing) {
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(out, l)) << gen.error();
-  EXPECT_EQ(result(), R"(  {
-    bool tint_hlsl_is_first_1 = true;
-    for(;;) {
-      if (!tint_hlsl_is_first_1) {
-        return;
-      }
-      tint_hlsl_is_first_1 = false;
-
-      discard;
+  EXPECT_EQ(result(), R"(  while (true) {
+    discard;
+    {
+      return;
     }
   }
 )");
@@ -89,26 +84,16 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_LoopNestedWithContinuing) {
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(out, outer)) << gen.error();
-  EXPECT_EQ(result(), R"(  {
-    bool tint_hlsl_is_first_1 = true;
-    for(;;) {
-      if (!tint_hlsl_is_first_1) {
-        lhs = rhs;
-      }
-      tint_hlsl_is_first_1 = false;
-
+  EXPECT_EQ(result(), R"(  while (true) {
+    while (true) {
+      discard;
       {
-        bool tint_hlsl_is_first_2 = true;
-        for(;;) {
-          if (!tint_hlsl_is_first_2) {
-            return;
-          }
-          tint_hlsl_is_first_2 = false;
-
-          discard;
-        }
+        return;
       }
     }
+    {
+      lhs = rhs;
+    }
   }
 )");
 }
@@ -153,18 +138,11 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_LoopWithVarUsedInContinuing) {
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(out, outer)) << gen.error();
-  EXPECT_EQ(result(), R"(  {
-    bool tint_hlsl_is_first_1 = true;
-    float lhs;
-    float other;
-    for(;;) {
-      if (!tint_hlsl_is_first_1) {
-        lhs = rhs;
-      }
-      tint_hlsl_is_first_1 = false;
-
-      lhs = 2.400000095f;
-      other = 0.0f;
+  EXPECT_EQ(result(), R"(  while (true) {
+    float lhs = 2.400000095f;
+    float other = 0.0f;
+    {
+      lhs = rhs;
     }
   }
 )");

test/bug/tint/744.wgsl

@@ -0,0 +1,29 @@
+[[block]] struct Uniforms {
+    aShape : vec2<u32>;
+    bShape : vec2<u32>;
+    outShape : vec2<u32>;
+};
+[[block]] struct Matrix {
+    numbers: array<u32>;
+};
+
+[[group(0), binding(0)]] var<storage> firstMatrix : [[access(read)]] Matrix;
+[[group(0), binding(1)]] var<storage> secondMatrix : [[access(read)]] Matrix;
+[[group(0), binding(2)]] var<storage> resultMatrix : [[access(write)]] Matrix;
+[[group(0), binding(3)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(2,2,1)]]
+fn main([[builtin(global_invocation_id)]] global_id  : vec3<u32>) {
+    let resultCell : vec2<u32> = vec2<u32>(global_id.y, global_id.x);
+    let dimInner : u32 = uniforms.aShape.y;
+    let dimOutter: u32 = uniforms.outShape.y;
+    var result : u32 = 0u;
+    for (var i : u32 = 0u; i < dimInner; i = i + 1u) {
+        let a : u32 = i + resultCell.x * dimInner;
+        let b : u32 = resultCell.y + i * dimOutter;
+        result = result + firstMatrix.numbers[a] * secondMatrix.numbers[b];
+    }
+
+    let index : u32 = resultCell.y + resultCell.x * dimOutter;
+    resultMatrix.numbers[index] = result;
+}

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

@@ -0,0 +1,41 @@
+struct Uniforms {
+  uint2 aShape;
+  uint2 bShape;
+  uint2 outShape;
+};
+struct tint_symbol_1 {
+  uint3 global_id : SV_DispatchThreadID;
+};
+
+ConstantBuffer<Uniforms> uniforms : register(b3, space0);
+
+ByteAddressBuffer firstMatrix : register(t0, space0);
+ByteAddressBuffer secondMatrix : register(t1, space0);
+RWByteAddressBuffer resultMatrix : register(u2, space0);
+
+[numthreads(2, 2, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint3 global_id = tint_symbol.global_id;
+  const uint2 resultCell = uint2(global_id.y, global_id.x);
+  const uint dimInner = uniforms.aShape.y;
+  const uint dimOutter = uniforms.outShape.y;
+  uint result = 0u;
+  {
+    uint i = 0u;
+    while (true) {
+      if (!((i < dimInner))) {
+        break;
+      }
+      const uint a = (i + (resultCell.x * dimInner));
+      const uint b = (resultCell.y + (i * dimOutter));
+      result = (result + (firstMatrix.Load((4u * a)) * secondMatrix.Load((4u * b))));
+      {
+        i = (i + 1u);
+      }
+    }
+  }
+  const uint index = (resultCell.y + (resultCell.x * dimOutter));
+  resultMatrix.Store((4u * index), asuint(result));
+  return;
+}
+

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

@@ -0,0 +1,47 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ packed_uint2 aShape;
+  /* 0x0008 */ packed_uint2 bShape;
+  /* 0x0010 */ packed_uint2 outShape;
+};
+struct Matrix {
+  /* 0x0000 */ uint numbers[1];
+};
+struct tint_symbol_2 {
+  uint3 global_id [[thread_position_in_grid]];
+};
+
+kernel void tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) {
+  uint3 const global_id = tint_symbol_1.global_id;
+  uint2 const resultCell = uint2(global_id.y, global_id.x);
+  uint const dimInner = uniforms.aShape.y;
+  uint const dimOutter = uniforms.outShape.y;
+  uint result = 0u;
+  {
+    uint i = 0u;
+    {
+      bool tint_msl_is_first_1 = true;
+      uint const a;
+      uint const b;
+      for(;;) {
+        if (!tint_msl_is_first_1) {
+          i = (i + 1u);
+        }
+        tint_msl_is_first_1 = false;
+
+        if (!((i < dimInner))) {
+          break;
+        }
+        a = (i + (resultCell.x * dimInner));
+        b = (resultCell.y + (i * dimOutter));
+        result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
+      }
+    }
+  }
+  uint const index = (resultCell.y + (resultCell.x * dimOutter));
+  resultMatrix.numbers[index] = result;
+  return;
+}
+

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

@@ -0,0 +1,127 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 71
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol
+               OpExecutionMode %main LocalSize 2 2 1
+               OpName %Matrix "Matrix"
+               OpMemberName %Matrix 0 "numbers"
+               OpName %firstMatrix "firstMatrix"
+               OpName %secondMatrix "secondMatrix"
+               OpName %resultMatrix "resultMatrix"
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "aShape"
+               OpMemberName %Uniforms 1 "bShape"
+               OpMemberName %Uniforms 2 "outShape"
+               OpName %uniforms "uniforms"
+               OpName %tint_symbol "tint_symbol"
+               OpName %main "main"
+               OpName %result "result"
+               OpName %i "i"
+               OpDecorate %Matrix Block
+               OpMemberDecorate %Matrix 0 Offset 0
+               OpDecorate %_runtimearr_uint ArrayStride 4
+               OpDecorate %firstMatrix NonWritable
+               OpDecorate %firstMatrix DescriptorSet 0
+               OpDecorate %firstMatrix Binding 0
+               OpDecorate %secondMatrix NonWritable
+               OpDecorate %secondMatrix DescriptorSet 0
+               OpDecorate %secondMatrix Binding 1
+               OpDecorate %resultMatrix NonReadable
+               OpDecorate %resultMatrix DescriptorSet 0
+               OpDecorate %resultMatrix Binding 2
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 8
+               OpMemberDecorate %Uniforms 2 Offset 16
+               OpDecorate %uniforms DescriptorSet 0
+               OpDecorate %uniforms Binding 3
+               OpDecorate %tint_symbol BuiltIn GlobalInvocationId
+       %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+     %Matrix = OpTypeStruct %_runtimearr_uint
+%_ptr_StorageBuffer_Matrix = OpTypePointer StorageBuffer %Matrix
+%firstMatrix = OpVariable %_ptr_StorageBuffer_Matrix StorageBuffer
+%secondMatrix = OpVariable %_ptr_StorageBuffer_Matrix StorageBuffer
+%resultMatrix = OpVariable %_ptr_StorageBuffer_Matrix StorageBuffer
+     %v2uint = OpTypeVector %uint 2
+   %Uniforms = OpTypeStruct %v2uint %v2uint %v2uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%tint_symbol = OpVariable %_ptr_Input_v3uint Input
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+     %uint_1 = OpConstant %uint 1
+%_ptr_Input_uint = OpTypePointer Input %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_2 = OpConstant %uint 2
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %35 = OpConstantNull %uint
+       %bool = OpTypeBool
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+       %main = OpFunction %void None %15
+         %18 = OpLabel
+     %result = OpVariable %_ptr_Function_uint Function %35
+          %i = OpVariable %_ptr_Function_uint Function %35
+         %21 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_1
+         %22 = OpLoad %uint %21
+         %24 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_0
+         %25 = OpLoad %uint %24
+         %26 = OpCompositeConstruct %v2uint %22 %25
+         %28 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
+         %29 = OpLoad %uint %28
+         %31 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_2 %uint_1
+         %32 = OpLoad %uint %31
+               OpStore %result %uint_0
+               OpStore %i %uint_0
+               OpBranch %37
+         %37 = OpLabel
+               OpLoopMerge %38 %39 None
+               OpBranch %40
+         %40 = OpLabel
+         %42 = OpLoad %uint %i
+         %43 = OpULessThan %bool %42 %29
+         %41 = OpLogicalNot %bool %43
+               OpSelectionMerge %45 None
+               OpBranchConditional %41 %46 %45
+         %46 = OpLabel
+               OpBranch %38
+         %45 = OpLabel
+         %47 = OpLoad %uint %i
+         %48 = OpCompositeExtract %uint %26 0
+         %49 = OpIMul %uint %48 %29
+         %50 = OpIAdd %uint %47 %49
+         %51 = OpCompositeExtract %uint %26 1
+         %52 = OpLoad %uint %i
+         %53 = OpIMul %uint %52 %32
+         %54 = OpIAdd %uint %51 %53
+         %55 = OpLoad %uint %result
+         %57 = OpAccessChain %_ptr_StorageBuffer_uint %firstMatrix %uint_0 %50
+         %58 = OpLoad %uint %57
+         %59 = OpAccessChain %_ptr_StorageBuffer_uint %secondMatrix %uint_0 %54
+         %60 = OpLoad %uint %59
+         %61 = OpIMul %uint %58 %60
+         %62 = OpIAdd %uint %55 %61
+               OpStore %result %62
+               OpBranch %39
+         %39 = OpLabel
+         %63 = OpLoad %uint %i
+         %64 = OpIAdd %uint %63 %uint_1
+               OpStore %i %64
+               OpBranch %37
+         %38 = OpLabel
+         %65 = OpCompositeExtract %uint %26 1
+         %66 = OpCompositeExtract %uint %26 0
+         %67 = OpIMul %uint %66 %32
+         %68 = OpIAdd %uint %65 %67
+         %69 = OpAccessChain %_ptr_StorageBuffer_uint %resultMatrix %uint_0 %68
+         %70 = OpLoad %uint %result
+               OpStore %69 %70
+               OpReturn
+               OpFunctionEnd

test/bug/tint/744.wgsl.expected.wgsl

@@ -0,0 +1,44 @@
+[[block]]
+struct Uniforms {
+  aShape : vec2<u32>;
+  bShape : vec2<u32>;
+  outShape : vec2<u32>;
+};
+
+[[block]]
+struct Matrix {
+  numbers : array<u32>;
+};
+
+[[group(0), binding(0)]] var<storage> firstMatrix : [[access(read)]] Matrix;
+
+[[group(0), binding(1)]] var<storage> secondMatrix : [[access(read)]] Matrix;
+
+[[group(0), binding(2)]] var<storage> resultMatrix : [[access(write)]] Matrix;
+
+[[group(0), binding(3)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(2, 2, 1)]]
+fn main([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
+  let resultCell : vec2<u32> = vec2<u32>(global_id.y, global_id.x);
+  let dimInner : u32 = uniforms.aShape.y;
+  let dimOutter : u32 = uniforms.outShape.y;
+  var result : u32 = 0u;
+  {
+    var i : u32 = 0u;
+    loop {
+      if (!((i < dimInner))) {
+        break;
+      }
+      let a : u32 = (i + (resultCell.x * dimInner));
+      let b : u32 = (resultCell.y + (i * dimOutter));
+      result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
+
+      continuing {
+        i = (i + 1u);
+      }
+    }
+  }
+  let index : u32 = (resultCell.y + (resultCell.x * dimOutter));
+  resultMatrix.numbers[index] = result;
+}

test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl

@@ -1 +1,7 @@
-SKIP: Failed to generate: error: pointers not supported in HLSL
+[numthreads(1, 1, 1)]
+void main() {
+  int i = 123;
+  const int use = (i + 1);
+  return;
+}
+

test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl

@@ -1 +1,8 @@
-SKIP: Failed to generate: error: pointers not supported in HLSL
+static int i = 123;
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int use = (i + 1);
+  return;
+}
+

test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl

@@ -1 +1,12 @@
-SKIP: Failed to generate: error: pointers not supported in HLSL
+struct S {
+  int a;
+};
+
+ConstantBuffer<S> v : register(b0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int use = (v.a + 1);
+  return;
+}
+

test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl

@@ -1 +1,9 @@
-SKIP: Failed to generate: error: pointers not supported in HLSL
+groupshared int i;
+
+[numthreads(1, 1, 1)]
+void main() {
+  i = 123;
+  const int use = (i + 1);
+  return;
+}
+

test/samples/compute_boids.wgsl.expected.hlsl

@@ -31,8 +31,8 @@ tint_symbol_2 vert_main(tint_symbol_1 tint_symbol) {
   const float2 a_particlePos = tint_symbol.a_particlePos;
   const float2 a_particleVel = tint_symbol.a_particleVel;
   const float2 a_pos = tint_symbol.a_pos;
-  float angle = -(  atan2(a_particleVel.x, a_particleVel.y));
-  float2 pos = float2(((a_pos.x *   cos(angle)) - (a_pos.y *   sin(angle))), ((a_pos.x *   sin(angle)) + (a_pos.y *   cos(angle))));
+  float angle = -(atan2(a_particleVel.x, a_particleVel.y));
+  float2 pos = float2(((a_pos.x * cos(angle)) - (a_pos.y * sin(angle))), ((a_pos.x * sin(angle)) + (a_pos.y * cos(angle))));
   const tint_symbol_2 tint_symbol_8 = {float4((pos + a_particlePos), 0.0f, 1.0f)};
   return tint_symbol_8;
 }
@@ -60,33 +60,31 @@ void comp_main(tint_symbol_5 tint_symbol_4) {
   float2 vel = float2(0.0f, 0.0f);
   {
     uint i = 0u;
-    {
-      bool tint_hlsl_is_first_1 = true;
-      for(;;) {
-        if (!tint_hlsl_is_first_1) {
+    while (true) {
+      if (!((i < 5u))) {
+        break;
+      }
+      if ((i == index)) {
+        {
           i = (i + 1u);
         }
-        tint_hlsl_is_first_1 = false;
-
-        if (!((i < 5u))) {
-          break;
-        }
-        if ((i == index)) {
-          continue;
-        }
-        pos = asfloat(particlesA.Load2((16u * i))).xy;
-        vel = asfloat(particlesA.Load2(((16u * i) + 8u))).xy;
-        if ((        distance(pos, vPos) < params.rule1Distance)) {
-          cMass = (cMass + pos);
-          cMassCount = (cMassCount + 1);
-        }
-        if ((        distance(pos, vPos) < params.rule2Distance)) {
-          colVel = (colVel - (pos - vPos));
-        }
-        if ((        distance(pos, vPos) < params.rule3Distance)) {
-          cVel = (cVel + vel);
-          cVelCount = (cVelCount + 1);
-        }
+        continue;
+      }
+      pos = asfloat(particlesA.Load2((16u * i))).xy;
+      vel = asfloat(particlesA.Load2(((16u * i) + 8u))).xy;
+      if ((distance(pos, vPos) < params.rule1Distance)) {
+        cMass = (cMass + pos);
+        cMassCount = (cMassCount + 1);
+      }
+      if ((distance(pos, vPos) < params.rule2Distance)) {
+        colVel = (colVel - (pos - vPos));
+      }
+      if ((distance(pos, vPos) < params.rule3Distance)) {
+        cVel = (cVel + vel);
+        cVelCount = (cVelCount + 1);
+      }
+      {
+        i = (i + 1u);
       }
     }
   }
@@ -97,7 +95,7 @@ void comp_main(tint_symbol_5 tint_symbol_4) {
     cVel = (cVel / float2(float(cVelCount), float(cVelCount)));
   }
   vVel = (((vVel + (cMass * params.rule1Scale)) + (colVel * params.rule2Scale)) + (cVel * params.rule3Scale));
-  vVel = (  normalize(vVel) *   clamp(  length(vVel), 0.0f, 0.100000001f));
+  vVel = (normalize(vVel) * clamp(length(vVel), 0.0f, 0.100000001f));
   vPos = (vPos + (vVel * params.deltaT));
   if ((vPos.x < -1.0f)) {
     vPos.x = 1.0f;