[msl] Add polyfills for degrees() and radians()
Use the generic polyfill transform to replace these with
multiplications.
Bug: 42251016
Change-Id: I1e12cdb392e147656b620b82cbf1d82509b16c5d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/195215
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/core/ir/transform/builtin_polyfill.cc b/src/tint/lang/core/ir/transform/builtin_polyfill.cc
index 6547d36..b3ea20f 100644
--- a/src/tint/lang/core/ir/transform/builtin_polyfill.cc
+++ b/src/tint/lang/core/ir/transform/builtin_polyfill.cc
@@ -41,6 +41,12 @@
namespace {
+/// Constant value used to polyfill the radians() builtin.
+static constexpr double kDegToRad = 0.017453292519943295474;
+
+/// Constant value used to polyfill the degrees() builtin.
+static constexpr double kRadToDeg = 57.295779513082322865;
+
/// PIMPL state for the transform.
struct State {
/// The polyfill config.
@@ -81,6 +87,11 @@
worklist.Push(builtin);
}
break;
+ case core::BuiltinFn::kDegrees:
+ if (config.degrees) {
+ worklist.Push(builtin);
+ }
+ break;
case core::BuiltinFn::kExtractBits:
if (config.extract_bits != BuiltinPolyfillLevel::kNone) {
worklist.Push(builtin);
@@ -101,6 +112,11 @@
worklist.Push(builtin);
}
break;
+ case core::BuiltinFn::kRadians:
+ if (config.radians) {
+ worklist.Push(builtin);
+ }
+ break;
case core::BuiltinFn::kSaturate:
if (config.saturate) {
worklist.Push(builtin);
@@ -157,6 +173,9 @@
case core::BuiltinFn::kCountTrailingZeros:
CountTrailingZeros(builtin);
break;
+ case core::BuiltinFn::kDegrees:
+ Degrees(builtin);
+ break;
case core::BuiltinFn::kExtractBits:
ExtractBits(builtin);
break;
@@ -169,6 +188,9 @@
case core::BuiltinFn::kInsertBits:
InsertBits(builtin);
break;
+ case core::BuiltinFn::kRadians:
+ Radians(builtin);
+ break;
case core::BuiltinFn::kSaturate:
Saturate(builtin);
break;
@@ -367,6 +389,24 @@
call->Destroy();
}
+ /// Polyfill an `degrees()` builtin call.
+ /// @param call the builtin call instruction
+ void Degrees(ir::CoreBuiltinCall* call) {
+ auto* arg = call->Args()[0];
+ auto* type = arg->Type()->DeepestElement();
+ ir::Value* value = nullptr;
+ if (type->Is<core::type::F16>()) {
+ value = b.Constant(f16(kRadToDeg));
+ } else if (type->Is<core::type::F32>()) {
+ value = b.Constant(f32(kRadToDeg));
+ }
+ b.InsertBefore(call, [&] {
+ auto* mul = b.Multiply(arg->Type(), arg, value);
+ mul->SetResults(Vector{call->DetachResult()});
+ });
+ call->Destroy();
+ }
+
/// Polyfill an `extractBits()` builtin call.
/// @param call the builtin call instruction
void ExtractBits(ir::CoreBuiltinCall* call) {
@@ -543,6 +583,24 @@
}
}
+ /// Polyfill an `radians()` builtin call.
+ /// @param call the builtin call instruction
+ void Radians(ir::CoreBuiltinCall* call) {
+ auto* arg = call->Args()[0];
+ auto* type = arg->Type()->DeepestElement();
+ ir::Value* value = nullptr;
+ if (type->Is<core::type::F16>()) {
+ value = b.Constant(f16(kDegToRad));
+ } else if (type->Is<core::type::F32>()) {
+ value = b.Constant(f32(kDegToRad));
+ }
+ b.InsertBefore(call, [&] {
+ auto* mul = b.Multiply(arg->Type(), arg, value);
+ mul->SetResults(Vector{call->DetachResult()});
+ });
+ call->Destroy();
+ }
+
/// Polyfill a `saturate()` builtin call.
/// @param call the builtin call instruction
void Saturate(ir::CoreBuiltinCall* call) {
diff --git a/src/tint/lang/core/ir/transform/builtin_polyfill.h b/src/tint/lang/core/ir/transform/builtin_polyfill.h
index 9795a9c..482af64 100644
--- a/src/tint/lang/core/ir/transform/builtin_polyfill.h
+++ b/src/tint/lang/core/ir/transform/builtin_polyfill.h
@@ -58,6 +58,8 @@
bool count_leading_zeros = false;
/// Should `countTrailingZeros()` be polyfilled?
bool count_trailing_zeros = false;
+ /// Should `degrees()` be polyfilled?
+ bool degrees = false;
/// How should `extractBits()` be polyfilled?
BuiltinPolyfillLevel extract_bits = BuiltinPolyfillLevel::kNone;
/// Should `firstLeadingBit()` be polyfilled?
@@ -66,6 +68,8 @@
bool first_trailing_bit = false;
/// How should `insertBits()` be polyfilled?
BuiltinPolyfillLevel insert_bits = BuiltinPolyfillLevel::kNone;
+ /// Should `radians()` be polyfilled?
+ bool radians = false;
/// Should `saturate()` be polyfilled?
bool saturate = false;
/// Should `textureSampleBaseClampToEdge()` be polyfilled for texture_2d<f32> textures?
diff --git a/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc b/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc
index 0e1dd0a..7e02718 100644
--- a/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc
+++ b/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc
@@ -609,6 +609,133 @@
EXPECT_EQ(expect, str());
}
+TEST_F(IR_BuiltinPolyfillTest, Degrees_NoPolyfill) {
+ Build(core::BuiltinFn::kDegrees, ty.f32(), Vector{ty.f32()});
+ auto* src = R"(
+%foo = func(%arg:f32):f32 {
+ $B1: {
+ %result:f32 = degrees %arg
+ ret %result
+ }
+}
+)";
+ auto* expect = src;
+
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ config.degrees = false;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinPolyfillTest, Degrees_F32) {
+ Build(core::BuiltinFn::kDegrees, ty.f32(), Vector{ty.f32()});
+ auto* src = R"(
+%foo = func(%arg:f32):f32 {
+ $B1: {
+ %result:f32 = degrees %arg
+ ret %result
+ }
+}
+)";
+ auto* expect = R"(
+%foo = func(%arg:f32):f32 {
+ $B1: {
+ %result:f32 = mul %arg, 57.295780181884765625f
+ ret %result
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ config.degrees = true;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinPolyfillTest, Degrees_F16) {
+ Build(core::BuiltinFn::kDegrees, ty.f16(), Vector{ty.f16()});
+ auto* src = R"(
+%foo = func(%arg:f16):f16 {
+ $B1: {
+ %result:f16 = degrees %arg
+ ret %result
+ }
+}
+)";
+ auto* expect = R"(
+%foo = func(%arg:f16):f16 {
+ $B1: {
+ %result:f16 = mul %arg, 57.28125h
+ ret %result
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ config.degrees = true;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinPolyfillTest, Degrees_Vec2F32) {
+ Build(core::BuiltinFn::kDegrees, ty.vec2<f32>(), Vector{ty.vec2<f32>()});
+ auto* src = R"(
+%foo = func(%arg:vec2<f32>):vec2<f32> {
+ $B1: {
+ %result:vec2<f32> = degrees %arg
+ ret %result
+ }
+}
+)";
+ auto* expect = R"(
+%foo = func(%arg:vec2<f32>):vec2<f32> {
+ $B1: {
+ %result:vec2<f32> = mul %arg, 57.295780181884765625f
+ ret %result
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ config.degrees = true;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinPolyfillTest, Degrees_Vec4F16) {
+ Build(core::BuiltinFn::kDegrees, ty.vec4<f16>(), Vector{ty.vec4<f16>()});
+ auto* src = R"(
+%foo = func(%arg:vec4<f16>):vec4<f16> {
+ $B1: {
+ %result:vec4<f16> = degrees %arg
+ ret %result
+ }
+}
+)";
+ auto* expect = R"(
+%foo = func(%arg:vec4<f16>):vec4<f16> {
+ $B1: {
+ %result:vec4<f16> = mul %arg, 57.28125h
+ ret %result
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ config.degrees = true;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(IR_BuiltinPolyfillTest, ExtractBits_NoPolyfill) {
Build(core::BuiltinFn::kExtractBits, ty.u32(), Vector{ty.u32(), ty.u32(), ty.u32()});
auto* src = R"(
@@ -1378,6 +1505,133 @@
EXPECT_EQ(expect, str());
}
+TEST_F(IR_BuiltinPolyfillTest, Radians_NoPolyfill) {
+ Build(core::BuiltinFn::kRadians, ty.f32(), Vector{ty.f32()});
+ auto* src = R"(
+%foo = func(%arg:f32):f32 {
+ $B1: {
+ %result:f32 = radians %arg
+ ret %result
+ }
+}
+)";
+ auto* expect = src;
+
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ config.radians = false;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinPolyfillTest, Radians_F32) {
+ Build(core::BuiltinFn::kRadians, ty.f32(), Vector{ty.f32()});
+ auto* src = R"(
+%foo = func(%arg:f32):f32 {
+ $B1: {
+ %result:f32 = radians %arg
+ ret %result
+ }
+}
+)";
+ auto* expect = R"(
+%foo = func(%arg:f32):f32 {
+ $B1: {
+ %result:f32 = mul %arg, 0.01745329238474369049f
+ ret %result
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ config.radians = true;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinPolyfillTest, Radians_F16) {
+ Build(core::BuiltinFn::kRadians, ty.f16(), Vector{ty.f16()});
+ auto* src = R"(
+%foo = func(%arg:f16):f16 {
+ $B1: {
+ %result:f16 = radians %arg
+ ret %result
+ }
+}
+)";
+ auto* expect = R"(
+%foo = func(%arg:f16):f16 {
+ $B1: {
+ %result:f16 = mul %arg, 0.0174407958984375h
+ ret %result
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ config.radians = true;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinPolyfillTest, Radians_Vec2F32) {
+ Build(core::BuiltinFn::kRadians, ty.vec2<f32>(), Vector{ty.vec2<f32>()});
+ auto* src = R"(
+%foo = func(%arg:vec2<f32>):vec2<f32> {
+ $B1: {
+ %result:vec2<f32> = radians %arg
+ ret %result
+ }
+}
+)";
+ auto* expect = R"(
+%foo = func(%arg:vec2<f32>):vec2<f32> {
+ $B1: {
+ %result:vec2<f32> = mul %arg, 0.01745329238474369049f
+ ret %result
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ config.radians = true;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinPolyfillTest, Radians_Vec4F16) {
+ Build(core::BuiltinFn::kRadians, ty.vec4<f16>(), Vector{ty.vec4<f16>()});
+ auto* src = R"(
+%foo = func(%arg:vec4<f16>):vec4<f16> {
+ $B1: {
+ %result:vec4<f16> = radians %arg
+ ret %result
+ }
+}
+)";
+ auto* expect = R"(
+%foo = func(%arg:vec4<f16>):vec4<f16> {
+ $B1: {
+ %result:vec4<f16> = mul %arg, 0.0174407958984375h
+ ret %result
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ config.radians = true;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(IR_BuiltinPolyfillTest, TextureSampleBaseClampToEdge_2d_f32) {
auto* texture_ty =
ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k2d, ty.f32());
diff --git a/src/tint/lang/msl/writer/raise/raise.cc b/src/tint/lang/msl/writer/raise/raise.cc
index ce9743c..4f08006 100644
--- a/src/tint/lang/msl/writer/raise/raise.cc
+++ b/src/tint/lang/msl/writer/raise/raise.cc
@@ -79,10 +79,12 @@
{
core::ir::transform::BuiltinPolyfillConfig core_polyfills{};
core_polyfills.clamp_int = true;
+ core_polyfills.degrees = true;
core_polyfills.extract_bits = core::ir::transform::BuiltinPolyfillLevel::kClampOrRangeCheck;
core_polyfills.first_leading_bit = true;
core_polyfills.first_trailing_bit = true;
core_polyfills.insert_bits = core::ir::transform::BuiltinPolyfillLevel::kClampOrRangeCheck;
+ core_polyfills.radians = true;
core_polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
RUN_TRANSFORM(core::ir::transform::BuiltinPolyfill, core_polyfills);
}
diff --git a/test/tint/builtins/degrees.spvasm.expected.ir.msl b/test/tint/builtins/degrees.spvasm.expected.ir.msl
index 52f1705..fa188f7 100644
--- a/test/tint/builtins/degrees.spvasm.expected.ir.msl
+++ b/test/tint/builtins/degrees.spvasm.expected.ir.msl
@@ -1,9 +1,13 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: degrees
-********************************************************************
-* 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 main_1() {
+ float a = 0.0f;
+ float b = 0.0f;
+ a = 42.0f;
+ b = (a * 57.295780181884765625f);
+}
+
+kernel void tint_symbol() {
+ main_1();
+}
diff --git a/test/tint/builtins/gen/var/degrees/0d170c.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/degrees/0d170c.wgsl.expected.ir.msl
index f6516f9..0a7036e 100644
--- a/test/tint/builtins/gen/var/degrees/0d170c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/degrees/0d170c.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device float4* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ float4 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float4 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+float4 degrees_0d170c() {
+ float4 arg_0 = float4(1.0f);
+ float4 res = (arg_0 * 57.295780181884765625f);
+ return res;
}
-%degrees_0d170c = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<f32>, read_write> = var, vec4<f32>(1.0f)
- %4:vec4<f32> = load %arg_0
- %5:vec4<f32> = degrees %4
- %res:ptr<function, vec4<f32>, read_write> = var, %5
- %7:vec4<f32> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %degrees_0d170c
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %degrees_0d170c
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %degrees_0d170c
- ret
- }
+fragment void fragment_main(device float4* 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) = degrees_0d170c();
}
-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. *
-********************************************************************
+kernel void compute_main(device float4* 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) = degrees_0d170c();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = degrees_0d170c();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/degrees/1ad5df.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/degrees/1ad5df.wgsl.expected.ir.msl
index 74cda5e..a087ea3 100644
--- a/test/tint/builtins/gen/var/degrees/1ad5df.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/degrees/1ad5df.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec2<f32>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device float2* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ float2 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float2 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+float2 degrees_1ad5df() {
+ float2 arg_0 = float2(1.0f);
+ float2 res = (arg_0 * 57.295780181884765625f);
+ return res;
}
-%degrees_1ad5df = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %4:vec2<f32> = load %arg_0
- %5:vec2<f32> = degrees %4
- %res:ptr<function, vec2<f32>, read_write> = var, %5
- %7:vec2<f32> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %degrees_1ad5df
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %degrees_1ad5df
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %degrees_1ad5df
- ret
- }
+fragment void fragment_main(device float2* 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) = degrees_1ad5df();
}
-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. *
-********************************************************************
+kernel void compute_main(device float2* 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) = degrees_1ad5df();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = degrees_1ad5df();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/degrees/2af623.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/degrees/2af623.wgsl.expected.ir.msl
index f9f1fb8..7bc712f 100644
--- a/test/tint/builtins/gen/var/degrees/2af623.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/degrees/2af623.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec3<f32>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device float3* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ float3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float3 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+float3 degrees_2af623() {
+ float3 arg_0 = float3(1.0f);
+ float3 res = (arg_0 * 57.295780181884765625f);
+ return res;
}
-%degrees_2af623 = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %4:vec3<f32> = load %arg_0
- %5:vec3<f32> = degrees %4
- %res:ptr<function, vec3<f32>, read_write> = var, %5
- %7:vec3<f32> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %degrees_2af623
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %degrees_2af623
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %degrees_2af623
- ret
- }
+fragment void fragment_main(device float3* 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) = degrees_2af623();
}
-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. *
-********************************************************************
+kernel void compute_main(device float3* 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) = degrees_2af623();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = degrees_2af623();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/degrees/3055d3.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/degrees/3055d3.wgsl.expected.ir.msl
index 036c9f4..b4a1509 100644
--- a/test/tint/builtins/gen/var/degrees/3055d3.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/degrees/3055d3.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec4<f16>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device half4* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ half4 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half4 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+half4 degrees_3055d3() {
+ half4 arg_0 = half4(1.0h);
+ half4 res = (arg_0 * 57.28125h);
+ return res;
}
-%degrees_3055d3 = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<f16>, read_write> = var, vec4<f16>(1.0h)
- %4:vec4<f16> = load %arg_0
- %5:vec4<f16> = degrees %4
- %res:ptr<function, vec4<f16>, read_write> = var, %5
- %7:vec4<f16> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %degrees_3055d3
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %degrees_3055d3
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %degrees_3055d3
- ret
- }
+fragment void fragment_main(device half4* 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) = degrees_3055d3();
}
-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. *
-********************************************************************
+kernel void compute_main(device half4* 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) = degrees_3055d3();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = degrees_3055d3();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/degrees/51f705.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/degrees/51f705.wgsl.expected.ir.msl
index 3e931a4..06b3bd7 100644
--- a/test/tint/builtins/gen/var/degrees/51f705.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/degrees/51f705.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../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)
+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]];
+};
+
+float degrees_51f705() {
+ float arg_0 = 1.0f;
+ float res = (arg_0 * 57.295780181884765625f);
+ return res;
}
-%degrees_51f705 = func():void {
- $B2: {
- %arg_0:ptr<function, f32, read_write> = var, 1.0f
- %4:f32 = load %arg_0
- %5:f32 = degrees %4
- %res:ptr<function, f32, read_write> = var, %5
- %7:f32 = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %degrees_51f705
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %degrees_51f705
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %degrees_51f705
- 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) = degrees_51f705();
}
-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. *
-********************************************************************
+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) = degrees_51f705();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = degrees_51f705();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/degrees/5e9805.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/degrees/5e9805.wgsl.expected.ir.msl
index ea9f15f..5759239 100644
--- a/test/tint/builtins/gen/var/degrees/5e9805.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/degrees/5e9805.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../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)
+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]];
+};
+
+half degrees_5e9805() {
+ half arg_0 = 1.0h;
+ half res = (arg_0 * 57.28125h);
+ return res;
}
-%degrees_5e9805 = func():void {
- $B2: {
- %arg_0:ptr<function, f16, read_write> = var, 1.0h
- %4:f16 = load %arg_0
- %5:f16 = degrees %4
- %res:ptr<function, f16, read_write> = var, %5
- %7:f16 = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %degrees_5e9805
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %degrees_5e9805
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %degrees_5e9805
- 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) = degrees_5e9805();
}
-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. *
-********************************************************************
+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) = degrees_5e9805();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = degrees_5e9805();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/degrees/dfe8f4.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/degrees/dfe8f4.wgsl.expected.ir.msl
index 4c30577..b74da9d 100644
--- a/test/tint/builtins/gen/var/degrees/dfe8f4.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/degrees/dfe8f4.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec3<f16>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device half3* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ half3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half3 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+half3 degrees_dfe8f4() {
+ half3 arg_0 = half3(1.0h);
+ half3 res = (arg_0 * 57.28125h);
+ return res;
}
-%degrees_dfe8f4 = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<f16>, read_write> = var, vec3<f16>(1.0h)
- %4:vec3<f16> = load %arg_0
- %5:vec3<f16> = degrees %4
- %res:ptr<function, vec3<f16>, read_write> = var, %5
- %7:vec3<f16> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %degrees_dfe8f4
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %degrees_dfe8f4
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %degrees_dfe8f4
- ret
- }
+fragment void fragment_main(device half3* 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) = degrees_dfe8f4();
}
-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. *
-********************************************************************
+kernel void compute_main(device half3* 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) = degrees_dfe8f4();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = degrees_dfe8f4();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/degrees/f59715.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/degrees/f59715.wgsl.expected.ir.msl
index 994b8d3..efe0c34 100644
--- a/test/tint/builtins/gen/var/degrees/f59715.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/degrees/f59715.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec2<f16>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device half2* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ half2 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half2 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+half2 degrees_f59715() {
+ half2 arg_0 = half2(1.0h);
+ half2 res = (arg_0 * 57.28125h);
+ return res;
}
-%degrees_f59715 = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<f16>, read_write> = var, vec2<f16>(1.0h)
- %4:vec2<f16> = load %arg_0
- %5:vec2<f16> = degrees %4
- %res:ptr<function, vec2<f16>, read_write> = var, %5
- %7:vec2<f16> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %degrees_f59715
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %degrees_f59715
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %degrees_f59715
- ret
- }
+fragment void fragment_main(device half2* 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) = degrees_f59715();
}
-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. *
-********************************************************************
+kernel void compute_main(device half2* 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) = degrees_f59715();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = degrees_f59715();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/radians/09b7fc.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/radians/09b7fc.wgsl.expected.ir.msl
index 1b0ccd9..eb84642 100644
--- a/test/tint/builtins/gen/var/radians/09b7fc.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/radians/09b7fc.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device float4* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ float4 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float4 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+float4 radians_09b7fc() {
+ float4 arg_0 = float4(1.0f);
+ float4 res = (arg_0 * 0.01745329238474369049f);
+ return res;
}
-%radians_09b7fc = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<f32>, read_write> = var, vec4<f32>(1.0f)
- %4:vec4<f32> = load %arg_0
- %5:vec4<f32> = radians %4
- %res:ptr<function, vec4<f32>, read_write> = var, %5
- %7:vec4<f32> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %radians_09b7fc
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %radians_09b7fc
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %radians_09b7fc
- ret
- }
+fragment void fragment_main(device float4* 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) = radians_09b7fc();
}
-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. *
-********************************************************************
+kernel void compute_main(device float4* 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) = radians_09b7fc();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = radians_09b7fc();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/radians/208fd9.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/radians/208fd9.wgsl.expected.ir.msl
index 5d51491..d830a38 100644
--- a/test/tint/builtins/gen/var/radians/208fd9.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/radians/208fd9.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../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)
+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]];
+};
+
+half radians_208fd9() {
+ half arg_0 = 1.0h;
+ half res = (arg_0 * 0.0174407958984375h);
+ return res;
}
-%radians_208fd9 = func():void {
- $B2: {
- %arg_0:ptr<function, f16, read_write> = var, 1.0h
- %4:f16 = load %arg_0
- %5:f16 = radians %4
- %res:ptr<function, f16, read_write> = var, %5
- %7:f16 = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %radians_208fd9
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %radians_208fd9
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %radians_208fd9
- 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) = radians_208fd9();
}
-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. *
-********************************************************************
+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) = radians_208fd9();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = radians_208fd9();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/radians/44f20b.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/radians/44f20b.wgsl.expected.ir.msl
index 0d55368..e354d2e 100644
--- a/test/tint/builtins/gen/var/radians/44f20b.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/radians/44f20b.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec4<f16>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device half4* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ half4 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half4 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+half4 radians_44f20b() {
+ half4 arg_0 = half4(1.0h);
+ half4 res = (arg_0 * 0.0174407958984375h);
+ return res;
}
-%radians_44f20b = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<f16>, read_write> = var, vec4<f16>(1.0h)
- %4:vec4<f16> = load %arg_0
- %5:vec4<f16> = radians %4
- %res:ptr<function, vec4<f16>, read_write> = var, %5
- %7:vec4<f16> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %radians_44f20b
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %radians_44f20b
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %radians_44f20b
- ret
- }
+fragment void fragment_main(device half4* 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) = radians_44f20b();
}
-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. *
-********************************************************************
+kernel void compute_main(device half4* 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) = radians_44f20b();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = radians_44f20b();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/radians/61687a.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/radians/61687a.wgsl.expected.ir.msl
index caa2be9..66f007f 100644
--- a/test/tint/builtins/gen/var/radians/61687a.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/radians/61687a.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec2<f32>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device float2* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ float2 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float2 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+float2 radians_61687a() {
+ float2 arg_0 = float2(1.0f);
+ float2 res = (arg_0 * 0.01745329238474369049f);
+ return res;
}
-%radians_61687a = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %4:vec2<f32> = load %arg_0
- %5:vec2<f32> = radians %4
- %res:ptr<function, vec2<f32>, read_write> = var, %5
- %7:vec2<f32> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %radians_61687a
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %radians_61687a
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %radians_61687a
- ret
- }
+fragment void fragment_main(device float2* 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) = radians_61687a();
}
-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. *
-********************************************************************
+kernel void compute_main(device float2* 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) = radians_61687a();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = radians_61687a();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/radians/6b0ff2.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/radians/6b0ff2.wgsl.expected.ir.msl
index ea20885..95b527d 100644
--- a/test/tint/builtins/gen/var/radians/6b0ff2.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/radians/6b0ff2.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../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)
+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]];
+};
+
+float radians_6b0ff2() {
+ float arg_0 = 1.0f;
+ float res = (arg_0 * 0.01745329238474369049f);
+ return res;
}
-%radians_6b0ff2 = func():void {
- $B2: {
- %arg_0:ptr<function, f32, read_write> = var, 1.0f
- %4:f32 = load %arg_0
- %5:f32 = radians %4
- %res:ptr<function, f32, read_write> = var, %5
- %7:f32 = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %radians_6b0ff2
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %radians_6b0ff2
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %radians_6b0ff2
- 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) = radians_6b0ff2();
}
-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. *
-********************************************************************
+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) = radians_6b0ff2();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = radians_6b0ff2();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/radians/7ea4c7.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/radians/7ea4c7.wgsl.expected.ir.msl
index 5f9034c..dbea81c 100644
--- a/test/tint/builtins/gen/var/radians/7ea4c7.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/radians/7ea4c7.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec3<f16>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device half3* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ half3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half3 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+half3 radians_7ea4c7() {
+ half3 arg_0 = half3(1.0h);
+ half3 res = (arg_0 * 0.0174407958984375h);
+ return res;
}
-%radians_7ea4c7 = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<f16>, read_write> = var, vec3<f16>(1.0h)
- %4:vec3<f16> = load %arg_0
- %5:vec3<f16> = radians %4
- %res:ptr<function, vec3<f16>, read_write> = var, %5
- %7:vec3<f16> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %radians_7ea4c7
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %radians_7ea4c7
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %radians_7ea4c7
- ret
- }
+fragment void fragment_main(device half3* 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) = radians_7ea4c7();
}
-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. *
-********************************************************************
+kernel void compute_main(device half3* 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) = radians_7ea4c7();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = radians_7ea4c7();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/radians/f96258.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/radians/f96258.wgsl.expected.ir.msl
index d0772e5..da70df2 100644
--- a/test/tint/builtins/gen/var/radians/f96258.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/radians/f96258.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec3<f32>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device float3* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ float3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float3 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+float3 radians_f96258() {
+ float3 arg_0 = float3(1.0f);
+ float3 res = (arg_0 * 0.01745329238474369049f);
+ return res;
}
-%radians_f96258 = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %4:vec3<f32> = load %arg_0
- %5:vec3<f32> = radians %4
- %res:ptr<function, vec3<f32>, read_write> = var, %5
- %7:vec3<f32> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %radians_f96258
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %radians_f96258
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %radians_f96258
- ret
- }
+fragment void fragment_main(device float3* 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) = radians_f96258();
}
-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. *
-********************************************************************
+kernel void compute_main(device float3* 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) = radians_f96258();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = radians_f96258();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/gen/var/radians/fbacf0.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/radians/fbacf0.wgsl.expected.ir.msl
index f5ae88a..4983022 100644
--- a/test/tint/builtins/gen/var/radians/fbacf0.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/radians/fbacf0.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec2<f16>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device half2* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ half2 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half2 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+half2 radians_fbacf0() {
+ half2 arg_0 = half2(1.0h);
+ half2 res = (arg_0 * 0.0174407958984375h);
+ return res;
}
-%radians_fbacf0 = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<f16>, read_write> = var, vec2<f16>(1.0h)
- %4:vec2<f16> = load %arg_0
- %5:vec2<f16> = radians %4
- %res:ptr<function, vec2<f16>, read_write> = var, %5
- %7:vec2<f16> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %radians_fbacf0
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %radians_fbacf0
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %radians_fbacf0
- ret
- }
+fragment void fragment_main(device half2* 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) = radians_fbacf0();
}
-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. *
-********************************************************************
+kernel void compute_main(device half2* 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) = radians_fbacf0();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = radians_fbacf0();
+ return out;
+}
+
+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};
+}
diff --git a/test/tint/builtins/radians.spvasm.expected.ir.msl b/test/tint/builtins/radians.spvasm.expected.ir.msl
index 8220f22..d01b84b 100644
--- a/test/tint/builtins/radians.spvasm.expected.ir.msl
+++ b/test/tint/builtins/radians.spvasm.expected.ir.msl
@@ -1,9 +1,13 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: radians
-********************************************************************
-* 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 main_1() {
+ float a = 0.0f;
+ float b = 0.0f;
+ a = 42.0f;
+ b = (a * 0.01745329238474369049f);
+}
+
+kernel void tint_symbol() {
+ main_1();
+}
diff --git a/test/tint/builtins/repeated_use.wgsl.expected.ir.msl b/test/tint/builtins/repeated_use.wgsl.expected.ir.msl
index 52f1705..a901cf9 100644
--- a/test/tint/builtins/repeated_use.wgsl.expected.ir.msl
+++ b/test/tint/builtins/repeated_use.wgsl.expected.ir.msl
@@ -1,9 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: degrees
-********************************************************************
-* 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 tint_symbol() {
+ float4 const va = float4(0.0f);
+ float4 const a = (va * 57.295780181884765625f);
+ float4 const vb = float4(1.0f);
+ float4 const b = (vb * 57.295780181884765625f);
+ float4 const vc = float4(1.0f, 2.0f, 3.0f, 4.0f);
+ float4 const c = (vc * 57.295780181884765625f);
+ float3 const vd = float3(0.0f);
+ float3 const d = (vd * 57.295780181884765625f);
+ float3 const ve = float3(1.0f);
+ float3 const e = (ve * 57.295780181884765625f);
+ float3 const vf = float3(1.0f, 2.0f, 3.0f);
+ float3 const f = (vf * 57.295780181884765625f);
+ float2 const vg = float2(0.0f);
+ float2 const g = (vg * 57.295780181884765625f);
+ float2 const vh = float2(1.0f);
+ float2 const h = (vh * 57.295780181884765625f);
+ float2 const vi = float2(1.0f, 2.0f);
+ float2 const i = (vi * 57.295780181884765625f);
+ float const vj = 1.0f;
+ float const j = (vj * 57.295780181884765625f);
+ float const vk = 2.0f;
+ float const k = (vk * 57.295780181884765625f);
+ float const vl = 3.0f;
+ float const l = (vl * 57.295780181884765625f);
+}
