Rename 'intrinsic' to 'builtin'

This matches the term used in the WGSL spec.

Change-Id: I4603332b828450c126ef806f1064ed54f372013f
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/78787
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

test/builtins/arrayLength/complex_via_let.wgsl

@@ -0,0 +1,12 @@
+struct S {
+    a : array<i32>;
+};
+
+@group(0) @binding(0) var<storage, read> G : S;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+    let p = &G;
+    let p2 = &((*p).a);
+    let l1 : u32 = arrayLength(p2);
+}

test/builtins/arrayLength/complex_via_let.wgsl.expected.glsl

@@ -0,0 +1,27 @@
+SKIP: FAILED
+
+#version 310 es
+precision mediump float;
+
+layout(binding = 0) buffer S_1 {
+  int a[];
+} G;
+void tint_symbol() {
+  uint tint_symbol_2 = 0u;
+  G.GetDimensions(tint_symbol_2);
+  uint tint_symbol_3 = ((tint_symbol_2 - 0u) / 4u);
+  uint l1 = tint_symbol_3;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}
+Error parsing GLSL shader:
+ERROR: 0:9: 'GetDimensions' : no such field in structure 
+ERROR: 0:9: '' : compilation terminated 
+ERROR: 2 compilation errors.  No code generated.
+
+
+

test/builtins/arrayLength/complex_via_let.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = ((tint_symbol_1 - 0u) / 4u);
+  const uint l1 = tint_symbol_2;
+  return;
+}

test/builtins/arrayLength/complex_via_let.wgsl.expected.msl

@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+struct S {
+  int a[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
+  return;
+}
+

test/builtins/arrayLength/complex_via_let.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+          %S = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %G = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/builtins/arrayLength/complex_via_let.wgsl.expected.wgsl

@@ -0,0 +1,12 @@
+struct S {
+  a : array<i32>;
+}
+
+@group(0) @binding(0) var<storage, read> G : S;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let p = &(G);
+  let p2 = &((*(p)).a);
+  let l1 : u32 = arrayLength(p2);
+}

test/builtins/arrayLength/complex_via_let_no_struct.wgsl

@@ -0,0 +1,8 @@
+@group(0) @binding(0) var<storage, read> G : array<i32>;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+    let p = &G;
+    let p2 = &(*p);
+    let l1 : u32 = arrayLength(p2);
+}

test/builtins/arrayLength/complex_via_let_no_struct.wgsl.expected.glsl

@@ -0,0 +1,27 @@
+SKIP: FAILED
+
+#version 310 es
+precision mediump float;
+
+layout(binding = 0) buffer G_block_1 {
+  int inner[];
+} G;
+void tint_symbol() {
+  uint tint_symbol_2 = 0u;
+  G.inner.GetDimensions(tint_symbol_2);
+  uint tint_symbol_3 = (tint_symbol_2 / 4u);
+  uint l1 = tint_symbol_3;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}
+Error parsing GLSL shader:
+ERROR: 0:9: '.' : cannot apply to an array: GetDimensions
+ERROR: 0:9: '' : compilation terminated 
+ERROR: 2 compilation errors.  No code generated.
+
+
+

test/builtins/arrayLength/complex_via_let_no_struct.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+  const uint l1 = tint_symbol_2;
+  return;
+}

test/builtins/arrayLength/complex_via_let_no_struct.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+  return;
+}
+

test/builtins/arrayLength/complex_via_let_no_struct.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %G_block "G_block"
+               OpMemberName %G_block 0 "inner"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %G_block Block
+               OpMemberDecorate %G_block 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+    %G_block = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_G_block = OpTypePointer StorageBuffer %G_block
+          %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/builtins/arrayLength/complex_via_let_no_struct.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+@group(0) @binding(0) var<storage, read> G : array<i32>;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let p = &(G);
+  let p2 = &(*(p));
+  let l1 : u32 = arrayLength(p2);
+}

test/builtins/arrayLength/deprecated.wgsl

@@ -0,0 +1,13 @@
+struct S {
+    a : array<i32>;
+};
+
+@group(0) @binding(0) var<storage, read> G : S;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+    let l1 : u32 = arrayLength(&G.a);
+
+    let p = &G.a;
+    let l2 : u32 = arrayLength(p);
+}

test/builtins/arrayLength/deprecated.wgsl.expected.glsl

@@ -0,0 +1,28 @@
+SKIP: FAILED
+
+#version 310 es
+precision mediump float;
+
+layout(binding = 0) buffer S_1 {
+  int a[];
+} G;
+void tint_symbol() {
+  uint tint_symbol_2 = 0u;
+  G.GetDimensions(tint_symbol_2);
+  uint tint_symbol_3 = ((tint_symbol_2 - 0u) / 4u);
+  uint l1 = tint_symbol_3;
+  uint l2 = tint_symbol_3;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}
+Error parsing GLSL shader:
+ERROR: 0:9: 'GetDimensions' : no such field in structure 
+ERROR: 0:9: '' : compilation terminated 
+ERROR: 2 compilation errors.  No code generated.
+
+
+

test/builtins/arrayLength/deprecated.wgsl.expected.hlsl

@@ -0,0 +1,11 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = ((tint_symbol_1 - 0u) / 4u);
+  const uint l1 = tint_symbol_2;
+  const uint l2 = tint_symbol_2;
+  return;
+}

test/builtins/arrayLength/deprecated.wgsl.expected.msl

@@ -0,0 +1,16 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+struct S {
+  int a[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
+  uint const l2 = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
+  return;
+}
+

test/builtins/arrayLength/deprecated.wgsl.expected.spvasm

@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 13
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+          %S = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %G = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+         %12 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/builtins/arrayLength/deprecated.wgsl.expected.wgsl

@@ -0,0 +1,12 @@
+struct S {
+  a : array<i32>;
+}
+
+@group(0) @binding(0) var<storage, read> G : S;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let l1 : u32 = arrayLength(&(G.a));
+  let p = &(G.a);
+  let l2 : u32 = arrayLength(p);
+}

test/builtins/arrayLength/simple.wgsl

@@ -0,0 +1,10 @@
+struct S {
+    a : array<i32>;
+};
+
+@group(0) @binding(0) var<storage, read> G : S;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+    let l1 : u32 = arrayLength(&G.a);
+}

test/builtins/arrayLength/simple.wgsl.expected.glsl

@@ -0,0 +1,27 @@
+SKIP: FAILED
+
+#version 310 es
+precision mediump float;
+
+layout(binding = 0) buffer S_1 {
+  int a[];
+} G;
+void tint_symbol() {
+  uint tint_symbol_2 = 0u;
+  G.GetDimensions(tint_symbol_2);
+  uint tint_symbol_3 = ((tint_symbol_2 - 0u) / 4u);
+  uint l1 = tint_symbol_3;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}
+Error parsing GLSL shader:
+ERROR: 0:9: 'GetDimensions' : no such field in structure 
+ERROR: 0:9: '' : compilation terminated 
+ERROR: 2 compilation errors.  No code generated.
+
+
+

test/builtins/arrayLength/simple.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = ((tint_symbol_1 - 0u) / 4u);
+  const uint l1 = tint_symbol_2;
+  return;
+}

test/builtins/arrayLength/simple.wgsl.expected.msl

@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+struct S {
+  int a[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
+  return;
+}
+

test/builtins/arrayLength/simple.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+          %S = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %G = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/builtins/arrayLength/simple.wgsl.expected.wgsl

@@ -0,0 +1,10 @@
+struct S {
+  a : array<i32>;
+}
+
+@group(0) @binding(0) var<storage, read> G : S;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let l1 : u32 = arrayLength(&(G.a));
+}

test/builtins/arrayLength/simple_no_struct.wgsl

@@ -0,0 +1,6 @@
+@group(0) @binding(0) var<storage, read> G : array<i32>;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+    let l1 : u32 = arrayLength(&G);
+}

test/builtins/arrayLength/simple_no_struct.wgsl.expected.glsl

@@ -0,0 +1,27 @@
+SKIP: FAILED
+
+#version 310 es
+precision mediump float;
+
+layout(binding = 0) buffer G_block_1 {
+  int inner[];
+} G;
+void tint_symbol() {
+  uint tint_symbol_2 = 0u;
+  G.inner.GetDimensions(tint_symbol_2);
+  uint tint_symbol_3 = (tint_symbol_2 / 4u);
+  uint l1 = tint_symbol_3;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}
+Error parsing GLSL shader:
+ERROR: 0:9: '.' : cannot apply to an array: GetDimensions
+ERROR: 0:9: '' : compilation terminated 
+ERROR: 2 compilation errors.  No code generated.
+
+
+

test/builtins/arrayLength/simple_no_struct.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+  const uint l1 = tint_symbol_2;
+  return;
+}

test/builtins/arrayLength/simple_no_struct.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+  return;
+}
+

test/builtins/arrayLength/simple_no_struct.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %G_block "G_block"
+               OpMemberName %G_block 0 "inner"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %G_block Block
+               OpMemberDecorate %G_block 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+    %G_block = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_G_block = OpTypePointer StorageBuffer %G_block
+          %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/builtins/arrayLength/simple_no_struct.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+@group(0) @binding(0) var<storage, read> G : array<i32>;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let l1 : u32 = arrayLength(&(G));
+}

test/builtins/arrayLength/via_let.wgsl

@@ -0,0 +1,12 @@
+struct S {
+    a : array<i32>;
+};
+
+@group(0) @binding(0) var<storage, read> G : S;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+    let p = &G.a;
+    let p2 = p;
+    let l1 : u32 = arrayLength(p2);
+}

test/builtins/arrayLength/via_let.wgsl.expected.glsl

@@ -0,0 +1,27 @@
+SKIP: FAILED
+
+#version 310 es
+precision mediump float;
+
+layout(binding = 0) buffer S_1 {
+  int a[];
+} G;
+void tint_symbol() {
+  uint tint_symbol_2 = 0u;
+  G.GetDimensions(tint_symbol_2);
+  uint tint_symbol_3 = ((tint_symbol_2 - 0u) / 4u);
+  uint l1 = tint_symbol_3;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}
+Error parsing GLSL shader:
+ERROR: 0:9: 'GetDimensions' : no such field in structure 
+ERROR: 0:9: '' : compilation terminated 
+ERROR: 2 compilation errors.  No code generated.
+
+
+

test/builtins/arrayLength/via_let.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = ((tint_symbol_1 - 0u) / 4u);
+  const uint l1 = tint_symbol_2;
+  return;
+}

test/builtins/arrayLength/via_let.wgsl.expected.msl

@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+struct S {
+  int a[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
+  return;
+}
+

test/builtins/arrayLength/via_let.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+          %S = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %G = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/builtins/arrayLength/via_let.wgsl.expected.wgsl

@@ -0,0 +1,12 @@
+struct S {
+  a : array<i32>;
+}
+
+@group(0) @binding(0) var<storage, read> G : S;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let p = &(G.a);
+  let p2 = p;
+  let l1 : u32 = arrayLength(p2);
+}

test/builtins/arrayLength/via_let_complex.wgsl

@@ -0,0 +1,13 @@
+struct S {
+    a : array<i32>;
+};
+
+@group(0) @binding(0) var<storage, read> G : S;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let p = &*&G;
+  let p2 = &*p;
+  let p3 = &((*p).a);
+  let l1 : u32 = arrayLength(&*p3);
+}

test/builtins/arrayLength/via_let_complex.wgsl.expected.glsl

@@ -0,0 +1,27 @@
+SKIP: FAILED
+
+#version 310 es
+precision mediump float;
+
+layout(binding = 0) buffer S_1 {
+  int a[];
+} G;
+void tint_symbol() {
+  uint tint_symbol_2 = 0u;
+  G.GetDimensions(tint_symbol_2);
+  uint tint_symbol_3 = ((tint_symbol_2 - 0u) / 4u);
+  uint l1 = tint_symbol_3;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}
+Error parsing GLSL shader:
+ERROR: 0:9: 'GetDimensions' : no such field in structure 
+ERROR: 0:9: '' : compilation terminated 
+ERROR: 2 compilation errors.  No code generated.
+
+
+

test/builtins/arrayLength/via_let_complex.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = ((tint_symbol_1 - 0u) / 4u);
+  const uint l1 = tint_symbol_2;
+  return;
+}

test/builtins/arrayLength/via_let_complex.wgsl.expected.msl

@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+struct S {
+  int a[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
+  return;
+}
+

test/builtins/arrayLength/via_let_complex.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+          %S = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %G = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/builtins/arrayLength/via_let_complex.wgsl.expected.wgsl

@@ -0,0 +1,13 @@
+struct S {
+  a : array<i32>;
+}
+
+@group(0) @binding(0) var<storage, read> G : S;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let p = &(*(&(G)));
+  let p2 = &(*(p));
+  let p3 = &((*(p)).a);
+  let l1 : u32 = arrayLength(&(*(p3)));
+}

test/builtins/arrayLength/via_let_complex_no_struct.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> G : array<i32>;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let p = &*&G;
+  let p2 = &*p;
+  let p3 = &(*p);
+  let l1 : u32 = arrayLength(&*p3);
+}

test/builtins/arrayLength/via_let_complex_no_struct.wgsl.expected.glsl

@@ -0,0 +1,27 @@
+SKIP: FAILED
+
+#version 310 es
+precision mediump float;
+
+layout(binding = 0) buffer G_block_1 {
+  int inner[];
+} G;
+void tint_symbol() {
+  uint tint_symbol_2 = 0u;
+  G.inner.GetDimensions(tint_symbol_2);
+  uint tint_symbol_3 = (tint_symbol_2 / 4u);
+  uint l1 = tint_symbol_3;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}
+Error parsing GLSL shader:
+ERROR: 0:9: '.' : cannot apply to an array: GetDimensions
+ERROR: 0:9: '' : compilation terminated 
+ERROR: 2 compilation errors.  No code generated.
+
+
+

test/builtins/arrayLength/via_let_complex_no_struct.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+  const uint l1 = tint_symbol_2;
+  return;
+}

test/builtins/arrayLength/via_let_complex_no_struct.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+  return;
+}
+

test/builtins/arrayLength/via_let_complex_no_struct.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %G_block "G_block"
+               OpMemberName %G_block 0 "inner"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %G_block Block
+               OpMemberDecorate %G_block 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+    %G_block = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_G_block = OpTypePointer StorageBuffer %G_block
+          %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/builtins/arrayLength/via_let_complex_no_struct.wgsl.expected.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> G : array<i32>;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let p = &(*(&(G)));
+  let p2 = &(*(p));
+  let p3 = &(*(p));
+  let l1 : u32 = arrayLength(&(*(p3)));
+}

test/builtins/arrayLength/via_let_no_struct.wgsl

@@ -0,0 +1,8 @@
+@group(0) @binding(0) var<storage, read> G : array<i32>;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+    let p = &G;
+    let p2 = p;
+    let l1 : u32 = arrayLength(p2);
+}

test/builtins/arrayLength/via_let_no_struct.wgsl.expected.glsl

@@ -0,0 +1,27 @@
+SKIP: FAILED
+
+#version 310 es
+precision mediump float;
+
+layout(binding = 0) buffer G_block_1 {
+  int inner[];
+} G;
+void tint_symbol() {
+  uint tint_symbol_2 = 0u;
+  G.inner.GetDimensions(tint_symbol_2);
+  uint tint_symbol_3 = (tint_symbol_2 / 4u);
+  uint l1 = tint_symbol_3;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}
+Error parsing GLSL shader:
+ERROR: 0:9: '.' : cannot apply to an array: GetDimensions
+ERROR: 0:9: '' : compilation terminated 
+ERROR: 2 compilation errors.  No code generated.
+
+
+

test/builtins/arrayLength/via_let_no_struct.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+  const uint l1 = tint_symbol_2;
+  return;
+}

test/builtins/arrayLength/via_let_no_struct.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+  return;
+}
+

test/builtins/arrayLength/via_let_no_struct.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %G_block "G_block"
+               OpMemberName %G_block 0 "inner"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %G_block Block
+               OpMemberDecorate %G_block 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+    %G_block = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_G_block = OpTypePointer StorageBuffer %G_block
+          %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/builtins/arrayLength/via_let_no_struct.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+@group(0) @binding(0) var<storage, read> G : array<i32>;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let p = &(G);
+  let p2 = p;
+  let l1 : u32 = arrayLength(p2);
+}