[msl] Emit min and max builtins
Bug: 42251016
Change-Id: I71d0428674a26830264d708070c6eea1323d9536
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/188347
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 35f1519..bfc968e 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -875,6 +875,8 @@
case core::BuiltinFn::kLdexp:
case core::BuiltinFn::kLog2:
case core::BuiltinFn::kLog:
+ case core::BuiltinFn::kMax:
+ case core::BuiltinFn::kMin:
case core::BuiltinFn::kMix:
case core::BuiltinFn::kNormalize:
case core::BuiltinFn::kPow:
diff --git a/test/tint/bug/tint/1332.wgsl.expected.ir.msl b/test/tint/bug/tint/1332.wgsl.expected.ir.msl
index bbc0093..3c72b34 100644
--- a/test/tint/bug/tint/1332.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1332.wgsl.expected.ir.msl
@@ -1,9 +1,7 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: max
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+kernel void compute_main() {
+ float const a = 1.24000000953674316406f;
+ float b = max(a, 1.17549435e-38f);
+}
diff --git a/test/tint/bug/tint/1563.wgsl.expected.ir.msl b/test/tint/bug/tint/1563.wgsl.expected.ir.msl
index 40ec97c..c12e4fc 100644
--- a/test/tint/bug/tint/1563.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1563.wgsl.expected.ir.msl
@@ -1,9 +1,10 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: min
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+float foo() {
+ int const oob = 99;
+ float const b = float4(0.0f)[min(uint(oob), 3u)];
+ float4 v = 0.0f;
+ v[min(uint(oob), 3u)] = b;
+ return b;
+}
diff --git a/test/tint/builtins/extractBits/scalar/i32.spvasm.expected.ir.msl b/test/tint/builtins/extractBits/scalar/i32.spvasm.expected.ir.msl
index 40ec97c..499f249 100644
--- a/test/tint/builtins/extractBits/scalar/i32.spvasm.expected.ir.msl
+++ b/test/tint/builtins/extractBits/scalar/i32.spvasm.expected.ir.msl
@@ -1,9 +1,15 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: min
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+void f_1() {
+ int v = 0;
+ uint offset_1 = 0u;
+ uint count = 0u;
+ int const v_1 = v;
+ uint const v_2 = count;
+ uint const v_3 = min(offset_1, 32u);
+ int const x_14 = extract_bits(v_1, v_3, min(v_2, (32u - v_3)));
+}
+kernel void f() {
+ f_1();
+}
diff --git a/test/tint/builtins/extractBits/scalar/u32.spvasm.expected.ir.msl b/test/tint/builtins/extractBits/scalar/u32.spvasm.expected.ir.msl
index 40ec97c..5cf9afb 100644
--- a/test/tint/builtins/extractBits/scalar/u32.spvasm.expected.ir.msl
+++ b/test/tint/builtins/extractBits/scalar/u32.spvasm.expected.ir.msl
@@ -1,9 +1,15 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: min
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+void f_1() {
+ uint v = 0u;
+ uint offset_1 = 0u;
+ uint count = 0u;
+ uint const v_1 = v;
+ uint const v_2 = count;
+ uint const v_3 = min(offset_1, 32u);
+ uint const x_11 = extract_bits(v_1, v_3, min(v_2, (32u - v_3)));
+}
+kernel void f() {
+ f_1();
+}
diff --git a/test/tint/builtins/extractBits/vec3/i32.spvasm.expected.ir.msl b/test/tint/builtins/extractBits/vec3/i32.spvasm.expected.ir.msl
index 40ec97c..4c0a0cf 100644
--- a/test/tint/builtins/extractBits/vec3/i32.spvasm.expected.ir.msl
+++ b/test/tint/builtins/extractBits/vec3/i32.spvasm.expected.ir.msl
@@ -1,9 +1,15 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: min
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+void f_1() {
+ int3 v = int3(0);
+ uint offset_1 = 0u;
+ uint count = 0u;
+ int3 const v_1 = v;
+ uint const v_2 = count;
+ uint const v_3 = min(offset_1, 32u);
+ int3 const x_15 = extract_bits(v_1, v_3, min(v_2, (32u - v_3)));
+}
+kernel void f() {
+ f_1();
+}
diff --git a/test/tint/builtins/extractBits/vec3/u32.spvasm.expected.ir.msl b/test/tint/builtins/extractBits/vec3/u32.spvasm.expected.ir.msl
index 40ec97c..680caac 100644
--- a/test/tint/builtins/extractBits/vec3/u32.spvasm.expected.ir.msl
+++ b/test/tint/builtins/extractBits/vec3/u32.spvasm.expected.ir.msl
@@ -1,9 +1,15 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: min
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+void f_1() {
+ uint3 v = uint3(0u);
+ uint offset_1 = 0u;
+ uint count = 0u;
+ uint3 const v_1 = v;
+ uint const v_2 = count;
+ uint const v_3 = min(offset_1, 32u);
+ uint3 const x_14 = extract_bits(v_1, v_3, min(v_2, (32u - v_3)));
+}
+kernel void f() {
+ f_1();
+}
diff --git a/test/tint/builtins/insertBits/scalar/i32.spvasm.expected.ir.msl b/test/tint/builtins/insertBits/scalar/i32.spvasm.expected.ir.msl
index 40ec97c..3ccaa47 100644
--- a/test/tint/builtins/insertBits/scalar/i32.spvasm.expected.ir.msl
+++ b/test/tint/builtins/insertBits/scalar/i32.spvasm.expected.ir.msl
@@ -1,9 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: min
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+void f_1() {
+ int v = 0;
+ int n = 0;
+ uint offset_1 = 0u;
+ uint count = 0u;
+ int const v_1 = v;
+ int const v_2 = n;
+ uint const v_3 = count;
+ uint const v_4 = min(offset_1, 32u);
+ int const x_15 = insert_bits(v_1, v_2, v_4, min(v_3, (32u - v_4)));
+}
+kernel void f() {
+ f_1();
+}
diff --git a/test/tint/builtins/insertBits/scalar/u32.spvasm.expected.ir.msl b/test/tint/builtins/insertBits/scalar/u32.spvasm.expected.ir.msl
index 40ec97c..b247b3b 100644
--- a/test/tint/builtins/insertBits/scalar/u32.spvasm.expected.ir.msl
+++ b/test/tint/builtins/insertBits/scalar/u32.spvasm.expected.ir.msl
@@ -1,9 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: min
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+void f_1() {
+ uint v = 0u;
+ uint n = 0u;
+ uint offset_1 = 0u;
+ uint count = 0u;
+ uint const v_1 = v;
+ uint const v_2 = n;
+ uint const v_3 = count;
+ uint const v_4 = min(offset_1, 32u);
+ uint const x_12 = insert_bits(v_1, v_2, v_4, min(v_3, (32u - v_4)));
+}
+kernel void f() {
+ f_1();
+}
diff --git a/test/tint/builtins/insertBits/vec3/i32.spvasm.expected.ir.msl b/test/tint/builtins/insertBits/vec3/i32.spvasm.expected.ir.msl
index 40ec97c..f7e8a00 100644
--- a/test/tint/builtins/insertBits/vec3/i32.spvasm.expected.ir.msl
+++ b/test/tint/builtins/insertBits/vec3/i32.spvasm.expected.ir.msl
@@ -1,9 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: min
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+void f_1() {
+ int3 v = int3(0);
+ int3 n = int3(0);
+ uint offset_1 = 0u;
+ uint count = 0u;
+ int3 const v_1 = v;
+ int3 const v_2 = n;
+ uint const v_3 = count;
+ uint const v_4 = min(offset_1, 32u);
+ int3 const x_16 = insert_bits(v_1, v_2, v_4, min(v_3, (32u - v_4)));
+}
+kernel void f() {
+ f_1();
+}
diff --git a/test/tint/builtins/insertBits/vec3/u32.spvasm.expected.ir.msl b/test/tint/builtins/insertBits/vec3/u32.spvasm.expected.ir.msl
index 40ec97c..ef958e4 100644
--- a/test/tint/builtins/insertBits/vec3/u32.spvasm.expected.ir.msl
+++ b/test/tint/builtins/insertBits/vec3/u32.spvasm.expected.ir.msl
@@ -1,9 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: min
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+void f_1() {
+ uint3 v = uint3(0u);
+ uint3 n = uint3(0u);
+ uint offset_1 = 0u;
+ uint count = 0u;
+ uint3 const v_1 = v;
+ uint3 const v_2 = n;
+ uint const v_3 = count;
+ uint const v_4 = min(offset_1, 32u);
+ uint3 const x_15 = insert_bits(v_1, v_2, v_4, min(v_3, (32u - v_4)));
+}
+kernel void f() {
+ f_1();
+}
diff --git a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_int_uclamp/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_int_uclamp/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index 55ef712..0000000
--- a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_int_uclamp/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,53 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(4) {
- field0:array<i32> @offset(0)
-}
-
-$B1: { # root
- %x_3:ptr<private, vec3<u32>, read_write> = var
- %x_6:ptr<storage, S, read_write> = var @binding_point(0, 0)
- %x_7:ptr<storage, S, read_write> = var @binding_point(0, 1)
- %x_8:ptr<storage, S, read_write> = var @binding_point(0, 2)
- %x_9:ptr<storage, S, read_write> = var @binding_point(0, 3)
-}
-
-%main_1 = func():void {
- $B2: {
- %7:u32 = load_vector_element %x_3, 0u
- %x_26:u32 = let %7
- %9:ptr<storage, i32, read_write> = access %x_9, 0u, %x_26
- %10:ptr<storage, i32, read_write> = access %x_6, 0u, %x_26
- %11:i32 = load %10
- %12:u32 = bitcast %11
- %13:u32 = let %12
- %14:ptr<storage, i32, read_write> = access %x_7, 0u, %x_26
- %15:i32 = load %14
- %16:u32 = bitcast %15
- %17:u32 = let %16
- %18:ptr<storage, i32, read_write> = access %x_8, 0u, %x_26
- %19:i32 = load %18
- %20:u32 = bitcast %19
- %21:u32 = let %20
- %22:u32 = max %13, %17
- %23:u32 = min %22, %21
- %24:i32 = bitcast %23
- store %9, %24
- ret
- }
-}
-%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%x_3_param:vec3<u32> [@global_invocation_id]):void {
- $B3: {
- store %x_3, %x_3_param
- %27:void = call %main_1
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_int_uclamp/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_int_uclamp/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index ada8f10..0000000
--- a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_int_uclamp/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,56 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(4) {
- field0:array<i32> @offset(0)
-}
-
-$B1: { # root
- %x_3:ptr<private, vec3<u32>, read_write> = var
- %x_6:ptr<storage, S, read_write> = var @binding_point(0, 0)
- %x_7:ptr<storage, S, read_write> = var @binding_point(0, 1)
- %x_8:ptr<storage, S, read_write> = var @binding_point(0, 2)
- %x_9:ptr<storage, S, read_write> = var @binding_point(0, 3)
-}
-
-%main_1 = func():void {
- $B2: {
- %7:u32 = load_vector_element %x_3, 0u
- %x_26:u32 = let %7
- %9:ptr<storage, i32, read_write> = access %x_6, 0u, %x_26
- %10:i32 = load %9
- %x_28:i32 = let %10
- %12:ptr<storage, i32, read_write> = access %x_7, 0u, %x_26
- %13:i32 = load %12
- %x_30:i32 = let %13
- %15:ptr<storage, i32, read_write> = access %x_8, 0u, %x_26
- %16:i32 = load %15
- %x_32:i32 = let %16
- %18:ptr<storage, i32, read_write> = access %x_9, 0u, %x_26
- %19:u32 = bitcast %x_28
- %20:u32 = let %19
- %21:u32 = bitcast %x_30
- %22:u32 = let %21
- %23:u32 = bitcast %x_32
- %24:u32 = let %23
- %25:u32 = max %20, %22
- %26:u32 = min %25, %24
- %27:i32 = bitcast %26
- store %18, %27
- ret
- }
-}
-%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%x_3_param:vec3<u32> [@global_invocation_id]):void {
- $B3: {
- store %x_3, %x_3_param
- %30:void = call %main_1
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_sclamp/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_sclamp/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index ec196d3..0000000
--- a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_sclamp/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,53 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(4) {
- field0:array<u32> @offset(0)
-}
-
-$B1: { # root
- %x_3:ptr<private, vec3<u32>, read_write> = var
- %x_6:ptr<storage, S, read_write> = var @binding_point(0, 0)
- %x_7:ptr<storage, S, read_write> = var @binding_point(0, 1)
- %x_8:ptr<storage, S, read_write> = var @binding_point(0, 2)
- %x_9:ptr<storage, S, read_write> = var @binding_point(0, 3)
-}
-
-%main_1 = func():void {
- $B2: {
- %7:u32 = load_vector_element %x_3, 0u
- %x_23:u32 = let %7
- %9:ptr<storage, u32, read_write> = access %x_9, 0u, %x_23
- %10:ptr<storage, u32, read_write> = access %x_6, 0u, %x_23
- %11:u32 = load %10
- %12:i32 = bitcast %11
- %13:i32 = let %12
- %14:ptr<storage, u32, read_write> = access %x_7, 0u, %x_23
- %15:u32 = load %14
- %16:i32 = bitcast %15
- %17:i32 = let %16
- %18:ptr<storage, u32, read_write> = access %x_8, 0u, %x_23
- %19:u32 = load %18
- %20:i32 = bitcast %19
- %21:i32 = let %20
- %22:i32 = max %13, %17
- %23:i32 = min %22, %21
- %24:u32 = bitcast %23
- store %9, %24
- ret
- }
-}
-%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%x_3_param:vec3<u32> [@global_invocation_id]):void {
- $B3: {
- store %x_3, %x_3_param
- %27:void = call %main_1
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_sclamp/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_sclamp/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index d11bca9..0000000
--- a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_sclamp/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,56 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(4) {
- field0:array<u32> @offset(0)
-}
-
-$B1: { # root
- %x_3:ptr<private, vec3<u32>, read_write> = var
- %x_6:ptr<storage, S, read_write> = var @binding_point(0, 0)
- %x_7:ptr<storage, S, read_write> = var @binding_point(0, 1)
- %x_8:ptr<storage, S, read_write> = var @binding_point(0, 2)
- %x_9:ptr<storage, S, read_write> = var @binding_point(0, 3)
-}
-
-%main_1 = func():void {
- $B2: {
- %7:u32 = load_vector_element %x_3, 0u
- %x_23:u32 = let %7
- %9:ptr<storage, u32, read_write> = access %x_6, 0u, %x_23
- %10:u32 = load %9
- %x_25:u32 = let %10
- %12:ptr<storage, u32, read_write> = access %x_7, 0u, %x_23
- %13:u32 = load %12
- %x_27:u32 = let %13
- %15:ptr<storage, u32, read_write> = access %x_8, 0u, %x_23
- %16:u32 = load %15
- %x_29:u32 = let %16
- %18:ptr<storage, u32, read_write> = access %x_9, 0u, %x_23
- %19:i32 = bitcast %x_25
- %20:i32 = let %19
- %21:i32 = bitcast %x_27
- %22:i32 = let %21
- %23:i32 = bitcast %x_29
- %24:i32 = let %23
- %25:i32 = max %20, %22
- %26:i32 = min %25, %24
- %27:u32 = bitcast %26
- store %18, %27
- ret
- }
-}
-%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%x_3_param:vec3<u32> [@global_invocation_id]):void {
- $B3: {
- store %x_3, %x_3_param
- %30:void = call %main_1
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_smax/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_smax/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index c9ca4ee..0000000
--- a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_smax/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,46 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(4) {
- field0:array<u32> @offset(0)
-}
-
-$B1: { # root
- %x_3:ptr<private, vec3<u32>, read_write> = var
- %x_6:ptr<storage, S, read_write> = var @binding_point(0, 0)
- %x_7:ptr<storage, S, read_write> = var @binding_point(0, 1)
- %x_8:ptr<storage, S, read_write> = var @binding_point(0, 2)
-}
-
-%main_1 = func():void {
- $B2: {
- %6:u32 = load_vector_element %x_3, 0u
- %x_21:u32 = let %6
- %8:ptr<storage, u32, read_write> = access %x_8, 0u, %x_21
- %9:ptr<storage, u32, read_write> = access %x_6, 0u, %x_21
- %10:u32 = load %9
- %11:i32 = bitcast %10
- %12:i32 = let %11
- %13:ptr<storage, u32, read_write> = access %x_7, 0u, %x_21
- %14:u32 = load %13
- %15:i32 = bitcast %14
- %16:i32 = max %12, %15
- %17:u32 = bitcast %16
- store %8, %17
- ret
- }
-}
-%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%x_3_param:vec3<u32> [@global_invocation_id]):void {
- $B3: {
- store %x_3, %x_3_param
- %20:void = call %main_1
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_smax/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_smax/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index cd3896b..0000000
--- a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_smax/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,48 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(4) {
- field0:array<u32> @offset(0)
-}
-
-$B1: { # root
- %x_3:ptr<private, vec3<u32>, read_write> = var
- %x_6:ptr<storage, S, read_write> = var @binding_point(0, 0)
- %x_7:ptr<storage, S, read_write> = var @binding_point(0, 1)
- %x_8:ptr<storage, S, read_write> = var @binding_point(0, 2)
-}
-
-%main_1 = func():void {
- $B2: {
- %6:u32 = load_vector_element %x_3, 0u
- %x_21:u32 = let %6
- %8:ptr<storage, u32, read_write> = access %x_6, 0u, %x_21
- %9:u32 = load %8
- %x_23:u32 = let %9
- %11:ptr<storage, u32, read_write> = access %x_7, 0u, %x_21
- %12:u32 = load %11
- %x_25:u32 = let %12
- %14:ptr<storage, u32, read_write> = access %x_8, 0u, %x_21
- %15:i32 = bitcast %x_23
- %16:i32 = let %15
- %17:i32 = bitcast %x_25
- %18:i32 = max %16, %17
- %19:u32 = bitcast %18
- store %14, %19
- ret
- }
-}
-%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%x_3_param:vec3<u32> [@global_invocation_id]):void {
- $B3: {
- store %x_3, %x_3_param
- %22:void = call %main_1
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_smin/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_smin/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index 4c56349..0000000
--- a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_smin/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,46 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(4) {
- field0:array<u32> @offset(0)
-}
-
-$B1: { # root
- %x_3:ptr<private, vec3<u32>, read_write> = var
- %x_6:ptr<storage, S, read_write> = var @binding_point(0, 0)
- %x_7:ptr<storage, S, read_write> = var @binding_point(0, 1)
- %x_8:ptr<storage, S, read_write> = var @binding_point(0, 2)
-}
-
-%main_1 = func():void {
- $B2: {
- %6:u32 = load_vector_element %x_3, 0u
- %x_21:u32 = let %6
- %8:ptr<storage, u32, read_write> = access %x_8, 0u, %x_21
- %9:ptr<storage, u32, read_write> = access %x_6, 0u, %x_21
- %10:u32 = load %9
- %11:i32 = bitcast %10
- %12:i32 = let %11
- %13:ptr<storage, u32, read_write> = access %x_7, 0u, %x_21
- %14:u32 = load %13
- %15:i32 = bitcast %14
- %16:i32 = min %12, %15
- %17:u32 = bitcast %16
- store %8, %17
- ret
- }
-}
-%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%x_3_param:vec3<u32> [@global_invocation_id]):void {
- $B3: {
- store %x_3, %x_3_param
- %20:void = call %main_1
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_smin/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_smin/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index 749865a..0000000
--- a/test/tint/vk-gl-cts/spirv_assembly/instruction/compute/signed_op/glsl_uint_smin/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,48 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(4) {
- field0:array<u32> @offset(0)
-}
-
-$B1: { # root
- %x_3:ptr<private, vec3<u32>, read_write> = var
- %x_6:ptr<storage, S, read_write> = var @binding_point(0, 0)
- %x_7:ptr<storage, S, read_write> = var @binding_point(0, 1)
- %x_8:ptr<storage, S, read_write> = var @binding_point(0, 2)
-}
-
-%main_1 = func():void {
- $B2: {
- %6:u32 = load_vector_element %x_3, 0u
- %x_21:u32 = let %6
- %8:ptr<storage, u32, read_write> = access %x_6, 0u, %x_21
- %9:u32 = load %8
- %x_23:u32 = let %9
- %11:ptr<storage, u32, read_write> = access %x_7, 0u, %x_21
- %12:u32 = load %11
- %x_25:u32 = let %12
- %14:ptr<storage, u32, read_write> = access %x_8, 0u, %x_21
- %15:i32 = bitcast %x_23
- %16:i32 = let %15
- %17:i32 = bitcast %x_25
- %18:i32 = min %16, %17
- %19:u32 = bitcast %18
- store %14, %19
- ret
- }
-}
-%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%x_3_param:vec3<u32> [@global_invocation_id]):void {
- $B3: {
- store %x_3, %x_3_param
- %22:void = call %main_1
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
