tint: Add abstract overload of modf

Required fleshing out resolver handling of abstract structures.

Bug: tint:1768

Change-Id: Icb2983ac9bf37d4e260735cfdba5ac5e7dfe3ee2
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/111400
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
Ben Clayton 2022-11-23 15:00:20 +00:00 committed by Dawn LUCI CQ
parent cd9873f8f0
commit 31b7fca82c
76 changed files with 711 additions and 175 deletions

View File

@ -504,8 +504,8 @@ fn log2<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
fn mix<T: f32_f16>(T, T, T) -> T
fn mix<N: num, T: f32_f16>(vec<N, T>, vec<N, T>, vec<N, T>) -> vec<N, T>
fn mix<N: num, T: f32_f16>(vec<N, T>, vec<N, T>, T) -> vec<N, T>
@const fn modf<T: f32_f16>(@test_value(-1.5) T) -> __modf_result<T>
@const fn modf<N: num, T: f32_f16>(@test_value(-1.5) vec<N, T>) -> __modf_result_vec<N, T>
@const fn modf<T: fa_f32_f16>(@test_value(-1.5) T) -> __modf_result<T>
@const fn modf<N: num, T: fa_f32_f16>(@test_value(-1.5) vec<N, T>) -> __modf_result_vec<N, T>
fn normalize<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
@const fn pack2x16float(vec2<f32>) -> u32
@const fn pack2x16snorm(vec2<f32>) -> u32

View File

@ -1374,8 +1374,8 @@ TEST_F(ResolverBuiltinFloatTest, Modf_Error_FirstParamInt) {
R"(error: no matching call to modf(i32, ptr<workgroup, f32, read_write>)
2 candidate functions:
modf(T) -> __modf_result_T where: T is f32 or f16
modf(vecN<T>) -> __modf_result_vecN_T where: T is f32 or f16
modf(T) -> __modf_result_T where: T is abstract-float, f32 or f16
modf(vecN<T>) -> __modf_result_vecN_T where: T is abstract-float, f32 or f16
)");
}
@ -1390,8 +1390,8 @@ TEST_F(ResolverBuiltinFloatTest, Modf_Error_SecondParamIntPtr) {
R"(error: no matching call to modf(f32, ptr<workgroup, i32, read_write>)
2 candidate functions:
modf(T) -> __modf_result_T where: T is f32 or f16
modf(vecN<T>) -> __modf_result_vecN_T where: T is f32 or f16
modf(T) -> __modf_result_T where: T is abstract-float, f32 or f16
modf(vecN<T>) -> __modf_result_vecN_T where: T is abstract-float, f32 or f16
)");
}
@ -1404,8 +1404,8 @@ TEST_F(ResolverBuiltinFloatTest, Modf_Error_SecondParamNotAPointer) {
EXPECT_EQ(r()->error(), R"(error: no matching call to modf(f32, f32)
2 candidate functions:
modf(T) -> __modf_result_T where: T is f32 or f16
modf(vecN<T>) -> __modf_result_vecN_T where: T is f32 or f16
modf(T) -> __modf_result_T where: T is abstract-float, f32 or f16
modf(vecN<T>) -> __modf_result_vecN_T where: T is abstract-float, f32 or f16
)");
}
@ -1420,8 +1420,8 @@ TEST_F(ResolverBuiltinFloatTest, Modf_Error_VectorSizesDontMatch) {
R"(error: no matching call to modf(vec2<f32>, ptr<workgroup, vec4<f32>, read_write>)
2 candidate functions:
modf(T) -> __modf_result_T where: T is f32 or f16
modf(vecN<T>) -> __modf_result_vecN_T where: T is f32 or f16
modf(T) -> __modf_result_T where: T is abstract-float, f32 or f16
modf(vecN<T>) -> __modf_result_vecN_T where: T is abstract-float, f32 or f16
)");
}

View File

@ -403,13 +403,26 @@ struct Composite : ImplConstant {
const sem::Type* target_ty,
const Source& source) const override {
// Convert each of the composite element types.
auto* el_ty = sem::Type::ElementOf(target_ty);
utils::Vector<const sem::Constant*, 4> conv_els;
conv_els.Reserve(elements.Length());
std::function<const sem::Type*(size_t idx)> target_el_ty;
if (auto* str = target_ty->As<sem::Struct>()) {
if (str->Members().size() != elements.Length()) {
TINT_ICE(Resolver, builder.Diagnostics())
<< "const-eval conversion of structure has mismatched element counts";
return utils::Failure;
}
target_el_ty = [str](size_t idx) { return str->Members()[idx]->Type(); };
} else {
auto* el_ty = sem::Type::ElementOf(target_ty);
target_el_ty = [el_ty](size_t) { return el_ty; };
}
for (auto* el : elements) {
// Note: This file is the only place where `sem::Constant`s are created, so this
// Note: This file is the only place where `sem::Constant`s are created, so the
// static_cast is safe.
auto conv_el = static_cast<const ImplConstant*>(el)->Convert(builder, el_ty, source);
auto conv_el = static_cast<const ImplConstant*>(el)->Convert(
builder, target_el_ty(conv_els.Length()), source);
if (!conv_el) {
return utils::Failure;
}

View File

@ -1339,7 +1339,8 @@ INSTANTIATE_TEST_SUITE_P( //
Modf,
ResolverConstEvalBuiltinTest,
testing::Combine(testing::Values(sem::BuiltinType::kModf),
testing::ValuesIn(Concat(ModfCases<f32>(), //
testing::ValuesIn(Concat(ModfCases<AFloat>(), //
ModfCases<f32>(), //
ModfCases<f16>()))));
std::vector<Case> Pack4x8snormCases() {

View File

@ -796,20 +796,20 @@ struct NameAndType {
std::string name;
const sem::Type* type;
};
const sem::Struct* build_struct(MatchState& state,
sem::Struct* build_struct(ProgramBuilder& b,
std::string name,
std::initializer_list<NameAndType> member_names_and_types) {
uint32_t offset = 0;
uint32_t max_align = 0;
sem::StructMemberList members;
for (auto& m : member_names_and_types) {
uint32_t align = m.type->Align();
uint32_t align = std::max<uint32_t>(m.type->Align(), 1);
uint32_t size = m.type->Size();
offset = utils::RoundUp(align, offset);
max_align = std::max(max_align, align);
members.emplace_back(state.builder.create<sem::StructMember>(
members.emplace_back(b.create<sem::StructMember>(
/* declaration */ nullptr,
/* name */ state.builder.Sym(m.name),
/* name */ b.Sym(m.name),
/* type */ m.type,
/* index */ static_cast<uint32_t>(members.size()),
/* offset */ offset,
@ -820,9 +820,9 @@ const sem::Struct* build_struct(MatchState& state,
}
uint32_t size_without_padding = offset;
uint32_t size_with_padding = utils::RoundUp(max_align, offset);
return state.builder.create<sem::Struct>(
return b.create<sem::Struct>(
/* declaration */ nullptr,
/* name */ state.builder.Sym(name),
/* name */ b.Sym(name),
/* members */ members,
/* align */ max_align,
/* size */ size_with_padding,
@ -830,23 +830,58 @@ const sem::Struct* build_struct(MatchState& state,
}
const sem::Struct* build_modf_result(MatchState& state, const sem::Type* el) {
std::string display_name;
if (el->Is<sem::F16>()) {
display_name = "__modf_result_f16";
} else {
display_name = "__modf_result";
}
return build_struct(state, display_name, {{"fract", el}, {"whole", el}});
auto build_f32 = [&] {
auto* ty = state.builder.create<sem::F32>();
return build_struct(state.builder, "__modf_result", {{"fract", ty}, {"whole", ty}});
};
auto build_f16 = [&] {
auto* ty = state.builder.create<sem::F16>();
return build_struct(state.builder, "__modf_result_f16", {{"fract", ty}, {"whole", ty}});
};
return Switch(
el, //
[&](const sem::F32*) { return build_f32(); }, //
[&](const sem::F16*) { return build_f16(); }, //
[&](const sem::AbstractFloat*) {
auto* abstract = build_struct(state.builder, "__modf_result_abstract",
{{"fract", el}, {"whole", el}});
abstract->SetConcreteTypes(utils::Vector{build_f32(), build_f16()});
return abstract;
},
[&](Default) {
TINT_ICE(Resolver, state.builder.Diagnostics())
<< "unhandled modf type: " << state.builder.FriendlyName(el);
return nullptr;
});
}
const sem::Struct* build_modf_result_vec(MatchState& state, Number& n, const sem::Type* el) {
std::string display_name;
if (el->Is<sem::F16>()) {
display_name = "__modf_result_vec" + std::to_string(n.Value()) + "_f16";
} else {
display_name = "__modf_result_vec" + std::to_string(n.Value());
}
auto prefix = "__modf_result_vec" + std::to_string(n.Value());
auto build_f32 = [&] {
auto* vec = state.builder.create<sem::Vector>(state.builder.create<sem::F32>(), n.Value());
return build_struct(state.builder, prefix, {{"fract", vec}, {"whole", vec}});
};
auto build_f16 = [&] {
auto* vec = state.builder.create<sem::Vector>(state.builder.create<sem::F16>(), n.Value());
return build_struct(state.builder, prefix + "_f16", {{"fract", vec}, {"whole", vec}});
};
return Switch(
el, //
[&](const sem::F32*) { return build_f32(); }, //
[&](const sem::F16*) { return build_f16(); }, //
[&](const sem::AbstractFloat*) {
auto* vec = state.builder.create<sem::Vector>(el, n.Value());
return build_struct(state, display_name, {{"fract", vec}, {"whole", vec}});
auto* abstract =
build_struct(state.builder, prefix + "_abstract", {{"fract", vec}, {"whole", vec}});
abstract->SetConcreteTypes(utils::Vector{build_f32(), build_f16()});
return abstract;
},
[&](Default) {
TINT_ICE(Resolver, state.builder.Diagnostics())
<< "unhandled modf type: " << state.builder.FriendlyName(el);
return nullptr;
});
}
const sem::Struct* build_frexp_result(MatchState& state, const sem::Type* el) {
std::string display_name;
@ -856,7 +891,7 @@ const sem::Struct* build_frexp_result(MatchState& state, const sem::Type* el) {
display_name = "__frexp_result";
}
auto* i32 = state.builder.create<sem::I32>();
return build_struct(state, display_name, {{"fract", el}, {"exp", i32}});
return build_struct(state.builder, display_name, {{"fract", el}, {"exp", i32}});
}
const sem::Struct* build_frexp_result_vec(MatchState& state, Number& n, const sem::Type* el) {
std::string display_name;
@ -867,11 +902,12 @@ const sem::Struct* build_frexp_result_vec(MatchState& state, Number& n, const se
}
auto* vec = state.builder.create<sem::Vector>(el, n.Value());
auto* vec_i32 = state.builder.create<sem::Vector>(state.builder.create<sem::I32>(), n.Value());
return build_struct(state, display_name, {{"fract", vec}, {"exp", vec_i32}});
return build_struct(state.builder, display_name, {{"fract", vec}, {"exp", vec_i32}});
}
const sem::Struct* build_atomic_compare_exchange_result(MatchState& state, const sem::Type* ty) {
return build_struct(
state, "__atomic_compare_exchange_result" + ty->FriendlyName(state.builder.Symbols()),
state.builder,
"__atomic_compare_exchange_result" + ty->FriendlyName(state.builder.Symbols()),
{{"old_value", const_cast<sem::Type*>(ty)},
{"exchanged", state.builder.create<sem::Bool>()}});
}

View File

@ -148,7 +148,6 @@ class IntrinsicTable {
/// `sem::EvaluationStage::kRuntime`, then only overloads with concrete argument types
/// will be considered, as all abstract-numerics will have been materialized
/// after shader creation time (sem::EvaluationStage::kConstant).
/// @param source the source of the call
/// @return a sem::TypeInitializer, sem::TypeConversion or nullptr if nothing matched
virtual InitOrConv Lookup(InitConvIntrinsic type,

View File

@ -12846,7 +12846,7 @@ constexpr OverloadInfo kOverloads[] = {
/* num parameters */ 1,
/* num template types */ 1,
/* num template numbers */ 0,
/* template types */ &kTemplateTypes[26],
/* template types */ &kTemplateTypes[23],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[878],
/* return matcher indices */ &kMatcherIndices[106],
@ -12858,7 +12858,7 @@ constexpr OverloadInfo kOverloads[] = {
/* num parameters */ 1,
/* num template types */ 1,
/* num template numbers */ 1,
/* template types */ &kTemplateTypes[26],
/* template types */ &kTemplateTypes[23],
/* template numbers */ &kTemplateNumbers[4],
/* parameters */ &kParameters[879],
/* return matcher indices */ &kMatcherIndices[45],
@ -14351,8 +14351,8 @@ constexpr IntrinsicInfo kBuiltins[] = {
},
{
/* [53] */
/* fn modf<T : f32_f16>(@test_value(-1.5) T) -> __modf_result<T> */
/* fn modf<N : num, T : f32_f16>(@test_value(-1.5) vec<N, T>) -> __modf_result_vec<N, T> */
/* fn modf<T : fa_f32_f16>(@test_value(-1.5) T) -> __modf_result<T> */
/* fn modf<N : num, T : fa_f32_f16>(@test_value(-1.5) vec<N, T>) -> __modf_result_vec<N, T> */
/* num overloads */ 2,
/* overloads */ &kOverloads[377],
},

View File

@ -1236,5 +1236,91 @@ TEST_F(MaterializeAbstractNumericToUnrelatedType, AIntToStructLetInit) {
} // namespace materialize_abstract_numeric_to_unrelated_type
////////////////////////////////////////////////////////////////////////////////
// Materialization tests for builtin-returned abstract structures
// These are too bespoke to slot into the more general materialization tests above
////////////////////////////////////////////////////////////////////////////////
namespace materialize_abstract_structure {
using MaterializeAbstractStructure = resolver::ResolverTest;
TEST_F(MaterializeAbstractStructure, Modf_Scalar_DefaultType) {
// var v = modf(1);
auto* call = Call("modf", 1_a);
WrapInFunction(Decl(Var("v", call)));
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get(call);
ASSERT_TRUE(sem->Is<sem::Materialize>());
auto* materialize = sem->As<sem::Materialize>();
ASSERT_TRUE(materialize->Type()->Is<sem::Struct>());
auto* concrete_str = materialize->Type()->As<sem::Struct>();
ASSERT_TRUE(concrete_str->Members()[0]->Type()->Is<sem::F32>());
ASSERT_TRUE(materialize->Expr()->Type()->Is<sem::Struct>());
auto* abstract_str = materialize->Expr()->Type()->As<sem::Struct>();
ASSERT_TRUE(abstract_str->Members()[0]->Type()->Is<sem::AbstractFloat>());
}
TEST_F(MaterializeAbstractStructure, Modf_Vector_DefaultType) {
// var v = modf(vec2(1));
auto* call = Call("modf", Construct(ty.vec2(nullptr), 1_a));
WrapInFunction(Decl(Var("v", call)));
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get(call);
ASSERT_TRUE(sem->Is<sem::Materialize>());
auto* materialize = sem->As<sem::Materialize>();
ASSERT_TRUE(materialize->Type()->Is<sem::Struct>());
auto* concrete_str = materialize->Type()->As<sem::Struct>();
ASSERT_TRUE(concrete_str->Members()[0]->Type()->Is<sem::Vector>());
ASSERT_TRUE(concrete_str->Members()[0]->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
ASSERT_TRUE(materialize->Expr()->Type()->Is<sem::Struct>());
auto* abstract_str = materialize->Expr()->Type()->As<sem::Struct>();
ASSERT_TRUE(abstract_str->Members()[0]->Type()->Is<sem::Vector>());
ASSERT_TRUE(
abstract_str->Members()[0]->Type()->As<sem::Vector>()->type()->Is<sem::AbstractFloat>());
}
TEST_F(MaterializeAbstractStructure, Modf_Scalar_ExplicitType) {
// var v = modf(1_h); // v is __modf_result_f16
// v = modf(1); // __modf_result_f16 <- __modf_result_abstract
Enable(ast::Extension::kF16);
auto* call = Call("modf", 1_a);
WrapInFunction(Decl(Var("v", Call("modf", 1_h))), //
Assign("v", call));
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get(call);
ASSERT_TRUE(sem->Is<sem::Materialize>());
auto* materialize = sem->As<sem::Materialize>();
ASSERT_TRUE(materialize->Type()->Is<sem::Struct>());
auto* concrete_str = materialize->Type()->As<sem::Struct>();
ASSERT_TRUE(concrete_str->Members()[0]->Type()->Is<sem::F16>());
ASSERT_TRUE(materialize->Expr()->Type()->Is<sem::Struct>());
auto* abstract_str = materialize->Expr()->Type()->As<sem::Struct>();
ASSERT_TRUE(abstract_str->Members()[0]->Type()->Is<sem::AbstractFloat>());
}
TEST_F(MaterializeAbstractStructure, Modf_Vector_ExplicitType) {
// var v = modf(vec2(1_h)); // v is __modf_result_vec2_f16
// v = modf(vec2(1)); // __modf_result_vec2_f16 <- __modf_result_vec2_abstract
Enable(ast::Extension::kF16);
auto* call = Call("modf", Construct(ty.vec2(nullptr), 1_a));
WrapInFunction(Decl(Var("v", Call("modf", Construct(ty.vec2(nullptr), 1_h)))),
Assign("v", call));
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get(call);
ASSERT_TRUE(sem->Is<sem::Materialize>());
auto* materialize = sem->As<sem::Materialize>();
ASSERT_TRUE(materialize->Type()->Is<sem::Struct>());
auto* concrete_str = materialize->Type()->As<sem::Struct>();
ASSERT_TRUE(concrete_str->Members()[0]->Type()->Is<sem::Vector>());
ASSERT_TRUE(concrete_str->Members()[0]->Type()->As<sem::Vector>()->type()->Is<sem::F16>());
ASSERT_TRUE(materialize->Expr()->Type()->Is<sem::Struct>());
auto* abstract_str = materialize->Expr()->Type()->As<sem::Struct>();
ASSERT_TRUE(abstract_str->Members()[0]->Type()->Is<sem::Vector>());
ASSERT_TRUE(
abstract_str->Members()[0]->Type()->As<sem::Vector>()->type()->Is<sem::AbstractFloat>());
}
} // namespace materialize_abstract_structure
} // namespace
} // namespace tint::resolver

View File

@ -1726,6 +1726,12 @@ const sem::Type* Resolver::ConcreteType(const sem::Type* ty,
return Array(source, source, el_ty, a->Count(), /* explicit_stride */ 0);
}
return nullptr;
},
[&](const sem::Struct* s) -> const sem::Type* {
if (auto& tys = s->ConcreteTypes(); !tys.IsEmpty()) {
return target_ty ? target_ty : tys[0];
}
return nullptr;
});
}
@ -1938,8 +1944,8 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
bool has_side_effects =
std::any_of(args.begin(), args.end(), [](auto* e) { return e->HasSideEffects(); });
// ct_init_or_conv is a helper for building either a sem::TypeInitializer or sem::TypeConversion
// call for a InitConvIntrinsic with an optional template argument type.
// ct_init_or_conv is a helper for building either a sem::TypeInitializer or
// sem::TypeConversion call for a InitConvIntrinsic with an optional template argument type.
auto ct_init_or_conv = [&](InitConvIntrinsic ty, const sem::Type* template_arg) -> sem::Call* {
auto arg_tys = utils::Transform(args, [](auto* arg) { return arg->Type(); });
auto ctor_or_conv =
@ -1987,8 +1993,8 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
if (!value) {
// Constant evaluation failed.
// Can happen for expressions that will fail validation (later).
// Use the kRuntime EvaluationStage, as kConstant will trigger an assertion in the
// sem::Expression initializer, which checks that kConstant is paired with a
// Use the kRuntime EvaluationStage, as kConstant will trigger an assertion in
// the sem::Expression initializer, which checks that kConstant is paired with a
// constant value.
stage = sem::EvaluationStage::kRuntime;
}
@ -1998,8 +2004,8 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
current_statement_, value, has_side_effects);
};
// ty_init_or_conv is a helper for building either a sem::TypeInitializer or sem::TypeConversion
// call for the given semantic type.
// ty_init_or_conv is a helper for building either a sem::TypeInitializer or
// sem::TypeConversion call for the given semantic type.
auto ty_init_or_conv = [&](const sem::Type* ty) {
return Switch(
ty, //
@ -2076,7 +2082,8 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
});
};
// ast::CallExpression has a target which is either an ast::Type or an ast::IdentifierExpression
// ast::CallExpression has a target which is either an ast::Type or an
// ast::IdentifierExpression
sem::Call* call = nullptr;
if (expr->target.type) {
// ast::CallExpression has an ast::Type as the target.
@ -2188,7 +2195,8 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
});
} else {
// ast::CallExpression has an ast::IdentifierExpression as the target.
// This call is either a function call, builtin call, type initializer or type conversion.
// This call is either a function call, builtin call, type initializer or type
// conversion.
auto* ident = expr->target.name;
Mark(ident);
auto* resolved = sem_.ResolvedSymbol(ident);
@ -2197,8 +2205,8 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
[&](sem::Type* ty) {
// A type initializer or conversions.
// Note: Unlike the code path where we're resolving the call target from an
// ast::Type, all types must already have the element type explicitly specified, so
// there's no need to infer element types.
// ast::Type, all types must already have the element type explicitly specified,
// so there's no need to infer element types.
return ty_init_or_conv(ty);
},
[&](sem::Function* func) { return FunctionCall(expr, func, args, arg_behaviors); },
@ -2264,7 +2272,8 @@ sem::Call* Resolver::BuiltinCall(const ast::CallExpression* expr,
AddWarning("use of deprecated builtin", expr->source);
}
// If the builtin is @const, and all arguments have constant values, evaluate the builtin now.
// If the builtin is @const, and all arguments have constant values, evaluate the builtin
// now.
auto stage = sem::EarliestStage(arg_stage, builtin.sem->Stage());
const sem::Constant* value = nullptr;
if (stage == sem::EvaluationStage::kConstant) {
@ -2894,7 +2903,8 @@ sem::Array* Resolver::Array(const ast::Array* arr) {
}
}
// Track the pipeline-overridable constants that are transitively referenced by this array type.
// Track the pipeline-overridable constants that are transitively referenced by this array
// type.
for (auto* var : transitively_referenced_overrides) {
out->AddTransitivelyReferencedOverride(var);
}
@ -2955,8 +2965,9 @@ bool Resolver::ArrayAttributes(utils::VectorRef<const ast::Attribute*> attribute
Mark(attr);
if (auto* sd = attr->As<ast::StrideAttribute>()) {
// If the element type is not plain, then el_ty->Align() may be 0, in which case we
// could get a DBZ in ArrayStrideAttribute(). In this case, validation will error about
// the invalid array element type (which is tested later), so this is just a seatbelt.
// could get a DBZ in ArrayStrideAttribute(). In this case, validation will error
// about the invalid array element type (which is tested later), so this is just a
// seatbelt.
if (IsPlain(el_ty)) {
explicit_stride = sd->stride;
if (!validator_.ArrayStrideAttribute(sd, el_ty->Size(), el_ty->Align())) {

View File

@ -152,6 +152,14 @@ class Struct final : public Castable<Struct, Type> {
/// including size and alignment information.
std::string Layout(const tint::SymbolTable& symbols) const;
/// @param concrete the conversion-rank ordered concrete versions of this abstract structure.
void SetConcreteTypes(utils::VectorRef<const Struct*> concrete) { concrete_types_ = concrete; }
/// @returns the conversion-rank ordered concrete versions of this abstract structure, or an
/// empty vector if this structure is not abstract.
/// @note only structures returned by builtins may be abstract (e.g. modf, frexp)
const utils::Vector<const Struct*, 2>& ConcreteTypes() const { return concrete_types_; }
private:
ast::Struct const* const declaration_;
const Symbol name_;
@ -161,6 +169,7 @@ class Struct final : public Castable<Struct, Type> {
const uint32_t size_no_padding_;
std::unordered_set<ast::AddressSpace> address_space_usage_;
std::unordered_set<PipelineStageUsage> pipeline_stage_uses_;
utils::Vector<const Struct*, 2> concrete_types_;
};
/// StructMember holds the semantic information for structure members.

View File

@ -238,6 +238,15 @@ uint32_t Type::ConversionRank(const Type* from, const Type* to) {
}
return kNoConversion;
},
[&](const Struct* from_str) {
auto& concrete_tys = from_str->ConcreteTypes();
for (size_t i = 0; i < concrete_tys.Length(); i++) {
if (concrete_tys[i] == to) {
return static_cast<uint32_t>(i + 1);
}
}
return kNoConversion;
},
[&](Default) { return kNoConversion; });
}

View File

@ -44,6 +44,22 @@ struct TypeTest : public TestHelper {
const sem::Matrix* mat4x3_af = create<Matrix>(vec3_af, 4u);
const sem::Reference* ref_u32 =
create<Reference>(u32, ast::AddressSpace::kPrivate, ast::Access::kReadWrite);
const sem::Struct* str_f32 = create<Struct>(nullptr,
Sym("str_f32"),
StructMemberList{
create<StructMember>(
/* declaration */ nullptr,
/* name */ Sym("x"),
/* type */ f32,
/* index */ 0u,
/* offset */ 0u,
/* align */ 4u,
/* size */ 4u,
/* location */ std::nullopt),
},
/* align*/ 4u,
/* size*/ 4u,
/* size_no_padding*/ 4u);
const sem::Struct* str_f16 = create<Struct>(nullptr,
Sym("str_f16"),
StructMemberList{
@ -60,7 +76,7 @@ struct TypeTest : public TestHelper {
/* align*/ 4u,
/* size*/ 4u,
/* size_no_padding*/ 4u);
const sem::Struct* str_af = create<Struct>(nullptr,
sem::Struct* str_af = create<Struct>(nullptr,
Sym("str_af"),
StructMemberList{
create<StructMember>(
@ -139,6 +155,8 @@ struct TypeTest : public TestHelper {
/* size */ 5u * 4u,
/* stride */ 5u * 4u,
/* implicit_stride */ 5u * 4u);
TypeTest() { str_af->SetConcreteTypes(utils::Vector{str_f32, str_f16}); }
};
TEST_F(TypeTest, ConversionRank) {
@ -178,6 +196,8 @@ TEST_F(TypeTest, ConversionRank) {
EXPECT_EQ(Type::ConversionRank(ai, af), 5u);
EXPECT_EQ(Type::ConversionRank(ai, f32), 6u);
EXPECT_EQ(Type::ConversionRank(ai, f16), 7u);
EXPECT_EQ(Type::ConversionRank(str_af, str_f32), 1u);
EXPECT_EQ(Type::ConversionRank(str_af, str_f16), 2u);
EXPECT_EQ(Type::ConversionRank(i32, f32), Type::kNoConversion);
EXPECT_EQ(Type::ConversionRank(f32, u32), Type::kNoConversion);
@ -199,6 +219,10 @@ TEST_F(TypeTest, ConversionRank) {
EXPECT_EQ(Type::ConversionRank(af, ai), Type::kNoConversion);
EXPECT_EQ(Type::ConversionRank(f32, ai), Type::kNoConversion);
EXPECT_EQ(Type::ConversionRank(f16, ai), Type::kNoConversion);
EXPECT_EQ(Type::ConversionRank(str_f32, str_f16), Type::kNoConversion);
EXPECT_EQ(Type::ConversionRank(str_f16, str_f32), Type::kNoConversion);
EXPECT_EQ(Type::ConversionRank(str_f32, str_af), Type::kNoConversion);
EXPECT_EQ(Type::ConversionRank(str_f16, str_af), Type::kNoConversion);
}
TEST_F(TypeTest, ElementOf) {

View File

@ -0,0 +1,43 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn modf(vec<3, fa>) -> __modf_result_vec<3, fa>
fn modf_68d8ee() {
var res = modf(vec3(-1.5));
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
modf_68d8ee();
return vec4<f32>();
}
@fragment
fn fragment_main() {
modf_68d8ee();
}
@compute @workgroup_size(1)
fn compute_main() {
modf_68d8ee();
}

View File

@ -0,0 +1,43 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn modf(vec<2, fa>) -> __modf_result_vec<2, fa>
fn modf_732aa6() {
var res = modf(vec2(-1.5));
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
modf_732aa6();
return vec4<f32>();
}
@fragment
fn fragment_main() {
modf_732aa6();
}
@compute @workgroup_size(1)
fn compute_main() {
modf_732aa6();
}

View File

@ -0,0 +1,43 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn modf(fa) -> __modf_result<fa>
fn modf_c15f48() {
var res = modf(-1.5);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
modf_c15f48();
return vec4<f32>();
}
@fragment
fn fragment_main() {
modf_c15f48();
}
@compute @workgroup_size(1)
fn compute_main() {
modf_c15f48();
}

View File

@ -0,0 +1,43 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn modf(vec<4, fa>) -> __modf_result_vec<4, fa>
fn modf_f3d1f9() {
var res = modf(vec4(-1.5));
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
modf_f3d1f9();
return vec4<f32>();
}
@fragment
fn fragment_main() {
modf_f3d1f9();
}
@compute @workgroup_size(1)
fn compute_main() {
modf_f3d1f9();
}

View File

@ -0,0 +1,44 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn modf(vec<3, fa>) -> __modf_result_vec<3, fa>
fn modf_68d8ee() {
const arg_0 = vec3(-1.5);
var res = modf(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
modf_68d8ee();
return vec4<f32>();
}
@fragment
fn fragment_main() {
modf_68d8ee();
}
@compute @workgroup_size(1)
fn compute_main() {
modf_68d8ee();
}

View File

@ -0,0 +1,44 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn modf(vec<2, fa>) -> __modf_result_vec<2, fa>
fn modf_732aa6() {
const arg_0 = vec2(-1.5);
var res = modf(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
modf_732aa6();
return vec4<f32>();
}
@fragment
fn fragment_main() {
modf_732aa6();
}
@compute @workgroup_size(1)
fn compute_main() {
modf_732aa6();
}

View File

@ -0,0 +1,44 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn modf(fa) -> __modf_result<fa>
fn modf_c15f48() {
const arg_0 = -1.5;
var res = modf(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
modf_c15f48();
return vec4<f32>();
}
@fragment
fn fragment_main() {
modf_c15f48();
}
@compute @workgroup_size(1)
fn compute_main() {
modf_c15f48();
}

View File

@ -0,0 +1,44 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn modf(vec<4, fa>) -> __modf_result_vec<4, fa>
fn modf_f3d1f9() {
const arg_0 = vec4(-1.5);
var res = modf(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
modf_f3d1f9();
return vec4<f32>();
}
@fragment
fn fragment_main() {
modf_f3d1f9();
}
@compute @workgroup_size(1)
fn compute_main() {
modf_f3d1f9();
}

View File

@ -1,6 +1,6 @@
@compute @workgroup_size(1)
fn main() {
const in = 1.23;
const in = 1.25;
let res = modf(in);
let fract : f32 = res.fract;
let whole : f32 = res.whole;

View File

@ -4,7 +4,7 @@ struct modf_result {
};
[numthreads(1, 1, 1)]
void main() {
const modf_result res = {0.230000019f, 1.0f};
const modf_result res = {0.25f, 1.0f};
const float fract = res.fract;
const float whole = res.whole;
return;

View File

@ -4,7 +4,7 @@ struct modf_result {
};
[numthreads(1, 1, 1)]
void main() {
const modf_result res = {0.230000019f, 1.0f};
const modf_result res = {0.25f, 1.0f};
const float fract = res.fract;
const float whole = res.whole;
return;

View File

@ -7,7 +7,7 @@ struct modf_result {
void tint_symbol() {
modf_result res = modf_result(0.230000019f, 1.0f);
modf_result res = modf_result(0.25f, 1.0f);
float tint_symbol_2 = res.fract;
float whole = res.whole;
}

View File

@ -7,7 +7,7 @@ struct modf_result {
float whole;
};
kernel void tint_symbol() {
modf_result const res = modf_result{.fract=0.230000019f, .whole=1.0f};
modf_result const res = modf_result{.fract=0.25f, .whole=1.0f};
float const fract = res.fract;
float const whole = res.whole;
return;

View File

@ -17,9 +17,9 @@
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%__modf_result = OpTypeStruct %float %float
%float_0_230000019 = OpConstant %float 0.230000019
%float_0_25 = OpConstant %float 0.25
%float_1 = OpConstant %float 1
%9 = OpConstantComposite %__modf_result %float_0_230000019 %float_1
%9 = OpConstantComposite %__modf_result %float_0_25 %float_1
%main = OpFunction %void None %1
%4 = OpLabel
%10 = OpCompositeExtract %float %9 0

View File

@ -1,6 +1,6 @@
@compute @workgroup_size(1)
fn main() {
const in = 1.23;
const in = 1.25;
let res = modf(in);
let fract : f32 = res.fract;
let whole : f32 = res.whole;

View File

@ -1,6 +1,6 @@
@compute @workgroup_size(1)
fn main() {
const in = 1.23;
const in = 1.25;
const res = modf(in);
let fract : f32 = res.fract;
let whole : f32 = res.whole;

View File

@ -4,9 +4,9 @@ struct modf_result {
};
[numthreads(1, 1, 1)]
void main() {
const modf_result tint_symbol_1 = {0.230000019f, 1.0f};
const modf_result tint_symbol_1 = {0.25f, 1.0f};
const float fract = tint_symbol_1.fract;
const modf_result tint_symbol_2 = {0.230000019f, 1.0f};
const modf_result tint_symbol_2 = {0.25f, 1.0f};
const float whole = tint_symbol_2.whole;
return;
}

View File

@ -4,9 +4,9 @@ struct modf_result {
};
[numthreads(1, 1, 1)]
void main() {
const modf_result tint_symbol_1 = {0.230000019f, 1.0f};
const modf_result tint_symbol_1 = {0.25f, 1.0f};
const float fract = tint_symbol_1.fract;
const modf_result tint_symbol_2 = {0.230000019f, 1.0f};
const modf_result tint_symbol_2 = {0.25f, 1.0f};
const float whole = tint_symbol_2.whole;
return;
}

View File

@ -7,9 +7,9 @@ struct modf_result {
void tint_symbol() {
modf_result tint_symbol_3 = modf_result(0.230000019f, 1.0f);
modf_result tint_symbol_3 = modf_result(0.25f, 1.0f);
float tint_symbol_2 = tint_symbol_3.fract;
modf_result tint_symbol_4 = modf_result(0.230000019f, 1.0f);
modf_result tint_symbol_4 = modf_result(0.25f, 1.0f);
float whole = tint_symbol_4.whole;
}

View File

@ -7,9 +7,9 @@ struct modf_result {
float whole;
};
kernel void tint_symbol() {
modf_result const tint_symbol_1 = modf_result{.fract=0.230000019f, .whole=1.0f};
modf_result const tint_symbol_1 = modf_result{.fract=0.25f, .whole=1.0f};
float const fract = tint_symbol_1.fract;
modf_result const tint_symbol_2 = modf_result{.fract=0.230000019f, .whole=1.0f};
modf_result const tint_symbol_2 = modf_result{.fract=0.25f, .whole=1.0f};
float const whole = tint_symbol_2.whole;
return;
}

View File

@ -11,7 +11,7 @@
%void = OpTypeVoid
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%float_0_230000019 = OpConstant %float 0.230000019
%float_0_25 = OpConstant %float 0.25
%float_1 = OpConstant %float 1
%main = OpFunction %void None %1
%4 = OpLabel

View File

@ -1,6 +1,6 @@
@compute @workgroup_size(1)
fn main() {
const in = 1.23;
const in = 1.25;
const res = modf(in);
let fract : f32 = res.fract;
let whole : f32 = res.whole;

View File

@ -1,7 +1,7 @@
@compute @workgroup_size(1)
fn main() {
const const_in = 1.23;
let runtime_in = 1.23;
const const_in = 1.25;
let runtime_in = 1.25;
var res = modf(const_in);
res = modf(runtime_in);
res = modf(const_in);

View File

@ -10,10 +10,10 @@ modf_result tint_modf(float param_0) {
[numthreads(1, 1, 1)]
void main() {
const float runtime_in = 1.230000019f;
modf_result res = {0.230000019f, 1.0f};
const float runtime_in = 1.25f;
modf_result res = {0.25f, 1.0f};
res = tint_modf(runtime_in);
const modf_result c = {0.230000019f, 1.0f};
const modf_result c = {0.25f, 1.0f};
res = c;
const float fract = res.fract;
const float whole = res.whole;

View File

@ -10,10 +10,10 @@ modf_result tint_modf(float param_0) {
[numthreads(1, 1, 1)]
void main() {
const float runtime_in = 1.230000019f;
modf_result res = {0.230000019f, 1.0f};
const float runtime_in = 1.25f;
modf_result res = {0.25f, 1.0f};
res = tint_modf(runtime_in);
const modf_result c = {0.230000019f, 1.0f};
const modf_result c = {0.25f, 1.0f};
res = c;
const float fract = res.fract;
const float whole = res.whole;

View File

@ -13,10 +13,10 @@ modf_result tint_modf(float param_0) {
void tint_symbol() {
float runtime_in = 1.230000019f;
modf_result res = modf_result(0.230000019f, 1.0f);
float runtime_in = 1.25f;
modf_result res = modf_result(0.25f, 1.0f);
res = tint_modf(runtime_in);
res = modf_result(0.230000019f, 1.0f);
res = modf_result(0.25f, 1.0f);
float tint_symbol_1 = res.fract;
float whole = res.whole;
}

View File

@ -13,10 +13,10 @@ modf_result tint_modf(float param_0) {
}
kernel void tint_symbol() {
float const runtime_in = 1.230000019f;
modf_result res = modf_result{.fract=0.230000019f, .whole=1.0f};
float const runtime_in = 1.25f;
modf_result res = modf_result{.fract=0.25f, .whole=1.0f};
res = tint_modf(runtime_in);
res = modf_result{.fract=0.230000019f, .whole=1.0f};
res = modf_result{.fract=0.25f, .whole=1.0f};
float const fract = res.fract;
float const whole = res.whole;
return;

View File

@ -18,11 +18,11 @@
%void = OpTypeVoid
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%float_1_23000002 = OpConstant %float 1.23000002
%float_1_25 = OpConstant %float 1.25
%__modf_result = OpTypeStruct %float %float
%float_0_230000019 = OpConstant %float 0.230000019
%float_0_25 = OpConstant %float 0.25
%float_1 = OpConstant %float 1
%10 = OpConstantComposite %__modf_result %float_0_230000019 %float_1
%10 = OpConstantComposite %__modf_result %float_0_25 %float_1
%_ptr_Function___modf_result = OpTypePointer Function %__modf_result
%13 = OpConstantNull %__modf_result
%uint = OpTypeInt 32 0
@ -33,7 +33,7 @@
%4 = OpLabel
%res = OpVariable %_ptr_Function___modf_result Function %13
OpStore %res %10
%14 = OpExtInst %__modf_result %15 ModfStruct %float_1_23000002
%14 = OpExtInst %__modf_result %15 ModfStruct %float_1_25
OpStore %res %14
OpStore %res %10
%19 = OpAccessChain %_ptr_Function_float %res %uint_0

View File

@ -1,7 +1,7 @@
@compute @workgroup_size(1)
fn main() {
const const_in = 1.23;
let runtime_in = 1.23;
const const_in = 1.25;
let runtime_in = 1.25;
var res = modf(const_in);
res = modf(runtime_in);
res = modf(const_in);

View File

@ -1,6 +1,6 @@
@compute @workgroup_size(1)
fn main() {
let in = 1.23;
let in = 1.25;
let res = modf(in);
let fract : f32 = res.fract;
let whole : f32 = res.whole;

View File

@ -10,7 +10,7 @@ modf_result tint_modf(float param_0) {
[numthreads(1, 1, 1)]
void main() {
const float tint_symbol = 1.230000019f;
const float tint_symbol = 1.25f;
const modf_result res = tint_modf(tint_symbol);
const float fract = res.fract;
const float whole = res.whole;

View File

@ -10,7 +10,7 @@ modf_result tint_modf(float param_0) {
[numthreads(1, 1, 1)]
void main() {
const float tint_symbol = 1.230000019f;
const float tint_symbol = 1.25f;
const modf_result res = tint_modf(tint_symbol);
const float fract = res.fract;
const float whole = res.whole;

View File

@ -13,7 +13,7 @@ modf_result tint_modf(float param_0) {
void tint_symbol() {
float tint_symbol_1 = 1.230000019f;
float tint_symbol_1 = 1.25f;
modf_result res = tint_modf(tint_symbol_1);
float tint_symbol_2 = res.fract;
float whole = res.whole;

View File

@ -13,7 +13,7 @@ modf_result tint_modf(float param_0) {
}
kernel void tint_symbol() {
float const in = 1.230000019f;
float const in = 1.25f;
modf_result const res = tint_modf(in);
float const fract = res.fract;
float const whole = res.whole;

View File

@ -17,11 +17,11 @@
%void = OpTypeVoid
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%float_1_23000002 = OpConstant %float 1.23000002
%float_1_25 = OpConstant %float 1.25
%__modf_result = OpTypeStruct %float %float
%main = OpFunction %void None %1
%4 = OpLabel
%7 = OpExtInst %__modf_result %9 ModfStruct %float_1_23000002
%7 = OpExtInst %__modf_result %9 ModfStruct %float_1_25
%10 = OpCompositeExtract %float %7 0
%11 = OpCompositeExtract %float %7 1
OpReturn

View File

@ -1,6 +1,6 @@
@compute @workgroup_size(1)
fn main() {
let in = 1.23;
let in = 1.25;
let res = modf(in);
let fract : f32 = res.fract;
let whole : f32 = res.whole;

View File

@ -1,6 +1,6 @@
@compute @workgroup_size(1)
fn main() {
const in = vec2(1.23, 3.45);
const in = vec2(1.25, 3.75);
let res = modf(in);
let fract : vec2<f32> = res.fract;
let whole : vec2<f32> = res.whole;

View File

@ -4,7 +4,7 @@ struct modf_result_vec2 {
};
[numthreads(1, 1, 1)]
void main() {
const modf_result_vec2 res = {float2(0.230000019f, 0.450000048f), float2(1.0f, 3.0f)};
const modf_result_vec2 res = {float2(0.25f, 0.75f), float2(1.0f, 3.0f)};
const float2 fract = res.fract;
const float2 whole = res.whole;
return;

View File

@ -4,7 +4,7 @@ struct modf_result_vec2 {
};
[numthreads(1, 1, 1)]
void main() {
const modf_result_vec2 res = {float2(0.230000019f, 0.450000048f), float2(1.0f, 3.0f)};
const modf_result_vec2 res = {float2(0.25f, 0.75f), float2(1.0f, 3.0f)};
const float2 fract = res.fract;
const float2 whole = res.whole;
return;

View File

@ -7,7 +7,7 @@ struct modf_result_vec2 {
void tint_symbol() {
modf_result_vec2 res = modf_result_vec2(vec2(0.230000019f, 0.450000048f), vec2(1.0f, 3.0f));
modf_result_vec2 res = modf_result_vec2(vec2(0.25f, 0.75f), vec2(1.0f, 3.0f));
vec2 tint_symbol_2 = res.fract;
vec2 whole = res.whole;
}

View File

@ -7,7 +7,7 @@ struct modf_result_vec2 {
float2 whole;
};
kernel void tint_symbol() {
modf_result_vec2 const res = modf_result_vec2{.fract=float2(0.230000019f, 0.450000048f), .whole=float2(1.0f, 3.0f)};
modf_result_vec2 const res = modf_result_vec2{.fract=float2(0.25f, 0.75f), .whole=float2(1.0f, 3.0f)};
float2 const fract = res.fract;
float2 const whole = res.whole;
return;

View File

@ -18,9 +18,9 @@
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%__modf_result_vec2 = OpTypeStruct %v2float %v2float
%float_0_230000019 = OpConstant %float 0.230000019
%float_0_450000048 = OpConstant %float 0.450000048
%10 = OpConstantComposite %v2float %float_0_230000019 %float_0_450000048
%float_0_25 = OpConstant %float 0.25
%float_0_75 = OpConstant %float 0.75
%10 = OpConstantComposite %v2float %float_0_25 %float_0_75
%float_1 = OpConstant %float 1
%float_3 = OpConstant %float 3
%13 = OpConstantComposite %v2float %float_1 %float_3

View File

@ -1,6 +1,6 @@
@compute @workgroup_size(1)
fn main() {
const in = vec2(1.23, 3.45);
const in = vec2(1.25, 3.75);
let res = modf(in);
let fract : vec2<f32> = res.fract;
let whole : vec2<f32> = res.whole;

View File

@ -1,6 +1,6 @@
@compute @workgroup_size(1)
fn main() {
const in = vec2(1.23, 3.45);
const in = vec2(1.25, 3.75);
const res = modf(in);
let fract : vec2<f32> = res.fract;
let whole : vec2<f32> = res.whole;

View File

@ -4,9 +4,9 @@ struct modf_result_vec2 {
};
[numthreads(1, 1, 1)]
void main() {
const modf_result_vec2 tint_symbol_1 = {float2(0.230000019f, 0.450000048f), float2(1.0f, 3.0f)};
const modf_result_vec2 tint_symbol_1 = {float2(0.25f, 0.75f), float2(1.0f, 3.0f)};
const float2 fract = tint_symbol_1.fract;
const modf_result_vec2 tint_symbol_2 = {float2(0.230000019f, 0.450000048f), float2(1.0f, 3.0f)};
const modf_result_vec2 tint_symbol_2 = {float2(0.25f, 0.75f), float2(1.0f, 3.0f)};
const float2 whole = tint_symbol_2.whole;
return;
}

View File

@ -4,9 +4,9 @@ struct modf_result_vec2 {
};
[numthreads(1, 1, 1)]
void main() {
const modf_result_vec2 tint_symbol_1 = {float2(0.230000019f, 0.450000048f), float2(1.0f, 3.0f)};
const modf_result_vec2 tint_symbol_1 = {float2(0.25f, 0.75f), float2(1.0f, 3.0f)};
const float2 fract = tint_symbol_1.fract;
const modf_result_vec2 tint_symbol_2 = {float2(0.230000019f, 0.450000048f), float2(1.0f, 3.0f)};
const modf_result_vec2 tint_symbol_2 = {float2(0.25f, 0.75f), float2(1.0f, 3.0f)};
const float2 whole = tint_symbol_2.whole;
return;
}

View File

@ -7,9 +7,9 @@ struct modf_result_vec2 {
void tint_symbol() {
modf_result_vec2 tint_symbol_3 = modf_result_vec2(vec2(0.230000019f, 0.450000048f), vec2(1.0f, 3.0f));
modf_result_vec2 tint_symbol_3 = modf_result_vec2(vec2(0.25f, 0.75f), vec2(1.0f, 3.0f));
vec2 tint_symbol_2 = tint_symbol_3.fract;
modf_result_vec2 tint_symbol_4 = modf_result_vec2(vec2(0.230000019f, 0.450000048f), vec2(1.0f, 3.0f));
modf_result_vec2 tint_symbol_4 = modf_result_vec2(vec2(0.25f, 0.75f), vec2(1.0f, 3.0f));
vec2 whole = tint_symbol_4.whole;
}

View File

@ -7,9 +7,9 @@ struct modf_result_vec2 {
float2 whole;
};
kernel void tint_symbol() {
modf_result_vec2 const tint_symbol_1 = modf_result_vec2{.fract=float2(0.230000019f, 0.450000048f), .whole=float2(1.0f, 3.0f)};
modf_result_vec2 const tint_symbol_1 = modf_result_vec2{.fract=float2(0.25f, 0.75f), .whole=float2(1.0f, 3.0f)};
float2 const fract = tint_symbol_1.fract;
modf_result_vec2 const tint_symbol_2 = modf_result_vec2{.fract=float2(0.230000019f, 0.450000048f), .whole=float2(1.0f, 3.0f)};
modf_result_vec2 const tint_symbol_2 = modf_result_vec2{.fract=float2(0.25f, 0.75f), .whole=float2(1.0f, 3.0f)};
float2 const whole = tint_symbol_2.whole;
return;
}

View File

@ -12,9 +12,9 @@
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%float_0_230000019 = OpConstant %float 0.230000019
%float_0_450000048 = OpConstant %float 0.450000048
%9 = OpConstantComposite %v2float %float_0_230000019 %float_0_450000048
%float_0_25 = OpConstant %float 0.25
%float_0_75 = OpConstant %float 0.75
%9 = OpConstantComposite %v2float %float_0_25 %float_0_75
%float_1 = OpConstant %float 1
%float_3 = OpConstant %float 3
%12 = OpConstantComposite %v2float %float_1 %float_3

View File

@ -1,6 +1,6 @@
@compute @workgroup_size(1)
fn main() {
const in = vec2(1.23, 3.45);
const in = vec2(1.25, 3.75);
const res = modf(in);
let fract : vec2<f32> = res.fract;
let whole : vec2<f32> = res.whole;

View File

@ -1,7 +1,7 @@
@compute @workgroup_size(1)
fn main() {
const const_in = vec2(1.23, 3.45);
let runtime_in = vec2(1.23, 3.45);
const const_in = vec2(1.25, 3.75);
let runtime_in = vec2(1.25, 3.75);
var res = modf(const_in);
res = modf(runtime_in);
res = modf(const_in);

View File

@ -10,10 +10,10 @@ modf_result_vec2 tint_modf(float2 param_0) {
[numthreads(1, 1, 1)]
void main() {
const float2 runtime_in = float2(1.230000019f, 3.450000048f);
modf_result_vec2 res = {float2(0.230000019f, 0.450000048f), float2(1.0f, 3.0f)};
const float2 runtime_in = float2(1.25f, 3.75f);
modf_result_vec2 res = {float2(0.25f, 0.75f), float2(1.0f, 3.0f)};
res = tint_modf(runtime_in);
const modf_result_vec2 c = {float2(0.230000019f, 0.450000048f), float2(1.0f, 3.0f)};
const modf_result_vec2 c = {float2(0.25f, 0.75f), float2(1.0f, 3.0f)};
res = c;
const float2 fract = res.fract;
const float2 whole = res.whole;

View File

@ -10,10 +10,10 @@ modf_result_vec2 tint_modf(float2 param_0) {
[numthreads(1, 1, 1)]
void main() {
const float2 runtime_in = float2(1.230000019f, 3.450000048f);
modf_result_vec2 res = {float2(0.230000019f, 0.450000048f), float2(1.0f, 3.0f)};
const float2 runtime_in = float2(1.25f, 3.75f);
modf_result_vec2 res = {float2(0.25f, 0.75f), float2(1.0f, 3.0f)};
res = tint_modf(runtime_in);
const modf_result_vec2 c = {float2(0.230000019f, 0.450000048f), float2(1.0f, 3.0f)};
const modf_result_vec2 c = {float2(0.25f, 0.75f), float2(1.0f, 3.0f)};
res = c;
const float2 fract = res.fract;
const float2 whole = res.whole;

View File

@ -13,10 +13,10 @@ modf_result_vec2 tint_modf(vec2 param_0) {
void tint_symbol() {
vec2 runtime_in = vec2(1.230000019f, 3.450000048f);
modf_result_vec2 res = modf_result_vec2(vec2(0.230000019f, 0.450000048f), vec2(1.0f, 3.0f));
vec2 runtime_in = vec2(1.25f, 3.75f);
modf_result_vec2 res = modf_result_vec2(vec2(0.25f, 0.75f), vec2(1.0f, 3.0f));
res = tint_modf(runtime_in);
res = modf_result_vec2(vec2(0.230000019f, 0.450000048f), vec2(1.0f, 3.0f));
res = modf_result_vec2(vec2(0.25f, 0.75f), vec2(1.0f, 3.0f));
vec2 tint_symbol_1 = res.fract;
vec2 whole = res.whole;
}

View File

@ -13,10 +13,10 @@ modf_result_vec2 tint_modf(float2 param_0) {
}
kernel void tint_symbol() {
float2 const runtime_in = float2(1.230000019f, 3.450000048f);
modf_result_vec2 res = modf_result_vec2{.fract=float2(0.230000019f, 0.450000048f), .whole=float2(1.0f, 3.0f)};
float2 const runtime_in = float2(1.25f, 3.75f);
modf_result_vec2 res = modf_result_vec2{.fract=float2(0.25f, 0.75f), .whole=float2(1.0f, 3.0f)};
res = tint_modf(runtime_in);
res = modf_result_vec2{.fract=float2(0.230000019f, 0.450000048f), .whole=float2(1.0f, 3.0f)};
res = modf_result_vec2{.fract=float2(0.25f, 0.75f), .whole=float2(1.0f, 3.0f)};
float2 const fract = res.fract;
float2 const whole = res.whole;
return;

View File

@ -19,13 +19,13 @@
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%float_1_23000002 = OpConstant %float 1.23000002
%float_3_45000005 = OpConstant %float 3.45000005
%9 = OpConstantComposite %v2float %float_1_23000002 %float_3_45000005
%float_1_25 = OpConstant %float 1.25
%float_3_75 = OpConstant %float 3.75
%9 = OpConstantComposite %v2float %float_1_25 %float_3_75
%__modf_result_vec2 = OpTypeStruct %v2float %v2float
%float_0_230000019 = OpConstant %float 0.230000019
%float_0_450000048 = OpConstant %float 0.450000048
%13 = OpConstantComposite %v2float %float_0_230000019 %float_0_450000048
%float_0_25 = OpConstant %float 0.25
%float_0_75 = OpConstant %float 0.75
%13 = OpConstantComposite %v2float %float_0_25 %float_0_75
%float_1 = OpConstant %float 1
%float_3 = OpConstant %float 3
%16 = OpConstantComposite %v2float %float_1 %float_3

View File

@ -1,7 +1,7 @@
@compute @workgroup_size(1)
fn main() {
const const_in = vec2(1.23, 3.45);
let runtime_in = vec2(1.23, 3.45);
const const_in = vec2(1.25, 3.75);
let runtime_in = vec2(1.25, 3.75);
var res = modf(const_in);
res = modf(runtime_in);
res = modf(const_in);

View File

@ -1,6 +1,6 @@
@compute @workgroup_size(1)
fn main() {
let in = vec2(1.23, 3.45);
let in = vec2(1.25, 3.75);
let res = modf(in);
let fract : vec2<f32> = res.fract;
let whole : vec2<f32> = res.whole;

View File

@ -10,7 +10,7 @@ modf_result_vec2 tint_modf(float2 param_0) {
[numthreads(1, 1, 1)]
void main() {
const float2 tint_symbol = float2(1.230000019f, 3.450000048f);
const float2 tint_symbol = float2(1.25f, 3.75f);
const modf_result_vec2 res = tint_modf(tint_symbol);
const float2 fract = res.fract;
const float2 whole = res.whole;

View File

@ -10,7 +10,7 @@ modf_result_vec2 tint_modf(float2 param_0) {
[numthreads(1, 1, 1)]
void main() {
const float2 tint_symbol = float2(1.230000019f, 3.450000048f);
const float2 tint_symbol = float2(1.25f, 3.75f);
const modf_result_vec2 res = tint_modf(tint_symbol);
const float2 fract = res.fract;
const float2 whole = res.whole;

View File

@ -13,7 +13,7 @@ modf_result_vec2 tint_modf(vec2 param_0) {
void tint_symbol() {
vec2 tint_symbol_1 = vec2(1.230000019f, 3.450000048f);
vec2 tint_symbol_1 = vec2(1.25f, 3.75f);
modf_result_vec2 res = tint_modf(tint_symbol_1);
vec2 tint_symbol_2 = res.fract;
vec2 whole = res.whole;

View File

@ -13,7 +13,7 @@ modf_result_vec2 tint_modf(float2 param_0) {
}
kernel void tint_symbol() {
float2 const in = float2(1.230000019f, 3.450000048f);
float2 const in = float2(1.25f, 3.75f);
modf_result_vec2 const res = tint_modf(in);
float2 const fract = res.fract;
float2 const whole = res.whole;

View File

@ -18,9 +18,9 @@
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%float_1_23000002 = OpConstant %float 1.23000002
%float_3_45000005 = OpConstant %float 3.45000005
%9 = OpConstantComposite %v2float %float_1_23000002 %float_3_45000005
%float_1_25 = OpConstant %float 1.25
%float_3_75 = OpConstant %float 3.75
%9 = OpConstantComposite %v2float %float_1_25 %float_3_75
%__modf_result_vec2 = OpTypeStruct %v2float %v2float
%main = OpFunction %void None %1
%4 = OpLabel

View File

@ -1,6 +1,6 @@
@compute @workgroup_size(1)
fn main() {
let in = vec2(1.23, 3.45);
let in = vec2(1.25, 3.75);
let res = modf(in);
let fract : vec2<f32> = res.fract;
let whole : vec2<f32> = res.whole;