[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); +}