[msl] Add polyfill for distance builtin
If the argument is scalar, we need to use `abs(a - b)`.
Add an MSL builtin function for distance that only supports vectors.
Bug: 42251016
Change-Id: I96b83cc20caacac9866059c454c4dc5e692b4c2c
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/193921
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/msl/builtin_fn.cc b/src/tint/lang/msl/builtin_fn.cc
index e345a4d..8c88765 100644
--- a/src/tint/lang/msl/builtin_fn.cc
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -90,6 +90,8 @@
return "sample_compare";
case BuiltinFn::kWrite:
return "write";
+ case BuiltinFn::kDistance:
+ return "distance";
case BuiltinFn::kLength:
return "length";
case BuiltinFn::kThreadgroupBarrier:
diff --git a/src/tint/lang/msl/builtin_fn.h b/src/tint/lang/msl/builtin_fn.h
index 04b22b1..0dac6ca 100644
--- a/src/tint/lang/msl/builtin_fn.h
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -71,6 +71,7 @@
kSample,
kSampleCompare,
kWrite,
+ kDistance,
kLength,
kThreadgroupBarrier,
kNone,
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index eb585fd..7978441 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -2600,14 +2600,19 @@
},
{
/* [300] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(9),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(93),
},
{
/* [301] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(93),
},
+ {
+ /* [302] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(9),
+ },
};
static_assert(ParameterIndex::CanIndex(kParameters),
@@ -3552,7 +3557,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(300),
+ /* parameters */ ParameterIndex(302),
/* return_matcher_indices */ MatcherIndicesIndex(59),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4333,7 +4338,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(300),
+ /* parameters */ ParameterIndex(302),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4461,16 +4466,27 @@
{
/* [160] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
+ /* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(11),
- /* parameters */ ParameterIndex(301),
+ /* parameters */ ParameterIndex(300),
/* return_matcher_indices */ MatcherIndicesIndex(1),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
/* [161] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 2,
+ /* templates */ TemplateIndex(11),
+ /* parameters */ ParameterIndex(300),
+ /* return_matcher_indices */ MatcherIndicesIndex(1),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [162] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4775,16 +4791,22 @@
},
{
/* [24] */
- /* fn length[N : num, T : f32_f16](vec<N, T>) -> T */
+ /* fn distance[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(160),
},
{
/* [25] */
- /* fn threadgroup_barrier(u32) */
+ /* fn length[N : num, T : f32_f16](vec<N, T>) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(161),
},
+ {
+ /* [26] */
+ /* fn threadgroup_barrier(u32) */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(162),
+ },
};
// clang-format on
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
index bb569df..f720c99 100644
--- a/src/tint/lang/msl/msl.def
+++ b/src/tint/lang/msl/msl.def
@@ -321,6 +321,7 @@
@member_function fn write[A: iu32](texture: texture_storage_2d_array<u32_texel_format, writable>, value: vec4<u32>, coords: vec2<u32>, array_index: A)
@member_function fn write(texture: texture_storage_3d<u32_texel_format, writable>, value: vec4<u32>, coords: vec3<u32>)
+fn distance[N: num, T: f32_f16](vec<N, T>, vec<N, T>) -> T
fn length[N: num, T: f32_f16](vec<N, T>) -> T
@stage("compute") fn threadgroup_barrier(u32)
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index 2939e83..a406bee 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -93,6 +93,7 @@
case core::BuiltinFn::kAtomicStore:
case core::BuiltinFn::kAtomicSub:
case core::BuiltinFn::kAtomicXor:
+ case core::BuiltinFn::kDistance:
case core::BuiltinFn::kLength:
case core::BuiltinFn::kTextureDimensions:
case core::BuiltinFn::kTextureGather:
@@ -159,6 +160,9 @@
break;
// Geometric builtins.
+ case core::BuiltinFn::kDistance:
+ Distance(builtin);
+ break;
case core::BuiltinFn::kLength:
Length(builtin);
break;
@@ -282,6 +286,24 @@
builtin->Destroy();
}
+ /// Polyfill a distance call if necessary.
+ /// @param builtin the builtin call instruction
+ void Distance(core::ir::CoreBuiltinCall* builtin) {
+ b.InsertBefore(builtin, [&] {
+ auto* arg0 = builtin->Args()[0];
+ auto* arg1 = builtin->Args()[1];
+ if (arg0->Type()->Is<core::type::Scalar>()) {
+ // Calls to `distance` with a scalar argument are replaced with `abs(a - b)`.
+ auto* sub = b.Subtract(builtin->Result(0)->Type(), arg0, arg1);
+ b.CallWithResult(builtin->DetachResult(), core::BuiltinFn::kAbs, sub);
+ } else {
+ b.CallWithResult<msl::ir::BuiltinCall>(builtin->DetachResult(),
+ msl::BuiltinFn::kDistance, arg0, arg1);
+ }
+ });
+ builtin->Destroy();
+ }
+
/// Polyfill a length call if necessary.
/// @param builtin the builtin call instruction
void Length(core::ir::CoreBuiltinCall* builtin) {
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
index cff9ff5..ed89002 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -758,6 +758,75 @@
EXPECT_EQ(expect, str());
}
+TEST_F(MslWriter_BuiltinPolyfillTest, Distance_Scalar) {
+ auto* value0 = b.FunctionParam<f32>("value0");
+ auto* value1 = b.FunctionParam<f32>("value1");
+ auto* func = b.Function("foo", ty.f32());
+ func->SetParams({value0, value1});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call<f32>(core::BuiltinFn::kDistance, value0, value1);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%value0:f32, %value1:f32):f32 {
+ $B1: {
+ %4:f32 = distance %value0, %value1
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%value0:f32, %value1:f32):f32 {
+ $B1: {
+ %4:f32 = sub %value0, %value1
+ %5:f32 = abs %4
+ ret %5
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, Distance_Vector) {
+ auto* value0 = b.FunctionParam<vec4<f32>>("value0");
+ auto* value1 = b.FunctionParam<vec4<f32>>("value1");
+ auto* func = b.Function("foo", ty.f32());
+ func->SetParams({value0, value1});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call<f32>(core::BuiltinFn::kDistance, value0, value1);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%value0:vec4<f32>, %value1:vec4<f32>):f32 {
+ $B1: {
+ %4:f32 = distance %value0, %value1
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%value0:vec4<f32>, %value1:vec4<f32>):f32 {
+ $B1: {
+ %4:f32 = msl.distance %value0, %value1
+ ret %4
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(MslWriter_BuiltinPolyfillTest, Length_Scalar) {
auto* value = b.FunctionParam<f32>("value");
auto* func = b.Function("foo", ty.f32());
diff --git a/test/tint/bug/tint/1061.spvasm.expected.ir.msl b/test/tint/bug/tint/1061.spvasm.expected.ir.msl
index 61fdca7..cd584c6 100644
--- a/test/tint/bug/tint/1061.spvasm.expected.ir.msl
+++ b/test/tint/bug/tint/1061.spvasm.expected.ir.msl
@@ -1,67 +1,39 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct buf0 {
+ float4 r;
+};
+struct tint_module_vars_struct {
+ const constant buf0* x_7;
+ thread float4* x_GLF_color;
+};
+struct main_out {
+ float4 x_GLF_color_1;
+};
+struct tint_symbol_outputs {
+ float4 main_out_x_GLF_color_1 [[color(0)]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(16) {
- r:vec4<f32> @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_7:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %f:ptr<function, f32, read_write> = var
- %v:ptr<function, vec4<f32>, read_write> = var
- store %f, 1.0f
- %6:f32 = load %f
- %7:f32 = sin %6
- %8:f32 = let %7
- %9:f32 = load %f
- %10:f32 = cos %9
- %11:f32 = let %10
- %12:f32 = load %f
- %13:f32 = exp2 %12
- %14:f32 = let %13
- %15:f32 = load %f
- %16:f32 = log %15
- %17:vec4<f32> = construct %8, %11, %14, %16
- store %v, %17
- %18:vec4<f32> = load %v
- %19:ptr<uniform, vec4<f32>, read> = access %x_7, 0u
- %20:vec4<f32> = load %19
- %21:f32 = distance %18, %20
- %22:bool = lt %21, 0.10000000149011611938f
- if %22 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- store %x_GLF_color, vec4<f32>(1.0f, 0.0f, 0.0f, 1.0f)
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
+void main_1(tint_module_vars_struct tint_module_vars) {
+ float f = 0.0f;
+ float4 v = 0.0f;
+ f = 1.0f;
+ float const v_1 = sin(f);
+ float const v_2 = cos(f);
+ float const v_3 = exp2(f);
+ v = float4(v_1, v_2, v_3, log(f));
+ if ((distance(v, (*tint_module_vars.x_7).r) < 0.10000000149011611938f)) {
+ (*tint_module_vars.x_GLF_color) = float4(1.0f, 0.0f, 0.0f, 1.0f);
+ } else {
+ (*tint_module_vars.x_GLF_color) = float4(0.0f);
}
}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %24:void = call %main_1
- %25:vec4<f32> = load %x_GLF_color
- %26:main_out = construct %25
- ret %26
- }
+main_out tint_symbol_inner(tint_module_vars_struct tint_module_vars) {
+ main_1(tint_module_vars);
+ return main_out{.x_GLF_color_1=(*tint_module_vars.x_GLF_color)};
}
-
-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. *
-********************************************************************
+fragment tint_symbol_outputs tint_symbol(const constant buf0* x_7 [[buffer(0)]]) {
+ thread float4 x_GLF_color = 0.0f;
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.x_7=x_7, .x_GLF_color=(&x_GLF_color)};
+ return tint_symbol_outputs{.main_out_x_GLF_color_1=tint_symbol_inner(tint_module_vars).x_GLF_color_1};
+}
diff --git a/test/tint/builtins/gen/var/distance/0657d4.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/distance/0657d4.wgsl.expected.ir.msl
index 9039640..553e62a 100644
--- a/test/tint/builtins/gen/var/distance/0657d4.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/distance/0657d4.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device float* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ float prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+float distance_0657d4() {
+ float3 arg_0 = float3(1.0f);
+ float3 arg_1 = float3(1.0f);
+ float res = distance(arg_0, arg_1);
+ return res;
}
-
-%distance_0657d4 = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %arg_1:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %5:vec3<f32> = load %arg_0
- %6:vec3<f32> = load %arg_1
- %7:f32 = distance %5, %6
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_0657d4();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %distance_0657d4
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_0657d4();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %distance_0657d4
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = distance_0657d4();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %distance_0657d4
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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/builtins/gen/var/distance/7272f3.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/distance/7272f3.wgsl.expected.ir.msl
index 4011b03..98fe780 100644
--- a/test/tint/builtins/gen/var/distance/7272f3.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/distance/7272f3.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device half* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ half prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f16, read_write> = var @binding_point(2, 0)
+half distance_7272f3() {
+ half4 arg_0 = half4(1.0h);
+ half4 arg_1 = half4(1.0h);
+ half res = distance(arg_0, arg_1);
+ return res;
}
-
-%distance_7272f3 = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<f16>, read_write> = var, vec4<f16>(1.0h)
- %arg_1:ptr<function, vec4<f16>, read_write> = var, vec4<f16>(1.0h)
- %5:vec4<f16> = load %arg_0
- %6:vec4<f16> = load %arg_1
- %7:f16 = distance %5, %6
- %res:ptr<function, f16, read_write> = var, %7
- %9:f16 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_7272f3();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %distance_7272f3
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_7272f3();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %distance_7272f3
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = distance_7272f3();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %distance_7272f3
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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/builtins/gen/var/distance/7d201f.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/distance/7d201f.wgsl.expected.ir.msl
index 191e2a4..18b2205 100644
--- a/test/tint/builtins/gen/var/distance/7d201f.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/distance/7d201f.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device half* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ half prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f16, read_write> = var @binding_point(2, 0)
+half distance_7d201f() {
+ half arg_0 = 1.0h;
+ half arg_1 = 1.0h;
+ half res = abs((arg_0 - arg_1));
+ return res;
}
-
-%distance_7d201f = func():void {
- $B2: {
- %arg_0:ptr<function, f16, read_write> = var, 1.0h
- %arg_1:ptr<function, f16, read_write> = var, 1.0h
- %5:f16 = load %arg_0
- %6:f16 = load %arg_1
- %7:f16 = distance %5, %6
- %res:ptr<function, f16, read_write> = var, %7
- %9:f16 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_7d201f();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %distance_7d201f
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_7d201f();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %distance_7d201f
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = distance_7d201f();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %distance_7d201f
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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/builtins/gen/var/distance/892a5d.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/distance/892a5d.wgsl.expected.ir.msl
index 747068c..7cac0cb 100644
--- a/test/tint/builtins/gen/var/distance/892a5d.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/distance/892a5d.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device half* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ half prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f16, read_write> = var @binding_point(2, 0)
+half distance_892a5d() {
+ half2 arg_0 = half2(1.0h);
+ half2 arg_1 = half2(1.0h);
+ half res = distance(arg_0, arg_1);
+ return res;
}
-
-%distance_892a5d = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<f16>, read_write> = var, vec2<f16>(1.0h)
- %arg_1:ptr<function, vec2<f16>, read_write> = var, vec2<f16>(1.0h)
- %5:vec2<f16> = load %arg_0
- %6:vec2<f16> = load %arg_1
- %7:f16 = distance %5, %6
- %res:ptr<function, f16, read_write> = var, %7
- %9:f16 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_892a5d();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %distance_892a5d
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_892a5d();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %distance_892a5d
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = distance_892a5d();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %distance_892a5d
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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/builtins/gen/var/distance/928fa0.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/distance/928fa0.wgsl.expected.ir.msl
index d82d9bc..cadcd1c 100644
--- a/test/tint/builtins/gen/var/distance/928fa0.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/distance/928fa0.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device half* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ half prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f16, read_write> = var @binding_point(2, 0)
+half distance_928fa0() {
+ half3 arg_0 = half3(1.0h);
+ half3 arg_1 = half3(1.0h);
+ half res = distance(arg_0, arg_1);
+ return res;
}
-
-%distance_928fa0 = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<f16>, read_write> = var, vec3<f16>(1.0h)
- %arg_1:ptr<function, vec3<f16>, read_write> = var, vec3<f16>(1.0h)
- %5:vec3<f16> = load %arg_0
- %6:vec3<f16> = load %arg_1
- %7:f16 = distance %5, %6
- %res:ptr<function, f16, read_write> = var, %7
- %9:f16 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_928fa0();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %distance_928fa0
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_928fa0();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %distance_928fa0
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = distance_928fa0();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %distance_928fa0
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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/builtins/gen/var/distance/9646ea.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/distance/9646ea.wgsl.expected.ir.msl
index 1a8b1cc..7974ebc 100644
--- a/test/tint/builtins/gen/var/distance/9646ea.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/distance/9646ea.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device float* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ float prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+float distance_9646ea() {
+ float4 arg_0 = float4(1.0f);
+ float4 arg_1 = float4(1.0f);
+ float res = distance(arg_0, arg_1);
+ return res;
}
-
-%distance_9646ea = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<f32>, read_write> = var, vec4<f32>(1.0f)
- %arg_1:ptr<function, vec4<f32>, read_write> = var, vec4<f32>(1.0f)
- %5:vec4<f32> = load %arg_0
- %6:vec4<f32> = load %arg_1
- %7:f32 = distance %5, %6
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_9646ea();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %distance_9646ea
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_9646ea();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %distance_9646ea
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = distance_9646ea();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %distance_9646ea
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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/builtins/gen/var/distance/aa4055.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/distance/aa4055.wgsl.expected.ir.msl
index a672973..71348d0 100644
--- a/test/tint/builtins/gen/var/distance/aa4055.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/distance/aa4055.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device float* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ float prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+float distance_aa4055() {
+ float2 arg_0 = float2(1.0f);
+ float2 arg_1 = float2(1.0f);
+ float res = distance(arg_0, arg_1);
+ return res;
}
-
-%distance_aa4055 = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %arg_1:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %5:vec2<f32> = load %arg_0
- %6:vec2<f32> = load %arg_1
- %7:f32 = distance %5, %6
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_aa4055();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %distance_aa4055
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_aa4055();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %distance_aa4055
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = distance_aa4055();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %distance_aa4055
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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/builtins/gen/var/distance/cfed73.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/distance/cfed73.wgsl.expected.ir.msl
index e238d18..50a492b 100644
--- a/test/tint/builtins/gen/var/distance/cfed73.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/distance/cfed73.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device float* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ float prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+float distance_cfed73() {
+ float arg_0 = 1.0f;
+ float arg_1 = 1.0f;
+ float res = abs((arg_0 - arg_1));
+ return res;
}
-
-%distance_cfed73 = func():void {
- $B2: {
- %arg_0:ptr<function, f32, read_write> = var, 1.0f
- %arg_1:ptr<function, f32, read_write> = var, 1.0f
- %5:f32 = load %arg_0
- %6:f32 = load %arg_1
- %7:f32 = distance %5, %6
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_cfed73();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %distance_cfed73
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = distance_cfed73();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %distance_cfed73
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = distance_cfed73();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %distance_cfed73
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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/samples/compute_boids.wgsl.expected.ir.msl b/test/tint/samples/compute_boids.wgsl.expected.ir.msl
index 6a0272f..cbf9c06 100644
--- a/test/tint/samples/compute_boids.wgsl.expected.ir.msl
+++ b/test/tint/samples/compute_boids.wgsl.expected.ir.msl
@@ -1,309 +1,147 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct SimParams {
+ float deltaT;
+ float rule1Distance;
+ float rule2Distance;
+ float rule3Distance;
+ float rule1Scale;
+ float rule2Scale;
+ float rule3Scale;
+};
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: SimParams = struct @align(4) {
- deltaT:f32 @offset(0)
- rule1Distance:f32 @offset(4)
- rule2Distance:f32 @offset(8)
- rule3Distance:f32 @offset(12)
- rule1Scale:f32 @offset(16)
- rule2Scale:f32 @offset(20)
- rule3Scale:f32 @offset(24)
+struct Particle {
+ float2 pos;
+ float2 vel;
+};
+struct Particles {
+ tint_array<Particle, 5> particles;
+};
+struct tint_module_vars_struct {
+ const constant SimParams* params;
+ device Particles* particlesA;
+ device Particles* particlesB;
+};
+struct vert_main_outputs {
+ float4 tint_symbol [[position]];
+};
+struct vert_main_inputs {
+ float2 a_particlePos [[attribute(0)]];
+ float2 a_particleVel [[attribute(1)]];
+ float2 a_pos [[attribute(2)]];
+};
+struct frag_main_outputs {
+ float4 tint_symbol_1 [[color(0)]];
+};
+
+float4 vert_main_inner(float2 a_particlePos, float2 a_particleVel, float2 a_pos) {
+ float angle = -(atan2(a_particleVel[0u], a_particleVel[1u]));
+ float const v = (a_pos[0u] * cos(angle));
+ float const v_1 = (v - (a_pos[1u] * sin(angle)));
+ float const v_2 = (a_pos[0u] * sin(angle));
+ float2 pos = float2(v_1, (v_2 + (a_pos[1u] * cos(angle))));
+ return float4((pos + a_particlePos), 0.0f, 1.0f);
}
-
-Particle = struct @align(8) {
- pos:vec2<f32> @offset(0)
- vel:vec2<f32> @offset(8)
+float4 frag_main_inner() {
+ return float4(1.0f);
}
-
-Particles = struct @align(8) {
- particles:array<Particle, 5> @offset(0)
-}
-
-$B1: { # root
- %params:ptr<uniform, SimParams, read> = var @binding_point(0, 0)
- %particlesA:ptr<storage, Particles, read_write> = var @binding_point(0, 1)
- %particlesB:ptr<storage, Particles, read_write> = var @binding_point(0, 2)
-}
-
-%vert_main = @vertex func(%a_particlePos:vec2<f32> [@location(0)], %a_particleVel:vec2<f32> [@location(1)], %a_pos:vec2<f32> [@location(2)]):vec4<f32> [@position] {
- $B2: {
- %8:f32 = access %a_particleVel, 0u
- %9:f32 = access %a_particleVel, 1u
- %10:f32 = atan2 %8, %9
- %11:f32 = negation %10
- %angle:ptr<function, f32, read_write> = var, %11
- %13:f32 = access %a_pos, 0u
- %14:f32 = load %angle
- %15:f32 = cos %14
- %16:f32 = mul %13, %15
- %17:f32 = let %16
- %18:f32 = access %a_pos, 1u
- %19:f32 = load %angle
- %20:f32 = sin %19
- %21:f32 = mul %18, %20
- %22:f32 = sub %17, %21
- %23:f32 = let %22
- %24:f32 = access %a_pos, 0u
- %25:f32 = load %angle
- %26:f32 = sin %25
- %27:f32 = mul %24, %26
- %28:f32 = let %27
- %29:f32 = access %a_pos, 1u
- %30:f32 = load %angle
- %31:f32 = cos %30
- %32:f32 = mul %29, %31
- %33:f32 = add %28, %32
- %34:vec2<f32> = construct %23, %33
- %pos:ptr<function, vec2<f32>, read_write> = var, %34
- %36:vec2<f32> = load %pos
- %37:vec2<f32> = add %36, %a_particlePos
- %38:vec4<f32> = construct %37, 0.0f, 1.0f
- ret %38
+void comp_main_inner(uint3 gl_GlobalInvocationID, tint_module_vars_struct tint_module_vars) {
+ uint index = gl_GlobalInvocationID[0u];
+ if ((index >= 5u)) {
+ return;
}
-}
-%frag_main = @fragment func():vec4<f32> [@location(0)] {
- $B3: {
- ret vec4<f32>(1.0f)
+ float2 vPos = (*tint_module_vars.particlesA).particles[index].pos;
+ float2 vVel = (*tint_module_vars.particlesA).particles[index].vel;
+ float2 cMass = float2(0.0f);
+ float2 cVel = float2(0.0f);
+ float2 colVel = float2(0.0f);
+ int cMassCount = 0;
+ int cVelCount = 0;
+ float2 pos = 0.0f;
+ float2 vel = 0.0f;
+ {
+ uint i = 0u;
+ while(true) {
+ if ((i < 5u)) {
+ } else {
+ break;
+ }
+ if ((i == index)) {
+ {
+ i = (i + 1u);
+ }
+ continue;
+ }
+ pos = (*tint_module_vars.particlesA).particles[i].pos.xy;
+ vel = (*tint_module_vars.particlesA).particles[i].vel.xy;
+ float const v_3 = distance(pos, vPos);
+ if ((v_3 < (*tint_module_vars.params).rule1Distance)) {
+ cMass = (cMass + pos);
+ cMassCount = (cMassCount + 1);
+ }
+ float const v_4 = distance(pos, vPos);
+ if ((v_4 < (*tint_module_vars.params).rule2Distance)) {
+ colVel = (colVel - (pos - vPos));
+ }
+ float const v_5 = distance(pos, vPos);
+ if ((v_5 < (*tint_module_vars.params).rule3Distance)) {
+ cVel = (cVel + vel);
+ cVelCount = (cVelCount + 1);
+ }
+ {
+ i = (i + 1u);
+ }
+ continue;
+ }
}
-}
-%comp_main = @compute @workgroup_size(1, 1, 1) func(%gl_GlobalInvocationID:vec3<u32> [@global_invocation_id]):void {
- $B4: {
- %42:u32 = access %gl_GlobalInvocationID, 0u
- %index:ptr<function, u32, read_write> = var, %42
- %44:u32 = load %index
- %45:bool = gte %44, 5u
- if %45 [t: $B5] { # if_1
- $B5: { # true
- ret
- }
- }
- %46:u32 = load %index
- %47:ptr<storage, vec2<f32>, read_write> = access %particlesA, 0u, %46, 0u
- %48:vec2<f32> = load %47
- %vPos:ptr<function, vec2<f32>, read_write> = var, %48
- %50:u32 = load %index
- %51:ptr<storage, vec2<f32>, read_write> = access %particlesA, 0u, %50, 1u
- %52:vec2<f32> = load %51
- %vVel:ptr<function, vec2<f32>, read_write> = var, %52
- %cMass:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(0.0f)
- %cVel:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(0.0f)
- %colVel:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(0.0f)
- %cMassCount:ptr<function, i32, read_write> = var, 0i
- %cVelCount:ptr<function, i32, read_write> = var, 0i
- %pos_1:ptr<function, vec2<f32>, read_write> = var # %pos_1: 'pos'
- %vel:ptr<function, vec2<f32>, read_write> = var
- loop [i: $B6, b: $B7, c: $B8] { # loop_1
- $B6: { # initializer
- %i:ptr<function, u32, read_write> = var, 0u
- next_iteration # -> $B7
- }
- $B7: { # body
- %62:u32 = load %i
- %63:bool = lt %62, 5u
- if %63 [t: $B9, f: $B10] { # if_2
- $B9: { # true
- exit_if # if_2
- }
- $B10: { # false
- exit_loop # loop_1
- }
- }
- %64:u32 = load %i
- %65:u32 = load %index
- %66:bool = eq %64, %65
- if %66 [t: $B11] { # if_3
- $B11: { # true
- continue # -> $B8
- }
- }
- %67:u32 = load %i
- %68:ptr<storage, vec2<f32>, read_write> = access %particlesA, 0u, %67, 0u
- %69:vec2<f32> = load %68
- %70:vec2<f32> = swizzle %69, xy
- store %pos_1, %70
- %71:u32 = load %i
- %72:ptr<storage, vec2<f32>, read_write> = access %particlesA, 0u, %71, 1u
- %73:vec2<f32> = load %72
- %74:vec2<f32> = swizzle %73, xy
- store %vel, %74
- %75:vec2<f32> = load %pos_1
- %76:vec2<f32> = load %vPos
- %77:f32 = distance %75, %76
- %78:ptr<uniform, f32, read> = access %params, 1u
- %79:f32 = load %78
- %80:bool = lt %77, %79
- if %80 [t: $B12] { # if_4
- $B12: { # true
- %81:vec2<f32> = load %cMass
- %82:vec2<f32> = load %pos_1
- %83:vec2<f32> = add %81, %82
- store %cMass, %83
- %84:i32 = load %cMassCount
- %85:i32 = add %84, 1i
- store %cMassCount, %85
- exit_if # if_4
- }
- }
- %86:vec2<f32> = load %pos_1
- %87:vec2<f32> = load %vPos
- %88:f32 = distance %86, %87
- %89:ptr<uniform, f32, read> = access %params, 2u
- %90:f32 = load %89
- %91:bool = lt %88, %90
- if %91 [t: $B13] { # if_5
- $B13: { # true
- %92:vec2<f32> = load %colVel
- %93:vec2<f32> = load %pos_1
- %94:vec2<f32> = load %vPos
- %95:vec2<f32> = sub %93, %94
- %96:vec2<f32> = sub %92, %95
- store %colVel, %96
- exit_if # if_5
- }
- }
- %97:vec2<f32> = load %pos_1
- %98:vec2<f32> = load %vPos
- %99:f32 = distance %97, %98
- %100:ptr<uniform, f32, read> = access %params, 3u
- %101:f32 = load %100
- %102:bool = lt %99, %101
- if %102 [t: $B14] { # if_6
- $B14: { # true
- %103:vec2<f32> = load %cVel
- %104:vec2<f32> = load %vel
- %105:vec2<f32> = add %103, %104
- store %cVel, %105
- %106:i32 = load %cVelCount
- %107:i32 = add %106, 1i
- store %cVelCount, %107
- exit_if # if_6
- }
- }
- continue # -> $B8
- }
- $B8: { # continuing
- %108:u32 = load %i
- %109:u32 = add %108, 1u
- store %i, %109
- next_iteration # -> $B7
- }
- }
- %110:i32 = load %cMassCount
- %111:bool = gt %110, 0i
- if %111 [t: $B15] { # if_7
- $B15: { # true
- %112:vec2<f32> = load %cMass
- %113:vec2<f32> = let %112
- %114:i32 = load %cMassCount
- %115:f32 = convert %114
- %116:f32 = let %115
- %117:i32 = load %cMassCount
- %118:f32 = convert %117
- %119:vec2<f32> = construct %116, %118
- %120:vec2<f32> = div %113, %119
- %121:vec2<f32> = load %vPos
- %122:vec2<f32> = sub %120, %121
- store %cMass, %122
- exit_if # if_7
- }
- }
- %123:i32 = load %cVelCount
- %124:bool = gt %123, 0i
- if %124 [t: $B16] { # if_8
- $B16: { # true
- %125:vec2<f32> = load %cVel
- %126:vec2<f32> = let %125
- %127:i32 = load %cVelCount
- %128:f32 = convert %127
- %129:f32 = let %128
- %130:i32 = load %cVelCount
- %131:f32 = convert %130
- %132:vec2<f32> = construct %129, %131
- %133:vec2<f32> = div %126, %132
- store %cVel, %133
- exit_if # if_8
- }
- }
- %134:vec2<f32> = load %vVel
- %135:vec2<f32> = load %cMass
- %136:ptr<uniform, f32, read> = access %params, 4u
- %137:f32 = load %136
- %138:vec2<f32> = mul %135, %137
- %139:vec2<f32> = add %134, %138
- %140:vec2<f32> = load %colVel
- %141:ptr<uniform, f32, read> = access %params, 5u
- %142:f32 = load %141
- %143:vec2<f32> = mul %140, %142
- %144:vec2<f32> = add %139, %143
- %145:vec2<f32> = load %cVel
- %146:ptr<uniform, f32, read> = access %params, 6u
- %147:f32 = load %146
- %148:vec2<f32> = mul %145, %147
- %149:vec2<f32> = add %144, %148
- store %vVel, %149
- %150:vec2<f32> = load %vVel
- %151:vec2<f32> = normalize %150
- %152:vec2<f32> = let %151
- %153:vec2<f32> = load %vVel
- %154:f32 = length %153
- %155:f32 = clamp %154, 0.0f, 0.10000000149011611938f
- %156:vec2<f32> = mul %152, %155
- store %vVel, %156
- %157:vec2<f32> = load %vPos
- %158:vec2<f32> = load %vVel
- %159:ptr<uniform, f32, read> = access %params, 0u
- %160:f32 = load %159
- %161:vec2<f32> = mul %158, %160
- %162:vec2<f32> = add %157, %161
- store %vPos, %162
- %163:f32 = load_vector_element %vPos, 0u
- %164:bool = lt %163, -1.0f
- if %164 [t: $B17] { # if_9
- $B17: { # true
- store_vector_element %vPos, 0u, 1.0f
- exit_if # if_9
- }
- }
- %165:f32 = load_vector_element %vPos, 0u
- %166:bool = gt %165, 1.0f
- if %166 [t: $B18] { # if_10
- $B18: { # true
- store_vector_element %vPos, 0u, -1.0f
- exit_if # if_10
- }
- }
- %167:f32 = load_vector_element %vPos, 1u
- %168:bool = lt %167, -1.0f
- if %168 [t: $B19] { # if_11
- $B19: { # true
- store_vector_element %vPos, 1u, 1.0f
- exit_if # if_11
- }
- }
- %169:f32 = load_vector_element %vPos, 1u
- %170:bool = gt %169, 1.0f
- if %170 [t: $B20] { # if_12
- $B20: { # true
- store_vector_element %vPos, 1u, -1.0f
- exit_if # if_12
- }
- }
- %171:u32 = load %index
- %172:ptr<storage, vec2<f32>, read_write> = access %particlesB, 0u, %171, 0u
- %173:vec2<f32> = load %vPos
- store %172, %173
- %174:u32 = load %index
- %175:ptr<storage, vec2<f32>, read_write> = access %particlesB, 0u, %174, 1u
- %176:vec2<f32> = load %vVel
- store %175, %176
- ret
+ if ((cMassCount > 0)) {
+ float2 const v_6 = cMass;
+ float const v_7 = float(cMassCount);
+ float2 const v_8 = (v_6 / float2(v_7, float(cMassCount)));
+ cMass = (v_8 - vPos);
}
+ if ((cVelCount > 0)) {
+ float2 const v_9 = cVel;
+ float const v_10 = float(cVelCount);
+ cVel = (v_9 / float2(v_10, float(cVelCount)));
+ }
+ vVel = (((vVel + (cMass * (*tint_module_vars.params).rule1Scale)) + (colVel * (*tint_module_vars.params).rule2Scale)) + (cVel * (*tint_module_vars.params).rule3Scale));
+ float2 const v_11 = normalize(vVel);
+ vVel = (v_11 * clamp(length(vVel), 0.0f, 0.10000000149011611938f));
+ vPos = (vPos + (vVel * (*tint_module_vars.params).deltaT));
+ if ((vPos[0u] < -1.0f)) {
+ vPos[0u] = 1.0f;
+ }
+ if ((vPos[0u] > 1.0f)) {
+ vPos[0u] = -1.0f;
+ }
+ if ((vPos[1u] < -1.0f)) {
+ vPos[1u] = 1.0f;
+ }
+ if ((vPos[1u] > 1.0f)) {
+ vPos[1u] = -1.0f;
+ }
+ (*tint_module_vars.particlesB).particles[index].pos = vPos;
+ (*tint_module_vars.particlesB).particles[index].vel = vVel;
}
-
-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. *
-********************************************************************
+vertex vert_main_outputs vert_main(vert_main_inputs inputs [[stage_in]]) {
+ return vert_main_outputs{.tint_symbol=vert_main_inner(inputs.a_particlePos, inputs.a_particleVel, inputs.a_pos)};
+}
+fragment frag_main_outputs frag_main() {
+ return frag_main_outputs{.tint_symbol_1=frag_main_inner()};
+}
+kernel void comp_main(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], const constant SimParams* params [[buffer(0)]], device Particles* particlesA [[buffer(1)]], device Particles* particlesB [[buffer(2)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.params=params, .particlesA=particlesA, .particlesB=particlesB};
+ comp_main_inner(gl_GlobalInvocationID, tint_module_vars);
+}
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-acos-ldexp/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-acos-ldexp/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index 3621034..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-acos-ldexp/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,59 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(4) {
- two:i32 @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_6:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %v:ptr<function, vec2<f32>, read_write> = var
- %d:ptr<function, f32, read_write> = var
- %6:ptr<uniform, i32, read> = access %x_6, 0u
- %7:i32 = load %6
- %8:vec2<i32> = construct %7, 3i
- %9:vec2<f32> = ldexp vec2<f32>(0.10000000149011611938f), %8
- %10:vec2<f32> = acos %9
- store %v, %10
- %11:vec2<f32> = load %v
- %12:f32 = distance %11, vec2<f32>(1.15927994251251220703f, 0.64349997043609619141f)
- store %d, %12
- %13:f32 = load %d
- %14:bool = lt %13, 0.00999999977648258209f
- if %14 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- store %x_GLF_color, vec4<f32>(1.0f, 0.0f, 0.0f, 1.0f)
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %16:void = call %main_1
- %17:vec4<f32> = load %x_GLF_color
- %18:main_out = construct %17
- ret %18
- }
-}
-
-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/graphicsfuzz/cov-apfloat-acos-ldexp/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-acos-ldexp/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index 4e0bafd..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-acos-ldexp/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,62 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(4) {
- two:i32 @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_6:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %v:ptr<function, vec2<f32>, read_write> = var
- %d:ptr<function, f32, read_write> = var
- %6:ptr<uniform, i32, read> = access %x_6, 0u
- %7:i32 = load %6
- %x_35:i32 = let %7
- %9:vec2<i32> = construct %x_35, 3i
- %10:vec2<f32> = ldexp vec2<f32>(0.10000000149011611938f), %9
- %11:vec2<f32> = acos %10
- store %v, %11
- %12:vec2<f32> = load %v
- %x_39:vec2<f32> = let %12
- %14:f32 = distance %x_39, vec2<f32>(1.15927994251251220703f, 0.64349997043609619141f)
- store %d, %14
- %15:f32 = load %d
- %x_41:f32 = let %15
- %17:bool = lt %x_41, 0.00999999977648258209f
- if %17 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- store %x_GLF_color, vec4<f32>(1.0f, 0.0f, 0.0f, 1.0f)
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %19:void = call %main_1
- %20:vec4<f32> = load %x_GLF_color
- %21:main_out = construct %20
- ret %21
- }
-}
-
-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/graphicsfuzz/cov-apfloat-tanh/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-tanh/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index 71c5ec1..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-tanh/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,9 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: distance
-********************************************************************
-* 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/graphicsfuzz/cov-apfloat-tanh/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-tanh/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index 71c5ec1..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-tanh/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,9 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: distance
-********************************************************************
-* 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/graphicsfuzz/cov-folding-rules-mix-uniform-weight/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-mix-uniform-weight/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index 9f6eba1..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-mix-uniform-weight/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,60 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(4) {
- one:f32 @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_6:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %v:ptr<function, vec3<f32>, read_write> = var
- %d:ptr<function, f32, read_write> = var
- %6:ptr<uniform, f32, read> = access %x_6, 0u
- %7:f32 = load %6
- %8:vec3<f32> = construct %7
- %9:vec3<f32> = mix vec3<f32>(5.0f, 8.0f, -12.19999980926513671875f), vec3<f32>(1.0f, 4.90000009536743164062f, -2.09999990463256835938f), %8
- store %v, %9
- %10:vec3<f32> = load %v
- %11:f32 = distance %10, vec3<f32>(1.0f, 4.90000009536743164062f, -2.09999990463256835938f)
- store %d, %11
- %12:f32 = load %d
- %13:bool = lt %12, 0.10000000149011611938f
- if %13 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- %14:f32 = load_vector_element %v, 0u
- %15:vec4<f32> = construct %14, 0.0f, 0.0f, 1.0f
- store %x_GLF_color, %15
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %17:void = call %main_1
- %18:vec4<f32> = load %x_GLF_color
- %19:main_out = construct %18
- ret %19
- }
-}
-
-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/graphicsfuzz/cov-folding-rules-mix-uniform-weight/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-mix-uniform-weight/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index 601ba78..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-mix-uniform-weight/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,64 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(4) {
- one:f32 @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_6:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %v:ptr<function, vec3<f32>, read_write> = var
- %d:ptr<function, f32, read_write> = var
- %6:ptr<uniform, f32, read> = access %x_6, 0u
- %7:f32 = load %6
- %x_36:f32 = let %7
- %9:vec3<f32> = construct %x_36, %x_36, %x_36
- %10:vec3<f32> = mix vec3<f32>(5.0f, 8.0f, -12.19999980926513671875f), vec3<f32>(1.0f, 4.90000009536743164062f, -2.09999990463256835938f), %9
- store %v, %10
- %11:vec3<f32> = load %v
- %x_39:vec3<f32> = let %11
- %13:f32 = distance %x_39, vec3<f32>(1.0f, 4.90000009536743164062f, -2.09999990463256835938f)
- store %d, %13
- %14:f32 = load %d
- %x_41:f32 = let %14
- %16:bool = lt %x_41, 0.10000000149011611938f
- if %16 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- %17:f32 = load_vector_element %v, 0u
- %x_47:f32 = let %17
- %19:vec4<f32> = construct %x_47, 0.0f, 0.0f, 1.0f
- store %x_GLF_color, %19
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %21:void = call %main_1
- %22:vec4<f32> = load %x_GLF_color
- %23:main_out = construct %22
- ret %23
- }
-}
-
-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/graphicsfuzz/cov-folding-rules-vec-mix-uniform/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-vec-mix-uniform/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index 56f034c..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-vec-mix-uniform/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,62 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(8) {
- zeroOne:vec2<f32> @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_6:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %v:ptr<function, vec2<f32>, read_write> = var
- %d:ptr<function, f32, read_write> = var
- %6:ptr<uniform, vec2<f32>, read> = access %x_6, 0u
- %7:vec2<f32> = load %6
- %8:vec2<f32> = mix vec2<f32>(2.0f, 3.0f), vec2<f32>(4.0f, 5.0f), %7
- store %v, %8
- %9:vec2<f32> = load %v
- %10:f32 = distance %9, vec2<f32>(2.0f, 5.0f)
- store %d, %10
- %11:f32 = load %d
- %12:bool = lt %11, 0.10000000149011611938f
- if %12 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- %13:f32 = load_vector_element %v, 0u
- %14:f32 = sub %13, 1.0f
- %15:f32 = load_vector_element %v, 1u
- %16:f32 = sub %15, 5.0f
- %17:vec4<f32> = construct %14, %16, 0.0f, 1.0f
- store %x_GLF_color, %17
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %19:void = call %main_1
- %20:vec4<f32> = load %x_GLF_color
- %21:main_out = construct %20
- ret %21
- }
-}
-
-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/graphicsfuzz/cov-folding-rules-vec-mix-uniform/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-vec-mix-uniform/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index a74ee08..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-vec-mix-uniform/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,67 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(8) {
- zeroOne:vec2<f32> @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_6:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %v:ptr<function, vec2<f32>, read_write> = var
- %d:ptr<function, f32, read_write> = var
- %6:ptr<uniform, vec2<f32>, read> = access %x_6, 0u
- %7:vec2<f32> = load %6
- %x_37:vec2<f32> = let %7
- %9:vec2<f32> = mix vec2<f32>(2.0f, 3.0f), vec2<f32>(4.0f, 5.0f), %x_37
- store %v, %9
- %10:vec2<f32> = load %v
- %x_39:vec2<f32> = let %10
- %12:f32 = distance %x_39, vec2<f32>(2.0f, 5.0f)
- store %d, %12
- %13:f32 = load %d
- %x_41:f32 = let %13
- %15:bool = lt %x_41, 0.10000000149011611938f
- if %15 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- %16:f32 = load_vector_element %v, 0u
- %x_47:f32 = let %16
- %18:f32 = load_vector_element %v, 1u
- %x_50:f32 = let %18
- %20:f32 = sub %x_47, 1.0f
- %21:f32 = sub %x_50, 5.0f
- %22:vec4<f32> = construct %20, %21, 0.0f, 1.0f
- store %x_GLF_color, %22
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %24:void = call %main_1
- %25:vec4<f32> = load %x_GLF_color
- %26:main_out = construct %25
- ret %26
- }
-}
-
-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/graphicsfuzz/cov-simplify-mul-identity/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-simplify-mul-identity/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index bf16665..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-simplify-mul-identity/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,67 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(4) {
- one:f32 @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_7:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %v:ptr<function, vec4<f32>, read_write> = var
- %res:ptr<function, vec4<f32>, read_write> = var
- store %v, vec4<f32>(8.3999996185302734375f, -864.66497802734375f, 945.41998291015625f, 1.0f)
- %6:ptr<uniform, f32, read> = access %x_7, 0u
- %7:f32 = load %6
- %x_31:f32 = let %7
- %9:vec4<f32> = construct %x_31, 0.0f, 0.0f, 0.0f
- %10:vec4<f32> = let %9
- %11:vec4<f32> = construct 0.0f, %x_31, 0.0f, 0.0f
- %12:vec4<f32> = let %11
- %13:vec4<f32> = construct 0.0f, 0.0f, %x_31, 0.0f
- %14:vec4<f32> = let %13
- %15:vec4<f32> = construct 0.0f, 0.0f, 0.0f, %x_31
- %16:mat4x4<f32> = construct %10, %12, %14, %15
- %17:vec4<f32> = load %v
- %18:vec4<f32> = mul %16, %17
- store %res, %18
- %19:vec4<f32> = load %v
- %20:vec4<f32> = load %res
- %21:f32 = distance %19, %20
- %22:bool = lt %21, 0.00999999977648258209f
- if %22 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- store %x_GLF_color, vec4<f32>(1.0f, 0.0f, 0.0f, 1.0f)
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %24:void = call %main_1
- %25:vec4<f32> = load %x_GLF_color
- %26:main_out = construct %25
- ret %26
- }
-}
-
-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/graphicsfuzz/cov-simplify-mul-identity/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-simplify-mul-identity/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index ea8cf6d..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-simplify-mul-identity/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,70 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(4) {
- one:f32 @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_7:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %v:ptr<function, vec4<f32>, read_write> = var
- %res:ptr<function, vec4<f32>, read_write> = var
- store %v, vec4<f32>(8.3999996185302734375f, -864.66497802734375f, 945.41998291015625f, 1.0f)
- %6:ptr<uniform, f32, read> = access %x_7, 0u
- %7:f32 = load %6
- %x_31:f32 = let %7
- %9:vec4<f32> = load %v
- %x_37:vec4<f32> = let %9
- %11:vec4<f32> = construct %x_31, 0.0f, 0.0f, 0.0f
- %12:vec4<f32> = let %11
- %13:vec4<f32> = construct 0.0f, %x_31, 0.0f, 0.0f
- %14:vec4<f32> = let %13
- %15:vec4<f32> = construct 0.0f, 0.0f, %x_31, 0.0f
- %16:vec4<f32> = let %15
- %17:vec4<f32> = construct 0.0f, 0.0f, 0.0f, %x_31
- %18:mat4x4<f32> = construct %12, %14, %16, %17
- %19:vec4<f32> = mul %18, %x_37
- store %res, %19
- %20:vec4<f32> = load %v
- %x_39:vec4<f32> = let %20
- %22:vec4<f32> = load %res
- %x_40:vec4<f32> = let %22
- %24:f32 = distance %x_39, %x_40
- %25:bool = lt %24, 0.00999999977648258209f
- if %25 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- store %x_GLF_color, vec4<f32>(1.0f, 0.0f, 0.0f, 1.0f)
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %27:void = call %main_1
- %28:vec4<f32> = load %x_GLF_color
- %29:main_out = construct %28
- ret %29
- }
-}
-
-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/graphicsfuzz/cov-struct-float-array-mix-uniform-vectors/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-struct-float-array-mix-uniform-vectors/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index 4d5c300..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-struct-float-array-mix-uniform-vectors/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,161 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: strided_arr = struct @align(4) {
- el:f32 @offset(0)
-}
-
-buf1 = struct @align(4) {
- x_GLF_uniform_float_values:array<strided_arr, 5> @offset(0)
-}
-
-buf2 = struct @align(8) {
- zeroVec:vec2<f32> @offset(0)
-}
-
-buf3 = struct @align(8) {
- oneVec:vec2<f32> @offset(0)
-}
-
-strided_arr_1 = struct @align(4) {
- el:i32 @offset(0)
-}
-
-buf0 = struct @align(4) {
- x_GLF_uniform_int_values:array<strided_arr_1, 2> @offset(0)
-}
-
-S = struct @align(4) {
- numbers:array<f32, 3> @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_7:ptr<uniform, buf1, read> = var @binding_point(0, 1)
- %x_9:ptr<uniform, buf2, read> = var @binding_point(0, 2)
- %x_12:ptr<uniform, buf3, read> = var @binding_point(0, 3)
- %x_15:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %obj:ptr<function, S, read_write> = var
- %a:ptr<function, f32, read_write> = var
- %x_49:ptr<function, vec2<f32>, read_write> = var
- %b:ptr<function, f32, read_write> = var
- %11:ptr<uniform, f32, read> = access %x_7, 0u, 3i, 0u
- %12:f32 = load %11
- %13:ptr<uniform, f32, read> = access %x_7, 0u, 2i, 0u
- %14:f32 = load %13
- %15:ptr<uniform, f32, read> = access %x_7, 0u, 4i, 0u
- %16:f32 = load %15
- %17:array<f32, 3> = construct %12, %14, %16
- %18:S = construct %17
- store %obj, %18
- %19:ptr<uniform, vec2<f32>, read> = access %x_9, 0u
- %20:f32 = load_vector_element %19, 0u
- %x_59:f32 = let %20
- %22:i32 = call %tint_f32_to_i32, %x_59
- %24:ptr<function, f32, read_write> = access %obj, 0u, %22
- %25:ptr<uniform, f32, read> = access %x_7, 0u, 0i, 0u
- %26:f32 = load %25
- store %24, %26
- %27:ptr<uniform, vec2<f32>, read> = access %x_9, 0u
- %28:f32 = load_vector_element %27, 0u
- %29:ptr<uniform, f32, read> = access %x_7, 0u, 0i, 0u
- %30:f32 = load %29
- %31:bool = gt %28, %30
- if %31 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- %32:ptr<uniform, vec2<f32>, read> = access %x_9, 0u
- %33:vec2<f32> = load %32
- store %x_49, %33
- exit_if # if_1
- }
- $B4: { # false
- %34:ptr<uniform, vec2<f32>, read> = access %x_12, 0u
- %35:vec2<f32> = load %34
- store %x_49, %35
- exit_if # if_1
- }
- }
- %36:f32 = load_vector_element %x_49, 1u
- store %a, %36
- %37:ptr<uniform, f32, read> = access %x_7, 0u, 0i, 0u
- %38:f32 = load %37
- %39:f32 = load %a
- %40:ptr<uniform, i32, read> = access %x_15, 0u, 0i, 0u
- %41:i32 = load %40
- %42:ptr<function, f32, read_write> = access %obj, 0u, %41
- %43:f32 = load %42
- %44:f32 = mix %38, %39, %43
- store %b, %44
- %45:f32 = load %b
- %46:ptr<uniform, f32, read> = access %x_7, 0u, 2i, 0u
- %47:f32 = load %46
- %48:f32 = distance %45, %47
- %49:ptr<uniform, f32, read> = access %x_7, 0u, 1i, 0u
- %50:f32 = load %49
- %51:bool = lt %48, %50
- if %51 [t: $B5, f: $B6] { # if_2
- $B5: { # true
- %52:ptr<uniform, i32, read> = access %x_15, 0u, 0i, 0u
- %53:i32 = load %52
- %54:f32 = convert %53
- %55:f32 = let %54
- %56:ptr<uniform, i32, read> = access %x_15, 0u, 1i, 0u
- %57:i32 = load %56
- %58:f32 = convert %57
- %59:f32 = let %58
- %60:ptr<uniform, i32, read> = access %x_15, 0u, 1i, 0u
- %61:i32 = load %60
- %62:f32 = convert %61
- %63:f32 = let %62
- %64:ptr<uniform, i32, read> = access %x_15, 0u, 0i, 0u
- %65:i32 = load %64
- %66:f32 = convert %65
- %67:vec4<f32> = construct %55, %59, %63, %66
- store %x_GLF_color, %67
- exit_if # if_2
- }
- $B6: { # false
- %68:ptr<uniform, i32, read> = access %x_15, 0u, 1i, 0u
- %69:i32 = load %68
- %70:f32 = convert %69
- %71:vec4<f32> = construct %70
- store %x_GLF_color, %71
- exit_if # if_2
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B7: {
- %73:void = call %main_1
- %74:vec4<f32> = load %x_GLF_color
- %75:main_out = construct %74
- ret %75
- }
-}
-%tint_f32_to_i32 = func(%value:f32):i32 {
- $B8: {
- %77:i32 = convert %value
- %78:bool = gte %value, -2147483648.0f
- %79:i32 = select -2147483648i, %77, %78
- %80:bool = lte %value, 2147483520.0f
- %81:i32 = select 2147483647i, %79, %80
- ret %81
- }
-}
-
-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/graphicsfuzz/cov-struct-float-array-mix-uniform-vectors/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-struct-float-array-mix-uniform-vectors/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index 4c0c448..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-struct-float-array-mix-uniform-vectors/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,183 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: strided_arr = struct @align(4) {
- el:f32 @offset(0)
-}
-
-buf1 = struct @align(4) {
- x_GLF_uniform_float_values:array<strided_arr, 5> @offset(0)
-}
-
-buf2 = struct @align(8) {
- zeroVec:vec2<f32> @offset(0)
-}
-
-buf3 = struct @align(8) {
- oneVec:vec2<f32> @offset(0)
-}
-
-strided_arr_1 = struct @align(4) {
- el:i32 @offset(0)
-}
-
-buf0 = struct @align(4) {
- x_GLF_uniform_int_values:array<strided_arr_1, 2> @offset(0)
-}
-
-S = struct @align(4) {
- numbers:array<f32, 3> @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_7:ptr<uniform, buf1, read> = var @binding_point(0, 1)
- %x_9:ptr<uniform, buf2, read> = var @binding_point(0, 2)
- %x_12:ptr<uniform, buf3, read> = var @binding_point(0, 3)
- %x_15:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %obj:ptr<function, S, read_write> = var
- %a:ptr<function, f32, read_write> = var
- %x_49:ptr<function, vec2<f32>, read_write> = var
- %b:ptr<function, f32, read_write> = var
- %11:ptr<uniform, f32, read> = access %x_7, 0u, 3i, 0u
- %12:f32 = load %11
- %x_51:f32 = let %12
- %14:ptr<uniform, f32, read> = access %x_7, 0u, 2i, 0u
- %15:f32 = load %14
- %x_53:f32 = let %15
- %17:ptr<uniform, f32, read> = access %x_7, 0u, 4i, 0u
- %18:f32 = load %17
- %x_55:f32 = let %18
- %20:array<f32, 3> = construct %x_51, %x_53, %x_55
- %21:S = construct %20
- store %obj, %21
- %22:ptr<uniform, vec2<f32>, read> = access %x_9, 0u
- %23:f32 = load_vector_element %22, 0u
- %x_59:f32 = let %23
- %25:ptr<uniform, f32, read> = access %x_7, 0u, 0i, 0u
- %26:f32 = load %25
- %x_62:f32 = let %26
- %28:i32 = call %tint_f32_to_i32, %x_59
- %30:ptr<function, f32, read_write> = access %obj, 0u, %28
- store %30, %x_62
- %31:ptr<uniform, vec2<f32>, read> = access %x_9, 0u
- %32:f32 = load_vector_element %31, 0u
- %x_65:f32 = let %32
- %34:ptr<uniform, f32, read> = access %x_7, 0u, 0i, 0u
- %35:f32 = load %34
- %x_67:f32 = let %35
- %37:bool = gt %x_65, %x_67
- if %37 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- %38:ptr<uniform, vec2<f32>, read> = access %x_9, 0u
- %39:vec2<f32> = load %38
- %x_73:vec2<f32> = let %39
- store %x_49, %x_73
- exit_if # if_1
- }
- $B4: { # false
- %41:ptr<uniform, vec2<f32>, read> = access %x_12, 0u
- %42:vec2<f32> = load %41
- %x_75:vec2<f32> = let %42
- store %x_49, %x_75
- exit_if # if_1
- }
- }
- %44:f32 = load_vector_element %x_49, 1u
- %x_77:f32 = let %44
- store %a, %x_77
- %46:ptr<uniform, f32, read> = access %x_7, 0u, 0i, 0u
- %47:f32 = load %46
- %x_79:f32 = let %47
- %49:f32 = load %a
- %x_80:f32 = let %49
- %51:ptr<uniform, i32, read> = access %x_15, 0u, 0i, 0u
- %52:i32 = load %51
- %x_82:i32 = let %52
- %54:ptr<function, f32, read_write> = access %obj, 0u, %x_82
- %55:f32 = load %54
- %x_84:f32 = let %55
- %57:f32 = mix %x_79, %x_80, %x_84
- store %b, %57
- %58:f32 = load %b
- %x_86:f32 = let %58
- %60:ptr<uniform, f32, read> = access %x_7, 0u, 2i, 0u
- %61:f32 = load %60
- %x_88:f32 = let %61
- %63:ptr<uniform, f32, read> = access %x_7, 0u, 1i, 0u
- %64:f32 = load %63
- %x_91:f32 = let %64
- %66:f32 = distance %x_86, %x_88
- %67:bool = lt %66, %x_91
- if %67 [t: $B5, f: $B6] { # if_2
- $B5: { # true
- %68:ptr<uniform, i32, read> = access %x_15, 0u, 0i, 0u
- %69:i32 = load %68
- %x_97:i32 = let %69
- %71:ptr<uniform, i32, read> = access %x_15, 0u, 1i, 0u
- %72:i32 = load %71
- %x_100:i32 = let %72
- %74:ptr<uniform, i32, read> = access %x_15, 0u, 1i, 0u
- %75:i32 = load %74
- %x_103:i32 = let %75
- %77:ptr<uniform, i32, read> = access %x_15, 0u, 0i, 0u
- %78:i32 = load %77
- %x_106:i32 = let %78
- %80:f32 = convert %x_97
- %81:f32 = let %80
- %82:f32 = convert %x_100
- %83:f32 = let %82
- %84:f32 = convert %x_103
- %85:f32 = let %84
- %86:f32 = convert %x_106
- %87:vec4<f32> = construct %81, %83, %85, %86
- store %x_GLF_color, %87
- exit_if # if_2
- }
- $B6: { # false
- %88:ptr<uniform, i32, read> = access %x_15, 0u, 1i, 0u
- %89:i32 = load %88
- %x_110:i32 = let %89
- %91:f32 = convert %x_110
- %x_111:f32 = let %91
- %93:vec4<f32> = construct %x_111, %x_111, %x_111, %x_111
- store %x_GLF_color, %93
- exit_if # if_2
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B7: {
- %95:void = call %main_1
- %96:vec4<f32> = load %x_GLF_color
- %97:main_out = construct %96
- ret %97
- }
-}
-%tint_f32_to_i32 = func(%value:f32):i32 {
- $B8: {
- %99:i32 = convert %value
- %100:bool = gte %value, -2147483648.0f
- %101:i32 = select -2147483648i, %99, %100
- %102:bool = lte %value, 2147483520.0f
- %103:i32 = select 2147483647i, %101, %102
- ret %103
- }
-}
-
-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/graphicsfuzz/cov-unpack-unorm-mix-always-one/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-unpack-unorm-mix-always-one/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index ae7455e..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-unpack-unorm-mix-always-one/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,129 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: strided_arr = struct @align(4) {
- el:u32 @offset(0)
-}
-
-buf0 = struct @align(4) {
- x_GLF_uniform_uint_values:array<strided_arr, 1> @offset(0)
-}
-
-strided_arr_1 = struct @align(4) {
- el:i32 @offset(0)
-}
-
-buf1 = struct @align(4) {
- x_GLF_uniform_int_values:array<strided_arr_1, 2> @offset(0)
-}
-
-strided_arr_2 = struct @align(4) {
- el:f32 @offset(0)
-}
-
-buf2 = struct @align(4) {
- x_GLF_uniform_float_values:array<strided_arr_2, 3> @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_6:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_8:ptr<uniform, buf1, read> = var @binding_point(0, 1)
- %x_10:ptr<uniform, buf2, read> = var @binding_point(0, 2)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %v:ptr<function, vec4<f32>, read_write> = var
- %7:ptr<uniform, u32, read> = access %x_6, 0u, 0i, 0u
- %8:u32 = load %7
- %9:u32 = let %8
- %10:ptr<uniform, u32, read> = access %x_6, 0u, 0i, 0u
- %11:u32 = load %10
- %12:u32 = select %11, 92382u, true
- %13:u32 = call %tint_div_u32, %9, %12
- %15:vec4<f32> = unpack4x8unorm %13
- store %v, %15
- %16:vec4<f32> = load %v
- %17:vec4<f32> = let %16
- %18:ptr<uniform, i32, read> = access %x_8, 0u, 0i, 0u
- %19:i32 = load %18
- %20:f32 = convert %19
- %21:f32 = let %20
- %22:ptr<uniform, i32, read> = access %x_8, 0u, 0i, 0u
- %23:i32 = load %22
- %24:f32 = convert %23
- %25:f32 = let %24
- %26:ptr<uniform, i32, read> = access %x_8, 0u, 0i, 0u
- %27:i32 = load %26
- %28:f32 = convert %27
- %29:ptr<uniform, f32, read> = access %x_10, 0u, 1i, 0u
- %30:f32 = load %29
- %31:ptr<uniform, f32, read> = access %x_10, 0u, 2i, 0u
- %32:f32 = load %31
- %33:f32 = div %30, %32
- %34:vec4<f32> = construct %21, %25, %28, %33
- %35:f32 = distance %17, %34
- %36:ptr<uniform, f32, read> = access %x_10, 0u, 0i, 0u
- %37:f32 = load %36
- %38:bool = lt %35, %37
- if %38 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- %39:ptr<uniform, i32, read> = access %x_8, 0u, 1i, 0u
- %40:i32 = load %39
- %41:f32 = convert %40
- %42:f32 = let %41
- %43:ptr<uniform, i32, read> = access %x_8, 0u, 0i, 0u
- %44:i32 = load %43
- %45:f32 = convert %44
- %46:f32 = let %45
- %47:ptr<uniform, i32, read> = access %x_8, 0u, 0i, 0u
- %48:i32 = load %47
- %49:f32 = convert %48
- %50:f32 = let %49
- %51:ptr<uniform, i32, read> = access %x_8, 0u, 1i, 0u
- %52:i32 = load %51
- %53:f32 = convert %52
- %54:vec4<f32> = construct %42, %46, %50, %53
- store %x_GLF_color, %54
- exit_if # if_1
- }
- $B4: { # false
- %55:ptr<uniform, i32, read> = access %x_8, 0u, 0i, 0u
- %56:i32 = load %55
- %57:f32 = convert %56
- %58:vec4<f32> = construct %57
- store %x_GLF_color, %58
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %60:void = call %main_1
- %61:vec4<f32> = load %x_GLF_color
- %62:main_out = construct %61
- ret %62
- }
-}
-%tint_div_u32 = func(%lhs:u32, %rhs:u32):u32 {
- $B6: {
- %65:bool = eq %rhs, 0u
- %66:u32 = select %rhs, 1u, %65
- %67:u32 = div %lhs, %66
- ret %67
- }
-}
-
-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/graphicsfuzz/cov-unpack-unorm-mix-always-one/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-unpack-unorm-mix-always-one/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index 5daef5c..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-unpack-unorm-mix-always-one/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,142 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: strided_arr = struct @align(4) {
- el:u32 @offset(0)
-}
-
-buf0 = struct @align(4) {
- x_GLF_uniform_uint_values:array<strided_arr, 1> @offset(0)
-}
-
-strided_arr_1 = struct @align(4) {
- el:i32 @offset(0)
-}
-
-buf1 = struct @align(4) {
- x_GLF_uniform_int_values:array<strided_arr_1, 2> @offset(0)
-}
-
-strided_arr_2 = struct @align(4) {
- el:f32 @offset(0)
-}
-
-buf2 = struct @align(4) {
- x_GLF_uniform_float_values:array<strided_arr_2, 3> @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_6:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_8:ptr<uniform, buf1, read> = var @binding_point(0, 1)
- %x_10:ptr<uniform, buf2, read> = var @binding_point(0, 2)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %v:ptr<function, vec4<f32>, read_write> = var
- %7:ptr<uniform, u32, read> = access %x_6, 0u, 0i, 0u
- %8:u32 = load %7
- %x_39:u32 = let %8
- %10:ptr<uniform, u32, read> = access %x_6, 0u, 0i, 0u
- %11:u32 = load %10
- %x_41:u32 = let %11
- %13:u32 = select %x_41, 92382u, true
- %14:u32 = call %tint_div_u32, %x_39, %13
- %16:vec4<f32> = unpack4x8unorm %14
- store %v, %16
- %17:vec4<f32> = load %v
- %x_45:vec4<f32> = let %17
- %19:ptr<uniform, i32, read> = access %x_8, 0u, 0i, 0u
- %20:i32 = load %19
- %x_47:i32 = let %20
- %22:ptr<uniform, i32, read> = access %x_8, 0u, 0i, 0u
- %23:i32 = load %22
- %x_50:i32 = let %23
- %25:ptr<uniform, i32, read> = access %x_8, 0u, 0i, 0u
- %26:i32 = load %25
- %x_53:i32 = let %26
- %28:ptr<uniform, f32, read> = access %x_10, 0u, 1i, 0u
- %29:f32 = load %28
- %x_56:f32 = let %29
- %31:ptr<uniform, f32, read> = access %x_10, 0u, 2i, 0u
- %32:f32 = load %31
- %x_58:f32 = let %32
- %34:ptr<uniform, f32, read> = access %x_10, 0u, 0i, 0u
- %35:f32 = load %34
- %x_63:f32 = let %35
- %37:f32 = convert %x_47
- %38:f32 = let %37
- %39:f32 = convert %x_50
- %40:f32 = let %39
- %41:f32 = convert %x_53
- %42:f32 = div %x_56, %x_58
- %43:vec4<f32> = construct %38, %40, %41, %42
- %44:f32 = distance %x_45, %43
- %45:bool = lt %44, %x_63
- if %45 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- %46:ptr<uniform, i32, read> = access %x_8, 0u, 1i, 0u
- %47:i32 = load %46
- %x_69:i32 = let %47
- %49:ptr<uniform, i32, read> = access %x_8, 0u, 0i, 0u
- %50:i32 = load %49
- %x_72:i32 = let %50
- %52:ptr<uniform, i32, read> = access %x_8, 0u, 0i, 0u
- %53:i32 = load %52
- %x_75:i32 = let %53
- %55:ptr<uniform, i32, read> = access %x_8, 0u, 1i, 0u
- %56:i32 = load %55
- %x_78:i32 = let %56
- %58:f32 = convert %x_69
- %59:f32 = let %58
- %60:f32 = convert %x_72
- %61:f32 = let %60
- %62:f32 = convert %x_75
- %63:f32 = let %62
- %64:f32 = convert %x_78
- %65:vec4<f32> = construct %59, %61, %63, %64
- store %x_GLF_color, %65
- exit_if # if_1
- }
- $B4: { # false
- %66:ptr<uniform, i32, read> = access %x_8, 0u, 0i, 0u
- %67:i32 = load %66
- %x_82:i32 = let %67
- %69:f32 = convert %x_82
- %x_83:f32 = let %69
- %71:vec4<f32> = construct %x_83, %x_83, %x_83, %x_83
- store %x_GLF_color, %71
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %73:void = call %main_1
- %74:vec4<f32> = load %x_GLF_color
- %75:main_out = construct %74
- ret %75
- }
-}
-%tint_div_u32 = func(%lhs:u32, %rhs:u32):u32 {
- $B6: {
- %78:bool = eq %rhs, 0u
- %79:u32 = select %rhs, 1u, %78
- %80:u32 = div %lhs, %79
- ret %80
- }
-}
-
-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. *
-********************************************************************
