[tint] Polyfill Smoothstep functions for all IR
We are extending the acceptable inputs for smoothstep to include the
cases where 'low' > 'high'.
For the remaining non-ir paths:
There will be no compiler errors for low > high but the result cannot
be expected to pass conformance.
I have added "smoothstep.wgsl" tint end2end test to cover this case.
Bug: 379909620
Change-Id: I963d3411d152c641c5f9761b7d835220ab7b334c
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/218175
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Peter McNeeley <petermcneeley@google.com>
diff --git a/src/dawn/tests/end2end/ShaderBuiltinPartialConstArgsErrorTests.cpp b/src/dawn/tests/end2end/ShaderBuiltinPartialConstArgsErrorTests.cpp
index cd2c005..bd7cf34 100644
--- a/src/dawn/tests/end2end/ShaderBuiltinPartialConstArgsErrorTests.cpp
+++ b/src/dawn/tests/end2end/ShaderBuiltinPartialConstArgsErrorTests.cpp
@@ -152,9 +152,8 @@
bool BadCaseForBuiltin() {
if (GetParam().mBuiltin == "smoothstep") {
- // The more case is bad because low can't be more than high.
// The equal case generates a divide by zero.
- return GetParam().mCompare != Compare::kLess;
+ return GetParam().mCompare == Compare::kEqual;
}
if (GetParam().mBuiltin == "clamp") {
return GetParam().mCompare == Compare::kMore;
diff --git a/src/tint/lang/core/constant/eval.cc b/src/tint/lang/core/constant/eval.cc
index 8462a3e..e255a57 100644
--- a/src/tint/lang/core/constant/eval.cc
+++ b/src/tint/lang/core/constant/eval.cc
@@ -3701,9 +3701,9 @@
auto create = [&](auto low, auto high, auto x) -> Eval::Result {
using NumberT = decltype(low);
- if (low >= high) {
- AddError(source) << "smoothstep called with 'low' (" << low
- << ") not less than 'high' (" << high << ")";
+ if (low == high) {
+ AddError(source) << "smoothstep called with 'low' (" << low << ") equal to 'high' ("
+ << high << ")";
if (!use_runtime_semantics_) {
return error;
}
diff --git a/src/tint/lang/core/constant/eval_builtin_test.cc b/src/tint/lang/core/constant/eval_builtin_test.cc
index 0eaf4d2..a0955c8 100644
--- a/src/tint/lang/core/constant/eval_builtin_test.cc
+++ b/src/tint/lang/core/constant/eval_builtin_test.cc
@@ -2591,9 +2591,9 @@
return "12:34 error: " + OverflowErrorMessage(a, op, b) + R"(
12:34 note: when calculating smoothstep)";
};
- auto error_low_ge_high = [](auto a, auto b) {
+ auto error_low_equal_high = [](auto a, auto b) {
StringStream ss;
- ss << "12:34 error: smoothstep called with 'low' (" << a << ") not less than 'high' (" << b
+ ss << "12:34 error: smoothstep called with 'low' (" << a << ") equal to 'high' (" << b
<< ")";
return ss.str();
};
@@ -2614,10 +2614,10 @@
// `high - low` underflows
E({T::Lowest(), T::Highest(), T(0)}, error_overflow(T::Highest(), "-", T::Lowest())),
// low == high
- E({T(0), T(0), T(0)}, error_low_ge_high(T(0), T(0))),
- E({T(1), T(1), T(5)}, error_low_ge_high(T(1), T(1))),
+ E({T(0), T(0), T(0)}, error_low_equal_high(T(0), T(0))),
+ E({T(1), T(1), T(5)}, error_low_equal_high(T(1), T(1))),
// low > high
- E({T(1), T(0), T(5)}, error_low_ge_high(T(1), T(0))),
+ C({T(1), T(0), T(5)}, T(0)),
};
}
INSTANTIATE_TEST_SUITE_P( //
diff --git a/src/tint/lang/core/ir/transform/builtin_polyfill.cc b/src/tint/lang/core/ir/transform/builtin_polyfill.cc
index 95dd6fe..da8c722 100644
--- a/src/tint/lang/core/ir/transform/builtin_polyfill.cc
+++ b/src/tint/lang/core/ir/transform/builtin_polyfill.cc
@@ -91,6 +91,9 @@
worklist.Push(builtin);
}
break;
+ case core::BuiltinFn::kSmoothstep:
+ worklist.Push(builtin);
+ break;
case core::BuiltinFn::kExtractBits:
if (config.extract_bits != BuiltinPolyfillLevel::kNone) {
worklist.Push(builtin);
@@ -192,6 +195,9 @@
case core::BuiltinFn::kDegrees:
Degrees(builtin);
break;
+ case core::BuiltinFn::kSmoothstep:
+ SmoothStep(builtin);
+ break;
case core::BuiltinFn::kExtractBits:
ExtractBits(builtin);
break;
@@ -407,6 +413,46 @@
call->Destroy();
}
+ /// Polyfill an `smoothStep()` builtin call.
+ /// @param call the builtin call instruction
+ void SmoothStep(ir::CoreBuiltinCall* call) {
+ auto* edge0_arg = call->Args()[0];
+ auto* edge1_arg = call->Args()[1];
+ auto* x_arg = call->Args()[2];
+ auto* type = x_arg->Type();
+ ir::Constant* zero = nullptr;
+ ir::Constant* one = nullptr;
+ ir::Constant* two = nullptr;
+ ir::Constant* three = nullptr;
+ if (type->DeepestElement()->Is<core::type::F32>()) {
+ zero = b.MatchWidth(0_f, type);
+ one = b.MatchWidth(1_f, type);
+ two = b.MatchWidth(2_f, type);
+ three = b.MatchWidth(3_f, type);
+ } else if (type->DeepestElement()->Is<core::type::F16>()) {
+ zero = b.MatchWidth(0_h, type);
+ one = b.MatchWidth(1_h, type);
+ two = b.MatchWidth(2_h, type);
+ three = b.MatchWidth(3_h, type);
+ }
+
+ b.InsertBefore(call, [&] {
+ auto* dividend = b.Subtract(type, x_arg, edge0_arg);
+ auto* divisor = b.Subtract(type, edge1_arg, edge0_arg);
+ auto* quotient = b.Divide(type, dividend, divisor);
+ auto* t_clamped = b.Call(type, core::BuiltinFn::kClamp, quotient, zero, one);
+
+ // Smoothstep is a well defined function.
+ // result = t * t * (3.0 - 2.0 * t);
+ auto* smooth_result =
+ b.Multiply(type, t_clamped,
+ b.Multiply(type, t_clamped,
+ b.Subtract(type, three, b.Multiply(type, two, t_clamped))));
+ smooth_result->SetResults(Vector{call->DetachResult()});
+ });
+ call->Destroy();
+ }
+
/// Polyfill an `extractBits()` builtin call.
/// @param call the builtin call instruction
void ExtractBits(ir::CoreBuiltinCall* call) {
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 7a4091b..163b943 100644
--- a/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc
+++ b/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc
@@ -188,6 +188,139 @@
EXPECT_EQ(expect, str());
}
+TEST_F(IR_BuiltinPolyfillTest, Smoothstep_F32) {
+ Build(core::BuiltinFn::kSmoothstep, ty.f32(), Vector{ty.f32(), ty.f32(), ty.f32()});
+ auto* src = R"(
+%foo = func(%arg:f32, %arg_1:f32, %arg_2:f32):f32 { # %arg_1: 'arg', %arg_2: 'arg'
+ $B1: {
+ %result:f32 = smoothstep %arg, %arg_1, %arg_2
+ ret %result
+ }
+}
+)";
+ auto* expect = R"(
+%foo = func(%arg:f32, %arg_1:f32, %arg_2:f32):f32 { # %arg_1: 'arg', %arg_2: 'arg'
+ $B1: {
+ %5:f32 = sub %arg_2, %arg
+ %6:f32 = sub %arg_1, %arg
+ %7:f32 = div %5, %6
+ %8:f32 = clamp %7, 0.0f, 1.0f
+ %9:f32 = mul 2.0f, %8
+ %10:f32 = sub 3.0f, %9
+ %11:f32 = mul %8, %10
+ %result:f32 = mul %8, %11
+ ret %result
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinPolyfillTest, Smoothstep_F16) {
+ Build(core::BuiltinFn::kSmoothstep, ty.f16(), Vector{ty.f16(), ty.f16(), ty.f16()});
+ auto* src = R"(
+%foo = func(%arg:f16, %arg_1:f16, %arg_2:f16):f16 { # %arg_1: 'arg', %arg_2: 'arg'
+ $B1: {
+ %result:f16 = smoothstep %arg, %arg_1, %arg_2
+ ret %result
+ }
+}
+)";
+ auto* expect = R"(
+%foo = func(%arg:f16, %arg_1:f16, %arg_2:f16):f16 { # %arg_1: 'arg', %arg_2: 'arg'
+ $B1: {
+ %5:f16 = sub %arg_2, %arg
+ %6:f16 = sub %arg_1, %arg
+ %7:f16 = div %5, %6
+ %8:f16 = clamp %7, 0.0h, 1.0h
+ %9:f16 = mul 2.0h, %8
+ %10:f16 = sub 3.0h, %9
+ %11:f16 = mul %8, %10
+ %result:f16 = mul %8, %11
+ ret %result
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinPolyfillTest, Smoothstep_Vec2F32) {
+ Build(core::BuiltinFn::kSmoothstep, ty.vec2<f32>(),
+ Vector{ty.vec2<f32>(), ty.vec2<f32>(), ty.vec2<f32>()});
+ auto* src = R"(
+%foo = func(%arg:vec2<f32>, %arg_1:vec2<f32>, %arg_2:vec2<f32>):vec2<f32> { # %arg_1: 'arg', %arg_2: 'arg'
+ $B1: {
+ %result:vec2<f32> = smoothstep %arg, %arg_1, %arg_2
+ ret %result
+ }
+}
+)";
+ auto* expect = R"(
+%foo = func(%arg:vec2<f32>, %arg_1:vec2<f32>, %arg_2:vec2<f32>):vec2<f32> { # %arg_1: 'arg', %arg_2: 'arg'
+ $B1: {
+ %5:vec2<f32> = sub %arg_2, %arg
+ %6:vec2<f32> = sub %arg_1, %arg
+ %7:vec2<f32> = div %5, %6
+ %8:vec2<f32> = clamp %7, vec2<f32>(0.0f), vec2<f32>(1.0f)
+ %9:vec2<f32> = mul vec2<f32>(2.0f), %8
+ %10:vec2<f32> = sub vec2<f32>(3.0f), %9
+ %11:vec2<f32> = mul %8, %10
+ %result:vec2<f32> = mul %8, %11
+ ret %result
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinPolyfillTest, Smoothstep_Vec4F16) {
+ Build(core::BuiltinFn::kSmoothstep, ty.vec4<f16>(),
+ Vector{ty.vec4<f16>(), ty.vec4<f16>(), ty.vec4<f16>()});
+ auto* src = R"(
+%foo = func(%arg:vec4<f16>, %arg_1:vec4<f16>, %arg_2:vec4<f16>):vec4<f16> { # %arg_1: 'arg', %arg_2: 'arg'
+ $B1: {
+ %result:vec4<f16> = smoothstep %arg, %arg_1, %arg_2
+ ret %result
+ }
+}
+)";
+ auto* expect = R"(
+%foo = func(%arg:vec4<f16>, %arg_1:vec4<f16>, %arg_2:vec4<f16>):vec4<f16> { # %arg_1: 'arg', %arg_2: 'arg'
+ $B1: {
+ %5:vec4<f16> = sub %arg_2, %arg
+ %6:vec4<f16> = sub %arg_1, %arg
+ %7:vec4<f16> = div %5, %6
+ %8:vec4<f16> = clamp %7, vec4<f16>(0.0h), vec4<f16>(1.0h)
+ %9:vec4<f16> = mul vec4<f16>(2.0h), %8
+ %10:vec4<f16> = sub vec4<f16>(3.0h), %9
+ %11:vec4<f16> = mul %8, %10
+ %result:vec4<f16> = mul %8, %11
+ ret %result
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ BuiltinPolyfillConfig config;
+ Run(BuiltinPolyfill, config);
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(IR_BuiltinPolyfillTest, CountLeadingZeros_NoPolyfill) {
Build(core::BuiltinFn::kCountLeadingZeros, ty.u32(), Vector{ty.u32()});
auto* src = R"(
diff --git a/src/tint/lang/spirv/writer/builtin_test.cc b/src/tint/lang/spirv/writer/builtin_test.cc
index 8e6942c..bc1a72e 100644
--- a/src/tint/lang/spirv/writer/builtin_test.cc
+++ b/src/tint/lang/spirv/writer/builtin_test.cc
@@ -1444,16 +1444,13 @@
ASSERT_TRUE(Generate()) << Error() << output_;
EXPECT_INST(params.spirv_inst);
}
-INSTANTIATE_TEST_SUITE_P(
- SpirvWriterTest,
- Builtin_3arg,
- testing::Values(BuiltinTestCase{kF32, core::BuiltinFn::kClamp, "NClamp"},
- BuiltinTestCase{kF32, core::BuiltinFn::kFma, "Fma"},
- BuiltinTestCase{kF16, core::BuiltinFn::kFma, "Fma"},
- BuiltinTestCase{kF32, core::BuiltinFn::kMix, "Mix"},
- BuiltinTestCase{kF16, core::BuiltinFn::kMix, "Mix"},
- BuiltinTestCase{kF32, core::BuiltinFn::kSmoothstep, "SmoothStep"},
- BuiltinTestCase{kF16, core::BuiltinFn::kSmoothstep, "SmoothStep"}));
+INSTANTIATE_TEST_SUITE_P(SpirvWriterTest,
+ Builtin_3arg,
+ testing::Values(BuiltinTestCase{kF32, core::BuiltinFn::kClamp, "NClamp"},
+ BuiltinTestCase{kF32, core::BuiltinFn::kFma, "Fma"},
+ BuiltinTestCase{kF16, core::BuiltinFn::kFma, "Fma"},
+ BuiltinTestCase{kF32, core::BuiltinFn::kMix, "Mix"},
+ BuiltinTestCase{kF16, core::BuiltinFn::kMix, "Mix"}));
TEST_F(SpirvWriterTest, Builtin_Clamp_Scalar_I32) {
auto* value = b.FunctionParam("value", ty.i32());
@@ -1557,6 +1554,57 @@
)");
}
+TEST_F(SpirvWriterTest, Builtin_Smoothstep_F32) {
+ auto* value = b.FunctionParam("value", ty.f32());
+ auto* low = b.FunctionParam("low", ty.f32());
+ auto* high = b.FunctionParam("high", ty.f32());
+ auto* func = b.Function("foo", ty.f32());
+ func->SetParams({value, low, high});
+
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(ty.f32(), core::BuiltinFn::kSmoothstep, value, low, high);
+ b.Return(func, result);
+ mod.SetName(result, "result");
+ });
+
+ ASSERT_TRUE(Generate()) << Error() << output_;
+ EXPECT_INST(R"(
+ %8 = OpFSub %float %high %value
+ %9 = OpFSub %float %low %value
+ %10 = OpFDiv %float %8 %9
+ %11 = OpExtInst %float %12 NClamp %10 %float_0 %float_1
+ %15 = OpFMul %float %float_2 %11
+ %17 = OpFSub %float %float_3 %15
+ %19 = OpFMul %float %11 %17
+ %result = OpFMul %float %11 %19
+)");
+}
+
+TEST_F(SpirvWriterTest, Builtin_Smoothstep_F16) {
+ auto* value = b.FunctionParam("value", ty.f16());
+ auto* low = b.FunctionParam("low", ty.f16());
+ auto* high = b.FunctionParam("high", ty.f16());
+ auto* func = b.Function("foo", ty.f16());
+ func->SetParams({value, low, high});
+
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(ty.f16(), core::BuiltinFn::kSmoothstep, value, low, high);
+ b.Return(func, result);
+ mod.SetName(result, "result");
+ });
+
+ ASSERT_TRUE(Generate()) << Error() << output_;
+ EXPECT_INST(R"(
+ %8 = OpFSub %half %high %value
+ %9 = OpFSub %half %low %value
+ %10 = OpFDiv %half %8 %9
+ %11 = OpExtInst %half %12 NClamp %10 %half_0x0p_0 %half_0x1p_0
+ %15 = OpFMul %half %half_0x1p_1 %11
+ %17 = OpFSub %half %half_0x1_8p_1 %15
+ %19 = OpFMul %half %11 %17
+ %result = OpFMul %half %11 %19
+)");
+}
TEST_F(SpirvWriterTest, Builtin_ExtractBits_Scalar_U32) {
auto* arg = b.FunctionParam("arg", ty.u32());
auto* offset = b.FunctionParam("offset", ty.u32());
diff --git a/src/tint/lang/wgsl/resolver/builtins_validation_test.cc b/src/tint/lang/wgsl/resolver/builtins_validation_test.cc
index 5d46e0a..3e38ed2 100644
--- a/src/tint/lang/wgsl/resolver/builtins_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/builtins_validation_test.cc
@@ -1669,7 +1669,7 @@
EXPECT_FALSE(r()->Resolve());
StringStream ss;
ss << "12:34 error: smoothstep called with 'low' (" << params.lowStr
- << ") not less than 'high' (" << params.highStr << ")";
+ << ") equal to 'high' (" << params.highStr << ")";
auto expect = ss.str();
EXPECT_EQ(r()->error(), expect);
}
@@ -1706,7 +1706,7 @@
EXPECT_FALSE(r()->Resolve());
StringStream ss;
ss << "12:34 error: smoothstep called with 'low' (" << params.lowStr
- << ") not less than 'high' (" << params.highStr << ")";
+ << ") equal to 'high' (" << params.highStr << ")";
auto expect = ss.str();
EXPECT_EQ(r()->error(), expect);
}
@@ -1738,36 +1738,36 @@
// AInt AInt
{DataType<f32>::AST, Mk(1_a), Mk(1_a), false, "1.0", "1.0"},
- {DataType<f32>::AST, Mk(1_a), Mk(0_a), false, "1.0", "0.0"},
+ {DataType<f32>::AST, Mk(1_a), Mk(0_a), true, "1.0", "0.0"},
// AFloat AInt
{DataType<f32>::AST, Mk(1.0_a), Mk(1_a), false, "1.0", "1.0"},
- {DataType<f32>::AST, Mk(1.0_a), Mk(0_a), false, "1.0", "0.0"},
+ {DataType<f32>::AST, Mk(1.0_a), Mk(0_a), true, "1.0", "0.0"},
// AInt AFloat
{DataType<f32>::AST, Mk(1_a), Mk(1.0_a), false, "1.0", "1.0"},
- {DataType<f32>::AST, Mk(1_a), Mk(0.0_a), false, "1.0", "0.0"},
+ {DataType<f32>::AST, Mk(1_a), Mk(0.0_a), true, "1.0", "0.0"},
// AFloat AFloat
{DataType<f32>::AST, Mk(1.0_a), Mk(1.0_a), false, "1.0", "1.0"},
- {DataType<f32>::AST, Mk(1.0_a), Mk(0.0_a), false, "1.0", "0.0"},
+ {DataType<f32>::AST, Mk(1.0_a), Mk(0.0_a), true, "1.0", "0.0"},
// AInt f32
{DataType<f32>::AST, Mk(1_a), Mk(1_f), false, "1.0", "1.0"},
- {DataType<f32>::AST, Mk(1_a), Mk(0_f), false, "1.0", "0.0"},
+ {DataType<f32>::AST, Mk(1_a), Mk(0_f), true, "1.0", "0.0"},
// f32 AInt
{DataType<f32>::AST, Mk(1_f), Mk(1_a), false, "1.0", "1.0"},
- {DataType<f32>::AST, Mk(1_f), Mk(0_a), false, "1.0", "0.0"},
+ {DataType<f32>::AST, Mk(1_f), Mk(0_a), true, "1.0", "0.0"},
// AFloat f32
{DataType<f32>::AST, Mk(1.0_a), Mk(1_f), false, "1.0", "1.0"},
- {DataType<f32>::AST, Mk(1.0_a), Mk(0_f), false, "1.0", "0.0"},
+ {DataType<f32>::AST, Mk(1.0_a), Mk(0_f), true, "1.0", "0.0"},
// f32 AFloat
{DataType<f32>::AST, Mk(1_a), Mk(1.0_a), false, "1.0", "1.0"},
- {DataType<f32>::AST, Mk(1_a), Mk(0.0_a), false, "1.0", "0.0"},
+ {DataType<f32>::AST, Mk(1_a), Mk(0.0_a), true, "1.0", "0.0"},
// f32 f32
{DataType<f32>::AST, Mk(1_f), Mk(1.0_f), false, "1.0", "1.0"},
- {DataType<f32>::AST, Mk(1_f), Mk(0.0_f), false, "1.0", "0.0"},
+ {DataType<f32>::AST, Mk(1_f), Mk(0.0_f), true, "1.0", "0.0"},
};
}
diff --git a/test/tint/bug/tint/379127084.wgsl.expected.glsl b/test/tint/bug/tint/379127084.wgsl.expected.glsl
index 563281e..d2138c1 100644
--- a/test/tint/bug/tint/379127084.wgsl.expected.glsl
+++ b/test/tint/bug/tint/379127084.wgsl.expected.glsl
@@ -83,7 +83,7 @@
vec4 _67_p = _66_j;
vec2 _skTemp5 = fract(_57_k);
vec2 _68_d = _skTemp5;
- vec2 _skTemp6 = smoothstep(vec2(0.0f), vec2(1.0f), _68_d);
+ vec2 _skTemp6 = (clamp(((_68_d - vec2(0.0f)) / (vec2(1.0f) - vec2(0.0f))), vec2(0.0f), vec2(1.0f)) * (clamp(((_68_d - vec2(0.0f)) / (vec2(1.0f) - vec2(0.0f))), vec2(0.0f), vec2(1.0f)) * (vec2(3.0f) - (vec2(2.0f) * clamp(((_68_d - vec2(0.0f)) / (vec2(1.0f) - vec2(0.0f))), vec2(0.0f), vec2(1.0f))))));
vec2 _69_e = _skTemp6;
vec4 _71_g = vec4(0.0f);
int _72_h = 0;
diff --git a/test/tint/bug/tint/379127084.wgsl.expected.ir.dxc.hlsl b/test/tint/bug/tint/379127084.wgsl.expected.ir.dxc.hlsl
index adb0840..ef82543 100644
--- a/test/tint/bug/tint/379127084.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/bug/tint/379127084.wgsl.expected.ir.dxc.hlsl
@@ -72,7 +72,7 @@
float4 _67_p = _66_j;
float2 _skTemp5 = frac(_57_k);
float2 _68_d = _skTemp5;
- float2 _skTemp6 = smoothstep((0.0f).xx, (1.0f).xx, _68_d);
+ float2 _skTemp6 = (clamp(((_68_d - (0.0f).xx) / ((1.0f).xx - (0.0f).xx)), (0.0f).xx, (1.0f).xx) * (clamp(((_68_d - (0.0f).xx) / ((1.0f).xx - (0.0f).xx)), (0.0f).xx, (1.0f).xx) * ((3.0f).xx - ((2.0f).xx * clamp(((_68_d - (0.0f).xx) / ((1.0f).xx - (0.0f).xx)), (0.0f).xx, (1.0f).xx)))));
float2 _69_e = _skTemp6;
float4 _71_g = (0.0f).xxxx;
int _72_h = int(0);
diff --git a/test/tint/bug/tint/379127084.wgsl.expected.ir.fxc.hlsl b/test/tint/bug/tint/379127084.wgsl.expected.ir.fxc.hlsl
index c5e0f24..b563655 100644
--- a/test/tint/bug/tint/379127084.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/bug/tint/379127084.wgsl.expected.ir.fxc.hlsl
@@ -72,7 +72,7 @@
float4 _67_p = _66_j;
float2 _skTemp5 = frac(_57_k);
float2 _68_d = _skTemp5;
- float2 _skTemp6 = smoothstep((0.0f).xx, (1.0f).xx, _68_d);
+ float2 _skTemp6 = (clamp(((_68_d - (0.0f).xx) / ((1.0f).xx - (0.0f).xx)), (0.0f).xx, (1.0f).xx) * (clamp(((_68_d - (0.0f).xx) / ((1.0f).xx - (0.0f).xx)), (0.0f).xx, (1.0f).xx) * ((3.0f).xx - ((2.0f).xx * clamp(((_68_d - (0.0f).xx) / ((1.0f).xx - (0.0f).xx)), (0.0f).xx, (1.0f).xx)))));
float2 _69_e = _skTemp6;
float4 _71_g = (0.0f).xxxx;
int _72_h = int(0);
diff --git a/test/tint/bug/tint/379127084.wgsl.expected.ir.msl b/test/tint/bug/tint/379127084.wgsl.expected.ir.msl
index f729348..47fc621 100644
--- a/test/tint/bug/tint/379127084.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/379127084.wgsl.expected.ir.msl
@@ -96,7 +96,7 @@
float4 const _67_p = _66_j;
float2 const _skTemp5 = fract(_57_k);
float2 const _68_d = _skTemp5;
- float2 const _skTemp6 = smoothstep(float2(0.0f), float2(1.0f), _68_d);
+ float2 const _skTemp6 = (clamp(((_68_d - float2(0.0f)) / (float2(1.0f) - float2(0.0f))), float2(0.0f), float2(1.0f)) * (clamp(((_68_d - float2(0.0f)) / (float2(1.0f) - float2(0.0f))), float2(0.0f), float2(1.0f)) * (float2(3.0f) - (float2(2.0f) * clamp(((_68_d - float2(0.0f)) / (float2(1.0f) - float2(0.0f))), float2(0.0f), float2(1.0f))))));
float2 const _69_e = _skTemp6;
float4 _71_g = 0.0f;
int _72_h = 0;
diff --git a/test/tint/bug/tint/379127084.wgsl.expected.spvasm b/test/tint/bug/tint/379127084.wgsl.expected.spvasm
index 19ab605..52e2ea3 100644
--- a/test/tint/bug/tint/379127084.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/379127084.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 600
+; Bound: 608
; Schema: 0
OpCapability Shader
%48 = OpExtInstImport "GLSL.std.450"
@@ -218,19 +218,21 @@
%float_256 = OpConstant %float 256
%201 = OpConstantComposite %v4float %float_0_00390625 %float_0_00390625 %float_0_00390625 %float_0_00390625
%206 = OpConstantNull %v2float
- %float_0_25 = OpConstant %float 0.25
%float_2 = OpConstant %float 2
+ %211 = OpConstantComposite %v2float %float_2 %float_2
+ %float_3 = OpConstant %float 3
+ %214 = OpConstantComposite %v2float %float_3 %float_3
+ %float_0_25 = OpConstant %float 0.25
%_ptr_Function_uint = OpTypePointer Function %uint
%int_1 = OpConstant %int 1
%int_4 = OpConstant %int 4
- %350 = OpConstantComposite %v2float %float_2 %float_2
- %372 = OpConstantComposite %v4float %float_0_5 %float_0_5 %float_0_5 %float_0_5
- %376 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+ %381 = OpConstantComposite %v4float %float_0_5 %float_0_5 %float_0_5 %float_0_5
+ %385 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
%v3float = OpTypeVector %float 3
%float_0_212599993 = OpConstant %float 0.212599993
%float_0_715200007 = OpConstant %float 0.715200007
%float_0_0722000003 = OpConstant %float 0.0722000003
- %388 = OpConstantComposite %v3float %float_0_212599993 %float_0_715200007 %float_0_0722000003
+ %397 = OpConstantComposite %v3float %float_0_212599993 %float_0_715200007 %float_0_0722000003
%float_0 = OpConstant %float 0
%uint_7 = OpConstant %uint 7
%float_n1 = OpConstant %float -1
@@ -243,14 +245,13 @@
%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
%uint_6 = OpConstant %uint 6
%float_0_333333343 = OpConstant %float 0.333333343
- %529 = OpConstantComposite %v3float %float_0 %float_0_666666687 %float_0_333333343
- %float_3 = OpConstant %float 3
- %540 = OpConstantNull %v3float
- %541 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+ %538 = OpConstantComposite %v3float %float_0 %float_0_666666687 %float_0_333333343
+ %548 = OpConstantNull %v3float
+ %549 = OpConstantComposite %v3float %float_1 %float_1 %float_1
%uint_8 = OpConstant %uint 8
- %586 = OpTypeFunction %FSOut %FSIn
- %589 = OpConstantNull %FSOut
- %593 = OpTypeFunction %void
+ %594 = OpTypeFunction %FSOut %FSIn
+ %597 = OpConstantNull %FSOut
+ %601 = OpTypeFunction %void
%_skslMain = OpFunction %void None %37
%_stageIn = OpFunctionParameter %FSIn
%_stageOut_root = OpFunctionParameter %_ptr_Function_FSOut
@@ -418,440 +419,447 @@
%_67_p = OpLoad %v4float %_66_j None
%203 = OpLoad %v2float %_57_k None
%_skTemp5 = OpExtInst %v2float %48 Fract %203
- %_skTemp6 = OpExtInst %v2float %48 SmoothStep %206 %126 %_skTemp5
+ %205 = OpFSub %v2float %_skTemp5 %206
+ %207 = OpFSub %v2float %126 %206
+ %208 = OpFDiv %v2float %205 %207
+ %209 = OpExtInst %v2float %48 NClamp %208 %206 %126
+ %210 = OpFMul %v2float %211 %209
+ %213 = OpFSub %v2float %214 %210
+ %216 = OpFMul %v2float %209 %213
+ %_skTemp6 = OpFMul %v2float %209 %216
OpStore %_72_h %int_0
- OpBranch %209
- %209 = OpLabel
- OpBranch %212
- %212 = OpLabel
- OpLoopMerge %213 %211 None
- OpBranch %210
- %210 = OpLabel
- %215 = OpLoad %v2uint %tint_loop_idx_0 None
- %216 = OpIEqual %v2bool %215 %95
- %217 = OpAll %bool %216
- OpSelectionMerge %218 None
- OpBranchConditional %217 %219 %218
- %219 = OpLabel
- OpBranch %213
- %218 = OpLabel
- %220 = OpLoad %int %_72_h None
- %221 = OpConvertSToF %float %220
- %222 = OpFAdd %float %221 %float_0_5
- %_73_i = OpFMul %float %222 %float_0_25
- %225 = OpLoad %20 %noiseSampler_1_Texture None
- %226 = OpLoad %17 %noiseSampler_1_Sampler None
- %227 = OpCompositeExtract %float %_67_p 0
- %228 = OpCompositeConstruct %v2float %227 %_73_i
- %229 = OpExtInst %float %48 NClamp %float_n0_474999994 %float_n16 %float_15_9899998
- %230 = OpSampledImage %162 %225 %226
- %_74_j = OpImageSampleImplicitLod %v4float %230 %228 Bias %229
- %232 = OpLoad %20 %noiseSampler_1_Texture None
- %233 = OpLoad %17 %noiseSampler_1_Sampler None
- %234 = OpCompositeExtract %float %_67_p 1
- %235 = OpCompositeConstruct %v2float %234 %_73_i
- %236 = OpExtInst %float %48 NClamp %float_n0_474999994 %float_n16 %float_15_9899998
- %237 = OpSampledImage %162 %232 %233
- %_75_k = OpImageSampleImplicitLod %v4float %237 %235 Bias %236
- %239 = OpLoad %20 %noiseSampler_1_Texture None
- %240 = OpLoad %17 %noiseSampler_1_Sampler None
- %241 = OpCompositeExtract %float %_67_p 3
- %242 = OpCompositeConstruct %v2float %241 %_73_i
- %243 = OpExtInst %float %48 NClamp %float_n0_474999994 %float_n16 %float_15_9899998
- %244 = OpSampledImage %162 %239 %240
- %_76_l = OpImageSampleImplicitLod %v4float %244 %242 Bias %243
- %246 = OpLoad %20 %noiseSampler_1_Texture None
- %247 = OpLoad %17 %noiseSampler_1_Sampler None
- %248 = OpCompositeExtract %float %_67_p 2
- %249 = OpCompositeConstruct %v2float %248 %_73_i
- %250 = OpExtInst %float %48 NClamp %float_n0_474999994 %float_n16 %float_15_9899998
- %251 = OpSampledImage %162 %246 %247
- %_77_m = OpImageSampleImplicitLod %v4float %251 %249 Bias %250
+ OpBranch %220
+ %220 = OpLabel
+ OpBranch %223
+ %223 = OpLabel
+ OpLoopMerge %224 %222 None
+ OpBranch %221
+ %221 = OpLabel
+ %226 = OpLoad %v2uint %tint_loop_idx_0 None
+ %227 = OpIEqual %v2bool %226 %95
+ %228 = OpAll %bool %227
+ OpSelectionMerge %229 None
+ OpBranchConditional %228 %230 %229
+ %230 = OpLabel
+ OpBranch %224
+ %229 = OpLabel
+ %231 = OpLoad %int %_72_h None
+ %232 = OpConvertSToF %float %231
+ %233 = OpFAdd %float %232 %float_0_5
+ %_73_i = OpFMul %float %233 %float_0_25
+ %236 = OpLoad %20 %noiseSampler_1_Texture None
+ %237 = OpLoad %17 %noiseSampler_1_Sampler None
+ %238 = OpCompositeExtract %float %_67_p 0
+ %239 = OpCompositeConstruct %v2float %238 %_73_i
+ %240 = OpExtInst %float %48 NClamp %float_n0_474999994 %float_n16 %float_15_9899998
+ %241 = OpSampledImage %162 %236 %237
+ %_74_j = OpImageSampleImplicitLod %v4float %241 %239 Bias %240
+ %243 = OpLoad %20 %noiseSampler_1_Texture None
+ %244 = OpLoad %17 %noiseSampler_1_Sampler None
+ %245 = OpCompositeExtract %float %_67_p 1
+ %246 = OpCompositeConstruct %v2float %245 %_73_i
+ %247 = OpExtInst %float %48 NClamp %float_n0_474999994 %float_n16 %float_15_9899998
+ %248 = OpSampledImage %162 %243 %244
+ %_75_k = OpImageSampleImplicitLod %v4float %248 %246 Bias %247
+ %250 = OpLoad %20 %noiseSampler_1_Texture None
+ %251 = OpLoad %17 %noiseSampler_1_Sampler None
+ %252 = OpCompositeExtract %float %_67_p 3
+ %253 = OpCompositeConstruct %v2float %252 %_73_i
+ %254 = OpExtInst %float %48 NClamp %float_n0_474999994 %float_n16 %float_15_9899998
+ %255 = OpSampledImage %162 %250 %251
+ %_76_l = OpImageSampleImplicitLod %v4float %255 %253 Bias %254
+ %257 = OpLoad %20 %noiseSampler_1_Texture None
+ %258 = OpLoad %17 %noiseSampler_1_Sampler None
+ %259 = OpCompositeExtract %float %_67_p 2
+ %260 = OpCompositeConstruct %v2float %259 %_73_i
+ %261 = OpExtInst %float %48 NClamp %float_n0_474999994 %float_n16 %float_15_9899998
+ %262 = OpSampledImage %162 %257 %258
+ %_77_m = OpImageSampleImplicitLod %v4float %262 %260 Bias %261
OpStore %_78_n %_skTemp5
- %254 = OpVectorShuffle %v2float %_74_j %_74_j 1 3
- %255 = OpVectorShuffle %v2float %_74_j %_74_j 0 2
- %256 = OpVectorTimesScalar %v2float %255 %float_0_00390625
- %257 = OpFAdd %v2float %254 %256
- %258 = OpVectorTimesScalar %v2float %257 %float_2
- %260 = OpCompositeConstruct %v2float %float_1 %float_1
- %261 = OpFSub %v2float %258 %260
- %262 = OpLoad %v2float %_78_n None
- %_skTemp7 = OpDot %float %261 %262
+ %265 = OpVectorShuffle %v2float %_74_j %_74_j 1 3
+ %266 = OpVectorShuffle %v2float %_74_j %_74_j 0 2
+ %267 = OpVectorTimesScalar %v2float %266 %float_0_00390625
+ %268 = OpFAdd %v2float %265 %267
+ %269 = OpVectorTimesScalar %v2float %268 %float_2
+ %270 = OpCompositeConstruct %v2float %float_1 %float_1
+ %271 = OpFSub %v2float %269 %270
+ %272 = OpLoad %v2float %_78_n None
+ %_skTemp7 = OpDot %float %271 %272
OpStore %_79_o %_skTemp7
- %265 = OpAccessChain %_ptr_Function_float %_78_n %uint_0
- %266 = OpLoad %float %265 None
- %267 = OpFSub %float %266 %float_1
- %268 = OpAccessChain %_ptr_Function_float %_78_n %uint_0
- OpStore %268 %267 None
- %269 = OpVectorShuffle %v2float %_75_k %_75_k 1 3
- %270 = OpVectorShuffle %v2float %_75_k %_75_k 0 2
- %271 = OpVectorTimesScalar %v2float %270 %float_0_00390625
- %272 = OpFAdd %v2float %269 %271
- %273 = OpVectorTimesScalar %v2float %272 %float_2
- %274 = OpCompositeConstruct %v2float %float_1 %float_1
- %275 = OpFSub %v2float %273 %274
- %276 = OpLoad %v2float %_78_n None
- %_skTemp8 = OpDot %float %275 %276
+ %275 = OpAccessChain %_ptr_Function_float %_78_n %uint_0
+ %276 = OpLoad %float %275 None
+ %277 = OpFSub %float %276 %float_1
+ %278 = OpAccessChain %_ptr_Function_float %_78_n %uint_0
+ OpStore %278 %277 None
+ %279 = OpVectorShuffle %v2float %_75_k %_75_k 1 3
+ %280 = OpVectorShuffle %v2float %_75_k %_75_k 0 2
+ %281 = OpVectorTimesScalar %v2float %280 %float_0_00390625
+ %282 = OpFAdd %v2float %279 %281
+ %283 = OpVectorTimesScalar %v2float %282 %float_2
+ %284 = OpCompositeConstruct %v2float %float_1 %float_1
+ %285 = OpFSub %v2float %283 %284
+ %286 = OpLoad %v2float %_78_n None
+ %_skTemp8 = OpDot %float %285 %286
OpStore %_80_p %_skTemp8
- %279 = OpLoad %float %_79_o None
- %280 = OpLoad %float %_80_p None
- %281 = OpCompositeExtract %float %_skTemp6 0
- %_skTemp9 = OpExtInst %float %48 FMix %279 %280 %281
- %283 = OpAccessChain %_ptr_Function_float %_78_n %uint_1
- %284 = OpLoad %float %283 None
- %285 = OpFSub %float %284 %float_1
- %286 = OpAccessChain %_ptr_Function_float %_78_n %uint_1
- OpStore %286 %285 None
- %287 = OpVectorShuffle %v2float %_76_l %_76_l 1 3
- %288 = OpVectorShuffle %v2float %_76_l %_76_l 0 2
- %289 = OpVectorTimesScalar %v2float %288 %float_0_00390625
- %290 = OpFAdd %v2float %287 %289
- %291 = OpVectorTimesScalar %v2float %290 %float_2
- %292 = OpCompositeConstruct %v2float %float_1 %float_1
- %293 = OpFSub %v2float %291 %292
- %294 = OpLoad %v2float %_78_n None
- %_skTemp10 = OpDot %float %293 %294
+ %289 = OpLoad %float %_79_o None
+ %290 = OpLoad %float %_80_p None
+ %291 = OpCompositeExtract %float %_skTemp6 0
+ %_skTemp9 = OpExtInst %float %48 FMix %289 %290 %291
+ %293 = OpAccessChain %_ptr_Function_float %_78_n %uint_1
+ %294 = OpLoad %float %293 None
+ %295 = OpFSub %float %294 %float_1
+ %296 = OpAccessChain %_ptr_Function_float %_78_n %uint_1
+ OpStore %296 %295 None
+ %297 = OpVectorShuffle %v2float %_76_l %_76_l 1 3
+ %298 = OpVectorShuffle %v2float %_76_l %_76_l 0 2
+ %299 = OpVectorTimesScalar %v2float %298 %float_0_00390625
+ %300 = OpFAdd %v2float %297 %299
+ %301 = OpVectorTimesScalar %v2float %300 %float_2
+ %302 = OpCompositeConstruct %v2float %float_1 %float_1
+ %303 = OpFSub %v2float %301 %302
+ %304 = OpLoad %v2float %_78_n None
+ %_skTemp10 = OpDot %float %303 %304
OpStore %_80_p %_skTemp10 None
- %296 = OpAccessChain %_ptr_Function_float %_78_n %uint_0
- %297 = OpLoad %float %296 None
- %298 = OpFAdd %float %297 %float_1
- %299 = OpAccessChain %_ptr_Function_float %_78_n %uint_0
- OpStore %299 %298 None
- %300 = OpVectorShuffle %v2float %_77_m %_77_m 1 3
- %301 = OpVectorShuffle %v2float %_77_m %_77_m 0 2
- %302 = OpVectorTimesScalar %v2float %301 %float_0_00390625
- %303 = OpFAdd %v2float %300 %302
- %304 = OpVectorTimesScalar %v2float %303 %float_2
- %305 = OpCompositeConstruct %v2float %float_1 %float_1
- %306 = OpFSub %v2float %304 %305
- %307 = OpLoad %v2float %_78_n None
- %_skTemp11 = OpDot %float %306 %307
+ %306 = OpAccessChain %_ptr_Function_float %_78_n %uint_0
+ %307 = OpLoad %float %306 None
+ %308 = OpFAdd %float %307 %float_1
+ %309 = OpAccessChain %_ptr_Function_float %_78_n %uint_0
+ OpStore %309 %308 None
+ %310 = OpVectorShuffle %v2float %_77_m %_77_m 1 3
+ %311 = OpVectorShuffle %v2float %_77_m %_77_m 0 2
+ %312 = OpVectorTimesScalar %v2float %311 %float_0_00390625
+ %313 = OpFAdd %v2float %310 %312
+ %314 = OpVectorTimesScalar %v2float %313 %float_2
+ %315 = OpCompositeConstruct %v2float %float_1 %float_1
+ %316 = OpFSub %v2float %314 %315
+ %317 = OpLoad %v2float %_78_n None
+ %_skTemp11 = OpDot %float %316 %317
OpStore %_79_o %_skTemp11 None
- %309 = OpLoad %float %_79_o None
- %310 = OpLoad %float %_80_p None
- %311 = OpCompositeExtract %float %_skTemp6 0
- %_skTemp12 = OpExtInst %float %48 FMix %309 %310 %311
- %313 = OpCompositeExtract %float %_skTemp6 1
- %_skTemp13 = OpExtInst %float %48 FMix %_skTemp9 %_skTemp12 %313
- %315 = OpLoad %int %_72_h None
- %316 = OpBitcast %uint %315
- %317 = OpExtInst %uint %48 UMin %316 %uint_3
- %318 = OpAccessChain %_ptr_Function_float %_71_g %317
- OpStore %318 %_skTemp13 None
- OpBranch %211
- %211 = OpLabel
- %319 = OpAccessChain %_ptr_Function_uint %tint_loop_idx_0 %uint_0
- %321 = OpLoad %uint %319 None
- %322 = OpIAdd %uint %321 %uint_1
- %323 = OpAccessChain %_ptr_Function_uint %tint_loop_idx_0 %uint_0
- OpStore %323 %322 None
- %324 = OpIEqual %bool %322 %uint_0
- %325 = OpSelect %uint %324 %uint_1 %uint_0
- %326 = OpAccessChain %_ptr_Function_uint %tint_loop_idx_0 %uint_1
- %327 = OpLoad %uint %326 None
- %328 = OpIAdd %uint %327 %325
- %329 = OpAccessChain %_ptr_Function_uint %tint_loop_idx_0 %uint_1
- OpStore %329 %328 None
- %330 = OpLoad %int %_72_h None
- %331 = OpIAdd %int %330 %int_1
- OpStore %_72_h %331 None
- %333 = OpLoad %int %_72_h None
- %334 = OpSGreaterThanEqual %bool %333 %int_4
- OpBranchConditional %334 %213 %212
- %213 = OpLabel
- %336 = OpLoad %v4float %_71_g None
- OpStore %_83_q %336
- %338 = OpINotEqual %bool %_56_d %int_0
- OpSelectionMerge %339 None
- OpBranchConditional %338 %340 %339
- %340 = OpLabel
- %341 = OpLoad %v4float %_83_q None
- %_skTemp14 = OpExtInst %v4float %48 FAbs %341
+ %319 = OpLoad %float %_79_o None
+ %320 = OpLoad %float %_80_p None
+ %321 = OpCompositeExtract %float %_skTemp6 0
+ %_skTemp12 = OpExtInst %float %48 FMix %319 %320 %321
+ %323 = OpCompositeExtract %float %_skTemp6 1
+ %_skTemp13 = OpExtInst %float %48 FMix %_skTemp9 %_skTemp12 %323
+ %325 = OpLoad %int %_72_h None
+ %326 = OpBitcast %uint %325
+ %327 = OpExtInst %uint %48 UMin %326 %uint_3
+ %328 = OpAccessChain %_ptr_Function_float %_71_g %327
+ OpStore %328 %_skTemp13 None
+ OpBranch %222
+ %222 = OpLabel
+ %329 = OpAccessChain %_ptr_Function_uint %tint_loop_idx_0 %uint_0
+ %331 = OpLoad %uint %329 None
+ %332 = OpIAdd %uint %331 %uint_1
+ %333 = OpAccessChain %_ptr_Function_uint %tint_loop_idx_0 %uint_0
+ OpStore %333 %332 None
+ %334 = OpIEqual %bool %332 %uint_0
+ %335 = OpSelect %uint %334 %uint_1 %uint_0
+ %336 = OpAccessChain %_ptr_Function_uint %tint_loop_idx_0 %uint_1
+ %337 = OpLoad %uint %336 None
+ %338 = OpIAdd %uint %337 %335
+ %339 = OpAccessChain %_ptr_Function_uint %tint_loop_idx_0 %uint_1
+ OpStore %339 %338 None
+ %340 = OpLoad %int %_72_h None
+ %341 = OpIAdd %int %340 %int_1
+ OpStore %_72_h %341 None
+ %343 = OpLoad %int %_72_h None
+ %344 = OpSGreaterThanEqual %bool %343 %int_4
+ OpBranchConditional %344 %224 %223
+ %224 = OpLabel
+ %346 = OpLoad %v4float %_71_g None
+ OpStore %_83_q %346
+ %348 = OpINotEqual %bool %_56_d %int_0
+ OpSelectionMerge %349 None
+ OpBranchConditional %348 %350 %349
+ %350 = OpLabel
+ %351 = OpLoad %v4float %_83_q None
+ %_skTemp14 = OpExtInst %v4float %48 FAbs %351
OpStore %_83_q %_skTemp14 None
- OpBranch %339
- %339 = OpLabel
- %343 = OpLoad %v4float %_58_l None
- %344 = OpLoad %v4float %_83_q None
- %345 = OpLoad %float %_60_n None
- %346 = OpVectorTimesScalar %v4float %344 %345
- %347 = OpFAdd %v4float %343 %346
- OpStore %_58_l %347 None
- %348 = OpLoad %v2float %_57_k None
- %349 = OpFMul %v2float %348 %350
- OpStore %_57_k %349 None
- %351 = OpLoad %float %_60_n None
- %352 = OpFMul %float %351 %float_0_5
- OpStore %_60_n %352 None
- %353 = OpLoad %v2float %_59_m None
- %354 = OpFMul %v2float %353 %350
- OpStore %_59_m %354 None
+ OpBranch %349
+ %349 = OpLabel
+ %353 = OpLoad %v4float %_58_l None
+ %354 = OpLoad %v4float %_83_q None
+ %355 = OpLoad %float %_60_n None
+ %356 = OpVectorTimesScalar %v4float %354 %355
+ %357 = OpFAdd %v4float %353 %356
+ OpStore %_58_l %357 None
+ %358 = OpLoad %v2float %_57_k None
+ %359 = OpFMul %v2float %358 %211
+ OpStore %_57_k %359 None
+ %360 = OpLoad %float %_60_n None
+ %361 = OpFMul %float %360 %float_0_5
+ OpStore %_60_n %361 None
+ %362 = OpLoad %v2float %_59_m None
+ %363 = OpFMul %v2float %362 %211
+ OpStore %_59_m %363 None
OpBranch %112
%114 = OpLabel
OpBranch %89
%112 = OpLabel
OpBranch %87
%87 = OpLabel
- %355 = OpAccessChain %_ptr_Function_uint %tint_loop_idx %uint_0
- %356 = OpLoad %uint %355 None
- %357 = OpIAdd %uint %356 %uint_1
- %358 = OpAccessChain %_ptr_Function_uint %tint_loop_idx %uint_0
- OpStore %358 %357 None
- %359 = OpIEqual %bool %357 %uint_0
- %360 = OpSelect %uint %359 %uint_1 %uint_0
- %361 = OpAccessChain %_ptr_Function_uint %tint_loop_idx %uint_1
- %362 = OpLoad %uint %361 None
- %363 = OpIAdd %uint %362 %360
- %364 = OpAccessChain %_ptr_Function_uint %tint_loop_idx %uint_1
- OpStore %364 %363 None
- %365 = OpLoad %int %_61_o None
- %366 = OpIAdd %int %365 %int_1
- OpStore %_61_o %366 None
+ %364 = OpAccessChain %_ptr_Function_uint %tint_loop_idx %uint_0
+ %365 = OpLoad %uint %364 None
+ %366 = OpIAdd %uint %365 %uint_1
+ %367 = OpAccessChain %_ptr_Function_uint %tint_loop_idx %uint_0
+ OpStore %367 %366 None
+ %368 = OpIEqual %bool %366 %uint_0
+ %369 = OpSelect %uint %368 %uint_1 %uint_0
+ %370 = OpAccessChain %_ptr_Function_uint %tint_loop_idx %uint_1
+ %371 = OpLoad %uint %370 None
+ %372 = OpIAdd %uint %371 %369
+ %373 = OpAccessChain %_ptr_Function_uint %tint_loop_idx %uint_1
+ OpStore %373 %372 None
+ %374 = OpLoad %int %_61_o None
+ %375 = OpIAdd %int %374 %int_1
+ OpStore %_61_o %375 None
OpBranch %88
%89 = OpLabel
- %367 = OpIEqual %bool %_56_d %int_0
- OpSelectionMerge %368 None
- OpBranchConditional %367 %369 %368
- %369 = OpLabel
- %370 = OpLoad %v4float %_58_l None
- %371 = OpFMul %v4float %370 %372
- %373 = OpFAdd %v4float %371 %372
- OpStore %_58_l %373 None
- OpBranch %368
- %368 = OpLabel
- %374 = OpLoad %v4float %_58_l None
- %_skTemp15 = OpExtInst %v4float %48 NClamp %374 %70 %376
+ %376 = OpIEqual %bool %_56_d %int_0
+ OpSelectionMerge %377 None
+ OpBranchConditional %376 %378 %377
+ %378 = OpLabel
+ %379 = OpLoad %v4float %_58_l None
+ %380 = OpFMul %v4float %379 %381
+ %382 = OpFAdd %v4float %380 %381
+ OpStore %_58_l %382 None
+ OpBranch %377
+ %377 = OpLabel
+ %383 = OpLoad %v4float %_58_l None
+ %_skTemp15 = OpExtInst %v4float %48 NClamp %383 %70 %385
OpStore %_58_l %_skTemp15 None
- %377 = OpLoad %v4float %_58_l None
- %378 = OpVectorShuffle %v3float %377 %377 0 1 2
- %380 = OpAccessChain %_ptr_Function_float %_58_l %uint_3
- %381 = OpLoad %float %380 None
- %382 = OpVectorTimesScalar %v3float %378 %381
- %383 = OpAccessChain %_ptr_Function_float %_58_l %uint_3
- %384 = OpLoad %float %383 None
- %385 = OpCompositeConstruct %v4float %382 %384
- %386 = OpVectorShuffle %v3float %385 %385 0 1 2
- %_skTemp16 = OpDot %float %388 %386
+ %386 = OpLoad %v4float %_58_l None
+ %387 = OpVectorShuffle %v3float %386 %386 0 1 2
+ %389 = OpAccessChain %_ptr_Function_float %_58_l %uint_3
+ %390 = OpLoad %float %389 None
+ %391 = OpVectorTimesScalar %v3float %387 %390
+ %392 = OpAccessChain %_ptr_Function_float %_58_l %uint_3
+ %393 = OpLoad %float %392 None
+ %394 = OpCompositeConstruct %v4float %391 %393
+ %395 = OpVectorShuffle %v3float %394 %394 0 1 2
+ %_skTemp16 = OpDot %float %397 %395
%_skTemp17 = OpExtInst %float %48 NClamp %_skTemp16 %float_0 %float_1
- %394 = OpCompositeConstruct %v4float %float_0 %float_0 %float_0 %_skTemp17
- OpStore %_84_a %394
- %396 = OpLoad %uint %shadingSsboIndex None
- %397 = OpAccessChain %_ptr_StorageBuffer__runtimearr_FSUniformData %_storage1 %uint_0
- %398 = OpArrayLength %uint %_storage1 0
- %399 = OpISub %uint %398 %uint_1
- %400 = OpExtInst %uint %48 UMin %396 %399
- %401 = OpAccessChain %_ptr_StorageBuffer_int %_storage1 %uint_0 %400 %uint_7
- %_85_d = OpLoad %int %401 None
- %404 = OpINotEqual %bool %_85_d %137
- OpSelectionMerge %405 None
- OpBranchConditional %404 %406 %407
- %406 = OpLabel
- %409 = OpAccessChain %_ptr_Function_float %_84_a %uint_1
- %410 = OpLoad %float %409 None
- %411 = OpAccessChain %_ptr_Function_float %_84_a %uint_2
- %412 = OpLoad %float %411 None
- %413 = OpFOrdLessThan %bool %410 %412
+ %403 = OpCompositeConstruct %v4float %float_0 %float_0 %float_0 %_skTemp17
+ OpStore %_84_a %403
+ %405 = OpLoad %uint %shadingSsboIndex None
+ %406 = OpAccessChain %_ptr_StorageBuffer__runtimearr_FSUniformData %_storage1 %uint_0
+ %407 = OpArrayLength %uint %_storage1 0
+ %408 = OpISub %uint %407 %uint_1
+ %409 = OpExtInst %uint %48 UMin %405 %408
+ %410 = OpAccessChain %_ptr_StorageBuffer_int %_storage1 %uint_0 %409 %uint_7
+ %_85_d = OpLoad %int %410 None
+ %413 = OpINotEqual %bool %_85_d %137
OpSelectionMerge %414 None
OpBranchConditional %413 %415 %416
%415 = OpLabel
- %417 = OpLoad %v4float %_84_a None
- %418 = OpVectorShuffle %v2float %417 %417 2 1
- %419 = OpCompositeConstruct %v4float %418 %float_n1 %float_0_666666687
- OpStore %_skTemp18 %419 None
- OpBranch %414
- %416 = OpLabel
- %422 = OpLoad %v4float %_84_a None
- %423 = OpVectorShuffle %v2float %422 %422 1 2
- %424 = OpCompositeConstruct %v4float %423 %float_0 %float_n0_333333343
- OpStore %_skTemp18 %424 None
- OpBranch %414
- %414 = OpLabel
+ %418 = OpAccessChain %_ptr_Function_float %_84_a %uint_1
+ %419 = OpLoad %float %418 None
+ %420 = OpAccessChain %_ptr_Function_float %_84_a %uint_2
+ %421 = OpLoad %float %420 None
+ %422 = OpFOrdLessThan %bool %419 %421
+ OpSelectionMerge %423 None
+ OpBranchConditional %422 %424 %425
+ %424 = OpLabel
+ %426 = OpLoad %v4float %_84_a None
+ %427 = OpVectorShuffle %v2float %426 %426 2 1
+ %428 = OpCompositeConstruct %v4float %427 %float_n1 %float_0_666666687
+ OpStore %_skTemp18 %428 None
+ OpBranch %423
+ %425 = OpLabel
+ %431 = OpLoad %v4float %_84_a None
+ %432 = OpVectorShuffle %v2float %431 %431 1 2
+ %433 = OpCompositeConstruct %v4float %432 %float_0 %float_n0_333333343
+ OpStore %_skTemp18 %433 None
+ OpBranch %423
+ %423 = OpLabel
%_86_e = OpLoad %v4float %_skTemp18 None
- %428 = OpAccessChain %_ptr_Function_float %_84_a %uint_0
- %429 = OpLoad %float %428 None
- %430 = OpCompositeExtract %float %_86_e 0
- %431 = OpFOrdLessThan %bool %429 %430
- OpSelectionMerge %432 None
- OpBranchConditional %431 %433 %434
- %433 = OpLabel
- %435 = OpCompositeExtract %float %_86_e 0
- %436 = OpAccessChain %_ptr_Function_float %_84_a %uint_0
- %437 = OpLoad %float %436 None
- %438 = OpVectorShuffle %v2float %_86_e %_86_e 1 3
- %439 = OpCompositeConstruct %v4float %435 %437 %438
- OpStore %_skTemp19 %439 None
- OpBranch %432
- %434 = OpLabel
- %440 = OpAccessChain %_ptr_Function_float %_84_a %uint_0
- %441 = OpLoad %float %440 None
- %442 = OpCompositeExtract %float %_86_e 0
- %443 = OpVectorShuffle %v2float %_86_e %_86_e 1 2
- %444 = OpCompositeConstruct %v4float %441 %442 %443
- OpStore %_skTemp19 %444 None
- OpBranch %432
- %432 = OpLabel
+ %437 = OpAccessChain %_ptr_Function_float %_84_a %uint_0
+ %438 = OpLoad %float %437 None
+ %439 = OpCompositeExtract %float %_86_e 0
+ %440 = OpFOrdLessThan %bool %438 %439
+ OpSelectionMerge %441 None
+ OpBranchConditional %440 %442 %443
+ %442 = OpLabel
+ %444 = OpCompositeExtract %float %_86_e 0
+ %445 = OpAccessChain %_ptr_Function_float %_84_a %uint_0
+ %446 = OpLoad %float %445 None
+ %447 = OpVectorShuffle %v2float %_86_e %_86_e 1 3
+ %448 = OpCompositeConstruct %v4float %444 %446 %447
+ OpStore %_skTemp19 %448 None
+ OpBranch %441
+ %443 = OpLabel
+ %449 = OpAccessChain %_ptr_Function_float %_84_a %uint_0
+ %450 = OpLoad %float %449 None
+ %451 = OpCompositeExtract %float %_86_e 0
+ %452 = OpVectorShuffle %v2float %_86_e %_86_e 1 2
+ %453 = OpCompositeConstruct %v4float %450 %451 %452
+ OpStore %_skTemp19 %453 None
+ OpBranch %441
+ %441 = OpLabel
%_87_f = OpLoad %v4float %_skTemp19 None
%_88_h = OpCompositeExtract %float %_87_f 0
- %447 = OpCompositeExtract %float %_87_f 1
- %448 = OpCompositeExtract %float %_87_f 2
- %_skTemp20 = OpExtInst %float %48 FMin %447 %448
+ %456 = OpCompositeExtract %float %_87_f 1
+ %457 = OpCompositeExtract %float %_87_f 2
+ %_skTemp20 = OpExtInst %float %48 FMin %456 %457
%_89_i = OpFSub %float %_88_h %_skTemp20
- %451 = OpFMul %float %_89_i %float_0_5
- %_90_j = OpFSub %float %_88_h %451
- %453 = OpCompositeExtract %float %_87_f 3
- %454 = OpCompositeExtract %float %_87_f 1
- %455 = OpCompositeExtract %float %_87_f 2
- %456 = OpFSub %float %454 %455
- %457 = OpFMul %float %_89_i %float_6
- %459 = OpFAdd %float %457 %float_9_99999975en05
- %461 = OpFDiv %float %456 %459
- %462 = OpFAdd %float %453 %461
- %_skTemp21 = OpExtInst %float %48 FAbs %462
- %464 = OpFMul %float %_90_j %float_2
- %465 = OpAccessChain %_ptr_Function_float %_84_a %uint_3
- %466 = OpLoad %float %465 None
- %467 = OpFSub %float %464 %466
- %_skTemp22 = OpExtInst %float %48 FAbs %467
- %469 = OpAccessChain %_ptr_Function_float %_84_a %uint_3
- %470 = OpLoad %float %469 None
- %471 = OpFAdd %float %470 %float_9_99999975en05
- %472 = OpFSub %float %471 %_skTemp22
- %_92_l = OpFDiv %float %_89_i %472
+ %460 = OpFMul %float %_89_i %float_0_5
+ %_90_j = OpFSub %float %_88_h %460
+ %462 = OpCompositeExtract %float %_87_f 3
+ %463 = OpCompositeExtract %float %_87_f 1
+ %464 = OpCompositeExtract %float %_87_f 2
+ %465 = OpFSub %float %463 %464
+ %466 = OpFMul %float %_89_i %float_6
+ %468 = OpFAdd %float %466 %float_9_99999975en05
+ %470 = OpFDiv %float %465 %468
+ %471 = OpFAdd %float %462 %470
+ %_skTemp21 = OpExtInst %float %48 FAbs %471
+ %473 = OpFMul %float %_90_j %float_2
%474 = OpAccessChain %_ptr_Function_float %_84_a %uint_3
%475 = OpLoad %float %474 None
- %476 = OpFAdd %float %475 %float_9_99999975en05
- %_93_m = OpFDiv %float %_90_j %476
+ %476 = OpFSub %float %473 %475
+ %_skTemp22 = OpExtInst %float %48 FAbs %476
%478 = OpAccessChain %_ptr_Function_float %_84_a %uint_3
%479 = OpLoad %float %478 None
- %480 = OpCompositeConstruct %v4float %_skTemp21 %_92_l %_93_m %479
- OpStore %_84_a %480 None
- OpBranch %405
- %407 = OpLabel
- %481 = OpAccessChain %_ptr_Function_float %_84_a %uint_3
- %482 = OpLoad %float %481 None
- %_skTemp23 = OpExtInst %float %48 FMax %482 %float_9_99999975en05
- %484 = OpLoad %v4float %_84_a None
- %485 = OpVectorShuffle %v3float %484 %484 0 1 2
- %486 = OpCompositeConstruct %v3float %_skTemp23 %_skTemp23 %_skTemp23
- %487 = OpFDiv %v3float %485 %486
- %488 = OpAccessChain %_ptr_Function_float %_84_a %uint_3
- %489 = OpLoad %float %488 None
- %490 = OpCompositeConstruct %v4float %487 %489
- OpStore %_84_a %490 None
- OpBranch %405
- %405 = OpLabel
- %491 = OpLoad %uint %shadingSsboIndex None
- %492 = OpAccessChain %_ptr_StorageBuffer__runtimearr_FSUniformData %_storage1 %uint_0
- %493 = OpArrayLength %uint %_storage1 0
- %494 = OpISub %uint %493 %uint_1
- %495 = OpExtInst %uint %48 UMin %491 %494
- %496 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %_storage1 %uint_0 %495 %uint_5
- %499 = OpLoad %mat4v4float %496 None
- %500 = OpLoad %v4float %_84_a None
- %501 = OpMatrixTimesVector %v4float %499 %500
- %502 = OpLoad %uint %shadingSsboIndex None
- %503 = OpAccessChain %_ptr_StorageBuffer__runtimearr_FSUniformData %_storage1 %uint_0
- %504 = OpArrayLength %uint %_storage1 0
- %505 = OpISub %uint %504 %uint_1
- %506 = OpExtInst %uint %48 UMin %502 %505
- %507 = OpAccessChain %_ptr_StorageBuffer_v4float %_storage1 %uint_0 %506 %uint_6
- %510 = OpLoad %v4float %507 None
- %511 = OpFAdd %v4float %501 %510
- OpStore %_94_f %511
- %513 = OpINotEqual %bool %_85_d %137
- OpSelectionMerge %514 None
- OpBranchConditional %513 %515 %516
- %515 = OpLabel
- %517 = OpAccessChain %_ptr_Function_float %_94_f %uint_2
- %518 = OpLoad %float %517 None
- %519 = OpFMul %float %float_2 %518
- %520 = OpFSub %float %519 %float_1
- %_skTemp24 = OpExtInst %float %48 FAbs %520
- %522 = OpFSub %float %float_1 %_skTemp24
- %523 = OpAccessChain %_ptr_Function_float %_94_f %uint_1
- %524 = OpLoad %float %523 None
- %_95_b = OpFMul %float %522 %524
- %526 = OpLoad %v4float %_94_f None
- %527 = OpVectorShuffle %v3float %526 %526 0 0 0
- %_96_c = OpFAdd %v3float %527 %529
+ %480 = OpFAdd %float %479 %float_9_99999975en05
+ %481 = OpFSub %float %480 %_skTemp22
+ %_92_l = OpFDiv %float %_89_i %481
+ %483 = OpAccessChain %_ptr_Function_float %_84_a %uint_3
+ %484 = OpLoad %float %483 None
+ %485 = OpFAdd %float %484 %float_9_99999975en05
+ %_93_m = OpFDiv %float %_90_j %485
+ %487 = OpAccessChain %_ptr_Function_float %_84_a %uint_3
+ %488 = OpLoad %float %487 None
+ %489 = OpCompositeConstruct %v4float %_skTemp21 %_92_l %_93_m %488
+ OpStore %_84_a %489 None
+ OpBranch %414
+ %416 = OpLabel
+ %490 = OpAccessChain %_ptr_Function_float %_84_a %uint_3
+ %491 = OpLoad %float %490 None
+ %_skTemp23 = OpExtInst %float %48 FMax %491 %float_9_99999975en05
+ %493 = OpLoad %v4float %_84_a None
+ %494 = OpVectorShuffle %v3float %493 %493 0 1 2
+ %495 = OpCompositeConstruct %v3float %_skTemp23 %_skTemp23 %_skTemp23
+ %496 = OpFDiv %v3float %494 %495
+ %497 = OpAccessChain %_ptr_Function_float %_84_a %uint_3
+ %498 = OpLoad %float %497 None
+ %499 = OpCompositeConstruct %v4float %496 %498
+ OpStore %_84_a %499 None
+ OpBranch %414
+ %414 = OpLabel
+ %500 = OpLoad %uint %shadingSsboIndex None
+ %501 = OpAccessChain %_ptr_StorageBuffer__runtimearr_FSUniformData %_storage1 %uint_0
+ %502 = OpArrayLength %uint %_storage1 0
+ %503 = OpISub %uint %502 %uint_1
+ %504 = OpExtInst %uint %48 UMin %500 %503
+ %505 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %_storage1 %uint_0 %504 %uint_5
+ %508 = OpLoad %mat4v4float %505 None
+ %509 = OpLoad %v4float %_84_a None
+ %510 = OpMatrixTimesVector %v4float %508 %509
+ %511 = OpLoad %uint %shadingSsboIndex None
+ %512 = OpAccessChain %_ptr_StorageBuffer__runtimearr_FSUniformData %_storage1 %uint_0
+ %513 = OpArrayLength %uint %_storage1 0
+ %514 = OpISub %uint %513 %uint_1
+ %515 = OpExtInst %uint %48 UMin %511 %514
+ %516 = OpAccessChain %_ptr_StorageBuffer_v4float %_storage1 %uint_0 %515 %uint_6
+ %519 = OpLoad %v4float %516 None
+ %520 = OpFAdd %v4float %510 %519
+ OpStore %_94_f %520
+ %522 = OpINotEqual %bool %_85_d %137
+ OpSelectionMerge %523 None
+ OpBranchConditional %522 %524 %525
+ %524 = OpLabel
+ %526 = OpAccessChain %_ptr_Function_float %_94_f %uint_2
+ %527 = OpLoad %float %526 None
+ %528 = OpFMul %float %float_2 %527
+ %529 = OpFSub %float %528 %float_1
+ %_skTemp24 = OpExtInst %float %48 FAbs %529
+ %531 = OpFSub %float %float_1 %_skTemp24
+ %532 = OpAccessChain %_ptr_Function_float %_94_f %uint_1
+ %533 = OpLoad %float %532 None
+ %_95_b = OpFMul %float %531 %533
+ %535 = OpLoad %v4float %_94_f None
+ %536 = OpVectorShuffle %v3float %535 %535 0 0 0
+ %_96_c = OpFAdd %v3float %536 %538
%_skTemp25 = OpExtInst %v3float %48 Fract %_96_c
- %532 = OpVectorTimesScalar %v3float %_skTemp25 %float_6
- %533 = OpCompositeConstruct %v3float %float_3 %float_3 %float_3
- %535 = OpFSub %v3float %532 %533
- %_skTemp26 = OpExtInst %v3float %48 FAbs %535
- %537 = OpCompositeConstruct %v3float %float_1 %float_1 %float_1
- %538 = OpFSub %v3float %_skTemp26 %537
- %_skTemp27 = OpExtInst %v3float %48 NClamp %538 %540 %541
- %542 = OpCompositeConstruct %v3float %float_0_5 %float_0_5 %float_0_5
- %543 = OpFSub %v3float %_skTemp27 %542
- %544 = OpVectorTimesScalar %v3float %543 %_95_b
- %545 = OpAccessChain %_ptr_Function_float %_94_f %uint_2
- %546 = OpLoad %float %545 None
- %547 = OpCompositeConstruct %v3float %546 %546 %546
- %548 = OpFAdd %v3float %544 %547
- %549 = OpAccessChain %_ptr_Function_float %_94_f %uint_3
- %550 = OpLoad %float %549 None
- %551 = OpVectorTimesScalar %v3float %548 %550
- %552 = OpAccessChain %_ptr_Function_float %_94_f %uint_3
- %553 = OpLoad %float %552 None
- %554 = OpCompositeConstruct %v4float %551 %553
- %_skTemp28 = OpExtInst %v4float %48 NClamp %554 %70 %376
+ %541 = OpVectorTimesScalar %v3float %_skTemp25 %float_6
+ %542 = OpCompositeConstruct %v3float %float_3 %float_3 %float_3
+ %543 = OpFSub %v3float %541 %542
+ %_skTemp26 = OpExtInst %v3float %48 FAbs %543
+ %545 = OpCompositeConstruct %v3float %float_1 %float_1 %float_1
+ %546 = OpFSub %v3float %_skTemp26 %545
+ %_skTemp27 = OpExtInst %v3float %48 NClamp %546 %548 %549
+ %550 = OpCompositeConstruct %v3float %float_0_5 %float_0_5 %float_0_5
+ %551 = OpFSub %v3float %_skTemp27 %550
+ %552 = OpVectorTimesScalar %v3float %551 %_95_b
+ %553 = OpAccessChain %_ptr_Function_float %_94_f %uint_2
+ %554 = OpLoad %float %553 None
+ %555 = OpCompositeConstruct %v3float %554 %554 %554
+ %556 = OpFAdd %v3float %552 %555
+ %557 = OpAccessChain %_ptr_Function_float %_94_f %uint_3
+ %558 = OpLoad %float %557 None
+ %559 = OpVectorTimesScalar %v3float %556 %558
+ %560 = OpAccessChain %_ptr_Function_float %_94_f %uint_3
+ %561 = OpLoad %float %560 None
+ %562 = OpCompositeConstruct %v4float %559 %561
+ %_skTemp28 = OpExtInst %v4float %48 NClamp %562 %70 %385
OpStore %_94_f %_skTemp28 None
- OpBranch %514
- %516 = OpLabel
- %556 = OpLoad %uint %shadingSsboIndex None
- %557 = OpAccessChain %_ptr_StorageBuffer__runtimearr_FSUniformData %_storage1 %uint_0
- %558 = OpArrayLength %uint %_storage1 0
- %559 = OpISub %uint %558 %uint_1
- %560 = OpExtInst %uint %48 UMin %556 %559
- %561 = OpAccessChain %_ptr_StorageBuffer_int %_storage1 %uint_0 %560 %uint_8
- %563 = OpLoad %int %561 None
- %564 = OpINotEqual %bool %563 %137
- OpSelectionMerge %565 None
- OpBranchConditional %564 %566 %567
- %566 = OpLabel
- %568 = OpLoad %v4float %_94_f None
- %_skTemp29 = OpExtInst %v4float %48 NClamp %568 %70 %376
+ OpBranch %523
+ %525 = OpLabel
+ %564 = OpLoad %uint %shadingSsboIndex None
+ %565 = OpAccessChain %_ptr_StorageBuffer__runtimearr_FSUniformData %_storage1 %uint_0
+ %566 = OpArrayLength %uint %_storage1 0
+ %567 = OpISub %uint %566 %uint_1
+ %568 = OpExtInst %uint %48 UMin %564 %567
+ %569 = OpAccessChain %_ptr_StorageBuffer_int %_storage1 %uint_0 %568 %uint_8
+ %571 = OpLoad %int %569 None
+ %572 = OpINotEqual %bool %571 %137
+ OpSelectionMerge %573 None
+ OpBranchConditional %572 %574 %575
+ %574 = OpLabel
+ %576 = OpLoad %v4float %_94_f None
+ %_skTemp29 = OpExtInst %v4float %48 NClamp %576 %70 %385
OpStore %_94_f %_skTemp29 None
- OpBranch %565
- %567 = OpLabel
- %570 = OpAccessChain %_ptr_Function_float %_94_f %uint_3
- %571 = OpLoad %float %570 None
- %_skTemp30 = OpExtInst %float %48 NClamp %571 %float_0 %float_1
- %573 = OpAccessChain %_ptr_Function_float %_94_f %uint_3
- OpStore %573 %_skTemp30 None
- OpBranch %565
- %565 = OpLabel
- %574 = OpLoad %v4float %_94_f None
- %575 = OpVectorShuffle %v3float %574 %574 0 1 2
- %576 = OpAccessChain %_ptr_Function_float %_94_f %uint_3
- %577 = OpLoad %float %576 None
- %578 = OpVectorTimesScalar %v3float %575 %577
- %579 = OpAccessChain %_ptr_Function_float %_94_f %uint_3
- %580 = OpLoad %float %579 None
- %581 = OpCompositeConstruct %v4float %578 %580
- OpStore %_94_f %581 None
- OpBranch %514
- %514 = OpLabel
+ OpBranch %573
+ %575 = OpLabel
+ %578 = OpAccessChain %_ptr_Function_float %_94_f %uint_3
+ %579 = OpLoad %float %578 None
+ %_skTemp30 = OpExtInst %float %48 NClamp %579 %float_0 %float_1
+ %581 = OpAccessChain %_ptr_Function_float %_94_f %uint_3
+ OpStore %581 %_skTemp30 None
+ OpBranch %573
+ %573 = OpLabel
+ %582 = OpLoad %v4float %_94_f None
+ %583 = OpVectorShuffle %v3float %582 %582 0 1 2
+ %584 = OpAccessChain %_ptr_Function_float %_94_f %uint_3
+ %585 = OpLoad %float %584 None
+ %586 = OpVectorTimesScalar %v3float %583 %585
+ %587 = OpAccessChain %_ptr_Function_float %_94_f %uint_3
+ %588 = OpLoad %float %587 None
+ %589 = OpCompositeConstruct %v4float %586 %588
+ OpStore %_94_f %589 None
+ OpBranch %523
+ %523 = OpLabel
%outColor_0 = OpLoad %v4float %_94_f None
- %583 = OpAccessChain %_ptr_Function_v4float %_stageOut_root %uint_0
- OpStore %583 %outColor_0 None
+ %591 = OpAccessChain %_ptr_Function_v4float %_stageOut_root %uint_0
+ OpStore %591 %outColor_0 None
OpReturn
OpFunctionEnd
- %main_inner = OpFunction %FSOut None %586
+ %main_inner = OpFunction %FSOut None %594
%_stageIn_0 = OpFunctionParameter %FSIn
- %587 = OpLabel
- %_stageOut = OpVariable %_ptr_Function_FSOut Function %589
- %590 = OpFunctionCall %void %_skslMain %_stageIn_0 %_stageOut
- %591 = OpLoad %FSOut %_stageOut None
- OpReturnValue %591
+ %595 = OpLabel
+ %_stageOut = OpVariable %_ptr_Function_FSOut Function %597
+ %598 = OpFunctionCall %void %_skslMain %_stageIn_0 %_stageOut
+ %599 = OpLoad %FSOut %_stageOut None
+ OpReturnValue %599
OpFunctionEnd
- %main = OpFunction %void None %593
- %594 = OpLabel
- %595 = OpLoad %v2uint %main_loc0_Input None
- %596 = OpLoad %v2float %main_loc1_Input None
- %597 = OpCompositeConstruct %FSIn %595 %596
- %598 = OpFunctionCall %FSOut %main_inner %597
- %599 = OpCompositeExtract %v4float %598 0
- OpStore %main_loc0_Output %599 None
+ %main = OpFunction %void None %601
+ %602 = OpLabel
+ %603 = OpLoad %v2uint %main_loc0_Input None
+ %604 = OpLoad %v2float %main_loc1_Input None
+ %605 = OpCompositeConstruct %FSIn %603 %604
+ %606 = OpFunctionCall %FSOut %main_inner %605
+ %607 = OpCompositeExtract %v4float %606 0
+ OpStore %main_loc0_Output %607 None
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.glsl b/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.glsl
index 5223f45..6ddc16e 100644
--- a/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.glsl
@@ -14,7 +14,9 @@
f16vec2 arg_0 = f16vec2(2.0hf);
f16vec2 arg_1 = f16vec2(4.0hf);
f16vec2 arg_2 = f16vec2(3.0hf);
- f16vec2 res = smoothstep(arg_0, arg_1, arg_2);
+ f16vec2 v_1 = arg_0;
+ f16vec2 v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), f16vec2(0.0hf), f16vec2(1.0hf));
+ f16vec2 res = (v_2 * (v_2 * (f16vec2(3.0hf) - (f16vec2(2.0hf) * v_2))));
return res;
}
void main() {
@@ -34,7 +36,9 @@
f16vec2 arg_0 = f16vec2(2.0hf);
f16vec2 arg_1 = f16vec2(4.0hf);
f16vec2 arg_2 = f16vec2(3.0hf);
- f16vec2 res = smoothstep(arg_0, arg_1, arg_2);
+ f16vec2 v_1 = arg_0;
+ f16vec2 v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), f16vec2(0.0hf), f16vec2(1.0hf));
+ f16vec2 res = (v_2 * (v_2 * (f16vec2(3.0hf) - (f16vec2(2.0hf) * v_2))));
return res;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -58,7 +62,9 @@
f16vec2 arg_0 = f16vec2(2.0hf);
f16vec2 arg_1 = f16vec2(4.0hf);
f16vec2 arg_2 = f16vec2(3.0hf);
- f16vec2 res = smoothstep(arg_0, arg_1, arg_2);
+ f16vec2 v = arg_0;
+ f16vec2 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), f16vec2(0.0hf), f16vec2(1.0hf));
+ f16vec2 res = (v_1 * (v_1 * (f16vec2(3.0hf) - (f16vec2(2.0hf) * v_1))));
return res;
}
VertexOutput vertex_main_inner() {
@@ -68,10 +74,10 @@
return tint_symbol;
}
void main() {
- VertexOutput v = vertex_main_inner();
- gl_Position = v.pos;
+ VertexOutput v_2 = vertex_main_inner();
+ gl_Position = v_2.pos;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
- vertex_main_loc0_Output = v.prevent_dce;
+ vertex_main_loc0_Output = v_2.prevent_dce;
gl_PointSize = 1.0f;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.ir.dxc.hlsl
index 7a08288..61267af 100644
--- a/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.ir.dxc.hlsl
@@ -7,7 +7,9 @@
vector<float16_t, 2> arg_0 = (float16_t(2.0h)).xx;
vector<float16_t, 2> arg_1 = (float16_t(4.0h)).xx;
vector<float16_t, 2> arg_2 = (float16_t(3.0h)).xx;
- vector<float16_t, 2> res = smoothstep(arg_0, arg_1, arg_2);
+ vector<float16_t, 2> v = arg_0;
+ vector<float16_t, 2> v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (float16_t(0.0h)).xx, (float16_t(1.0h)).xx);
+ vector<float16_t, 2> res = (v_1 * (v_1 * ((float16_t(3.0h)).xx - ((float16_t(2.0h)).xx * v_1))));
return res;
}
@@ -24,7 +26,9 @@
vector<float16_t, 2> arg_0 = (float16_t(2.0h)).xx;
vector<float16_t, 2> arg_1 = (float16_t(4.0h)).xx;
vector<float16_t, 2> arg_2 = (float16_t(3.0h)).xx;
- vector<float16_t, 2> res = smoothstep(arg_0, arg_1, arg_2);
+ vector<float16_t, 2> v = arg_0;
+ vector<float16_t, 2> v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (float16_t(0.0h)).xx, (float16_t(1.0h)).xx);
+ vector<float16_t, 2> res = (v_1 * (v_1 * ((float16_t(3.0h)).xx - ((float16_t(2.0h)).xx * v_1))));
return res;
}
@@ -51,7 +55,9 @@
vector<float16_t, 2> arg_0 = (float16_t(2.0h)).xx;
vector<float16_t, 2> arg_1 = (float16_t(4.0h)).xx;
vector<float16_t, 2> arg_2 = (float16_t(3.0h)).xx;
- vector<float16_t, 2> res = smoothstep(arg_0, arg_1, arg_2);
+ vector<float16_t, 2> v = arg_0;
+ vector<float16_t, 2> v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (float16_t(0.0h)).xx, (float16_t(1.0h)).xx);
+ vector<float16_t, 2> res = (v_1 * (v_1 * ((float16_t(3.0h)).xx - ((float16_t(2.0h)).xx * v_1))));
return res;
}
@@ -59,13 +65,13 @@
VertexOutput tint_symbol = (VertexOutput)0;
tint_symbol.pos = (0.0f).xxxx;
tint_symbol.prevent_dce = smoothstep_12c031();
- VertexOutput v = tint_symbol;
- return v;
+ VertexOutput v_2 = tint_symbol;
+ return v_2;
}
vertex_main_outputs vertex_main() {
- VertexOutput v_1 = vertex_main_inner();
- vertex_main_outputs v_2 = {v_1.prevent_dce, v_1.pos};
- return v_2;
+ VertexOutput v_3 = vertex_main_inner();
+ vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos};
+ return v_4;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.ir.msl
index c54ffb7..882805d 100644
--- a/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.ir.msl
@@ -12,7 +12,9 @@
half2 arg_0 = half2(2.0h);
half2 arg_1 = half2(4.0h);
half2 arg_2 = half2(3.0h);
- half2 res = smoothstep(arg_0, arg_1, arg_2);
+ half2 const v = arg_0;
+ half2 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), half2(0.0h), half2(1.0h));
+ half2 res = (v_1 * (v_1 * (half2(3.0h) - (half2(2.0h) * v_1))));
return res;
}
@@ -34,7 +36,9 @@
half2 arg_0 = half2(2.0h);
half2 arg_1 = half2(4.0h);
half2 arg_2 = half2(3.0h);
- half2 res = smoothstep(arg_0, arg_1, arg_2);
+ half2 const v = arg_0;
+ half2 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), half2(0.0h), half2(1.0h));
+ half2 res = (v_1 * (v_1 * (half2(3.0h) - (half2(2.0h) * v_1))));
return res;
}
@@ -62,7 +66,9 @@
half2 arg_0 = half2(2.0h);
half2 arg_1 = half2(4.0h);
half2 arg_2 = half2(3.0h);
- half2 res = smoothstep(arg_0, arg_1, arg_2);
+ half2 const v = arg_0;
+ half2 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), half2(0.0h), half2(1.0h));
+ half2 res = (v_1 * (v_1 * (half2(3.0h) - (half2(2.0h) * v_1))));
return res;
}
@@ -74,9 +80,9 @@
}
vertex vertex_main_outputs vertex_main() {
- VertexOutput const v = vertex_main_inner();
+ VertexOutput const v_2 = vertex_main_inner();
vertex_main_outputs tint_wrapper_result = {};
- tint_wrapper_result.VertexOutput_pos = v.pos;
- tint_wrapper_result.VertexOutput_prevent_dce = v.prevent_dce;
+ tint_wrapper_result.VertexOutput_pos = v_2.pos;
+ tint_wrapper_result.VertexOutput_prevent_dce = v_2.prevent_dce;
return tint_wrapper_result;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.spvasm b/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.spvasm
index 8222d80..8326d6a 100644
--- a/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/smoothstep/12c031.wgsl.expected.spvasm
@@ -4,13 +4,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 35
+; Bound: 45
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
- %23 = OpExtInstImport "GLSL.std.450"
+ %26 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %fragment_main "fragment_main"
OpExecutionMode %fragment_main OriginUpperLeft
@@ -40,8 +40,11 @@
%14 = OpConstantComposite %v2half %half_0x1p_2 %half_0x1p_2
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
%17 = OpConstantComposite %v2half %half_0x1_8p_1 %half_0x1_8p_1
+ %27 = OpConstantNull %v2half
+%half_0x1p_0 = OpConstant %half 0x1p+0
+ %28 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
%void = OpTypeVoid
- %28 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_StorageBuffer_v2half = OpTypePointer StorageBuffer %v2half
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -57,16 +60,23 @@
%19 = OpLoad %v2half %arg_0 None
%20 = OpLoad %v2half %arg_1 None
%21 = OpLoad %v2half %arg_2 None
- %22 = OpExtInst %v2half %23 SmoothStep %19 %20 %21
- OpStore %res %22
- %25 = OpLoad %v2half %res None
- OpReturnValue %25
+ %22 = OpFSub %v2half %21 %19
+ %23 = OpFSub %v2half %20 %19
+ %24 = OpFDiv %v2half %22 %23
+ %25 = OpExtInst %v2half %26 NClamp %24 %27 %28
+ %30 = OpFMul %v2half %11 %25
+ %31 = OpFSub %v2half %17 %30
+ %32 = OpFMul %v2half %25 %31
+ %33 = OpFMul %v2half %25 %32
+ OpStore %res %33
+ %35 = OpLoad %v2half %res None
+ OpReturnValue %35
OpFunctionEnd
-%fragment_main = OpFunction %void None %28
- %29 = OpLabel
- %30 = OpFunctionCall %v2half %smoothstep_12c031
- %31 = OpAccessChain %_ptr_StorageBuffer_v2half %1 %uint_0
- OpStore %31 %30 None
+%fragment_main = OpFunction %void None %38
+ %39 = OpLabel
+ %40 = OpFunctionCall %v2half %smoothstep_12c031
+ %41 = OpAccessChain %_ptr_StorageBuffer_v2half %1 %uint_0
+ OpStore %41 %40 None
OpReturn
OpFunctionEnd
;
@@ -75,13 +85,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 35
+; Bound: 45
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
- %23 = OpExtInstImport "GLSL.std.450"
+ %26 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %compute_main LocalSize 1 1 1
@@ -111,8 +121,11 @@
%14 = OpConstantComposite %v2half %half_0x1p_2 %half_0x1p_2
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
%17 = OpConstantComposite %v2half %half_0x1_8p_1 %half_0x1_8p_1
+ %27 = OpConstantNull %v2half
+%half_0x1p_0 = OpConstant %half 0x1p+0
+ %28 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
%void = OpTypeVoid
- %28 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_StorageBuffer_v2half = OpTypePointer StorageBuffer %v2half
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -128,16 +141,23 @@
%19 = OpLoad %v2half %arg_0 None
%20 = OpLoad %v2half %arg_1 None
%21 = OpLoad %v2half %arg_2 None
- %22 = OpExtInst %v2half %23 SmoothStep %19 %20 %21
- OpStore %res %22
- %25 = OpLoad %v2half %res None
- OpReturnValue %25
+ %22 = OpFSub %v2half %21 %19
+ %23 = OpFSub %v2half %20 %19
+ %24 = OpFDiv %v2half %22 %23
+ %25 = OpExtInst %v2half %26 NClamp %24 %27 %28
+ %30 = OpFMul %v2half %11 %25
+ %31 = OpFSub %v2half %17 %30
+ %32 = OpFMul %v2half %25 %31
+ %33 = OpFMul %v2half %25 %32
+ OpStore %res %33
+ %35 = OpLoad %v2half %res None
+ OpReturnValue %35
OpFunctionEnd
-%compute_main = OpFunction %void None %28
- %29 = OpLabel
- %30 = OpFunctionCall %v2half %smoothstep_12c031
- %31 = OpAccessChain %_ptr_StorageBuffer_v2half %1 %uint_0
- OpStore %31 %30 None
+%compute_main = OpFunction %void None %38
+ %39 = OpLabel
+ %40 = OpFunctionCall %v2half %smoothstep_12c031
+ %41 = OpAccessChain %_ptr_StorageBuffer_v2half %1 %uint_0
+ OpStore %41 %40 None
OpReturn
OpFunctionEnd
;
@@ -146,14 +166,14 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 55
+; Bound: 65
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
- %28 = OpExtInstImport "GLSL.std.450"
+ %31 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %vertex_main_position_Output %vertex_main_loc0_Output %vertex_main___point_size_Output
OpName %vertex_main_position_Output "vertex_main_position_Output"
@@ -194,17 +214,20 @@
%19 = OpConstantComposite %v2half %half_0x1p_2 %half_0x1p_2
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
%22 = OpConstantComposite %v2half %half_0x1_8p_1 %half_0x1_8p_1
+ %32 = OpConstantNull %v2half
+%half_0x1p_0 = OpConstant %half 0x1p+0
+ %33 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
%VertexOutput = OpTypeStruct %v4float %v2half
- %33 = OpTypeFunction %VertexOutput
+ %43 = OpTypeFunction %VertexOutput
%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
- %37 = OpConstantNull %VertexOutput
+ %47 = OpConstantNull %VertexOutput
%_ptr_Function_v4float = OpTypePointer Function %v4float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
- %42 = OpConstantNull %v4float
+ %52 = OpConstantNull %v4float
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
- %49 = OpTypeFunction %void
+ %59 = OpTypeFunction %void
%float_1 = OpConstant %float 1
%smoothstep_12c031 = OpFunction %v2half None %12
%13 = OpLabel
@@ -218,29 +241,36 @@
%24 = OpLoad %v2half %arg_0 None
%25 = OpLoad %v2half %arg_1 None
%26 = OpLoad %v2half %arg_2 None
- %27 = OpExtInst %v2half %28 SmoothStep %24 %25 %26
- OpStore %res %27
- %30 = OpLoad %v2half %res None
- OpReturnValue %30
+ %27 = OpFSub %v2half %26 %24
+ %28 = OpFSub %v2half %25 %24
+ %29 = OpFDiv %v2half %27 %28
+ %30 = OpExtInst %v2half %31 NClamp %29 %32 %33
+ %35 = OpFMul %v2half %16 %30
+ %36 = OpFSub %v2half %22 %35
+ %37 = OpFMul %v2half %30 %36
+ %38 = OpFMul %v2half %30 %37
+ OpStore %res %38
+ %40 = OpLoad %v2half %res None
+ OpReturnValue %40
OpFunctionEnd
-%vertex_main_inner = OpFunction %VertexOutput None %33
- %34 = OpLabel
- %out = OpVariable %_ptr_Function_VertexOutput Function %37
- %38 = OpAccessChain %_ptr_Function_v4float %out %uint_0
- OpStore %38 %42 None
- %43 = OpAccessChain %_ptr_Function_v2half %out %uint_1
- %45 = OpFunctionCall %v2half %smoothstep_12c031
- OpStore %43 %45 None
- %46 = OpLoad %VertexOutput %out None
- OpReturnValue %46
+%vertex_main_inner = OpFunction %VertexOutput None %43
+ %44 = OpLabel
+ %out = OpVariable %_ptr_Function_VertexOutput Function %47
+ %48 = OpAccessChain %_ptr_Function_v4float %out %uint_0
+ OpStore %48 %52 None
+ %53 = OpAccessChain %_ptr_Function_v2half %out %uint_1
+ %55 = OpFunctionCall %v2half %smoothstep_12c031
+ OpStore %53 %55 None
+ %56 = OpLoad %VertexOutput %out None
+ OpReturnValue %56
OpFunctionEnd
-%vertex_main = OpFunction %void None %49
- %50 = OpLabel
- %51 = OpFunctionCall %VertexOutput %vertex_main_inner
- %52 = OpCompositeExtract %v4float %51 0
- OpStore %vertex_main_position_Output %52 None
- %53 = OpCompositeExtract %v2half %51 1
- OpStore %vertex_main_loc0_Output %53 None
+%vertex_main = OpFunction %void None %59
+ %60 = OpLabel
+ %61 = OpFunctionCall %VertexOutput %vertex_main_inner
+ %62 = OpCompositeExtract %v4float %61 0
+ OpStore %vertex_main_position_Output %62 None
+ %63 = OpCompositeExtract %v2half %61 1
+ OpStore %vertex_main_loc0_Output %63 None
OpStore %vertex_main___point_size_Output %float_1 None
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.glsl b/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.glsl
index 8e853d1..ab200d2 100644
--- a/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.glsl
@@ -13,7 +13,9 @@
vec2 arg_0 = vec2(2.0f);
vec2 arg_1 = vec2(4.0f);
vec2 arg_2 = vec2(3.0f);
- vec2 res = smoothstep(arg_0, arg_1, arg_2);
+ vec2 v_1 = arg_0;
+ vec2 v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), vec2(0.0f), vec2(1.0f));
+ vec2 res = (v_2 * (v_2 * (vec2(3.0f) - (vec2(2.0f) * v_2))));
return res;
}
void main() {
@@ -32,7 +34,9 @@
vec2 arg_0 = vec2(2.0f);
vec2 arg_1 = vec2(4.0f);
vec2 arg_2 = vec2(3.0f);
- vec2 res = smoothstep(arg_0, arg_1, arg_2);
+ vec2 v_1 = arg_0;
+ vec2 v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), vec2(0.0f), vec2(1.0f));
+ vec2 res = (v_2 * (v_2 * (vec2(3.0f) - (vec2(2.0f) * v_2))));
return res;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -55,7 +59,9 @@
vec2 arg_0 = vec2(2.0f);
vec2 arg_1 = vec2(4.0f);
vec2 arg_2 = vec2(3.0f);
- vec2 res = smoothstep(arg_0, arg_1, arg_2);
+ vec2 v = arg_0;
+ vec2 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), vec2(0.0f), vec2(1.0f));
+ vec2 res = (v_1 * (v_1 * (vec2(3.0f) - (vec2(2.0f) * v_1))));
return res;
}
VertexOutput vertex_main_inner() {
@@ -65,10 +71,10 @@
return tint_symbol;
}
void main() {
- VertexOutput v = vertex_main_inner();
- gl_Position = v.pos;
+ VertexOutput v_2 = vertex_main_inner();
+ gl_Position = v_2.pos;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
- vertex_main_loc0_Output = v.prevent_dce;
+ vertex_main_loc0_Output = v_2.prevent_dce;
gl_PointSize = 1.0f;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.ir.dxc.hlsl
index 078c2f5..a6209ab 100644
--- a/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.ir.dxc.hlsl
@@ -7,7 +7,9 @@
float2 arg_0 = (2.0f).xx;
float2 arg_1 = (4.0f).xx;
float2 arg_2 = (3.0f).xx;
- float2 res = smoothstep(arg_0, arg_1, arg_2);
+ float2 v = arg_0;
+ float2 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xx, (1.0f).xx);
+ float2 res = (v_1 * (v_1 * ((3.0f).xx - ((2.0f).xx * v_1))));
return res;
}
@@ -24,7 +26,9 @@
float2 arg_0 = (2.0f).xx;
float2 arg_1 = (4.0f).xx;
float2 arg_2 = (3.0f).xx;
- float2 res = smoothstep(arg_0, arg_1, arg_2);
+ float2 v = arg_0;
+ float2 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xx, (1.0f).xx);
+ float2 res = (v_1 * (v_1 * ((3.0f).xx - ((2.0f).xx * v_1))));
return res;
}
@@ -51,7 +55,9 @@
float2 arg_0 = (2.0f).xx;
float2 arg_1 = (4.0f).xx;
float2 arg_2 = (3.0f).xx;
- float2 res = smoothstep(arg_0, arg_1, arg_2);
+ float2 v = arg_0;
+ float2 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xx, (1.0f).xx);
+ float2 res = (v_1 * (v_1 * ((3.0f).xx - ((2.0f).xx * v_1))));
return res;
}
@@ -59,13 +65,13 @@
VertexOutput tint_symbol = (VertexOutput)0;
tint_symbol.pos = (0.0f).xxxx;
tint_symbol.prevent_dce = smoothstep_392c19();
- VertexOutput v = tint_symbol;
- return v;
+ VertexOutput v_2 = tint_symbol;
+ return v_2;
}
vertex_main_outputs vertex_main() {
- VertexOutput v_1 = vertex_main_inner();
- vertex_main_outputs v_2 = {v_1.prevent_dce, v_1.pos};
- return v_2;
+ VertexOutput v_3 = vertex_main_inner();
+ vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos};
+ return v_4;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.ir.fxc.hlsl
index 078c2f5..a6209ab 100644
--- a/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.ir.fxc.hlsl
@@ -7,7 +7,9 @@
float2 arg_0 = (2.0f).xx;
float2 arg_1 = (4.0f).xx;
float2 arg_2 = (3.0f).xx;
- float2 res = smoothstep(arg_0, arg_1, arg_2);
+ float2 v = arg_0;
+ float2 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xx, (1.0f).xx);
+ float2 res = (v_1 * (v_1 * ((3.0f).xx - ((2.0f).xx * v_1))));
return res;
}
@@ -24,7 +26,9 @@
float2 arg_0 = (2.0f).xx;
float2 arg_1 = (4.0f).xx;
float2 arg_2 = (3.0f).xx;
- float2 res = smoothstep(arg_0, arg_1, arg_2);
+ float2 v = arg_0;
+ float2 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xx, (1.0f).xx);
+ float2 res = (v_1 * (v_1 * ((3.0f).xx - ((2.0f).xx * v_1))));
return res;
}
@@ -51,7 +55,9 @@
float2 arg_0 = (2.0f).xx;
float2 arg_1 = (4.0f).xx;
float2 arg_2 = (3.0f).xx;
- float2 res = smoothstep(arg_0, arg_1, arg_2);
+ float2 v = arg_0;
+ float2 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xx, (1.0f).xx);
+ float2 res = (v_1 * (v_1 * ((3.0f).xx - ((2.0f).xx * v_1))));
return res;
}
@@ -59,13 +65,13 @@
VertexOutput tint_symbol = (VertexOutput)0;
tint_symbol.pos = (0.0f).xxxx;
tint_symbol.prevent_dce = smoothstep_392c19();
- VertexOutput v = tint_symbol;
- return v;
+ VertexOutput v_2 = tint_symbol;
+ return v_2;
}
vertex_main_outputs vertex_main() {
- VertexOutput v_1 = vertex_main_inner();
- vertex_main_outputs v_2 = {v_1.prevent_dce, v_1.pos};
- return v_2;
+ VertexOutput v_3 = vertex_main_inner();
+ vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos};
+ return v_4;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.ir.msl
index 5d244bb..f637e45 100644
--- a/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.ir.msl
@@ -12,7 +12,9 @@
float2 arg_0 = float2(2.0f);
float2 arg_1 = float2(4.0f);
float2 arg_2 = float2(3.0f);
- float2 res = smoothstep(arg_0, arg_1, arg_2);
+ float2 const v = arg_0;
+ float2 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), float2(0.0f), float2(1.0f));
+ float2 res = (v_1 * (v_1 * (float2(3.0f) - (float2(2.0f) * v_1))));
return res;
}
@@ -34,7 +36,9 @@
float2 arg_0 = float2(2.0f);
float2 arg_1 = float2(4.0f);
float2 arg_2 = float2(3.0f);
- float2 res = smoothstep(arg_0, arg_1, arg_2);
+ float2 const v = arg_0;
+ float2 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), float2(0.0f), float2(1.0f));
+ float2 res = (v_1 * (v_1 * (float2(3.0f) - (float2(2.0f) * v_1))));
return res;
}
@@ -62,7 +66,9 @@
float2 arg_0 = float2(2.0f);
float2 arg_1 = float2(4.0f);
float2 arg_2 = float2(3.0f);
- float2 res = smoothstep(arg_0, arg_1, arg_2);
+ float2 const v = arg_0;
+ float2 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), float2(0.0f), float2(1.0f));
+ float2 res = (v_1 * (v_1 * (float2(3.0f) - (float2(2.0f) * v_1))));
return res;
}
@@ -74,9 +80,9 @@
}
vertex vertex_main_outputs vertex_main() {
- VertexOutput const v = vertex_main_inner();
+ VertexOutput const v_2 = vertex_main_inner();
vertex_main_outputs tint_wrapper_result = {};
- tint_wrapper_result.VertexOutput_pos = v.pos;
- tint_wrapper_result.VertexOutput_prevent_dce = v.prevent_dce;
+ tint_wrapper_result.VertexOutput_pos = v_2.pos;
+ tint_wrapper_result.VertexOutput_prevent_dce = v_2.prevent_dce;
return tint_wrapper_result;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.spvasm b/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.spvasm
index 7b730ba..f6ffc87 100644
--- a/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/smoothstep/392c19.wgsl.expected.spvasm
@@ -4,10 +4,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 35
+; Bound: 45
; Schema: 0
OpCapability Shader
- %23 = OpExtInstImport "GLSL.std.450"
+ %26 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %fragment_main "fragment_main"
OpExecutionMode %fragment_main OriginUpperLeft
@@ -37,8 +37,11 @@
%14 = OpConstantComposite %v2float %float_4 %float_4
%float_3 = OpConstant %float 3
%17 = OpConstantComposite %v2float %float_3 %float_3
+ %27 = OpConstantNull %v2float
+ %float_1 = OpConstant %float 1
+ %28 = OpConstantComposite %v2float %float_1 %float_1
%void = OpTypeVoid
- %28 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_StorageBuffer_v2float = OpTypePointer StorageBuffer %v2float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -54,16 +57,23 @@
%19 = OpLoad %v2float %arg_0 None
%20 = OpLoad %v2float %arg_1 None
%21 = OpLoad %v2float %arg_2 None
- %22 = OpExtInst %v2float %23 SmoothStep %19 %20 %21
- OpStore %res %22
- %25 = OpLoad %v2float %res None
- OpReturnValue %25
+ %22 = OpFSub %v2float %21 %19
+ %23 = OpFSub %v2float %20 %19
+ %24 = OpFDiv %v2float %22 %23
+ %25 = OpExtInst %v2float %26 NClamp %24 %27 %28
+ %30 = OpFMul %v2float %11 %25
+ %31 = OpFSub %v2float %17 %30
+ %32 = OpFMul %v2float %25 %31
+ %33 = OpFMul %v2float %25 %32
+ OpStore %res %33
+ %35 = OpLoad %v2float %res None
+ OpReturnValue %35
OpFunctionEnd
-%fragment_main = OpFunction %void None %28
- %29 = OpLabel
- %30 = OpFunctionCall %v2float %smoothstep_392c19
- %31 = OpAccessChain %_ptr_StorageBuffer_v2float %1 %uint_0
- OpStore %31 %30 None
+%fragment_main = OpFunction %void None %38
+ %39 = OpLabel
+ %40 = OpFunctionCall %v2float %smoothstep_392c19
+ %41 = OpAccessChain %_ptr_StorageBuffer_v2float %1 %uint_0
+ OpStore %41 %40 None
OpReturn
OpFunctionEnd
;
@@ -72,10 +82,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 35
+; Bound: 45
; Schema: 0
OpCapability Shader
- %23 = OpExtInstImport "GLSL.std.450"
+ %26 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %compute_main LocalSize 1 1 1
@@ -105,8 +115,11 @@
%14 = OpConstantComposite %v2float %float_4 %float_4
%float_3 = OpConstant %float 3
%17 = OpConstantComposite %v2float %float_3 %float_3
+ %27 = OpConstantNull %v2float
+ %float_1 = OpConstant %float 1
+ %28 = OpConstantComposite %v2float %float_1 %float_1
%void = OpTypeVoid
- %28 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_StorageBuffer_v2float = OpTypePointer StorageBuffer %v2float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -122,16 +135,23 @@
%19 = OpLoad %v2float %arg_0 None
%20 = OpLoad %v2float %arg_1 None
%21 = OpLoad %v2float %arg_2 None
- %22 = OpExtInst %v2float %23 SmoothStep %19 %20 %21
- OpStore %res %22
- %25 = OpLoad %v2float %res None
- OpReturnValue %25
+ %22 = OpFSub %v2float %21 %19
+ %23 = OpFSub %v2float %20 %19
+ %24 = OpFDiv %v2float %22 %23
+ %25 = OpExtInst %v2float %26 NClamp %24 %27 %28
+ %30 = OpFMul %v2float %11 %25
+ %31 = OpFSub %v2float %17 %30
+ %32 = OpFMul %v2float %25 %31
+ %33 = OpFMul %v2float %25 %32
+ OpStore %res %33
+ %35 = OpLoad %v2float %res None
+ OpReturnValue %35
OpFunctionEnd
-%compute_main = OpFunction %void None %28
- %29 = OpLabel
- %30 = OpFunctionCall %v2float %smoothstep_392c19
- %31 = OpAccessChain %_ptr_StorageBuffer_v2float %1 %uint_0
- OpStore %31 %30 None
+%compute_main = OpFunction %void None %38
+ %39 = OpLabel
+ %40 = OpFunctionCall %v2float %smoothstep_392c19
+ %41 = OpAccessChain %_ptr_StorageBuffer_v2float %1 %uint_0
+ OpStore %41 %40 None
OpReturn
OpFunctionEnd
;
@@ -140,10 +160,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 54
+; Bound: 63
; Schema: 0
OpCapability Shader
- %27 = OpExtInstImport "GLSL.std.450"
+ %30 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %vertex_main_position_Output %vertex_main_loc0_Output %vertex_main___point_size_Output
OpName %vertex_main_position_Output "vertex_main_position_Output"
@@ -183,18 +203,20 @@
%18 = OpConstantComposite %v2float %float_4 %float_4
%float_3 = OpConstant %float 3
%21 = OpConstantComposite %v2float %float_3 %float_3
+ %31 = OpConstantNull %v2float
+ %float_1 = OpConstant %float 1
+ %32 = OpConstantComposite %v2float %float_1 %float_1
%VertexOutput = OpTypeStruct %v4float %v2float
- %32 = OpTypeFunction %VertexOutput
+ %42 = OpTypeFunction %VertexOutput
%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
- %36 = OpConstantNull %VertexOutput
+ %46 = OpConstantNull %VertexOutput
%_ptr_Function_v4float = OpTypePointer Function %v4float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
- %41 = OpConstantNull %v4float
+ %51 = OpConstantNull %v4float
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
- %48 = OpTypeFunction %void
- %float_1 = OpConstant %float 1
+ %58 = OpTypeFunction %void
%smoothstep_392c19 = OpFunction %v2float None %11
%12 = OpLabel
%arg_0 = OpVariable %_ptr_Function_v2float Function
@@ -207,29 +229,36 @@
%23 = OpLoad %v2float %arg_0 None
%24 = OpLoad %v2float %arg_1 None
%25 = OpLoad %v2float %arg_2 None
- %26 = OpExtInst %v2float %27 SmoothStep %23 %24 %25
- OpStore %res %26
- %29 = OpLoad %v2float %res None
- OpReturnValue %29
+ %26 = OpFSub %v2float %25 %23
+ %27 = OpFSub %v2float %24 %23
+ %28 = OpFDiv %v2float %26 %27
+ %29 = OpExtInst %v2float %30 NClamp %28 %31 %32
+ %34 = OpFMul %v2float %15 %29
+ %35 = OpFSub %v2float %21 %34
+ %36 = OpFMul %v2float %29 %35
+ %37 = OpFMul %v2float %29 %36
+ OpStore %res %37
+ %39 = OpLoad %v2float %res None
+ OpReturnValue %39
OpFunctionEnd
-%vertex_main_inner = OpFunction %VertexOutput None %32
- %33 = OpLabel
- %out = OpVariable %_ptr_Function_VertexOutput Function %36
- %37 = OpAccessChain %_ptr_Function_v4float %out %uint_0
- OpStore %37 %41 None
- %42 = OpAccessChain %_ptr_Function_v2float %out %uint_1
- %44 = OpFunctionCall %v2float %smoothstep_392c19
- OpStore %42 %44 None
- %45 = OpLoad %VertexOutput %out None
- OpReturnValue %45
+%vertex_main_inner = OpFunction %VertexOutput None %42
+ %43 = OpLabel
+ %out = OpVariable %_ptr_Function_VertexOutput Function %46
+ %47 = OpAccessChain %_ptr_Function_v4float %out %uint_0
+ OpStore %47 %51 None
+ %52 = OpAccessChain %_ptr_Function_v2float %out %uint_1
+ %54 = OpFunctionCall %v2float %smoothstep_392c19
+ OpStore %52 %54 None
+ %55 = OpLoad %VertexOutput %out None
+ OpReturnValue %55
OpFunctionEnd
-%vertex_main = OpFunction %void None %48
- %49 = OpLabel
- %50 = OpFunctionCall %VertexOutput %vertex_main_inner
- %51 = OpCompositeExtract %v4float %50 0
- OpStore %vertex_main_position_Output %51 None
- %52 = OpCompositeExtract %v2float %50 1
- OpStore %vertex_main_loc0_Output %52 None
+%vertex_main = OpFunction %void None %58
+ %59 = OpLabel
+ %60 = OpFunctionCall %VertexOutput %vertex_main_inner
+ %61 = OpCompositeExtract %v4float %60 0
+ OpStore %vertex_main_position_Output %61 None
+ %62 = OpCompositeExtract %v2float %60 1
+ OpStore %vertex_main_loc0_Output %62 None
OpStore %vertex_main___point_size_Output %float_1 None
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.glsl b/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.glsl
index 239c4b1..98437fc 100644
--- a/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.glsl
@@ -13,7 +13,9 @@
vec4 arg_0 = vec4(2.0f);
vec4 arg_1 = vec4(4.0f);
vec4 arg_2 = vec4(3.0f);
- vec4 res = smoothstep(arg_0, arg_1, arg_2);
+ vec4 v_1 = arg_0;
+ vec4 v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), vec4(0.0f), vec4(1.0f));
+ vec4 res = (v_2 * (v_2 * (vec4(3.0f) - (vec4(2.0f) * v_2))));
return res;
}
void main() {
@@ -32,7 +34,9 @@
vec4 arg_0 = vec4(2.0f);
vec4 arg_1 = vec4(4.0f);
vec4 arg_2 = vec4(3.0f);
- vec4 res = smoothstep(arg_0, arg_1, arg_2);
+ vec4 v_1 = arg_0;
+ vec4 v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), vec4(0.0f), vec4(1.0f));
+ vec4 res = (v_2 * (v_2 * (vec4(3.0f) - (vec4(2.0f) * v_2))));
return res;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -55,7 +59,9 @@
vec4 arg_0 = vec4(2.0f);
vec4 arg_1 = vec4(4.0f);
vec4 arg_2 = vec4(3.0f);
- vec4 res = smoothstep(arg_0, arg_1, arg_2);
+ vec4 v = arg_0;
+ vec4 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), vec4(0.0f), vec4(1.0f));
+ vec4 res = (v_1 * (v_1 * (vec4(3.0f) - (vec4(2.0f) * v_1))));
return res;
}
VertexOutput vertex_main_inner() {
@@ -65,10 +71,10 @@
return tint_symbol;
}
void main() {
- VertexOutput v = vertex_main_inner();
- gl_Position = v.pos;
+ VertexOutput v_2 = vertex_main_inner();
+ gl_Position = v_2.pos;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
- vertex_main_loc0_Output = v.prevent_dce;
+ vertex_main_loc0_Output = v_2.prevent_dce;
gl_PointSize = 1.0f;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.ir.dxc.hlsl
index 44223b2..1aa1382 100644
--- a/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.ir.dxc.hlsl
@@ -7,7 +7,9 @@
float4 arg_0 = (2.0f).xxxx;
float4 arg_1 = (4.0f).xxxx;
float4 arg_2 = (3.0f).xxxx;
- float4 res = smoothstep(arg_0, arg_1, arg_2);
+ float4 v = arg_0;
+ float4 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xxxx, (1.0f).xxxx);
+ float4 res = (v_1 * (v_1 * ((3.0f).xxxx - ((2.0f).xxxx * v_1))));
return res;
}
@@ -24,7 +26,9 @@
float4 arg_0 = (2.0f).xxxx;
float4 arg_1 = (4.0f).xxxx;
float4 arg_2 = (3.0f).xxxx;
- float4 res = smoothstep(arg_0, arg_1, arg_2);
+ float4 v = arg_0;
+ float4 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xxxx, (1.0f).xxxx);
+ float4 res = (v_1 * (v_1 * ((3.0f).xxxx - ((2.0f).xxxx * v_1))));
return res;
}
@@ -51,7 +55,9 @@
float4 arg_0 = (2.0f).xxxx;
float4 arg_1 = (4.0f).xxxx;
float4 arg_2 = (3.0f).xxxx;
- float4 res = smoothstep(arg_0, arg_1, arg_2);
+ float4 v = arg_0;
+ float4 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xxxx, (1.0f).xxxx);
+ float4 res = (v_1 * (v_1 * ((3.0f).xxxx - ((2.0f).xxxx * v_1))));
return res;
}
@@ -59,13 +65,13 @@
VertexOutput tint_symbol = (VertexOutput)0;
tint_symbol.pos = (0.0f).xxxx;
tint_symbol.prevent_dce = smoothstep_40864c();
- VertexOutput v = tint_symbol;
- return v;
+ VertexOutput v_2 = tint_symbol;
+ return v_2;
}
vertex_main_outputs vertex_main() {
- VertexOutput v_1 = vertex_main_inner();
- vertex_main_outputs v_2 = {v_1.prevent_dce, v_1.pos};
- return v_2;
+ VertexOutput v_3 = vertex_main_inner();
+ vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos};
+ return v_4;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.ir.fxc.hlsl
index 44223b2..1aa1382 100644
--- a/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.ir.fxc.hlsl
@@ -7,7 +7,9 @@
float4 arg_0 = (2.0f).xxxx;
float4 arg_1 = (4.0f).xxxx;
float4 arg_2 = (3.0f).xxxx;
- float4 res = smoothstep(arg_0, arg_1, arg_2);
+ float4 v = arg_0;
+ float4 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xxxx, (1.0f).xxxx);
+ float4 res = (v_1 * (v_1 * ((3.0f).xxxx - ((2.0f).xxxx * v_1))));
return res;
}
@@ -24,7 +26,9 @@
float4 arg_0 = (2.0f).xxxx;
float4 arg_1 = (4.0f).xxxx;
float4 arg_2 = (3.0f).xxxx;
- float4 res = smoothstep(arg_0, arg_1, arg_2);
+ float4 v = arg_0;
+ float4 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xxxx, (1.0f).xxxx);
+ float4 res = (v_1 * (v_1 * ((3.0f).xxxx - ((2.0f).xxxx * v_1))));
return res;
}
@@ -51,7 +55,9 @@
float4 arg_0 = (2.0f).xxxx;
float4 arg_1 = (4.0f).xxxx;
float4 arg_2 = (3.0f).xxxx;
- float4 res = smoothstep(arg_0, arg_1, arg_2);
+ float4 v = arg_0;
+ float4 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xxxx, (1.0f).xxxx);
+ float4 res = (v_1 * (v_1 * ((3.0f).xxxx - ((2.0f).xxxx * v_1))));
return res;
}
@@ -59,13 +65,13 @@
VertexOutput tint_symbol = (VertexOutput)0;
tint_symbol.pos = (0.0f).xxxx;
tint_symbol.prevent_dce = smoothstep_40864c();
- VertexOutput v = tint_symbol;
- return v;
+ VertexOutput v_2 = tint_symbol;
+ return v_2;
}
vertex_main_outputs vertex_main() {
- VertexOutput v_1 = vertex_main_inner();
- vertex_main_outputs v_2 = {v_1.prevent_dce, v_1.pos};
- return v_2;
+ VertexOutput v_3 = vertex_main_inner();
+ vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos};
+ return v_4;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.ir.msl
index 5f844a3..38cded6 100644
--- a/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.ir.msl
@@ -12,7 +12,9 @@
float4 arg_0 = float4(2.0f);
float4 arg_1 = float4(4.0f);
float4 arg_2 = float4(3.0f);
- float4 res = smoothstep(arg_0, arg_1, arg_2);
+ float4 const v = arg_0;
+ float4 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), float4(0.0f), float4(1.0f));
+ float4 res = (v_1 * (v_1 * (float4(3.0f) - (float4(2.0f) * v_1))));
return res;
}
@@ -34,7 +36,9 @@
float4 arg_0 = float4(2.0f);
float4 arg_1 = float4(4.0f);
float4 arg_2 = float4(3.0f);
- float4 res = smoothstep(arg_0, arg_1, arg_2);
+ float4 const v = arg_0;
+ float4 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), float4(0.0f), float4(1.0f));
+ float4 res = (v_1 * (v_1 * (float4(3.0f) - (float4(2.0f) * v_1))));
return res;
}
@@ -62,7 +66,9 @@
float4 arg_0 = float4(2.0f);
float4 arg_1 = float4(4.0f);
float4 arg_2 = float4(3.0f);
- float4 res = smoothstep(arg_0, arg_1, arg_2);
+ float4 const v = arg_0;
+ float4 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), float4(0.0f), float4(1.0f));
+ float4 res = (v_1 * (v_1 * (float4(3.0f) - (float4(2.0f) * v_1))));
return res;
}
@@ -74,9 +80,9 @@
}
vertex vertex_main_outputs vertex_main() {
- VertexOutput const v = vertex_main_inner();
+ VertexOutput const v_2 = vertex_main_inner();
vertex_main_outputs tint_wrapper_result = {};
- tint_wrapper_result.VertexOutput_pos = v.pos;
- tint_wrapper_result.VertexOutput_prevent_dce = v.prevent_dce;
+ tint_wrapper_result.VertexOutput_pos = v_2.pos;
+ tint_wrapper_result.VertexOutput_prevent_dce = v_2.prevent_dce;
return tint_wrapper_result;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.spvasm b/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.spvasm
index 9bc965f..cc92228 100644
--- a/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/smoothstep/40864c.wgsl.expected.spvasm
@@ -4,10 +4,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 35
+; Bound: 45
; Schema: 0
OpCapability Shader
- %23 = OpExtInstImport "GLSL.std.450"
+ %26 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %fragment_main "fragment_main"
OpExecutionMode %fragment_main OriginUpperLeft
@@ -37,8 +37,11 @@
%14 = OpConstantComposite %v4float %float_4 %float_4 %float_4 %float_4
%float_3 = OpConstant %float 3
%17 = OpConstantComposite %v4float %float_3 %float_3 %float_3 %float_3
+ %27 = OpConstantNull %v4float
+ %float_1 = OpConstant %float 1
+ %28 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
%void = OpTypeVoid
- %28 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -54,16 +57,23 @@
%19 = OpLoad %v4float %arg_0 None
%20 = OpLoad %v4float %arg_1 None
%21 = OpLoad %v4float %arg_2 None
- %22 = OpExtInst %v4float %23 SmoothStep %19 %20 %21
- OpStore %res %22
- %25 = OpLoad %v4float %res None
- OpReturnValue %25
+ %22 = OpFSub %v4float %21 %19
+ %23 = OpFSub %v4float %20 %19
+ %24 = OpFDiv %v4float %22 %23
+ %25 = OpExtInst %v4float %26 NClamp %24 %27 %28
+ %30 = OpFMul %v4float %11 %25
+ %31 = OpFSub %v4float %17 %30
+ %32 = OpFMul %v4float %25 %31
+ %33 = OpFMul %v4float %25 %32
+ OpStore %res %33
+ %35 = OpLoad %v4float %res None
+ OpReturnValue %35
OpFunctionEnd
-%fragment_main = OpFunction %void None %28
- %29 = OpLabel
- %30 = OpFunctionCall %v4float %smoothstep_40864c
- %31 = OpAccessChain %_ptr_StorageBuffer_v4float %1 %uint_0
- OpStore %31 %30 None
+%fragment_main = OpFunction %void None %38
+ %39 = OpLabel
+ %40 = OpFunctionCall %v4float %smoothstep_40864c
+ %41 = OpAccessChain %_ptr_StorageBuffer_v4float %1 %uint_0
+ OpStore %41 %40 None
OpReturn
OpFunctionEnd
;
@@ -72,10 +82,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 35
+; Bound: 45
; Schema: 0
OpCapability Shader
- %23 = OpExtInstImport "GLSL.std.450"
+ %26 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %compute_main LocalSize 1 1 1
@@ -105,8 +115,11 @@
%14 = OpConstantComposite %v4float %float_4 %float_4 %float_4 %float_4
%float_3 = OpConstant %float 3
%17 = OpConstantComposite %v4float %float_3 %float_3 %float_3 %float_3
+ %27 = OpConstantNull %v4float
+ %float_1 = OpConstant %float 1
+ %28 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
%void = OpTypeVoid
- %28 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -122,16 +135,23 @@
%19 = OpLoad %v4float %arg_0 None
%20 = OpLoad %v4float %arg_1 None
%21 = OpLoad %v4float %arg_2 None
- %22 = OpExtInst %v4float %23 SmoothStep %19 %20 %21
- OpStore %res %22
- %25 = OpLoad %v4float %res None
- OpReturnValue %25
+ %22 = OpFSub %v4float %21 %19
+ %23 = OpFSub %v4float %20 %19
+ %24 = OpFDiv %v4float %22 %23
+ %25 = OpExtInst %v4float %26 NClamp %24 %27 %28
+ %30 = OpFMul %v4float %11 %25
+ %31 = OpFSub %v4float %17 %30
+ %32 = OpFMul %v4float %25 %31
+ %33 = OpFMul %v4float %25 %32
+ OpStore %res %33
+ %35 = OpLoad %v4float %res None
+ OpReturnValue %35
OpFunctionEnd
-%compute_main = OpFunction %void None %28
- %29 = OpLabel
- %30 = OpFunctionCall %v4float %smoothstep_40864c
- %31 = OpAccessChain %_ptr_StorageBuffer_v4float %1 %uint_0
- OpStore %31 %30 None
+%compute_main = OpFunction %void None %38
+ %39 = OpLabel
+ %40 = OpFunctionCall %v4float %smoothstep_40864c
+ %41 = OpAccessChain %_ptr_StorageBuffer_v4float %1 %uint_0
+ OpStore %41 %40 None
OpReturn
OpFunctionEnd
;
@@ -140,10 +160,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 51
+; Bound: 59
; Schema: 0
OpCapability Shader
- %25 = OpExtInstImport "GLSL.std.450"
+ %28 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %vertex_main_position_Output %vertex_main_loc0_Output %vertex_main___point_size_Output
OpName %vertex_main_position_Output "vertex_main_position_Output"
@@ -181,17 +201,18 @@
%16 = OpConstantComposite %v4float %float_4 %float_4 %float_4 %float_4
%float_3 = OpConstant %float 3
%19 = OpConstantComposite %v4float %float_3 %float_3 %float_3 %float_3
+ %29 = OpConstantNull %v4float
+ %float_1 = OpConstant %float 1
+ %30 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
%VertexOutput = OpTypeStruct %v4float %v4float
- %30 = OpTypeFunction %VertexOutput
+ %40 = OpTypeFunction %VertexOutput
%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
- %34 = OpConstantNull %VertexOutput
+ %44 = OpConstantNull %VertexOutput
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
- %38 = OpConstantNull %v4float
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
- %45 = OpTypeFunction %void
- %float_1 = OpConstant %float 1
+ %54 = OpTypeFunction %void
%smoothstep_40864c = OpFunction %v4float None %9
%10 = OpLabel
%arg_0 = OpVariable %_ptr_Function_v4float Function
@@ -204,29 +225,36 @@
%21 = OpLoad %v4float %arg_0 None
%22 = OpLoad %v4float %arg_1 None
%23 = OpLoad %v4float %arg_2 None
- %24 = OpExtInst %v4float %25 SmoothStep %21 %22 %23
- OpStore %res %24
- %27 = OpLoad %v4float %res None
- OpReturnValue %27
+ %24 = OpFSub %v4float %23 %21
+ %25 = OpFSub %v4float %22 %21
+ %26 = OpFDiv %v4float %24 %25
+ %27 = OpExtInst %v4float %28 NClamp %26 %29 %30
+ %32 = OpFMul %v4float %13 %27
+ %33 = OpFSub %v4float %19 %32
+ %34 = OpFMul %v4float %27 %33
+ %35 = OpFMul %v4float %27 %34
+ OpStore %res %35
+ %37 = OpLoad %v4float %res None
+ OpReturnValue %37
OpFunctionEnd
-%vertex_main_inner = OpFunction %VertexOutput None %30
- %31 = OpLabel
- %out = OpVariable %_ptr_Function_VertexOutput Function %34
- %35 = OpAccessChain %_ptr_Function_v4float %out %uint_0
- OpStore %35 %38 None
- %39 = OpAccessChain %_ptr_Function_v4float %out %uint_1
- %41 = OpFunctionCall %v4float %smoothstep_40864c
- OpStore %39 %41 None
- %42 = OpLoad %VertexOutput %out None
- OpReturnValue %42
+%vertex_main_inner = OpFunction %VertexOutput None %40
+ %41 = OpLabel
+ %out = OpVariable %_ptr_Function_VertexOutput Function %44
+ %45 = OpAccessChain %_ptr_Function_v4float %out %uint_0
+ OpStore %45 %29 None
+ %48 = OpAccessChain %_ptr_Function_v4float %out %uint_1
+ %50 = OpFunctionCall %v4float %smoothstep_40864c
+ OpStore %48 %50 None
+ %51 = OpLoad %VertexOutput %out None
+ OpReturnValue %51
OpFunctionEnd
-%vertex_main = OpFunction %void None %45
- %46 = OpLabel
- %47 = OpFunctionCall %VertexOutput %vertex_main_inner
- %48 = OpCompositeExtract %v4float %47 0
- OpStore %vertex_main_position_Output %48 None
- %49 = OpCompositeExtract %v4float %47 1
- OpStore %vertex_main_loc0_Output %49 None
+%vertex_main = OpFunction %void None %54
+ %55 = OpLabel
+ %56 = OpFunctionCall %VertexOutput %vertex_main_inner
+ %57 = OpCompositeExtract %v4float %56 0
+ OpStore %vertex_main_position_Output %57 None
+ %58 = OpCompositeExtract %v4float %56 1
+ OpStore %vertex_main_loc0_Output %58 None
OpStore %vertex_main___point_size_Output %float_1 None
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.glsl b/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.glsl
index dfa3b53..7dff434 100644
--- a/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.glsl
@@ -14,7 +14,9 @@
float16_t arg_0 = 2.0hf;
float16_t arg_1 = 4.0hf;
float16_t arg_2 = 3.0hf;
- float16_t res = smoothstep(arg_0, arg_1, arg_2);
+ float16_t v_1 = arg_0;
+ float16_t v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), 0.0hf, 1.0hf);
+ float16_t res = (v_2 * (v_2 * (3.0hf - (2.0hf * v_2))));
return res;
}
void main() {
@@ -34,7 +36,9 @@
float16_t arg_0 = 2.0hf;
float16_t arg_1 = 4.0hf;
float16_t arg_2 = 3.0hf;
- float16_t res = smoothstep(arg_0, arg_1, arg_2);
+ float16_t v_1 = arg_0;
+ float16_t v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), 0.0hf, 1.0hf);
+ float16_t res = (v_2 * (v_2 * (3.0hf - (2.0hf * v_2))));
return res;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -58,7 +62,9 @@
float16_t arg_0 = 2.0hf;
float16_t arg_1 = 4.0hf;
float16_t arg_2 = 3.0hf;
- float16_t res = smoothstep(arg_0, arg_1, arg_2);
+ float16_t v = arg_0;
+ float16_t v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0hf, 1.0hf);
+ float16_t res = (v_1 * (v_1 * (3.0hf - (2.0hf * v_1))));
return res;
}
VertexOutput vertex_main_inner() {
@@ -68,10 +74,10 @@
return tint_symbol;
}
void main() {
- VertexOutput v = vertex_main_inner();
- gl_Position = v.pos;
+ VertexOutput v_2 = vertex_main_inner();
+ gl_Position = v_2.pos;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
- vertex_main_loc0_Output = v.prevent_dce;
+ vertex_main_loc0_Output = v_2.prevent_dce;
gl_PointSize = 1.0f;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.ir.dxc.hlsl
index b99d216..e754813 100644
--- a/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.ir.dxc.hlsl
@@ -7,7 +7,9 @@
float16_t arg_0 = float16_t(2.0h);
float16_t arg_1 = float16_t(4.0h);
float16_t arg_2 = float16_t(3.0h);
- float16_t res = smoothstep(arg_0, arg_1, arg_2);
+ float16_t v = arg_0;
+ float16_t v_1 = clamp(((arg_2 - v) / (arg_1 - v)), float16_t(0.0h), float16_t(1.0h));
+ float16_t res = (v_1 * (v_1 * (float16_t(3.0h) - (float16_t(2.0h) * v_1))));
return res;
}
@@ -24,7 +26,9 @@
float16_t arg_0 = float16_t(2.0h);
float16_t arg_1 = float16_t(4.0h);
float16_t arg_2 = float16_t(3.0h);
- float16_t res = smoothstep(arg_0, arg_1, arg_2);
+ float16_t v = arg_0;
+ float16_t v_1 = clamp(((arg_2 - v) / (arg_1 - v)), float16_t(0.0h), float16_t(1.0h));
+ float16_t res = (v_1 * (v_1 * (float16_t(3.0h) - (float16_t(2.0h) * v_1))));
return res;
}
@@ -51,7 +55,9 @@
float16_t arg_0 = float16_t(2.0h);
float16_t arg_1 = float16_t(4.0h);
float16_t arg_2 = float16_t(3.0h);
- float16_t res = smoothstep(arg_0, arg_1, arg_2);
+ float16_t v = arg_0;
+ float16_t v_1 = clamp(((arg_2 - v) / (arg_1 - v)), float16_t(0.0h), float16_t(1.0h));
+ float16_t res = (v_1 * (v_1 * (float16_t(3.0h) - (float16_t(2.0h) * v_1))));
return res;
}
@@ -59,13 +65,13 @@
VertexOutput tint_symbol = (VertexOutput)0;
tint_symbol.pos = (0.0f).xxxx;
tint_symbol.prevent_dce = smoothstep_586e12();
- VertexOutput v = tint_symbol;
- return v;
+ VertexOutput v_2 = tint_symbol;
+ return v_2;
}
vertex_main_outputs vertex_main() {
- VertexOutput v_1 = vertex_main_inner();
- vertex_main_outputs v_2 = {v_1.prevent_dce, v_1.pos};
- return v_2;
+ VertexOutput v_3 = vertex_main_inner();
+ vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos};
+ return v_4;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.ir.msl
index e713a21..f08231e 100644
--- a/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.ir.msl
@@ -12,7 +12,9 @@
half arg_0 = 2.0h;
half arg_1 = 4.0h;
half arg_2 = 3.0h;
- half res = smoothstep(arg_0, arg_1, arg_2);
+ half const v = arg_0;
+ half const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0h, 1.0h);
+ half res = (v_1 * (v_1 * (3.0h - (2.0h * v_1))));
return res;
}
@@ -34,7 +36,9 @@
half arg_0 = 2.0h;
half arg_1 = 4.0h;
half arg_2 = 3.0h;
- half res = smoothstep(arg_0, arg_1, arg_2);
+ half const v = arg_0;
+ half const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0h, 1.0h);
+ half res = (v_1 * (v_1 * (3.0h - (2.0h * v_1))));
return res;
}
@@ -62,7 +66,9 @@
half arg_0 = 2.0h;
half arg_1 = 4.0h;
half arg_2 = 3.0h;
- half res = smoothstep(arg_0, arg_1, arg_2);
+ half const v = arg_0;
+ half const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0h, 1.0h);
+ half res = (v_1 * (v_1 * (3.0h - (2.0h * v_1))));
return res;
}
@@ -74,9 +80,9 @@
}
vertex vertex_main_outputs vertex_main() {
- VertexOutput const v = vertex_main_inner();
+ VertexOutput const v_2 = vertex_main_inner();
vertex_main_outputs tint_wrapper_result = {};
- tint_wrapper_result.VertexOutput_pos = v.pos;
- tint_wrapper_result.VertexOutput_prevent_dce = v.prevent_dce;
+ tint_wrapper_result.VertexOutput_pos = v_2.pos;
+ tint_wrapper_result.VertexOutput_prevent_dce = v_2.prevent_dce;
return tint_wrapper_result;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.spvasm b/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.spvasm
index d4e083e..ee7237c 100644
--- a/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/smoothstep/586e12.wgsl.expected.spvasm
@@ -4,13 +4,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 31
+; Bound: 40
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
- %19 = OpExtInstImport "GLSL.std.450"
+ %22 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %fragment_main "fragment_main"
OpExecutionMode %fragment_main OriginUpperLeft
@@ -36,8 +36,10 @@
%half_0x1p_1 = OpConstant %half 0x1p+1
%half_0x1p_2 = OpConstant %half 0x1p+2
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
+%half_0x0p_0 = OpConstant %half 0x0p+0
+%half_0x1p_0 = OpConstant %half 0x1p+0
%void = OpTypeVoid
- %24 = OpTypeFunction %void
+ %33 = OpTypeFunction %void
%_ptr_StorageBuffer_half = OpTypePointer StorageBuffer %half
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -53,16 +55,23 @@
%15 = OpLoad %half %arg_0 None
%16 = OpLoad %half %arg_1 None
%17 = OpLoad %half %arg_2 None
- %18 = OpExtInst %half %19 SmoothStep %15 %16 %17
- OpStore %res %18
- %21 = OpLoad %half %res None
- OpReturnValue %21
+ %18 = OpFSub %half %17 %15
+ %19 = OpFSub %half %16 %15
+ %20 = OpFDiv %half %18 %19
+ %21 = OpExtInst %half %22 NClamp %20 %half_0x0p_0 %half_0x1p_0
+ %25 = OpFMul %half %half_0x1p_1 %21
+ %26 = OpFSub %half %half_0x1_8p_1 %25
+ %27 = OpFMul %half %21 %26
+ %28 = OpFMul %half %21 %27
+ OpStore %res %28
+ %30 = OpLoad %half %res None
+ OpReturnValue %30
OpFunctionEnd
-%fragment_main = OpFunction %void None %24
- %25 = OpLabel
- %26 = OpFunctionCall %half %smoothstep_586e12
- %27 = OpAccessChain %_ptr_StorageBuffer_half %1 %uint_0
- OpStore %27 %26 None
+%fragment_main = OpFunction %void None %33
+ %34 = OpLabel
+ %35 = OpFunctionCall %half %smoothstep_586e12
+ %36 = OpAccessChain %_ptr_StorageBuffer_half %1 %uint_0
+ OpStore %36 %35 None
OpReturn
OpFunctionEnd
;
@@ -71,13 +80,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 31
+; Bound: 40
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
- %19 = OpExtInstImport "GLSL.std.450"
+ %22 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %compute_main LocalSize 1 1 1
@@ -103,8 +112,10 @@
%half_0x1p_1 = OpConstant %half 0x1p+1
%half_0x1p_2 = OpConstant %half 0x1p+2
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
+%half_0x0p_0 = OpConstant %half 0x0p+0
+%half_0x1p_0 = OpConstant %half 0x1p+0
%void = OpTypeVoid
- %24 = OpTypeFunction %void
+ %33 = OpTypeFunction %void
%_ptr_StorageBuffer_half = OpTypePointer StorageBuffer %half
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -120,16 +131,23 @@
%15 = OpLoad %half %arg_0 None
%16 = OpLoad %half %arg_1 None
%17 = OpLoad %half %arg_2 None
- %18 = OpExtInst %half %19 SmoothStep %15 %16 %17
- OpStore %res %18
- %21 = OpLoad %half %res None
- OpReturnValue %21
+ %18 = OpFSub %half %17 %15
+ %19 = OpFSub %half %16 %15
+ %20 = OpFDiv %half %18 %19
+ %21 = OpExtInst %half %22 NClamp %20 %half_0x0p_0 %half_0x1p_0
+ %25 = OpFMul %half %half_0x1p_1 %21
+ %26 = OpFSub %half %half_0x1_8p_1 %25
+ %27 = OpFMul %half %21 %26
+ %28 = OpFMul %half %21 %27
+ OpStore %res %28
+ %30 = OpLoad %half %res None
+ OpReturnValue %30
OpFunctionEnd
-%compute_main = OpFunction %void None %24
- %25 = OpLabel
- %26 = OpFunctionCall %half %smoothstep_586e12
- %27 = OpAccessChain %_ptr_StorageBuffer_half %1 %uint_0
- OpStore %27 %26 None
+%compute_main = OpFunction %void None %33
+ %34 = OpLabel
+ %35 = OpFunctionCall %half %smoothstep_586e12
+ %36 = OpAccessChain %_ptr_StorageBuffer_half %1 %uint_0
+ OpStore %36 %35 None
OpReturn
OpFunctionEnd
;
@@ -138,14 +156,14 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 51
+; Bound: 60
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
- %24 = OpExtInstImport "GLSL.std.450"
+ %27 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %vertex_main_position_Output %vertex_main_loc0_Output %vertex_main___point_size_Output
OpName %vertex_main_position_Output "vertex_main_position_Output"
@@ -182,17 +200,19 @@
%half_0x1p_1 = OpConstant %half 0x1p+1
%half_0x1p_2 = OpConstant %half 0x1p+2
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
+%half_0x0p_0 = OpConstant %half 0x0p+0
+%half_0x1p_0 = OpConstant %half 0x1p+0
%VertexOutput = OpTypeStruct %v4float %half
- %29 = OpTypeFunction %VertexOutput
+ %38 = OpTypeFunction %VertexOutput
%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
- %33 = OpConstantNull %VertexOutput
+ %42 = OpConstantNull %VertexOutput
%_ptr_Function_v4float = OpTypePointer Function %v4float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
- %38 = OpConstantNull %v4float
+ %47 = OpConstantNull %v4float
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
- %45 = OpTypeFunction %void
+ %54 = OpTypeFunction %void
%float_1 = OpConstant %float 1
%smoothstep_586e12 = OpFunction %half None %11
%12 = OpLabel
@@ -206,29 +226,36 @@
%20 = OpLoad %half %arg_0 None
%21 = OpLoad %half %arg_1 None
%22 = OpLoad %half %arg_2 None
- %23 = OpExtInst %half %24 SmoothStep %20 %21 %22
- OpStore %res %23
- %26 = OpLoad %half %res None
- OpReturnValue %26
+ %23 = OpFSub %half %22 %20
+ %24 = OpFSub %half %21 %20
+ %25 = OpFDiv %half %23 %24
+ %26 = OpExtInst %half %27 NClamp %25 %half_0x0p_0 %half_0x1p_0
+ %30 = OpFMul %half %half_0x1p_1 %26
+ %31 = OpFSub %half %half_0x1_8p_1 %30
+ %32 = OpFMul %half %26 %31
+ %33 = OpFMul %half %26 %32
+ OpStore %res %33
+ %35 = OpLoad %half %res None
+ OpReturnValue %35
OpFunctionEnd
-%vertex_main_inner = OpFunction %VertexOutput None %29
- %30 = OpLabel
- %out = OpVariable %_ptr_Function_VertexOutput Function %33
- %34 = OpAccessChain %_ptr_Function_v4float %out %uint_0
- OpStore %34 %38 None
- %39 = OpAccessChain %_ptr_Function_half %out %uint_1
- %41 = OpFunctionCall %half %smoothstep_586e12
- OpStore %39 %41 None
- %42 = OpLoad %VertexOutput %out None
- OpReturnValue %42
+%vertex_main_inner = OpFunction %VertexOutput None %38
+ %39 = OpLabel
+ %out = OpVariable %_ptr_Function_VertexOutput Function %42
+ %43 = OpAccessChain %_ptr_Function_v4float %out %uint_0
+ OpStore %43 %47 None
+ %48 = OpAccessChain %_ptr_Function_half %out %uint_1
+ %50 = OpFunctionCall %half %smoothstep_586e12
+ OpStore %48 %50 None
+ %51 = OpLoad %VertexOutput %out None
+ OpReturnValue %51
OpFunctionEnd
-%vertex_main = OpFunction %void None %45
- %46 = OpLabel
- %47 = OpFunctionCall %VertexOutput %vertex_main_inner
- %48 = OpCompositeExtract %v4float %47 0
- OpStore %vertex_main_position_Output %48 None
- %49 = OpCompositeExtract %half %47 1
- OpStore %vertex_main_loc0_Output %49 None
+%vertex_main = OpFunction %void None %54
+ %55 = OpLabel
+ %56 = OpFunctionCall %VertexOutput %vertex_main_inner
+ %57 = OpCompositeExtract %v4float %56 0
+ OpStore %vertex_main_position_Output %57 None
+ %58 = OpCompositeExtract %half %56 1
+ OpStore %vertex_main_loc0_Output %58 None
OpStore %vertex_main___point_size_Output %float_1 None
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.glsl b/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.glsl
index 1a8b2a6..be25200 100644
--- a/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.glsl
@@ -13,7 +13,9 @@
float arg_0 = 2.0f;
float arg_1 = 4.0f;
float arg_2 = 3.0f;
- float res = smoothstep(arg_0, arg_1, arg_2);
+ float v_1 = arg_0;
+ float v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), 0.0f, 1.0f);
+ float res = (v_2 * (v_2 * (3.0f - (2.0f * v_2))));
return res;
}
void main() {
@@ -32,7 +34,9 @@
float arg_0 = 2.0f;
float arg_1 = 4.0f;
float arg_2 = 3.0f;
- float res = smoothstep(arg_0, arg_1, arg_2);
+ float v_1 = arg_0;
+ float v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), 0.0f, 1.0f);
+ float res = (v_2 * (v_2 * (3.0f - (2.0f * v_2))));
return res;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -55,7 +59,9 @@
float arg_0 = 2.0f;
float arg_1 = 4.0f;
float arg_2 = 3.0f;
- float res = smoothstep(arg_0, arg_1, arg_2);
+ float v = arg_0;
+ float v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0f, 1.0f);
+ float res = (v_1 * (v_1 * (3.0f - (2.0f * v_1))));
return res;
}
VertexOutput vertex_main_inner() {
@@ -65,10 +71,10 @@
return tint_symbol;
}
void main() {
- VertexOutput v = vertex_main_inner();
- gl_Position = v.pos;
+ VertexOutput v_2 = vertex_main_inner();
+ gl_Position = v_2.pos;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
- vertex_main_loc0_Output = v.prevent_dce;
+ vertex_main_loc0_Output = v_2.prevent_dce;
gl_PointSize = 1.0f;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.ir.dxc.hlsl
index 4ad7f1b..056f690 100644
--- a/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.ir.dxc.hlsl
@@ -7,7 +7,9 @@
float arg_0 = 2.0f;
float arg_1 = 4.0f;
float arg_2 = 3.0f;
- float res = smoothstep(arg_0, arg_1, arg_2);
+ float v = arg_0;
+ float v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0f, 1.0f);
+ float res = (v_1 * (v_1 * (3.0f - (2.0f * v_1))));
return res;
}
@@ -24,7 +26,9 @@
float arg_0 = 2.0f;
float arg_1 = 4.0f;
float arg_2 = 3.0f;
- float res = smoothstep(arg_0, arg_1, arg_2);
+ float v = arg_0;
+ float v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0f, 1.0f);
+ float res = (v_1 * (v_1 * (3.0f - (2.0f * v_1))));
return res;
}
@@ -51,7 +55,9 @@
float arg_0 = 2.0f;
float arg_1 = 4.0f;
float arg_2 = 3.0f;
- float res = smoothstep(arg_0, arg_1, arg_2);
+ float v = arg_0;
+ float v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0f, 1.0f);
+ float res = (v_1 * (v_1 * (3.0f - (2.0f * v_1))));
return res;
}
@@ -59,13 +65,13 @@
VertexOutput tint_symbol = (VertexOutput)0;
tint_symbol.pos = (0.0f).xxxx;
tint_symbol.prevent_dce = smoothstep_6c4975();
- VertexOutput v = tint_symbol;
- return v;
+ VertexOutput v_2 = tint_symbol;
+ return v_2;
}
vertex_main_outputs vertex_main() {
- VertexOutput v_1 = vertex_main_inner();
- vertex_main_outputs v_2 = {v_1.prevent_dce, v_1.pos};
- return v_2;
+ VertexOutput v_3 = vertex_main_inner();
+ vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos};
+ return v_4;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.ir.fxc.hlsl
index 4ad7f1b..056f690 100644
--- a/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.ir.fxc.hlsl
@@ -7,7 +7,9 @@
float arg_0 = 2.0f;
float arg_1 = 4.0f;
float arg_2 = 3.0f;
- float res = smoothstep(arg_0, arg_1, arg_2);
+ float v = arg_0;
+ float v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0f, 1.0f);
+ float res = (v_1 * (v_1 * (3.0f - (2.0f * v_1))));
return res;
}
@@ -24,7 +26,9 @@
float arg_0 = 2.0f;
float arg_1 = 4.0f;
float arg_2 = 3.0f;
- float res = smoothstep(arg_0, arg_1, arg_2);
+ float v = arg_0;
+ float v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0f, 1.0f);
+ float res = (v_1 * (v_1 * (3.0f - (2.0f * v_1))));
return res;
}
@@ -51,7 +55,9 @@
float arg_0 = 2.0f;
float arg_1 = 4.0f;
float arg_2 = 3.0f;
- float res = smoothstep(arg_0, arg_1, arg_2);
+ float v = arg_0;
+ float v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0f, 1.0f);
+ float res = (v_1 * (v_1 * (3.0f - (2.0f * v_1))));
return res;
}
@@ -59,13 +65,13 @@
VertexOutput tint_symbol = (VertexOutput)0;
tint_symbol.pos = (0.0f).xxxx;
tint_symbol.prevent_dce = smoothstep_6c4975();
- VertexOutput v = tint_symbol;
- return v;
+ VertexOutput v_2 = tint_symbol;
+ return v_2;
}
vertex_main_outputs vertex_main() {
- VertexOutput v_1 = vertex_main_inner();
- vertex_main_outputs v_2 = {v_1.prevent_dce, v_1.pos};
- return v_2;
+ VertexOutput v_3 = vertex_main_inner();
+ vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos};
+ return v_4;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.ir.msl
index 8c990a4..d13344a 100644
--- a/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.ir.msl
@@ -12,7 +12,9 @@
float arg_0 = 2.0f;
float arg_1 = 4.0f;
float arg_2 = 3.0f;
- float res = smoothstep(arg_0, arg_1, arg_2);
+ float const v = arg_0;
+ float const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0f, 1.0f);
+ float res = (v_1 * (v_1 * (3.0f - (2.0f * v_1))));
return res;
}
@@ -34,7 +36,9 @@
float arg_0 = 2.0f;
float arg_1 = 4.0f;
float arg_2 = 3.0f;
- float res = smoothstep(arg_0, arg_1, arg_2);
+ float const v = arg_0;
+ float const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0f, 1.0f);
+ float res = (v_1 * (v_1 * (3.0f - (2.0f * v_1))));
return res;
}
@@ -62,7 +66,9 @@
float arg_0 = 2.0f;
float arg_1 = 4.0f;
float arg_2 = 3.0f;
- float res = smoothstep(arg_0, arg_1, arg_2);
+ float const v = arg_0;
+ float const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), 0.0f, 1.0f);
+ float res = (v_1 * (v_1 * (3.0f - (2.0f * v_1))));
return res;
}
@@ -74,9 +80,9 @@
}
vertex vertex_main_outputs vertex_main() {
- VertexOutput const v = vertex_main_inner();
+ VertexOutput const v_2 = vertex_main_inner();
vertex_main_outputs tint_wrapper_result = {};
- tint_wrapper_result.VertexOutput_pos = v.pos;
- tint_wrapper_result.VertexOutput_prevent_dce = v.prevent_dce;
+ tint_wrapper_result.VertexOutput_pos = v_2.pos;
+ tint_wrapper_result.VertexOutput_prevent_dce = v_2.prevent_dce;
return tint_wrapper_result;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.spvasm b/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.spvasm
index 2a25f1f..c493cf6 100644
--- a/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/smoothstep/6c4975.wgsl.expected.spvasm
@@ -4,10 +4,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 31
+; Bound: 40
; Schema: 0
OpCapability Shader
- %19 = OpExtInstImport "GLSL.std.450"
+ %22 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %fragment_main "fragment_main"
OpExecutionMode %fragment_main OriginUpperLeft
@@ -33,8 +33,10 @@
%float_2 = OpConstant %float 2
%float_4 = OpConstant %float 4
%float_3 = OpConstant %float 3
+ %float_0 = OpConstant %float 0
+ %float_1 = OpConstant %float 1
%void = OpTypeVoid
- %24 = OpTypeFunction %void
+ %33 = OpTypeFunction %void
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -50,16 +52,23 @@
%15 = OpLoad %float %arg_0 None
%16 = OpLoad %float %arg_1 None
%17 = OpLoad %float %arg_2 None
- %18 = OpExtInst %float %19 SmoothStep %15 %16 %17
- OpStore %res %18
- %21 = OpLoad %float %res None
- OpReturnValue %21
+ %18 = OpFSub %float %17 %15
+ %19 = OpFSub %float %16 %15
+ %20 = OpFDiv %float %18 %19
+ %21 = OpExtInst %float %22 NClamp %20 %float_0 %float_1
+ %25 = OpFMul %float %float_2 %21
+ %26 = OpFSub %float %float_3 %25
+ %27 = OpFMul %float %21 %26
+ %28 = OpFMul %float %21 %27
+ OpStore %res %28
+ %30 = OpLoad %float %res None
+ OpReturnValue %30
OpFunctionEnd
-%fragment_main = OpFunction %void None %24
- %25 = OpLabel
- %26 = OpFunctionCall %float %smoothstep_6c4975
- %27 = OpAccessChain %_ptr_StorageBuffer_float %1 %uint_0
- OpStore %27 %26 None
+%fragment_main = OpFunction %void None %33
+ %34 = OpLabel
+ %35 = OpFunctionCall %float %smoothstep_6c4975
+ %36 = OpAccessChain %_ptr_StorageBuffer_float %1 %uint_0
+ OpStore %36 %35 None
OpReturn
OpFunctionEnd
;
@@ -68,10 +77,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 31
+; Bound: 40
; Schema: 0
OpCapability Shader
- %19 = OpExtInstImport "GLSL.std.450"
+ %22 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %compute_main LocalSize 1 1 1
@@ -97,8 +106,10 @@
%float_2 = OpConstant %float 2
%float_4 = OpConstant %float 4
%float_3 = OpConstant %float 3
+ %float_0 = OpConstant %float 0
+ %float_1 = OpConstant %float 1
%void = OpTypeVoid
- %24 = OpTypeFunction %void
+ %33 = OpTypeFunction %void
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -114,16 +125,23 @@
%15 = OpLoad %float %arg_0 None
%16 = OpLoad %float %arg_1 None
%17 = OpLoad %float %arg_2 None
- %18 = OpExtInst %float %19 SmoothStep %15 %16 %17
- OpStore %res %18
- %21 = OpLoad %float %res None
- OpReturnValue %21
+ %18 = OpFSub %float %17 %15
+ %19 = OpFSub %float %16 %15
+ %20 = OpFDiv %float %18 %19
+ %21 = OpExtInst %float %22 NClamp %20 %float_0 %float_1
+ %25 = OpFMul %float %float_2 %21
+ %26 = OpFSub %float %float_3 %25
+ %27 = OpFMul %float %21 %26
+ %28 = OpFMul %float %21 %27
+ OpStore %res %28
+ %30 = OpLoad %float %res None
+ OpReturnValue %30
OpFunctionEnd
-%compute_main = OpFunction %void None %24
- %25 = OpLabel
- %26 = OpFunctionCall %float %smoothstep_6c4975
- %27 = OpAccessChain %_ptr_StorageBuffer_float %1 %uint_0
- OpStore %27 %26 None
+%compute_main = OpFunction %void None %33
+ %34 = OpLabel
+ %35 = OpFunctionCall %float %smoothstep_6c4975
+ %36 = OpAccessChain %_ptr_StorageBuffer_float %1 %uint_0
+ OpStore %36 %35 None
OpReturn
OpFunctionEnd
;
@@ -132,10 +150,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 49
+; Bound: 57
; Schema: 0
OpCapability Shader
- %22 = OpExtInstImport "GLSL.std.450"
+ %25 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %vertex_main_position_Output %vertex_main_loc0_Output %vertex_main___point_size_Output
OpName %vertex_main_position_Output "vertex_main_position_Output"
@@ -170,18 +188,19 @@
%float_2 = OpConstant %float 2
%float_4 = OpConstant %float 4
%float_3 = OpConstant %float 3
+ %float_0 = OpConstant %float 0
+ %float_1 = OpConstant %float 1
%VertexOutput = OpTypeStruct %v4float %float
- %27 = OpTypeFunction %VertexOutput
+ %36 = OpTypeFunction %VertexOutput
%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
- %31 = OpConstantNull %VertexOutput
+ %40 = OpConstantNull %VertexOutput
%_ptr_Function_v4float = OpTypePointer Function %v4float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
- %36 = OpConstantNull %v4float
+ %45 = OpConstantNull %v4float
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
- %43 = OpTypeFunction %void
- %float_1 = OpConstant %float 1
+ %52 = OpTypeFunction %void
%smoothstep_6c4975 = OpFunction %float None %9
%10 = OpLabel
%arg_0 = OpVariable %_ptr_Function_float Function
@@ -194,29 +213,36 @@
%18 = OpLoad %float %arg_0 None
%19 = OpLoad %float %arg_1 None
%20 = OpLoad %float %arg_2 None
- %21 = OpExtInst %float %22 SmoothStep %18 %19 %20
- OpStore %res %21
- %24 = OpLoad %float %res None
- OpReturnValue %24
+ %21 = OpFSub %float %20 %18
+ %22 = OpFSub %float %19 %18
+ %23 = OpFDiv %float %21 %22
+ %24 = OpExtInst %float %25 NClamp %23 %float_0 %float_1
+ %28 = OpFMul %float %float_2 %24
+ %29 = OpFSub %float %float_3 %28
+ %30 = OpFMul %float %24 %29
+ %31 = OpFMul %float %24 %30
+ OpStore %res %31
+ %33 = OpLoad %float %res None
+ OpReturnValue %33
OpFunctionEnd
-%vertex_main_inner = OpFunction %VertexOutput None %27
- %28 = OpLabel
- %out = OpVariable %_ptr_Function_VertexOutput Function %31
- %32 = OpAccessChain %_ptr_Function_v4float %out %uint_0
- OpStore %32 %36 None
- %37 = OpAccessChain %_ptr_Function_float %out %uint_1
- %39 = OpFunctionCall %float %smoothstep_6c4975
- OpStore %37 %39 None
- %40 = OpLoad %VertexOutput %out None
- OpReturnValue %40
+%vertex_main_inner = OpFunction %VertexOutput None %36
+ %37 = OpLabel
+ %out = OpVariable %_ptr_Function_VertexOutput Function %40
+ %41 = OpAccessChain %_ptr_Function_v4float %out %uint_0
+ OpStore %41 %45 None
+ %46 = OpAccessChain %_ptr_Function_float %out %uint_1
+ %48 = OpFunctionCall %float %smoothstep_6c4975
+ OpStore %46 %48 None
+ %49 = OpLoad %VertexOutput %out None
+ OpReturnValue %49
OpFunctionEnd
-%vertex_main = OpFunction %void None %43
- %44 = OpLabel
- %45 = OpFunctionCall %VertexOutput %vertex_main_inner
- %46 = OpCompositeExtract %v4float %45 0
- OpStore %vertex_main_position_Output %46 None
- %47 = OpCompositeExtract %float %45 1
- OpStore %vertex_main_loc0_Output %47 None
+%vertex_main = OpFunction %void None %52
+ %53 = OpLabel
+ %54 = OpFunctionCall %VertexOutput %vertex_main_inner
+ %55 = OpCompositeExtract %v4float %54 0
+ OpStore %vertex_main_position_Output %55 None
+ %56 = OpCompositeExtract %float %54 1
+ OpStore %vertex_main_loc0_Output %56 None
OpStore %vertex_main___point_size_Output %float_1 None
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.glsl b/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.glsl
index 1801249..f7c2bd8 100644
--- a/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.glsl
@@ -14,7 +14,9 @@
f16vec3 arg_0 = f16vec3(2.0hf);
f16vec3 arg_1 = f16vec3(4.0hf);
f16vec3 arg_2 = f16vec3(3.0hf);
- f16vec3 res = smoothstep(arg_0, arg_1, arg_2);
+ f16vec3 v_1 = arg_0;
+ f16vec3 v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), f16vec3(0.0hf), f16vec3(1.0hf));
+ f16vec3 res = (v_2 * (v_2 * (f16vec3(3.0hf) - (f16vec3(2.0hf) * v_2))));
return res;
}
void main() {
@@ -34,7 +36,9 @@
f16vec3 arg_0 = f16vec3(2.0hf);
f16vec3 arg_1 = f16vec3(4.0hf);
f16vec3 arg_2 = f16vec3(3.0hf);
- f16vec3 res = smoothstep(arg_0, arg_1, arg_2);
+ f16vec3 v_1 = arg_0;
+ f16vec3 v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), f16vec3(0.0hf), f16vec3(1.0hf));
+ f16vec3 res = (v_2 * (v_2 * (f16vec3(3.0hf) - (f16vec3(2.0hf) * v_2))));
return res;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -58,7 +62,9 @@
f16vec3 arg_0 = f16vec3(2.0hf);
f16vec3 arg_1 = f16vec3(4.0hf);
f16vec3 arg_2 = f16vec3(3.0hf);
- f16vec3 res = smoothstep(arg_0, arg_1, arg_2);
+ f16vec3 v = arg_0;
+ f16vec3 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), f16vec3(0.0hf), f16vec3(1.0hf));
+ f16vec3 res = (v_1 * (v_1 * (f16vec3(3.0hf) - (f16vec3(2.0hf) * v_1))));
return res;
}
VertexOutput vertex_main_inner() {
@@ -68,10 +74,10 @@
return tint_symbol;
}
void main() {
- VertexOutput v = vertex_main_inner();
- gl_Position = v.pos;
+ VertexOutput v_2 = vertex_main_inner();
+ gl_Position = v_2.pos;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
- vertex_main_loc0_Output = v.prevent_dce;
+ vertex_main_loc0_Output = v_2.prevent_dce;
gl_PointSize = 1.0f;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.ir.dxc.hlsl
index 8cae982..3c12dfc 100644
--- a/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.ir.dxc.hlsl
@@ -7,7 +7,9 @@
vector<float16_t, 3> arg_0 = (float16_t(2.0h)).xxx;
vector<float16_t, 3> arg_1 = (float16_t(4.0h)).xxx;
vector<float16_t, 3> arg_2 = (float16_t(3.0h)).xxx;
- vector<float16_t, 3> res = smoothstep(arg_0, arg_1, arg_2);
+ vector<float16_t, 3> v = arg_0;
+ vector<float16_t, 3> v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (float16_t(0.0h)).xxx, (float16_t(1.0h)).xxx);
+ vector<float16_t, 3> res = (v_1 * (v_1 * ((float16_t(3.0h)).xxx - ((float16_t(2.0h)).xxx * v_1))));
return res;
}
@@ -24,7 +26,9 @@
vector<float16_t, 3> arg_0 = (float16_t(2.0h)).xxx;
vector<float16_t, 3> arg_1 = (float16_t(4.0h)).xxx;
vector<float16_t, 3> arg_2 = (float16_t(3.0h)).xxx;
- vector<float16_t, 3> res = smoothstep(arg_0, arg_1, arg_2);
+ vector<float16_t, 3> v = arg_0;
+ vector<float16_t, 3> v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (float16_t(0.0h)).xxx, (float16_t(1.0h)).xxx);
+ vector<float16_t, 3> res = (v_1 * (v_1 * ((float16_t(3.0h)).xxx - ((float16_t(2.0h)).xxx * v_1))));
return res;
}
@@ -51,7 +55,9 @@
vector<float16_t, 3> arg_0 = (float16_t(2.0h)).xxx;
vector<float16_t, 3> arg_1 = (float16_t(4.0h)).xxx;
vector<float16_t, 3> arg_2 = (float16_t(3.0h)).xxx;
- vector<float16_t, 3> res = smoothstep(arg_0, arg_1, arg_2);
+ vector<float16_t, 3> v = arg_0;
+ vector<float16_t, 3> v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (float16_t(0.0h)).xxx, (float16_t(1.0h)).xxx);
+ vector<float16_t, 3> res = (v_1 * (v_1 * ((float16_t(3.0h)).xxx - ((float16_t(2.0h)).xxx * v_1))));
return res;
}
@@ -59,13 +65,13 @@
VertexOutput tint_symbol = (VertexOutput)0;
tint_symbol.pos = (0.0f).xxxx;
tint_symbol.prevent_dce = smoothstep_6e7a74();
- VertexOutput v = tint_symbol;
- return v;
+ VertexOutput v_2 = tint_symbol;
+ return v_2;
}
vertex_main_outputs vertex_main() {
- VertexOutput v_1 = vertex_main_inner();
- vertex_main_outputs v_2 = {v_1.prevent_dce, v_1.pos};
- return v_2;
+ VertexOutput v_3 = vertex_main_inner();
+ vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos};
+ return v_4;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.ir.msl
index dc8d71a..9ba8ced 100644
--- a/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.ir.msl
@@ -12,7 +12,9 @@
half3 arg_0 = half3(2.0h);
half3 arg_1 = half3(4.0h);
half3 arg_2 = half3(3.0h);
- half3 res = smoothstep(arg_0, arg_1, arg_2);
+ half3 const v = arg_0;
+ half3 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), half3(0.0h), half3(1.0h));
+ half3 res = (v_1 * (v_1 * (half3(3.0h) - (half3(2.0h) * v_1))));
return res;
}
@@ -34,7 +36,9 @@
half3 arg_0 = half3(2.0h);
half3 arg_1 = half3(4.0h);
half3 arg_2 = half3(3.0h);
- half3 res = smoothstep(arg_0, arg_1, arg_2);
+ half3 const v = arg_0;
+ half3 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), half3(0.0h), half3(1.0h));
+ half3 res = (v_1 * (v_1 * (half3(3.0h) - (half3(2.0h) * v_1))));
return res;
}
@@ -62,7 +66,9 @@
half3 arg_0 = half3(2.0h);
half3 arg_1 = half3(4.0h);
half3 arg_2 = half3(3.0h);
- half3 res = smoothstep(arg_0, arg_1, arg_2);
+ half3 const v = arg_0;
+ half3 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), half3(0.0h), half3(1.0h));
+ half3 res = (v_1 * (v_1 * (half3(3.0h) - (half3(2.0h) * v_1))));
return res;
}
@@ -74,9 +80,9 @@
}
vertex vertex_main_outputs vertex_main() {
- VertexOutput const v = vertex_main_inner();
+ VertexOutput const v_2 = vertex_main_inner();
vertex_main_outputs tint_wrapper_result = {};
- tint_wrapper_result.VertexOutput_pos = v.pos;
- tint_wrapper_result.VertexOutput_prevent_dce = v.prevent_dce;
+ tint_wrapper_result.VertexOutput_pos = v_2.pos;
+ tint_wrapper_result.VertexOutput_prevent_dce = v_2.prevent_dce;
return tint_wrapper_result;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.spvasm b/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.spvasm
index cfd0e38..9221c01 100644
--- a/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/smoothstep/6e7a74.wgsl.expected.spvasm
@@ -4,13 +4,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 35
+; Bound: 45
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
- %23 = OpExtInstImport "GLSL.std.450"
+ %26 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %fragment_main "fragment_main"
OpExecutionMode %fragment_main OriginUpperLeft
@@ -40,8 +40,11 @@
%14 = OpConstantComposite %v3half %half_0x1p_2 %half_0x1p_2 %half_0x1p_2
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
%17 = OpConstantComposite %v3half %half_0x1_8p_1 %half_0x1_8p_1 %half_0x1_8p_1
+ %27 = OpConstantNull %v3half
+%half_0x1p_0 = OpConstant %half 0x1p+0
+ %28 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
%void = OpTypeVoid
- %28 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_StorageBuffer_v3half = OpTypePointer StorageBuffer %v3half
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -57,16 +60,23 @@
%19 = OpLoad %v3half %arg_0 None
%20 = OpLoad %v3half %arg_1 None
%21 = OpLoad %v3half %arg_2 None
- %22 = OpExtInst %v3half %23 SmoothStep %19 %20 %21
- OpStore %res %22
- %25 = OpLoad %v3half %res None
- OpReturnValue %25
+ %22 = OpFSub %v3half %21 %19
+ %23 = OpFSub %v3half %20 %19
+ %24 = OpFDiv %v3half %22 %23
+ %25 = OpExtInst %v3half %26 NClamp %24 %27 %28
+ %30 = OpFMul %v3half %11 %25
+ %31 = OpFSub %v3half %17 %30
+ %32 = OpFMul %v3half %25 %31
+ %33 = OpFMul %v3half %25 %32
+ OpStore %res %33
+ %35 = OpLoad %v3half %res None
+ OpReturnValue %35
OpFunctionEnd
-%fragment_main = OpFunction %void None %28
- %29 = OpLabel
- %30 = OpFunctionCall %v3half %smoothstep_6e7a74
- %31 = OpAccessChain %_ptr_StorageBuffer_v3half %1 %uint_0
- OpStore %31 %30 None
+%fragment_main = OpFunction %void None %38
+ %39 = OpLabel
+ %40 = OpFunctionCall %v3half %smoothstep_6e7a74
+ %41 = OpAccessChain %_ptr_StorageBuffer_v3half %1 %uint_0
+ OpStore %41 %40 None
OpReturn
OpFunctionEnd
;
@@ -75,13 +85,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 35
+; Bound: 45
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
- %23 = OpExtInstImport "GLSL.std.450"
+ %26 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %compute_main LocalSize 1 1 1
@@ -111,8 +121,11 @@
%14 = OpConstantComposite %v3half %half_0x1p_2 %half_0x1p_2 %half_0x1p_2
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
%17 = OpConstantComposite %v3half %half_0x1_8p_1 %half_0x1_8p_1 %half_0x1_8p_1
+ %27 = OpConstantNull %v3half
+%half_0x1p_0 = OpConstant %half 0x1p+0
+ %28 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
%void = OpTypeVoid
- %28 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_StorageBuffer_v3half = OpTypePointer StorageBuffer %v3half
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -128,16 +141,23 @@
%19 = OpLoad %v3half %arg_0 None
%20 = OpLoad %v3half %arg_1 None
%21 = OpLoad %v3half %arg_2 None
- %22 = OpExtInst %v3half %23 SmoothStep %19 %20 %21
- OpStore %res %22
- %25 = OpLoad %v3half %res None
- OpReturnValue %25
+ %22 = OpFSub %v3half %21 %19
+ %23 = OpFSub %v3half %20 %19
+ %24 = OpFDiv %v3half %22 %23
+ %25 = OpExtInst %v3half %26 NClamp %24 %27 %28
+ %30 = OpFMul %v3half %11 %25
+ %31 = OpFSub %v3half %17 %30
+ %32 = OpFMul %v3half %25 %31
+ %33 = OpFMul %v3half %25 %32
+ OpStore %res %33
+ %35 = OpLoad %v3half %res None
+ OpReturnValue %35
OpFunctionEnd
-%compute_main = OpFunction %void None %28
- %29 = OpLabel
- %30 = OpFunctionCall %v3half %smoothstep_6e7a74
- %31 = OpAccessChain %_ptr_StorageBuffer_v3half %1 %uint_0
- OpStore %31 %30 None
+%compute_main = OpFunction %void None %38
+ %39 = OpLabel
+ %40 = OpFunctionCall %v3half %smoothstep_6e7a74
+ %41 = OpAccessChain %_ptr_StorageBuffer_v3half %1 %uint_0
+ OpStore %41 %40 None
OpReturn
OpFunctionEnd
;
@@ -146,14 +166,14 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 55
+; Bound: 65
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
- %28 = OpExtInstImport "GLSL.std.450"
+ %31 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %vertex_main_position_Output %vertex_main_loc0_Output %vertex_main___point_size_Output
OpName %vertex_main_position_Output "vertex_main_position_Output"
@@ -194,17 +214,20 @@
%19 = OpConstantComposite %v3half %half_0x1p_2 %half_0x1p_2 %half_0x1p_2
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
%22 = OpConstantComposite %v3half %half_0x1_8p_1 %half_0x1_8p_1 %half_0x1_8p_1
+ %32 = OpConstantNull %v3half
+%half_0x1p_0 = OpConstant %half 0x1p+0
+ %33 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
%VertexOutput = OpTypeStruct %v4float %v3half
- %33 = OpTypeFunction %VertexOutput
+ %43 = OpTypeFunction %VertexOutput
%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
- %37 = OpConstantNull %VertexOutput
+ %47 = OpConstantNull %VertexOutput
%_ptr_Function_v4float = OpTypePointer Function %v4float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
- %42 = OpConstantNull %v4float
+ %52 = OpConstantNull %v4float
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
- %49 = OpTypeFunction %void
+ %59 = OpTypeFunction %void
%float_1 = OpConstant %float 1
%smoothstep_6e7a74 = OpFunction %v3half None %12
%13 = OpLabel
@@ -218,29 +241,36 @@
%24 = OpLoad %v3half %arg_0 None
%25 = OpLoad %v3half %arg_1 None
%26 = OpLoad %v3half %arg_2 None
- %27 = OpExtInst %v3half %28 SmoothStep %24 %25 %26
- OpStore %res %27
- %30 = OpLoad %v3half %res None
- OpReturnValue %30
+ %27 = OpFSub %v3half %26 %24
+ %28 = OpFSub %v3half %25 %24
+ %29 = OpFDiv %v3half %27 %28
+ %30 = OpExtInst %v3half %31 NClamp %29 %32 %33
+ %35 = OpFMul %v3half %16 %30
+ %36 = OpFSub %v3half %22 %35
+ %37 = OpFMul %v3half %30 %36
+ %38 = OpFMul %v3half %30 %37
+ OpStore %res %38
+ %40 = OpLoad %v3half %res None
+ OpReturnValue %40
OpFunctionEnd
-%vertex_main_inner = OpFunction %VertexOutput None %33
- %34 = OpLabel
- %out = OpVariable %_ptr_Function_VertexOutput Function %37
- %38 = OpAccessChain %_ptr_Function_v4float %out %uint_0
- OpStore %38 %42 None
- %43 = OpAccessChain %_ptr_Function_v3half %out %uint_1
- %45 = OpFunctionCall %v3half %smoothstep_6e7a74
- OpStore %43 %45 None
- %46 = OpLoad %VertexOutput %out None
- OpReturnValue %46
+%vertex_main_inner = OpFunction %VertexOutput None %43
+ %44 = OpLabel
+ %out = OpVariable %_ptr_Function_VertexOutput Function %47
+ %48 = OpAccessChain %_ptr_Function_v4float %out %uint_0
+ OpStore %48 %52 None
+ %53 = OpAccessChain %_ptr_Function_v3half %out %uint_1
+ %55 = OpFunctionCall %v3half %smoothstep_6e7a74
+ OpStore %53 %55 None
+ %56 = OpLoad %VertexOutput %out None
+ OpReturnValue %56
OpFunctionEnd
-%vertex_main = OpFunction %void None %49
- %50 = OpLabel
- %51 = OpFunctionCall %VertexOutput %vertex_main_inner
- %52 = OpCompositeExtract %v4float %51 0
- OpStore %vertex_main_position_Output %52 None
- %53 = OpCompositeExtract %v3half %51 1
- OpStore %vertex_main_loc0_Output %53 None
+%vertex_main = OpFunction %void None %59
+ %60 = OpLabel
+ %61 = OpFunctionCall %VertexOutput %vertex_main_inner
+ %62 = OpCompositeExtract %v4float %61 0
+ OpStore %vertex_main_position_Output %62 None
+ %63 = OpCompositeExtract %v3half %61 1
+ OpStore %vertex_main_loc0_Output %63 None
OpStore %vertex_main___point_size_Output %float_1 None
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.glsl b/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.glsl
index 199026c..bd4a615 100644
--- a/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.glsl
@@ -13,7 +13,9 @@
vec3 arg_0 = vec3(2.0f);
vec3 arg_1 = vec3(4.0f);
vec3 arg_2 = vec3(3.0f);
- vec3 res = smoothstep(arg_0, arg_1, arg_2);
+ vec3 v_1 = arg_0;
+ vec3 v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), vec3(0.0f), vec3(1.0f));
+ vec3 res = (v_2 * (v_2 * (vec3(3.0f) - (vec3(2.0f) * v_2))));
return res;
}
void main() {
@@ -32,7 +34,9 @@
vec3 arg_0 = vec3(2.0f);
vec3 arg_1 = vec3(4.0f);
vec3 arg_2 = vec3(3.0f);
- vec3 res = smoothstep(arg_0, arg_1, arg_2);
+ vec3 v_1 = arg_0;
+ vec3 v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), vec3(0.0f), vec3(1.0f));
+ vec3 res = (v_2 * (v_2 * (vec3(3.0f) - (vec3(2.0f) * v_2))));
return res;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -55,7 +59,9 @@
vec3 arg_0 = vec3(2.0f);
vec3 arg_1 = vec3(4.0f);
vec3 arg_2 = vec3(3.0f);
- vec3 res = smoothstep(arg_0, arg_1, arg_2);
+ vec3 v = arg_0;
+ vec3 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), vec3(0.0f), vec3(1.0f));
+ vec3 res = (v_1 * (v_1 * (vec3(3.0f) - (vec3(2.0f) * v_1))));
return res;
}
VertexOutput vertex_main_inner() {
@@ -65,10 +71,10 @@
return tint_symbol;
}
void main() {
- VertexOutput v = vertex_main_inner();
- gl_Position = v.pos;
+ VertexOutput v_2 = vertex_main_inner();
+ gl_Position = v_2.pos;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
- vertex_main_loc0_Output = v.prevent_dce;
+ vertex_main_loc0_Output = v_2.prevent_dce;
gl_PointSize = 1.0f;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.ir.dxc.hlsl
index 26c059e..6f54ea0 100644
--- a/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.ir.dxc.hlsl
@@ -7,7 +7,9 @@
float3 arg_0 = (2.0f).xxx;
float3 arg_1 = (4.0f).xxx;
float3 arg_2 = (3.0f).xxx;
- float3 res = smoothstep(arg_0, arg_1, arg_2);
+ float3 v = arg_0;
+ float3 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xxx, (1.0f).xxx);
+ float3 res = (v_1 * (v_1 * ((3.0f).xxx - ((2.0f).xxx * v_1))));
return res;
}
@@ -24,7 +26,9 @@
float3 arg_0 = (2.0f).xxx;
float3 arg_1 = (4.0f).xxx;
float3 arg_2 = (3.0f).xxx;
- float3 res = smoothstep(arg_0, arg_1, arg_2);
+ float3 v = arg_0;
+ float3 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xxx, (1.0f).xxx);
+ float3 res = (v_1 * (v_1 * ((3.0f).xxx - ((2.0f).xxx * v_1))));
return res;
}
@@ -51,7 +55,9 @@
float3 arg_0 = (2.0f).xxx;
float3 arg_1 = (4.0f).xxx;
float3 arg_2 = (3.0f).xxx;
- float3 res = smoothstep(arg_0, arg_1, arg_2);
+ float3 v = arg_0;
+ float3 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xxx, (1.0f).xxx);
+ float3 res = (v_1 * (v_1 * ((3.0f).xxx - ((2.0f).xxx * v_1))));
return res;
}
@@ -59,13 +65,13 @@
VertexOutput tint_symbol = (VertexOutput)0;
tint_symbol.pos = (0.0f).xxxx;
tint_symbol.prevent_dce = smoothstep_aad1db();
- VertexOutput v = tint_symbol;
- return v;
+ VertexOutput v_2 = tint_symbol;
+ return v_2;
}
vertex_main_outputs vertex_main() {
- VertexOutput v_1 = vertex_main_inner();
- vertex_main_outputs v_2 = {v_1.prevent_dce, v_1.pos};
- return v_2;
+ VertexOutput v_3 = vertex_main_inner();
+ vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos};
+ return v_4;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.ir.fxc.hlsl
index 26c059e..6f54ea0 100644
--- a/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.ir.fxc.hlsl
@@ -7,7 +7,9 @@
float3 arg_0 = (2.0f).xxx;
float3 arg_1 = (4.0f).xxx;
float3 arg_2 = (3.0f).xxx;
- float3 res = smoothstep(arg_0, arg_1, arg_2);
+ float3 v = arg_0;
+ float3 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xxx, (1.0f).xxx);
+ float3 res = (v_1 * (v_1 * ((3.0f).xxx - ((2.0f).xxx * v_1))));
return res;
}
@@ -24,7 +26,9 @@
float3 arg_0 = (2.0f).xxx;
float3 arg_1 = (4.0f).xxx;
float3 arg_2 = (3.0f).xxx;
- float3 res = smoothstep(arg_0, arg_1, arg_2);
+ float3 v = arg_0;
+ float3 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xxx, (1.0f).xxx);
+ float3 res = (v_1 * (v_1 * ((3.0f).xxx - ((2.0f).xxx * v_1))));
return res;
}
@@ -51,7 +55,9 @@
float3 arg_0 = (2.0f).xxx;
float3 arg_1 = (4.0f).xxx;
float3 arg_2 = (3.0f).xxx;
- float3 res = smoothstep(arg_0, arg_1, arg_2);
+ float3 v = arg_0;
+ float3 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (0.0f).xxx, (1.0f).xxx);
+ float3 res = (v_1 * (v_1 * ((3.0f).xxx - ((2.0f).xxx * v_1))));
return res;
}
@@ -59,13 +65,13 @@
VertexOutput tint_symbol = (VertexOutput)0;
tint_symbol.pos = (0.0f).xxxx;
tint_symbol.prevent_dce = smoothstep_aad1db();
- VertexOutput v = tint_symbol;
- return v;
+ VertexOutput v_2 = tint_symbol;
+ return v_2;
}
vertex_main_outputs vertex_main() {
- VertexOutput v_1 = vertex_main_inner();
- vertex_main_outputs v_2 = {v_1.prevent_dce, v_1.pos};
- return v_2;
+ VertexOutput v_3 = vertex_main_inner();
+ vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos};
+ return v_4;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.ir.msl
index 9daf992..bf51c16 100644
--- a/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.ir.msl
@@ -12,7 +12,9 @@
float3 arg_0 = float3(2.0f);
float3 arg_1 = float3(4.0f);
float3 arg_2 = float3(3.0f);
- float3 res = smoothstep(arg_0, arg_1, arg_2);
+ float3 const v = arg_0;
+ float3 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), float3(0.0f), float3(1.0f));
+ float3 res = (v_1 * (v_1 * (float3(3.0f) - (float3(2.0f) * v_1))));
return res;
}
@@ -34,7 +36,9 @@
float3 arg_0 = float3(2.0f);
float3 arg_1 = float3(4.0f);
float3 arg_2 = float3(3.0f);
- float3 res = smoothstep(arg_0, arg_1, arg_2);
+ float3 const v = arg_0;
+ float3 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), float3(0.0f), float3(1.0f));
+ float3 res = (v_1 * (v_1 * (float3(3.0f) - (float3(2.0f) * v_1))));
return res;
}
@@ -62,7 +66,9 @@
float3 arg_0 = float3(2.0f);
float3 arg_1 = float3(4.0f);
float3 arg_2 = float3(3.0f);
- float3 res = smoothstep(arg_0, arg_1, arg_2);
+ float3 const v = arg_0;
+ float3 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), float3(0.0f), float3(1.0f));
+ float3 res = (v_1 * (v_1 * (float3(3.0f) - (float3(2.0f) * v_1))));
return res;
}
@@ -74,9 +80,9 @@
}
vertex vertex_main_outputs vertex_main() {
- VertexOutput const v = vertex_main_inner();
+ VertexOutput const v_2 = vertex_main_inner();
vertex_main_outputs tint_wrapper_result = {};
- tint_wrapper_result.VertexOutput_pos = v.pos;
- tint_wrapper_result.VertexOutput_prevent_dce = v.prevent_dce;
+ tint_wrapper_result.VertexOutput_pos = v_2.pos;
+ tint_wrapper_result.VertexOutput_prevent_dce = v_2.prevent_dce;
return tint_wrapper_result;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.spvasm b/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.spvasm
index 07adfc7..c7b7d58 100644
--- a/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/smoothstep/aad1db.wgsl.expected.spvasm
@@ -4,10 +4,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 35
+; Bound: 45
; Schema: 0
OpCapability Shader
- %23 = OpExtInstImport "GLSL.std.450"
+ %26 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %fragment_main "fragment_main"
OpExecutionMode %fragment_main OriginUpperLeft
@@ -37,8 +37,11 @@
%14 = OpConstantComposite %v3float %float_4 %float_4 %float_4
%float_3 = OpConstant %float 3
%17 = OpConstantComposite %v3float %float_3 %float_3 %float_3
+ %27 = OpConstantNull %v3float
+ %float_1 = OpConstant %float 1
+ %28 = OpConstantComposite %v3float %float_1 %float_1 %float_1
%void = OpTypeVoid
- %28 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -54,16 +57,23 @@
%19 = OpLoad %v3float %arg_0 None
%20 = OpLoad %v3float %arg_1 None
%21 = OpLoad %v3float %arg_2 None
- %22 = OpExtInst %v3float %23 SmoothStep %19 %20 %21
- OpStore %res %22
- %25 = OpLoad %v3float %res None
- OpReturnValue %25
+ %22 = OpFSub %v3float %21 %19
+ %23 = OpFSub %v3float %20 %19
+ %24 = OpFDiv %v3float %22 %23
+ %25 = OpExtInst %v3float %26 NClamp %24 %27 %28
+ %30 = OpFMul %v3float %11 %25
+ %31 = OpFSub %v3float %17 %30
+ %32 = OpFMul %v3float %25 %31
+ %33 = OpFMul %v3float %25 %32
+ OpStore %res %33
+ %35 = OpLoad %v3float %res None
+ OpReturnValue %35
OpFunctionEnd
-%fragment_main = OpFunction %void None %28
- %29 = OpLabel
- %30 = OpFunctionCall %v3float %smoothstep_aad1db
- %31 = OpAccessChain %_ptr_StorageBuffer_v3float %1 %uint_0
- OpStore %31 %30 None
+%fragment_main = OpFunction %void None %38
+ %39 = OpLabel
+ %40 = OpFunctionCall %v3float %smoothstep_aad1db
+ %41 = OpAccessChain %_ptr_StorageBuffer_v3float %1 %uint_0
+ OpStore %41 %40 None
OpReturn
OpFunctionEnd
;
@@ -72,10 +82,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 35
+; Bound: 45
; Schema: 0
OpCapability Shader
- %23 = OpExtInstImport "GLSL.std.450"
+ %26 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %compute_main LocalSize 1 1 1
@@ -105,8 +115,11 @@
%14 = OpConstantComposite %v3float %float_4 %float_4 %float_4
%float_3 = OpConstant %float 3
%17 = OpConstantComposite %v3float %float_3 %float_3 %float_3
+ %27 = OpConstantNull %v3float
+ %float_1 = OpConstant %float 1
+ %28 = OpConstantComposite %v3float %float_1 %float_1 %float_1
%void = OpTypeVoid
- %28 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -122,16 +135,23 @@
%19 = OpLoad %v3float %arg_0 None
%20 = OpLoad %v3float %arg_1 None
%21 = OpLoad %v3float %arg_2 None
- %22 = OpExtInst %v3float %23 SmoothStep %19 %20 %21
- OpStore %res %22
- %25 = OpLoad %v3float %res None
- OpReturnValue %25
+ %22 = OpFSub %v3float %21 %19
+ %23 = OpFSub %v3float %20 %19
+ %24 = OpFDiv %v3float %22 %23
+ %25 = OpExtInst %v3float %26 NClamp %24 %27 %28
+ %30 = OpFMul %v3float %11 %25
+ %31 = OpFSub %v3float %17 %30
+ %32 = OpFMul %v3float %25 %31
+ %33 = OpFMul %v3float %25 %32
+ OpStore %res %33
+ %35 = OpLoad %v3float %res None
+ OpReturnValue %35
OpFunctionEnd
-%compute_main = OpFunction %void None %28
- %29 = OpLabel
- %30 = OpFunctionCall %v3float %smoothstep_aad1db
- %31 = OpAccessChain %_ptr_StorageBuffer_v3float %1 %uint_0
- OpStore %31 %30 None
+%compute_main = OpFunction %void None %38
+ %39 = OpLabel
+ %40 = OpFunctionCall %v3float %smoothstep_aad1db
+ %41 = OpAccessChain %_ptr_StorageBuffer_v3float %1 %uint_0
+ OpStore %41 %40 None
OpReturn
OpFunctionEnd
;
@@ -140,10 +160,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 54
+; Bound: 63
; Schema: 0
OpCapability Shader
- %27 = OpExtInstImport "GLSL.std.450"
+ %30 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %vertex_main_position_Output %vertex_main_loc0_Output %vertex_main___point_size_Output
OpName %vertex_main_position_Output "vertex_main_position_Output"
@@ -183,18 +203,20 @@
%18 = OpConstantComposite %v3float %float_4 %float_4 %float_4
%float_3 = OpConstant %float 3
%21 = OpConstantComposite %v3float %float_3 %float_3 %float_3
+ %31 = OpConstantNull %v3float
+ %float_1 = OpConstant %float 1
+ %32 = OpConstantComposite %v3float %float_1 %float_1 %float_1
%VertexOutput = OpTypeStruct %v4float %v3float
- %32 = OpTypeFunction %VertexOutput
+ %42 = OpTypeFunction %VertexOutput
%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
- %36 = OpConstantNull %VertexOutput
+ %46 = OpConstantNull %VertexOutput
%_ptr_Function_v4float = OpTypePointer Function %v4float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
- %41 = OpConstantNull %v4float
+ %51 = OpConstantNull %v4float
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
- %48 = OpTypeFunction %void
- %float_1 = OpConstant %float 1
+ %58 = OpTypeFunction %void
%smoothstep_aad1db = OpFunction %v3float None %11
%12 = OpLabel
%arg_0 = OpVariable %_ptr_Function_v3float Function
@@ -207,29 +229,36 @@
%23 = OpLoad %v3float %arg_0 None
%24 = OpLoad %v3float %arg_1 None
%25 = OpLoad %v3float %arg_2 None
- %26 = OpExtInst %v3float %27 SmoothStep %23 %24 %25
- OpStore %res %26
- %29 = OpLoad %v3float %res None
- OpReturnValue %29
+ %26 = OpFSub %v3float %25 %23
+ %27 = OpFSub %v3float %24 %23
+ %28 = OpFDiv %v3float %26 %27
+ %29 = OpExtInst %v3float %30 NClamp %28 %31 %32
+ %34 = OpFMul %v3float %15 %29
+ %35 = OpFSub %v3float %21 %34
+ %36 = OpFMul %v3float %29 %35
+ %37 = OpFMul %v3float %29 %36
+ OpStore %res %37
+ %39 = OpLoad %v3float %res None
+ OpReturnValue %39
OpFunctionEnd
-%vertex_main_inner = OpFunction %VertexOutput None %32
- %33 = OpLabel
- %out = OpVariable %_ptr_Function_VertexOutput Function %36
- %37 = OpAccessChain %_ptr_Function_v4float %out %uint_0
- OpStore %37 %41 None
- %42 = OpAccessChain %_ptr_Function_v3float %out %uint_1
- %44 = OpFunctionCall %v3float %smoothstep_aad1db
- OpStore %42 %44 None
- %45 = OpLoad %VertexOutput %out None
- OpReturnValue %45
+%vertex_main_inner = OpFunction %VertexOutput None %42
+ %43 = OpLabel
+ %out = OpVariable %_ptr_Function_VertexOutput Function %46
+ %47 = OpAccessChain %_ptr_Function_v4float %out %uint_0
+ OpStore %47 %51 None
+ %52 = OpAccessChain %_ptr_Function_v3float %out %uint_1
+ %54 = OpFunctionCall %v3float %smoothstep_aad1db
+ OpStore %52 %54 None
+ %55 = OpLoad %VertexOutput %out None
+ OpReturnValue %55
OpFunctionEnd
-%vertex_main = OpFunction %void None %48
- %49 = OpLabel
- %50 = OpFunctionCall %VertexOutput %vertex_main_inner
- %51 = OpCompositeExtract %v4float %50 0
- OpStore %vertex_main_position_Output %51 None
- %52 = OpCompositeExtract %v3float %50 1
- OpStore %vertex_main_loc0_Output %52 None
+%vertex_main = OpFunction %void None %58
+ %59 = OpLabel
+ %60 = OpFunctionCall %VertexOutput %vertex_main_inner
+ %61 = OpCompositeExtract %v4float %60 0
+ OpStore %vertex_main_position_Output %61 None
+ %62 = OpCompositeExtract %v3float %60 1
+ OpStore %vertex_main_loc0_Output %62 None
OpStore %vertex_main___point_size_Output %float_1 None
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.glsl b/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.glsl
index d8ca294..b0bf8bb 100644
--- a/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.glsl
@@ -14,7 +14,9 @@
f16vec4 arg_0 = f16vec4(2.0hf);
f16vec4 arg_1 = f16vec4(4.0hf);
f16vec4 arg_2 = f16vec4(3.0hf);
- f16vec4 res = smoothstep(arg_0, arg_1, arg_2);
+ f16vec4 v_1 = arg_0;
+ f16vec4 v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), f16vec4(0.0hf), f16vec4(1.0hf));
+ f16vec4 res = (v_2 * (v_2 * (f16vec4(3.0hf) - (f16vec4(2.0hf) * v_2))));
return res;
}
void main() {
@@ -34,7 +36,9 @@
f16vec4 arg_0 = f16vec4(2.0hf);
f16vec4 arg_1 = f16vec4(4.0hf);
f16vec4 arg_2 = f16vec4(3.0hf);
- f16vec4 res = smoothstep(arg_0, arg_1, arg_2);
+ f16vec4 v_1 = arg_0;
+ f16vec4 v_2 = clamp(((arg_2 - v_1) / (arg_1 - v_1)), f16vec4(0.0hf), f16vec4(1.0hf));
+ f16vec4 res = (v_2 * (v_2 * (f16vec4(3.0hf) - (f16vec4(2.0hf) * v_2))));
return res;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -58,7 +62,9 @@
f16vec4 arg_0 = f16vec4(2.0hf);
f16vec4 arg_1 = f16vec4(4.0hf);
f16vec4 arg_2 = f16vec4(3.0hf);
- f16vec4 res = smoothstep(arg_0, arg_1, arg_2);
+ f16vec4 v = arg_0;
+ f16vec4 v_1 = clamp(((arg_2 - v) / (arg_1 - v)), f16vec4(0.0hf), f16vec4(1.0hf));
+ f16vec4 res = (v_1 * (v_1 * (f16vec4(3.0hf) - (f16vec4(2.0hf) * v_1))));
return res;
}
VertexOutput vertex_main_inner() {
@@ -68,10 +74,10 @@
return tint_symbol;
}
void main() {
- VertexOutput v = vertex_main_inner();
- gl_Position = v.pos;
+ VertexOutput v_2 = vertex_main_inner();
+ gl_Position = v_2.pos;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
- vertex_main_loc0_Output = v.prevent_dce;
+ vertex_main_loc0_Output = v_2.prevent_dce;
gl_PointSize = 1.0f;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.ir.dxc.hlsl
index 92da0c9..2ff8660 100644
--- a/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.ir.dxc.hlsl
@@ -7,7 +7,9 @@
vector<float16_t, 4> arg_0 = (float16_t(2.0h)).xxxx;
vector<float16_t, 4> arg_1 = (float16_t(4.0h)).xxxx;
vector<float16_t, 4> arg_2 = (float16_t(3.0h)).xxxx;
- vector<float16_t, 4> res = smoothstep(arg_0, arg_1, arg_2);
+ vector<float16_t, 4> v = arg_0;
+ vector<float16_t, 4> v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (float16_t(0.0h)).xxxx, (float16_t(1.0h)).xxxx);
+ vector<float16_t, 4> res = (v_1 * (v_1 * ((float16_t(3.0h)).xxxx - ((float16_t(2.0h)).xxxx * v_1))));
return res;
}
@@ -24,7 +26,9 @@
vector<float16_t, 4> arg_0 = (float16_t(2.0h)).xxxx;
vector<float16_t, 4> arg_1 = (float16_t(4.0h)).xxxx;
vector<float16_t, 4> arg_2 = (float16_t(3.0h)).xxxx;
- vector<float16_t, 4> res = smoothstep(arg_0, arg_1, arg_2);
+ vector<float16_t, 4> v = arg_0;
+ vector<float16_t, 4> v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (float16_t(0.0h)).xxxx, (float16_t(1.0h)).xxxx);
+ vector<float16_t, 4> res = (v_1 * (v_1 * ((float16_t(3.0h)).xxxx - ((float16_t(2.0h)).xxxx * v_1))));
return res;
}
@@ -51,7 +55,9 @@
vector<float16_t, 4> arg_0 = (float16_t(2.0h)).xxxx;
vector<float16_t, 4> arg_1 = (float16_t(4.0h)).xxxx;
vector<float16_t, 4> arg_2 = (float16_t(3.0h)).xxxx;
- vector<float16_t, 4> res = smoothstep(arg_0, arg_1, arg_2);
+ vector<float16_t, 4> v = arg_0;
+ vector<float16_t, 4> v_1 = clamp(((arg_2 - v) / (arg_1 - v)), (float16_t(0.0h)).xxxx, (float16_t(1.0h)).xxxx);
+ vector<float16_t, 4> res = (v_1 * (v_1 * ((float16_t(3.0h)).xxxx - ((float16_t(2.0h)).xxxx * v_1))));
return res;
}
@@ -59,13 +65,13 @@
VertexOutput tint_symbol = (VertexOutput)0;
tint_symbol.pos = (0.0f).xxxx;
tint_symbol.prevent_dce = smoothstep_c43ebd();
- VertexOutput v = tint_symbol;
- return v;
+ VertexOutput v_2 = tint_symbol;
+ return v_2;
}
vertex_main_outputs vertex_main() {
- VertexOutput v_1 = vertex_main_inner();
- vertex_main_outputs v_2 = {v_1.prevent_dce, v_1.pos};
- return v_2;
+ VertexOutput v_3 = vertex_main_inner();
+ vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos};
+ return v_4;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.ir.msl
index 766d3e5..5291a11 100644
--- a/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.ir.msl
@@ -12,7 +12,9 @@
half4 arg_0 = half4(2.0h);
half4 arg_1 = half4(4.0h);
half4 arg_2 = half4(3.0h);
- half4 res = smoothstep(arg_0, arg_1, arg_2);
+ half4 const v = arg_0;
+ half4 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), half4(0.0h), half4(1.0h));
+ half4 res = (v_1 * (v_1 * (half4(3.0h) - (half4(2.0h) * v_1))));
return res;
}
@@ -34,7 +36,9 @@
half4 arg_0 = half4(2.0h);
half4 arg_1 = half4(4.0h);
half4 arg_2 = half4(3.0h);
- half4 res = smoothstep(arg_0, arg_1, arg_2);
+ half4 const v = arg_0;
+ half4 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), half4(0.0h), half4(1.0h));
+ half4 res = (v_1 * (v_1 * (half4(3.0h) - (half4(2.0h) * v_1))));
return res;
}
@@ -62,7 +66,9 @@
half4 arg_0 = half4(2.0h);
half4 arg_1 = half4(4.0h);
half4 arg_2 = half4(3.0h);
- half4 res = smoothstep(arg_0, arg_1, arg_2);
+ half4 const v = arg_0;
+ half4 const v_1 = clamp(((arg_2 - v) / (arg_1 - v)), half4(0.0h), half4(1.0h));
+ half4 res = (v_1 * (v_1 * (half4(3.0h) - (half4(2.0h) * v_1))));
return res;
}
@@ -74,9 +80,9 @@
}
vertex vertex_main_outputs vertex_main() {
- VertexOutput const v = vertex_main_inner();
+ VertexOutput const v_2 = vertex_main_inner();
vertex_main_outputs tint_wrapper_result = {};
- tint_wrapper_result.VertexOutput_pos = v.pos;
- tint_wrapper_result.VertexOutput_prevent_dce = v.prevent_dce;
+ tint_wrapper_result.VertexOutput_pos = v_2.pos;
+ tint_wrapper_result.VertexOutput_prevent_dce = v_2.prevent_dce;
return tint_wrapper_result;
}
diff --git a/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.spvasm b/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.spvasm
index 6225a3c..90637e4 100644
--- a/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/smoothstep/c43ebd.wgsl.expected.spvasm
@@ -4,13 +4,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 35
+; Bound: 45
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
- %23 = OpExtInstImport "GLSL.std.450"
+ %26 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %fragment_main "fragment_main"
OpExecutionMode %fragment_main OriginUpperLeft
@@ -40,8 +40,11 @@
%14 = OpConstantComposite %v4half %half_0x1p_2 %half_0x1p_2 %half_0x1p_2 %half_0x1p_2
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
%17 = OpConstantComposite %v4half %half_0x1_8p_1 %half_0x1_8p_1 %half_0x1_8p_1 %half_0x1_8p_1
+ %27 = OpConstantNull %v4half
+%half_0x1p_0 = OpConstant %half 0x1p+0
+ %28 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
%void = OpTypeVoid
- %28 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_StorageBuffer_v4half = OpTypePointer StorageBuffer %v4half
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -57,16 +60,23 @@
%19 = OpLoad %v4half %arg_0 None
%20 = OpLoad %v4half %arg_1 None
%21 = OpLoad %v4half %arg_2 None
- %22 = OpExtInst %v4half %23 SmoothStep %19 %20 %21
- OpStore %res %22
- %25 = OpLoad %v4half %res None
- OpReturnValue %25
+ %22 = OpFSub %v4half %21 %19
+ %23 = OpFSub %v4half %20 %19
+ %24 = OpFDiv %v4half %22 %23
+ %25 = OpExtInst %v4half %26 NClamp %24 %27 %28
+ %30 = OpFMul %v4half %11 %25
+ %31 = OpFSub %v4half %17 %30
+ %32 = OpFMul %v4half %25 %31
+ %33 = OpFMul %v4half %25 %32
+ OpStore %res %33
+ %35 = OpLoad %v4half %res None
+ OpReturnValue %35
OpFunctionEnd
-%fragment_main = OpFunction %void None %28
- %29 = OpLabel
- %30 = OpFunctionCall %v4half %smoothstep_c43ebd
- %31 = OpAccessChain %_ptr_StorageBuffer_v4half %1 %uint_0
- OpStore %31 %30 None
+%fragment_main = OpFunction %void None %38
+ %39 = OpLabel
+ %40 = OpFunctionCall %v4half %smoothstep_c43ebd
+ %41 = OpAccessChain %_ptr_StorageBuffer_v4half %1 %uint_0
+ OpStore %41 %40 None
OpReturn
OpFunctionEnd
;
@@ -75,13 +85,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 35
+; Bound: 45
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
- %23 = OpExtInstImport "GLSL.std.450"
+ %26 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %compute_main LocalSize 1 1 1
@@ -111,8 +121,11 @@
%14 = OpConstantComposite %v4half %half_0x1p_2 %half_0x1p_2 %half_0x1p_2 %half_0x1p_2
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
%17 = OpConstantComposite %v4half %half_0x1_8p_1 %half_0x1_8p_1 %half_0x1_8p_1 %half_0x1_8p_1
+ %27 = OpConstantNull %v4half
+%half_0x1p_0 = OpConstant %half 0x1p+0
+ %28 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
%void = OpTypeVoid
- %28 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_StorageBuffer_v4half = OpTypePointer StorageBuffer %v4half
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
@@ -128,16 +141,23 @@
%19 = OpLoad %v4half %arg_0 None
%20 = OpLoad %v4half %arg_1 None
%21 = OpLoad %v4half %arg_2 None
- %22 = OpExtInst %v4half %23 SmoothStep %19 %20 %21
- OpStore %res %22
- %25 = OpLoad %v4half %res None
- OpReturnValue %25
+ %22 = OpFSub %v4half %21 %19
+ %23 = OpFSub %v4half %20 %19
+ %24 = OpFDiv %v4half %22 %23
+ %25 = OpExtInst %v4half %26 NClamp %24 %27 %28
+ %30 = OpFMul %v4half %11 %25
+ %31 = OpFSub %v4half %17 %30
+ %32 = OpFMul %v4half %25 %31
+ %33 = OpFMul %v4half %25 %32
+ OpStore %res %33
+ %35 = OpLoad %v4half %res None
+ OpReturnValue %35
OpFunctionEnd
-%compute_main = OpFunction %void None %28
- %29 = OpLabel
- %30 = OpFunctionCall %v4half %smoothstep_c43ebd
- %31 = OpAccessChain %_ptr_StorageBuffer_v4half %1 %uint_0
- OpStore %31 %30 None
+%compute_main = OpFunction %void None %38
+ %39 = OpLabel
+ %40 = OpFunctionCall %v4half %smoothstep_c43ebd
+ %41 = OpAccessChain %_ptr_StorageBuffer_v4half %1 %uint_0
+ OpStore %41 %40 None
OpReturn
OpFunctionEnd
;
@@ -146,14 +166,14 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
-; Bound: 55
+; Bound: 65
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
- %28 = OpExtInstImport "GLSL.std.450"
+ %31 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %vertex_main_position_Output %vertex_main_loc0_Output %vertex_main___point_size_Output
OpName %vertex_main_position_Output "vertex_main_position_Output"
@@ -194,17 +214,20 @@
%19 = OpConstantComposite %v4half %half_0x1p_2 %half_0x1p_2 %half_0x1p_2 %half_0x1p_2
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
%22 = OpConstantComposite %v4half %half_0x1_8p_1 %half_0x1_8p_1 %half_0x1_8p_1 %half_0x1_8p_1
+ %32 = OpConstantNull %v4half
+%half_0x1p_0 = OpConstant %half 0x1p+0
+ %33 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
%VertexOutput = OpTypeStruct %v4float %v4half
- %33 = OpTypeFunction %VertexOutput
+ %43 = OpTypeFunction %VertexOutput
%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
- %37 = OpConstantNull %VertexOutput
+ %47 = OpConstantNull %VertexOutput
%_ptr_Function_v4float = OpTypePointer Function %v4float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
- %42 = OpConstantNull %v4float
+ %52 = OpConstantNull %v4float
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
- %49 = OpTypeFunction %void
+ %59 = OpTypeFunction %void
%float_1 = OpConstant %float 1
%smoothstep_c43ebd = OpFunction %v4half None %12
%13 = OpLabel
@@ -218,29 +241,36 @@
%24 = OpLoad %v4half %arg_0 None
%25 = OpLoad %v4half %arg_1 None
%26 = OpLoad %v4half %arg_2 None
- %27 = OpExtInst %v4half %28 SmoothStep %24 %25 %26
- OpStore %res %27
- %30 = OpLoad %v4half %res None
- OpReturnValue %30
+ %27 = OpFSub %v4half %26 %24
+ %28 = OpFSub %v4half %25 %24
+ %29 = OpFDiv %v4half %27 %28
+ %30 = OpExtInst %v4half %31 NClamp %29 %32 %33
+ %35 = OpFMul %v4half %16 %30
+ %36 = OpFSub %v4half %22 %35
+ %37 = OpFMul %v4half %30 %36
+ %38 = OpFMul %v4half %30 %37
+ OpStore %res %38
+ %40 = OpLoad %v4half %res None
+ OpReturnValue %40
OpFunctionEnd
-%vertex_main_inner = OpFunction %VertexOutput None %33
- %34 = OpLabel
- %out = OpVariable %_ptr_Function_VertexOutput Function %37
- %38 = OpAccessChain %_ptr_Function_v4float %out %uint_0
- OpStore %38 %42 None
- %43 = OpAccessChain %_ptr_Function_v4half %out %uint_1
- %45 = OpFunctionCall %v4half %smoothstep_c43ebd
- OpStore %43 %45 None
- %46 = OpLoad %VertexOutput %out None
- OpReturnValue %46
+%vertex_main_inner = OpFunction %VertexOutput None %43
+ %44 = OpLabel
+ %out = OpVariable %_ptr_Function_VertexOutput Function %47
+ %48 = OpAccessChain %_ptr_Function_v4float %out %uint_0
+ OpStore %48 %52 None
+ %53 = OpAccessChain %_ptr_Function_v4half %out %uint_1
+ %55 = OpFunctionCall %v4half %smoothstep_c43ebd
+ OpStore %53 %55 None
+ %56 = OpLoad %VertexOutput %out None
+ OpReturnValue %56
OpFunctionEnd
-%vertex_main = OpFunction %void None %49
- %50 = OpLabel
- %51 = OpFunctionCall %VertexOutput %vertex_main_inner
- %52 = OpCompositeExtract %v4float %51 0
- OpStore %vertex_main_position_Output %52 None
- %53 = OpCompositeExtract %v4half %51 1
- OpStore %vertex_main_loc0_Output %53 None
+%vertex_main = OpFunction %void None %59
+ %60 = OpLabel
+ %61 = OpFunctionCall %VertexOutput %vertex_main_inner
+ %62 = OpCompositeExtract %v4float %61 0
+ OpStore %vertex_main_position_Output %62 None
+ %63 = OpCompositeExtract %v4half %61 1
+ OpStore %vertex_main_loc0_Output %63 None
OpStore %vertex_main___point_size_Output %float_1 None
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/smoothstep.wgsl b/test/tint/builtins/smoothstep.wgsl
new file mode 100644
index 0000000..566ac17
--- /dev/null
+++ b/test/tint/builtins/smoothstep.wgsl
@@ -0,0 +1,10 @@
+// Verifies that smoothstep with (effective) constant
+// values of low > high compiles on all backends
+
+@compute @workgroup_size(1)
+fn main() {
+ let low = 1.0;
+ let high = 0.0;
+ let x_val = 0.5;
+ let res = smoothstep(low,high, x_val);
+}
diff --git a/test/tint/builtins/smoothstep.wgsl.expected.dxc.hlsl b/test/tint/builtins/smoothstep.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..05337ff
--- /dev/null
+++ b/test/tint/builtins/smoothstep.wgsl.expected.dxc.hlsl
@@ -0,0 +1,8 @@
+[numthreads(1, 1, 1)]
+void main() {
+ float low = 1.0f;
+ float high = 0.0f;
+ float x_val = 0.5f;
+ float res = smoothstep(low, high, x_val);
+ return;
+}
diff --git a/test/tint/builtins/smoothstep.wgsl.expected.fxc.hlsl b/test/tint/builtins/smoothstep.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..05337ff
--- /dev/null
+++ b/test/tint/builtins/smoothstep.wgsl.expected.fxc.hlsl
@@ -0,0 +1,8 @@
+[numthreads(1, 1, 1)]
+void main() {
+ float low = 1.0f;
+ float high = 0.0f;
+ float x_val = 0.5f;
+ float res = smoothstep(low, high, x_val);
+ return;
+}
diff --git a/test/tint/builtins/smoothstep.wgsl.expected.glsl b/test/tint/builtins/smoothstep.wgsl.expected.glsl
new file mode 100644
index 0000000..3380ee6
--- /dev/null
+++ b/test/tint/builtins/smoothstep.wgsl.expected.glsl
@@ -0,0 +1,9 @@
+#version 310 es
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ float low = 1.0f;
+ float high = 0.0f;
+ float x_val = 0.5f;
+ float res = (clamp(((x_val - low) / (high - low)), 0.0f, 1.0f) * (clamp(((x_val - low) / (high - low)), 0.0f, 1.0f) * (3.0f - (2.0f * clamp(((x_val - low) / (high - low)), 0.0f, 1.0f)))));
+}
diff --git a/test/tint/builtins/smoothstep.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/smoothstep.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..5a08ab3
--- /dev/null
+++ b/test/tint/builtins/smoothstep.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,9 @@
+
+[numthreads(1, 1, 1)]
+void main() {
+ float low = 1.0f;
+ float high = 0.0f;
+ float x_val = 0.5f;
+ float res = (clamp(((x_val - low) / (high - low)), 0.0f, 1.0f) * (clamp(((x_val - low) / (high - low)), 0.0f, 1.0f) * (3.0f - (2.0f * clamp(((x_val - low) / (high - low)), 0.0f, 1.0f)))));
+}
+
diff --git a/test/tint/builtins/smoothstep.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/smoothstep.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..5a08ab3
--- /dev/null
+++ b/test/tint/builtins/smoothstep.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,9 @@
+
+[numthreads(1, 1, 1)]
+void main() {
+ float low = 1.0f;
+ float high = 0.0f;
+ float x_val = 0.5f;
+ float res = (clamp(((x_val - low) / (high - low)), 0.0f, 1.0f) * (clamp(((x_val - low) / (high - low)), 0.0f, 1.0f) * (3.0f - (2.0f * clamp(((x_val - low) / (high - low)), 0.0f, 1.0f)))));
+}
+
diff --git a/test/tint/builtins/smoothstep.wgsl.expected.ir.msl b/test/tint/builtins/smoothstep.wgsl.expected.ir.msl
new file mode 100644
index 0000000..6ca00e9
--- /dev/null
+++ b/test/tint/builtins/smoothstep.wgsl.expected.ir.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+using namespace metal;
+
+kernel void tint_symbol() {
+ float const low = 1.0f;
+ float const high = 0.0f;
+ float const x_val = 0.5f;
+ float const res = (clamp(((x_val - low) / (high - low)), 0.0f, 1.0f) * (clamp(((x_val - low) / (high - low)), 0.0f, 1.0f) * (3.0f - (2.0f * clamp(((x_val - low) / (high - low)), 0.0f, 1.0f)))));
+}
diff --git a/test/tint/builtins/smoothstep.wgsl.expected.msl b/test/tint/builtins/smoothstep.wgsl.expected.msl
new file mode 100644
index 0000000..c66e2ea
--- /dev/null
+++ b/test/tint/builtins/smoothstep.wgsl.expected.msl
@@ -0,0 +1,11 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+ float const low = 1.0f;
+ float const high = 0.0f;
+ float const x_val = 0.5f;
+ float const res = smoothstep(low, high, x_val);
+ return;
+}
+
diff --git a/test/tint/builtins/smoothstep.wgsl.expected.spvasm b/test/tint/builtins/smoothstep.wgsl.expected.spvasm
new file mode 100644
index 0000000..1652c5c
--- /dev/null
+++ b/test/tint/builtins/smoothstep.wgsl.expected.spvasm
@@ -0,0 +1,35 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 20
+; Schema: 0
+ OpCapability Shader
+ %13 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %low "low"
+ OpName %high "high"
+ OpName %x_val "x_val"
+ OpName %res "res"
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+ %low = OpConstant %float 1
+ %high = OpConstant %float 0
+ %x_val = OpConstant %float 0.5
+ %float_2 = OpConstant %float 2
+ %float_3 = OpConstant %float 3
+ %main = OpFunction %void None %3
+ %4 = OpLabel
+ %9 = OpFSub %float %x_val %low
+ %10 = OpFSub %float %high %low
+ %11 = OpFDiv %float %9 %10
+ %12 = OpExtInst %float %13 NClamp %11 %high %low
+ %14 = OpFMul %float %float_2 %12
+ %16 = OpFSub %float %float_3 %14
+ %18 = OpFMul %float %12 %16
+ %res = OpFMul %float %12 %18
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/smoothstep.wgsl.expected.wgsl b/test/tint/builtins/smoothstep.wgsl.expected.wgsl
new file mode 100644
index 0000000..a730121
--- /dev/null
+++ b/test/tint/builtins/smoothstep.wgsl.expected.wgsl
@@ -0,0 +1,7 @@
+@compute @workgroup_size(1)
+fn main() {
+ let low = 1.0;
+ let high = 0.0;
+ let x_val = 0.5;
+ let res = smoothstep(low, high, x_val);
+}
diff --git a/webgpu-cts/compat-expectations.txt b/webgpu-cts/compat-expectations.txt
index aefe74f..994b607 100644
--- a/webgpu-cts/compat-expectations.txt
+++ b/webgpu-cts/compat-expectations.txt
@@ -2067,6 +2067,9 @@
### This section represents tests which may require CTS changes.
+# Temporary Failures change in behavior of Smoothstep function
+crbug.com/379909620 webgpu:shader,validation,expression,call,builtin,smoothstep:* [ Failure ]
+
# Failures due to change in `@align()` validation.
crbug.com/375467276 webgpu:shader,execution,expression,access,structure,index:buffer_align:* [ Failure ]
diff --git a/webgpu-cts/expectations.txt b/webgpu-cts/expectations.txt
index 74a6489..f1f5579 100644
--- a/webgpu-cts/expectations.txt
+++ b/webgpu-cts/expectations.txt
@@ -1524,6 +1524,9 @@
crbug.com/dawn/366000875 webgpu:shader,validation,decl,var:var_access_mode_bad_other_template_contents:accessMode="read";prefix="storage,";suffix="," [ Failure ]
crbug.com/dawn/366000875 webgpu:shader,validation,decl,var:var_access_mode_bad_other_template_contents:accessMode="read_write";prefix="storage,";suffix="," [ Failure ]
+# Temporary Failures change in behavior of Smoothstep function
+crbug.com/379909620 webgpu:shader,validation,expression,call,builtin,smoothstep:* [ Failure ]
+
# Failures due to change in `@align()` validation.
crbug.com/375467276 webgpu:shader,execution,expression,access,structure,index:buffer_align:* [ Failure ]
crbug.com/375467276 webgpu:shader,validation,shader_io,align:* [ Failure ]
