transform/shader_io: Generate a wrapper function

This is a major reworking of this transform. The old transform code
was getting unwieldy, with part of the complication coming from the
handling of multiple return statements. By generating a wrapper
function instead, we can avoid a lot of this complexity.

The original entry point function is stripped of all shader IO
attributes (as well as `stage` and `workgroup_size`), but the body is
left unmodified. A new entry point wrapper function is introduced
which calls the original function, packing/unpacking the shader inputs
as necessary, and propagates the result to the corresponding shader
outputs.

The new code has been refactored to use a state object with the
different parts of the transform split into separate functions, which
makes it much more manageable.

Fixed: tint:1076
Bug: tint:920
Change-Id: I3490a0ea7a3509a4e198ce730e476516649d8d96
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60521
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

test/var/inferred/function-let.wgsl.expected.hlsl

@@ -47,7 +47,13 @@ struct tint_symbol {
   float4 value : SV_Target0;
 };
 
-tint_symbol main() {
-  const tint_symbol tint_symbol_3 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_3;
+float4 main_inner() {
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol main() {
+  const float4 inner_result = main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }

test/var/inferred/function-let.wgsl.expected.msl

@@ -51,8 +51,14 @@ void let_decls() {
   tint_array_wrapper const v15 = ret_MyArray();
 }
 
-fragment tint_symbol_1 tint_symbol() {
-  tint_symbol_1 const tint_symbol_4 = {.value=float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_4;
+float4 tint_symbol_inner() {
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+fragment tint_symbol_1 tint_symbol() {
+  float4 const inner_result = tint_symbol_inner();
+  tint_symbol_1 wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 

test/var/inferred/function-var.wgsl.expected.hlsl

@@ -47,7 +47,13 @@ struct tint_symbol {
   float4 value : SV_Target0;
 };
 
-tint_symbol main() {
-  const tint_symbol tint_symbol_3 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_3;
+float4 main_inner() {
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol main() {
+  const float4 inner_result = main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }

test/var/inferred/function-var.wgsl.expected.msl

@@ -51,8 +51,14 @@ void var_decls() {
   tint_array_wrapper v15 = ret_MyArray();
 }
 
-fragment tint_symbol_1 tint_symbol() {
-  tint_symbol_1 const tint_symbol_4 = {.value=float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_4;
+float4 tint_symbol_inner() {
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+fragment tint_symbol_1 tint_symbol() {
+  float4 const inner_result = tint_symbol_inner();
+  tint_symbol_1 wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 

test/var/inferred/global-let.wgsl.expected.hlsl

@@ -16,7 +16,13 @@ struct tint_symbol {
   float4 value : SV_Target0;
 };
 
-tint_symbol main() {
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+float4 main_inner() {
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol main() {
+  const float4 inner_result = main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }

test/var/inferred/global-let.wgsl.expected.msl

@@ -20,8 +20,14 @@ constant float3 v6 = float3(1.0f, 1.0f, 1.0f);
 constant float3x3 v7 = float3x3(float3(1.0f, 1.0f, 1.0f), float3(1.0f, 1.0f, 1.0f), float3(1.0f, 1.0f, 1.0f));
 constant MyStruct v8 = {};
 constant tint_array_wrapper v9 = {.arr={}};
-fragment tint_symbol_1 tint_symbol() {
-  tint_symbol_1 const tint_symbol_2 = {.value=float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_2;
+float4 tint_symbol_inner() {
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+fragment tint_symbol_1 tint_symbol() {
+  float4 const inner_result = tint_symbol_inner();
+  tint_symbol_1 wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 

test/var/initialization/workgroup/array.wgsl.expected.hlsl

@@ -4,9 +4,7 @@ struct tint_symbol_1 {
   uint local_invocation_index : SV_GroupIndex;
 };
 
-[numthreads(1, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint local_invocation_index = tint_symbol.local_invocation_index;
+void main_inner(uint local_invocation_index) {
   {
     for(uint idx = local_invocation_index; (idx < 3u); idx = (idx + 1u)) {
       const uint i = idx;
@@ -15,5 +13,10 @@ void main(tint_symbol_1 tint_symbol) {
   }
   GroupMemoryBarrierWithGroupSync();
   v;
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.local_invocation_index);
   return;
 }

test/var/initialization/workgroup/array.wgsl.expected.msl

@@ -5,14 +5,18 @@ struct tint_array_wrapper {
   int arr[3];
 };
 
-kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  threadgroup tint_array_wrapper tint_symbol_2;
+void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array_wrapper* const tint_symbol_1) {
   for(uint idx = local_invocation_index; (idx < 3u); idx = (idx + 1u)) {
     uint const i = idx;
-    tint_symbol_2.arr[i] = int();
+    (*(tint_symbol_1)).arr[i] = int();
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  (void) tint_symbol_2;
+  (void) *(tint_symbol_1);
+}
+
+kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup tint_array_wrapper tint_symbol_2;
+  tint_symbol_inner(local_invocation_index, &(tint_symbol_2));
   return;
 }
 

test/var/initialization/workgroup/matrix.wgsl.expected.hlsl

@@ -4,13 +4,16 @@ struct tint_symbol_1 {
   uint local_invocation_index : SV_GroupIndex;
 };
 
-[numthreads(1, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint local_invocation_index = tint_symbol.local_invocation_index;
+void main_inner(uint local_invocation_index) {
   {
     v = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
   }
   GroupMemoryBarrierWithGroupSync();
   v;
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.local_invocation_index);
   return;
 }

test/var/initialization/workgroup/scalar.wgsl.expected.hlsl

@@ -4,13 +4,16 @@ struct tint_symbol_1 {
   uint local_invocation_index : SV_GroupIndex;
 };
 
-[numthreads(1, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint local_invocation_index = tint_symbol.local_invocation_index;
+void main_inner(uint local_invocation_index) {
   {
     v = 0;
   }
   GroupMemoryBarrierWithGroupSync();
   v;
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.local_invocation_index);
   return;
 }

test/var/initialization/workgroup/scalar.wgsl.expected.msl

@@ -1,13 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  threadgroup int tint_symbol_2;
+void tint_symbol_inner(uint local_invocation_index, threadgroup int* const tint_symbol_1) {
   {
-    tint_symbol_2 = int();
+    *(tint_symbol_1) = int();
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  (void) tint_symbol_2;
+  (void) *(tint_symbol_1);
+}
+
+kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup int tint_symbol_2;
+  tint_symbol_inner(local_invocation_index, &(tint_symbol_2));
   return;
 }
 

test/var/initialization/workgroup/struct.wgsl.expected.hlsl

@@ -9,14 +9,17 @@ struct tint_symbol_1 {
   uint local_invocation_index : SV_GroupIndex;
 };
 
-[numthreads(1, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint local_invocation_index = tint_symbol.local_invocation_index;
+void main_inner(uint local_invocation_index) {
   {
     const S tint_symbol_2 = (S)0;
     v = tint_symbol_2;
   }
   GroupMemoryBarrierWithGroupSync();
   v;
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.local_invocation_index);
   return;
 }

test/var/initialization/workgroup/struct.wgsl.expected.msl

@@ -6,14 +6,18 @@ struct S {
   float b;
 };
 
-kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  threadgroup S tint_symbol_3;
+void tint_symbol_inner(uint local_invocation_index, threadgroup S* const tint_symbol_2) {
   {
-    S const tint_symbol_2 = {};
-    tint_symbol_3 = tint_symbol_2;
+    S const tint_symbol_1 = {};
+    *(tint_symbol_2) = tint_symbol_1;
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  (void) tint_symbol_3;
+  (void) *(tint_symbol_2);
+}
+
+kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup S tint_symbol_3;
+  tint_symbol_inner(local_invocation_index, &(tint_symbol_3));
   return;
 }
 

test/var/initialization/workgroup/vector.wgsl.expected.hlsl

@@ -4,13 +4,16 @@ struct tint_symbol_1 {
   uint local_invocation_index : SV_GroupIndex;
 };
 
-[numthreads(1, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint local_invocation_index = tint_symbol.local_invocation_index;
+void main_inner(uint local_invocation_index) {
   {
     v = int3(0, 0, 0);
   }
   GroupMemoryBarrierWithGroupSync();
   v;
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.local_invocation_index);
   return;
 }

test/var/initialization/workgroup/vector.wgsl.expected.msl

@@ -1,13 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  threadgroup int3 tint_symbol_2;
+void tint_symbol_inner(uint local_invocation_index, threadgroup int3* const tint_symbol_1) {
   {
-    tint_symbol_2 = int3();
+    *(tint_symbol_1) = int3();
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  (void) tint_symbol_2;
+  (void) *(tint_symbol_1);
+}
+
+kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup int3 tint_symbol_2;
+  tint_symbol_inner(local_invocation_index, &(tint_symbol_2));
   return;
 }
 

test/var/uses/workgroup.wgsl.expected.hlsl

@@ -29,15 +29,18 @@ struct tint_symbol_1 {
   uint local_invocation_index : SV_GroupIndex;
 };
 
-[numthreads(1, 1, 1)]
-void main1(tint_symbol_1 tint_symbol) {
-  const uint local_invocation_index = tint_symbol.local_invocation_index;
+void main1_inner(uint local_invocation_index) {
   {
     a = 0;
   }
   GroupMemoryBarrierWithGroupSync();
   a = 42;
   uses_a();
+}
+
+[numthreads(1, 1, 1)]
+void main1(tint_symbol_1 tint_symbol) {
+  main1_inner(tint_symbol.local_invocation_index);
   return;
 }
 
@@ -45,15 +48,18 @@ struct tint_symbol_3 {
   uint local_invocation_index_1 : SV_GroupIndex;
 };
 
-[numthreads(1, 1, 1)]
-void main2(tint_symbol_3 tint_symbol_2) {
-  const uint local_invocation_index_1 = tint_symbol_2.local_invocation_index_1;
+void main2_inner(uint local_invocation_index_1) {
   {
     b = 0;
   }
   GroupMemoryBarrierWithGroupSync();
   b = 7;
   uses_b();
+}
+
+[numthreads(1, 1, 1)]
+void main2(tint_symbol_3 tint_symbol_2) {
+  main2_inner(tint_symbol_2.local_invocation_index_1);
   return;
 }
 
@@ -61,9 +67,7 @@ struct tint_symbol_5 {
   uint local_invocation_index_2 : SV_GroupIndex;
 };
 
-[numthreads(1, 1, 1)]
-void main3(tint_symbol_5 tint_symbol_4) {
-  const uint local_invocation_index_2 = tint_symbol_4.local_invocation_index_2;
+void main3_inner(uint local_invocation_index_2) {
   {
     a = 0;
     b = 0;
@@ -71,6 +75,11 @@ void main3(tint_symbol_5 tint_symbol_4) {
   GroupMemoryBarrierWithGroupSync();
   outer();
   no_uses();
+}
+
+[numthreads(1, 1, 1)]
+void main3(tint_symbol_5 tint_symbol_4) {
+  main3_inner(tint_symbol_4.local_invocation_index_2);
   return;
 }
 

test/var/uses/workgroup.wgsl.expected.msl

@@ -1,61 +1,73 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void uses_a(threadgroup int* const tint_symbol_3) {
-  *(tint_symbol_3) = as_type<int>((as_type<uint>(*(tint_symbol_3)) + as_type<uint>(1)));
+void uses_a(threadgroup int* const tint_symbol) {
+  *(tint_symbol) = as_type<int>((as_type<uint>(*(tint_symbol)) + as_type<uint>(1)));
 }
 
-void uses_b(threadgroup int* const tint_symbol_4) {
-  *(tint_symbol_4) = as_type<int>((as_type<uint>(*(tint_symbol_4)) * as_type<uint>(2)));
+void uses_b(threadgroup int* const tint_symbol_1) {
+  *(tint_symbol_1) = as_type<int>((as_type<uint>(*(tint_symbol_1)) * as_type<uint>(2)));
 }
 
-void uses_a_and_b(threadgroup int* const tint_symbol_5, threadgroup int* const tint_symbol_6) {
-  *(tint_symbol_5) = *(tint_symbol_6);
+void uses_a_and_b(threadgroup int* const tint_symbol_2, threadgroup int* const tint_symbol_3) {
+  *(tint_symbol_2) = *(tint_symbol_3);
 }
 
 void no_uses() {
 }
 
-void outer(threadgroup int* const tint_symbol_7, threadgroup int* const tint_symbol_8) {
-  *(tint_symbol_7) = 0;
-  uses_a(tint_symbol_7);
-  uses_a_and_b(tint_symbol_8, tint_symbol_7);
-  uses_b(tint_symbol_8);
+void outer(threadgroup int* const tint_symbol_4, threadgroup int* const tint_symbol_5) {
+  *(tint_symbol_4) = 0;
+  uses_a(tint_symbol_4);
+  uses_a_and_b(tint_symbol_5, tint_symbol_4);
+  uses_b(tint_symbol_5);
   no_uses();
 }
 
+void main1_inner(uint local_invocation_index, threadgroup int* const tint_symbol_6) {
+  {
+    *(tint_symbol_6) = int();
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  *(tint_symbol_6) = 42;
+  uses_a(tint_symbol_6);
+}
+
 kernel void main1(uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  threadgroup int tint_symbol_9;
+  threadgroup int tint_symbol_7;
+  main1_inner(local_invocation_index, &(tint_symbol_7));
+  return;
+}
+
+void main2_inner(uint local_invocation_index_1, threadgroup int* const tint_symbol_8) {
   {
-    tint_symbol_9 = int();
+    *(tint_symbol_8) = int();
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  tint_symbol_9 = 42;
-  uses_a(&(tint_symbol_9));
-  return;
+  *(tint_symbol_8) = 7;
+  uses_b(tint_symbol_8);
 }
 
 kernel void main2(uint local_invocation_index_1 [[thread_index_in_threadgroup]]) {
-  threadgroup int tint_symbol_10;
-  {
-    tint_symbol_10 = int();
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  tint_symbol_10 = 7;
-  uses_b(&(tint_symbol_10));
+  threadgroup int tint_symbol_9;
+  main2_inner(local_invocation_index_1, &(tint_symbol_9));
   return;
 }
 
-kernel void main3(uint local_invocation_index_2 [[thread_index_in_threadgroup]]) {
-  threadgroup int tint_symbol_11;
-  threadgroup int tint_symbol_12;
+void main3_inner(uint local_invocation_index_2, threadgroup int* const tint_symbol_10, threadgroup int* const tint_symbol_11) {
   {
-    tint_symbol_11 = int();
-    tint_symbol_12 = int();
+    *(tint_symbol_10) = int();
+    *(tint_symbol_11) = int();
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  outer(&(tint_symbol_11), &(tint_symbol_12));
+  outer(tint_symbol_10, tint_symbol_11);
   no_uses();
+}
+
+kernel void main3(uint local_invocation_index_2 [[thread_index_in_threadgroup]]) {
+  threadgroup int tint_symbol_12;
+  threadgroup int tint_symbol_13;
+  main3_inner(local_invocation_index_2, &(tint_symbol_12), &(tint_symbol_13));
   return;
 }