tint: Implement const-eval of modf
Bug: tint:1581 Change-Id: I53151ebf43601cd6afcdd2ec91d0ff9c4e650ef3 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/111241 Reviewed-by: Antonio Maiorano <amaiorano@google.com> Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: Ben Clayton <bclayton@google.com>
This commit is contained in:
Ben Clayton
Dawn LUCI CQ
parent
92d858ac3c
commit
329dfd7cbd
|
|
@ -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>
|
||||
fn modf<T: f32_f16>(T) -> __modf_result<T>
|
||||
fn modf<N: num, T: f32_f16>(vec<N, T>) -> __modf_result_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>
|
||||
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
|
||||
|
|
|
|||
|
|
@ -2266,6 +2266,40 @@ ConstEval::Result ConstEval::min(const sem::Type* ty,
|
|||
return TransformElements(builder, ty, transform, args[0], args[1]);
|
||||
}
|
||||
|
||||
ConstEval::Result ConstEval::modf(const sem::Type* ty,
|
||||
utils::VectorRef<const sem::Constant*> args,
|
||||
const Source& source) {
|
||||
auto transform_fract = [&](const sem::Constant* c) {
|
||||
auto create = [&](auto e) {
|
||||
return CreateElement(builder, source, c->Type(),
|
||||
decltype(e)(e.value - std::trunc(e.value)));
|
||||
};
|
||||
return Dispatch_fa_f32_f16(create, c);
|
||||
};
|
||||
auto transform_whole = [&](const sem::Constant* c) {
|
||||
auto create = [&](auto e) {
|
||||
return CreateElement(builder, source, c->Type(), decltype(e)(std::trunc(e.value)));
|
||||
};
|
||||
return Dispatch_fa_f32_f16(create, c);
|
||||
};
|
||||
|
||||
utils::Vector<const sem::Constant*, 2> fields;
|
||||
|
||||
if (auto fract = TransformElements(builder, args[0]->Type(), transform_fract, args[0])) {
|
||||
fields.Push(fract.Get());
|
||||
} else {
|
||||
return utils::Failure;
|
||||
}
|
||||
|
||||
if (auto whole = TransformElements(builder, args[0]->Type(), transform_whole, args[0])) {
|
||||
fields.Push(whole.Get());
|
||||
} else {
|
||||
return utils::Failure;
|
||||
}
|
||||
|
||||
return CreateComposite(builder, ty, std::move(fields));
|
||||
}
|
||||
|
||||
ConstEval::Result ConstEval::pack2x16float(const sem::Type* ty,
|
||||
utils::VectorRef<const sem::Constant*> args,
|
||||
const Source& source) {
|
||||
|
|
|
|||
|
|
@ -628,6 +628,15 @@ class ConstEval {
|
|||
utils::VectorRef<const sem::Constant*> args,
|
||||
const Source& source);
|
||||
|
||||
/// modf builtin
|
||||
/// @param ty the expression type
|
||||
/// @param args the input arguments
|
||||
/// @param source the source location
|
||||
/// @return the result value, or null if the value cannot be calculated
|
||||
Result modf(const sem::Type* ty,
|
||||
utils::VectorRef<const sem::Constant*> args,
|
||||
const Source& source);
|
||||
|
||||
/// pack2x16float builtin
|
||||
/// @param ty the expression type
|
||||
/// @param args the input arguments
|
||||
|
|
|
|||
|
|
@ -98,6 +98,11 @@ static Case C(std::initializer_list<Types> args, Types result) {
|
|||
return Case{utils::Vector<Types, 8>{args}, utils::Vector<Types, 2>{std::move(result)}};
|
||||
}
|
||||
|
||||
/// Creates a Case with Values for args and result
|
||||
static Case C(std::initializer_list<Types> args, std::initializer_list<Types> results) {
|
||||
return Case{utils::Vector<Types, 8>{args}, utils::Vector<Types, 2>{results}};
|
||||
}
|
||||
|
||||
/// Convenience overload that creates a Case with just scalars
|
||||
static Case C(std::initializer_list<ScalarTypes> sargs, ScalarTypes sresult) {
|
||||
utils::Vector<Types, 8> args;
|
||||
|
|
@ -109,6 +114,20 @@ static Case C(std::initializer_list<ScalarTypes> sargs, ScalarTypes sresult) {
|
|||
return Case{std::move(args), utils::Vector<Types, 2>{std::move(result)}};
|
||||
}
|
||||
|
||||
/// Creates a Case with Values for args and result
|
||||
static Case C(std::initializer_list<ScalarTypes> sargs,
|
||||
std::initializer_list<ScalarTypes> sresults) {
|
||||
utils::Vector<Types, 8> args;
|
||||
for (auto& sa : sargs) {
|
||||
std::visit([&](auto&& v) { return args.Push(Val(v)); }, sa);
|
||||
}
|
||||
utils::Vector<Types, 2> results;
|
||||
for (auto& sa : sresults) {
|
||||
std::visit([&](auto&& v) { return results.Push(Val(v)); }, sa);
|
||||
}
|
||||
return Case{std::move(args), std::move(results)};
|
||||
}
|
||||
|
||||
/// Creates a Case with Values for args and expected error
|
||||
static Case E(std::initializer_list<Types> args, std::string err) {
|
||||
return Case{utils::Vector<Types, 8>{args}, std::move(err)};
|
||||
|
|
@ -1290,6 +1309,38 @@ INSTANTIATE_TEST_SUITE_P( //
|
|||
MinCases<AFloat>(),
|
||||
MinCases<f32>(),
|
||||
MinCases<f16>()))));
|
||||
template <typename T>
|
||||
std::vector<Case> ModfCases() {
|
||||
return {
|
||||
// Scalar tests
|
||||
// in fract whole
|
||||
C({T(0.0)}, {T(0.0), T(0.0)}), //
|
||||
C({T(1.0)}, {T(0.0), T(1.0)}), //
|
||||
C({T(2.0)}, {T(0.0), T(2.0)}), //
|
||||
C({T(1.5)}, {T(0.5), T(1.0)}), //
|
||||
C({T(4.25)}, {T(0.25), T(4.0)}), //
|
||||
C({T(-1.0)}, {T(0.0), T(-1.0)}), //
|
||||
C({T(-2.0)}, {T(0.0), T(-2.0)}), //
|
||||
C({T(-1.5)}, {T(-0.5), T(-1.0)}), //
|
||||
C({T(-4.25)}, {T(-0.25), T(-4.0)}), //
|
||||
C({T::Lowest()}, {T(0.0), T::Lowest()}), //
|
||||
C({T::Highest()}, {T(0.0), T::Highest()}), //
|
||||
|
||||
// Vector tests
|
||||
// in fract whole
|
||||
C({Vec(T(0.0), T(0.0))}, {Vec(T(0.0), T(0.0)), Vec(T(0.0), T(0.0))}),
|
||||
C({Vec(T(1.0), T(2.0))}, {Vec(T(0.0), T(0.0)), Vec(T(1), T(2))}),
|
||||
C({Vec(T(-2.0), T(1.0))}, {Vec(T(0.0), T(0.0)), Vec(T(-2), T(1))}),
|
||||
C({Vec(T(1.5), T(-2.25))}, {Vec(T(0.5), T(-0.25)), Vec(T(1.0), T(-2.0))}),
|
||||
C({Vec(T::Lowest(), T::Highest())}, {Vec(T(0.0), T(0.0)), Vec(T::Lowest(), T::Highest())}),
|
||||
};
|
||||
}
|
||||
INSTANTIATE_TEST_SUITE_P( //
|
||||
Modf,
|
||||
ResolverConstEvalBuiltinTest,
|
||||
testing::Combine(testing::Values(sem::BuiltinType::kModf),
|
||||
testing::ValuesIn(Concat(ModfCases<f32>(), //
|
||||
ModfCases<f16>()))));
|
||||
|
||||
std::vector<Case> Pack4x8snormCases() {
|
||||
return {
|
||||
|
|
|
|||
|
|
@ -12851,7 +12851,7 @@ constexpr OverloadInfo kOverloads[] = {
|
|||
/* parameters */ &kParameters[878],
|
||||
/* return matcher indices */ &kMatcherIndices[106],
|
||||
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
|
||||
/* const eval */ nullptr,
|
||||
/* const eval */ &ConstEval::modf,
|
||||
},
|
||||
{
|
||||
/* [378] */
|
||||
|
|
@ -12863,7 +12863,7 @@ constexpr OverloadInfo kOverloads[] = {
|
|||
/* parameters */ &kParameters[879],
|
||||
/* return matcher indices */ &kMatcherIndices[45],
|
||||
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
|
||||
/* const eval */ nullptr,
|
||||
/* const eval */ &ConstEval::modf,
|
||||
},
|
||||
{
|
||||
/* [379] */
|
||||
|
|
@ -14351,8 +14351,8 @@ constexpr IntrinsicInfo kBuiltins[] = {
|
|||
},
|
||||
{
|
||||
/* [53] */
|
||||
/* fn modf<T : f32_f16>(T) -> __modf_result<T> */
|
||||
/* fn modf<N : num, T : f32_f16>(vec<N, T>) -> __modf_result_vec<N, T> */
|
||||
/* 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> */
|
||||
/* num overloads */ 2,
|
||||
/* overloads */ &kOverloads[377],
|
||||
},
|
||||
|
|
|
|||
|
|
@ -530,9 +530,13 @@ bool Validator::AddressSpaceLayout(const sem::Variable* var,
|
|||
}
|
||||
|
||||
if (auto* str = var->Type()->UnwrapRef()->As<sem::Struct>()) {
|
||||
if (!AddressSpaceLayout(str, var->AddressSpace(), str->Declaration()->source, layouts)) {
|
||||
AddNote("see declaration of variable", var->Declaration()->source);
|
||||
return false;
|
||||
// Check the structure has a declaration. Builtins like modf() and frexp() return untypeable
|
||||
// structures, and so they have no declaration. Just skip validation for these.
|
||||
if (auto* str_decl = str->Declaration()) {
|
||||
if (!AddressSpaceLayout(str, var->AddressSpace(), str_decl->source, layouts)) {
|
||||
AddNote("see declaration of variable", var->Declaration()->source);
|
||||
return false;
|
||||
}
|
||||
}
|
||||
} else {
|
||||
Source source = var->Declaration()->source;
|
||||
|
|
|
|||
|
|
@ -939,9 +939,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
|
|||
return true;
|
||||
}
|
||||
case sem::BuiltinType::kAtomicCompareExchangeWeak: {
|
||||
// Emit the builtin return type unique to this overload. This does not
|
||||
// exist in the AST, so it will not be generated in Generate().
|
||||
if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
|
||||
if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
|
||||
return false;
|
||||
}
|
||||
|
||||
|
|
@ -2373,10 +2371,12 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan
|
|||
return true;
|
||||
},
|
||||
[&](const sem::Struct* s) {
|
||||
if (!EmitType(out, s, ast::AddressSpace::kNone, ast::Access::kUndefined, "")) {
|
||||
if (!EmitStructType(&helpers_, s)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
out << StructName(s);
|
||||
|
||||
ScopedParen sp(out);
|
||||
|
||||
for (size_t i = 0; i < s->Members().size(); i++) {
|
||||
|
|
@ -2998,6 +2998,11 @@ bool GeneratorImpl::EmitTypeAndName(std::ostream& out,
|
|||
}
|
||||
|
||||
bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
|
||||
auto it = emitted_structs_.emplace(str);
|
||||
if (!it.second) {
|
||||
return true;
|
||||
}
|
||||
|
||||
auto address_space_uses = str->AddressSpaceUsage();
|
||||
line(b) << "struct " << StructName(str) << " {";
|
||||
EmitStructMembers(b, str);
|
||||
|
|
@ -3007,14 +3012,6 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
|
|||
return true;
|
||||
}
|
||||
|
||||
bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
|
||||
auto it = emitted_structs_.emplace(str);
|
||||
if (!it.second) {
|
||||
return true;
|
||||
}
|
||||
return EmitStructType(buffer, str);
|
||||
}
|
||||
|
||||
bool GeneratorImpl::EmitStructMembers(TextBuffer* b, const sem::Struct* str) {
|
||||
ScopedIndent si(b);
|
||||
for (auto* mem : str->Members()) {
|
||||
|
|
|
|||
|
|
@ -430,17 +430,12 @@ class GeneratorImpl : public TextGenerator {
|
|||
ast::AddressSpace address_space,
|
||||
ast::Access access,
|
||||
const std::string& name);
|
||||
/// Handles generating a structure declaration
|
||||
/// Handles generating a structure declaration. If the structure has already been emitted, then
|
||||
/// this function will simply return `true` without emitting anything.
|
||||
/// @param buffer the text buffer that the type declaration will be written to
|
||||
/// @param ty the struct to generate
|
||||
/// @returns true if the struct is emitted
|
||||
bool EmitStructType(TextBuffer* buffer, const sem::Struct* ty);
|
||||
/// Handles generating a structure declaration only the first time called. Subsequent calls are
|
||||
/// a no-op and return true.
|
||||
/// @param buffer the text buffer that the type declaration will be written to
|
||||
/// @param ty the struct to generate
|
||||
/// @returns true if the struct is emitted
|
||||
bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty);
|
||||
/// Handles generating the members of a structure
|
||||
/// @param buffer the text buffer that the struct members will be written to
|
||||
/// @param ty the struct to generate
|
||||
|
|
|
|||
|
|
@ -416,9 +416,9 @@ TEST_F(GlslGeneratorImplTest_Builtin, FMA_f16) {
|
|||
EXPECT_EQ(out.str(), "((a) * (b) + (c))");
|
||||
}
|
||||
|
||||
TEST_F(GlslGeneratorImplTest_Builtin, Modf_Scalar_f32) {
|
||||
auto* call = Call("modf", 1_f);
|
||||
WrapInFunction(CallStmt(call));
|
||||
TEST_F(GlslGeneratorImplTest_Builtin, Runtime_Modf_Scalar_f32) {
|
||||
WrapInFunction(Decl(Let("f", Expr(1.5_f))), //
|
||||
Decl(Let("v", Call("modf", "f"))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
|
|
@ -438,7 +438,8 @@ modf_result tint_modf(float param_0) {
|
|||
|
||||
|
||||
void test_function() {
|
||||
tint_modf(1.0f);
|
||||
float f = 1.5f;
|
||||
modf_result v = tint_modf(f);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
@ -449,11 +450,11 @@ void main() {
|
|||
)");
|
||||
}
|
||||
|
||||
TEST_F(GlslGeneratorImplTest_Builtin, Modf_Scalar_f16) {
|
||||
TEST_F(GlslGeneratorImplTest_Builtin, Runtime_Modf_Scalar_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
auto* call = Call("modf", 1_h);
|
||||
WrapInFunction(CallStmt(call));
|
||||
WrapInFunction(Decl(Let("f", Expr(1.5_h))), //
|
||||
Decl(Let("v", Call("modf", "f"))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
|
|
@ -474,7 +475,8 @@ modf_result_f16 tint_modf(float16_t param_0) {
|
|||
|
||||
|
||||
void test_function() {
|
||||
tint_modf(1.0hf);
|
||||
float16_t f = 1.5hf;
|
||||
modf_result_f16 v = tint_modf(f);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
@ -485,9 +487,9 @@ void main() {
|
|||
)");
|
||||
}
|
||||
|
||||
TEST_F(GlslGeneratorImplTest_Builtin, Modf_Vector_f32) {
|
||||
auto* call = Call("modf", vec3<f32>());
|
||||
WrapInFunction(CallStmt(call));
|
||||
TEST_F(GlslGeneratorImplTest_Builtin, Runtime_Modf_Vector_f32) {
|
||||
WrapInFunction(Decl(Let("f", vec3<f32>(1.5_f, 2.5_f, 3.5_f))), //
|
||||
Decl(Let("v", Call("modf", "f"))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
|
|
@ -507,7 +509,8 @@ modf_result_vec3 tint_modf(vec3 param_0) {
|
|||
|
||||
|
||||
void test_function() {
|
||||
tint_modf(vec3(0.0f));
|
||||
vec3 f = vec3(1.5f, 2.5f, 3.5f);
|
||||
modf_result_vec3 v = tint_modf(f);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
@ -518,11 +521,11 @@ void main() {
|
|||
)");
|
||||
}
|
||||
|
||||
TEST_F(GlslGeneratorImplTest_Builtin, Modf_Vector_f16) {
|
||||
TEST_F(GlslGeneratorImplTest_Builtin, Runtime_Modf_Vector_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
auto* call = Call("modf", vec3<f16>());
|
||||
WrapInFunction(CallStmt(call));
|
||||
WrapInFunction(Decl(Let("f", vec3<f16>(1.5_h, 2.5_h, 3.5_h))), //
|
||||
Decl(Let("v", Call("modf", "f"))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
|
|
@ -543,7 +546,118 @@ modf_result_vec3_f16 tint_modf(f16vec3 param_0) {
|
|||
|
||||
|
||||
void test_function() {
|
||||
tint_modf(f16vec3(0.0hf));
|
||||
f16vec3 f = f16vec3(1.5hf, 2.5hf, 3.5hf);
|
||||
modf_result_vec3_f16 v = tint_modf(f);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
void main() {
|
||||
test_function();
|
||||
return;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(GlslGeneratorImplTest_Builtin, Const_Modf_Scalar_f32) {
|
||||
WrapInFunction(Decl(Let("v", Call("modf", 1.5_f))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(#version 310 es
|
||||
|
||||
struct modf_result {
|
||||
float fract;
|
||||
float whole;
|
||||
};
|
||||
|
||||
|
||||
void test_function() {
|
||||
modf_result v = modf_result(0.5f, 1.0f);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
void main() {
|
||||
test_function();
|
||||
return;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(GlslGeneratorImplTest_Builtin, Const_Modf_Scalar_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
WrapInFunction(Decl(Let("v", Call("modf", 1.5_h))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(#version 310 es
|
||||
#extension GL_AMD_gpu_shader_half_float : require
|
||||
|
||||
struct modf_result_f16 {
|
||||
float16_t fract;
|
||||
float16_t whole;
|
||||
};
|
||||
|
||||
|
||||
void test_function() {
|
||||
modf_result_f16 v = modf_result_f16(0.5hf, 1.0hf);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
void main() {
|
||||
test_function();
|
||||
return;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(GlslGeneratorImplTest_Builtin, Const_Modf_Vector_f32) {
|
||||
WrapInFunction(Decl(Let("v", Call("modf", vec3<f32>(1.5_f, 2.5_f, 3.5_f)))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(#version 310 es
|
||||
|
||||
struct modf_result_vec3 {
|
||||
vec3 fract;
|
||||
vec3 whole;
|
||||
};
|
||||
|
||||
|
||||
void test_function() {
|
||||
modf_result_vec3 v = modf_result_vec3(vec3(0.5f), vec3(1.0f, 2.0f, 3.0f));
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
void main() {
|
||||
test_function();
|
||||
return;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(GlslGeneratorImplTest_Builtin, Const_Modf_Vector_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
WrapInFunction(Decl(Let("v", Call("modf", vec3<f16>(1.5_h, 2.5_h, 3.5_h)))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(#version 310 es
|
||||
#extension GL_AMD_gpu_shader_half_float : require
|
||||
|
||||
struct modf_result_vec3_f16 {
|
||||
f16vec3 fract;
|
||||
f16vec3 whole;
|
||||
};
|
||||
|
||||
|
||||
void test_function() {
|
||||
modf_result_vec3_f16 v = modf_result_vec3_f16(f16vec3(0.5hf), f16vec3(1.0hf, 2.0hf, 3.0hf));
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -1651,9 +1651,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
|
|||
return true;
|
||||
}
|
||||
case sem::BuiltinType::kAtomicCompareExchangeWeak: {
|
||||
// Emit the builtin return type unique to this overload. This does not
|
||||
// exist in the AST, so it will not be generated in Generate().
|
||||
if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
|
||||
if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
|
||||
return false;
|
||||
}
|
||||
|
||||
|
|
@ -2506,7 +2504,7 @@ bool GeneratorImpl::EmitCase(const ast::SwitchStatement* s, size_t case_idx) {
|
|||
out << "default";
|
||||
} else {
|
||||
out << "case ";
|
||||
if (!EmitConstant(out, selector->Value())) {
|
||||
if (!EmitConstant(out, selector->Value(), /* is_variable_initializer */ false)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
|
@ -2552,7 +2550,13 @@ bool GeneratorImpl::EmitDiscard(const ast::DiscardStatement*) {
|
|||
bool GeneratorImpl::EmitExpression(std::ostream& out, const ast::Expression* expr) {
|
||||
if (auto* sem = builder_.Sem().Get(expr)) {
|
||||
if (auto* constant = sem->ConstantValue()) {
|
||||
return EmitConstant(out, constant);
|
||||
bool is_variable_initializer = false;
|
||||
if (auto* stmt = sem->Stmt()) {
|
||||
if (auto* decl = As<ast::VariableDeclStatement>(stmt->Declaration())) {
|
||||
is_variable_initializer = decl->variable->initializer == expr;
|
||||
}
|
||||
}
|
||||
return EmitConstant(out, constant, is_variable_initializer);
|
||||
}
|
||||
}
|
||||
return Switch(
|
||||
|
|
@ -3042,7 +3046,9 @@ bool GeneratorImpl::EmitEntryPointFunction(const ast::Function* func) {
|
|||
return true;
|
||||
}
|
||||
|
||||
bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constant) {
|
||||
bool GeneratorImpl::EmitConstant(std::ostream& out,
|
||||
const sem::Constant* constant,
|
||||
bool is_variable_initializer) {
|
||||
return Switch(
|
||||
constant->Type(), //
|
||||
[&](const sem::Bool*) {
|
||||
|
|
@ -3072,7 +3078,7 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan
|
|||
if (constant->AllEqual()) {
|
||||
{
|
||||
ScopedParen sp(out);
|
||||
if (!EmitConstant(out, constant->Index(0))) {
|
||||
if (!EmitConstant(out, constant->Index(0), is_variable_initializer)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
|
@ -3093,7 +3099,7 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan
|
|||
if (i > 0) {
|
||||
out << ", ";
|
||||
}
|
||||
if (!EmitConstant(out, constant->Index(i))) {
|
||||
if (!EmitConstant(out, constant->Index(i), is_variable_initializer)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
|
@ -3110,7 +3116,7 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan
|
|||
if (i > 0) {
|
||||
out << ", ";
|
||||
}
|
||||
if (!EmitConstant(out, constant->Index(i))) {
|
||||
if (!EmitConstant(out, constant->Index(i), is_variable_initializer)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
|
@ -3139,7 +3145,7 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan
|
|||
if (i > 0) {
|
||||
out << ", ";
|
||||
}
|
||||
if (!EmitConstant(out, constant->Index(i))) {
|
||||
if (!EmitConstant(out, constant->Index(i), is_variable_initializer)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
|
@ -3147,25 +3153,45 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan
|
|||
return true;
|
||||
},
|
||||
[&](const sem::Struct* s) {
|
||||
if (!EmitStructType(&helpers_, s)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (constant->AllZero()) {
|
||||
out << "(";
|
||||
if (!EmitType(out, s, ast::AddressSpace::kNone, ast::Access::kUndefined, "")) {
|
||||
return false;
|
||||
}
|
||||
out << ")0";
|
||||
out << "(" << StructName(s) << ")0";
|
||||
return true;
|
||||
}
|
||||
|
||||
out << "{";
|
||||
TINT_DEFER(out << "}");
|
||||
|
||||
for (size_t i = 0; i < s->Members().size(); i++) {
|
||||
if (i > 0) {
|
||||
out << ", ";
|
||||
auto emit_member_values = [&](std::ostream& o) {
|
||||
o << "{";
|
||||
for (size_t i = 0; i < s->Members().size(); i++) {
|
||||
if (i > 0) {
|
||||
o << ", ";
|
||||
}
|
||||
if (!EmitConstant(o, constant->Index(i), is_variable_initializer)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
if (!EmitConstant(out, constant->Index(i))) {
|
||||
o << "}";
|
||||
return true;
|
||||
};
|
||||
|
||||
if (is_variable_initializer) {
|
||||
if (!emit_member_values(out)) {
|
||||
return false;
|
||||
}
|
||||
} else {
|
||||
// HLSL requires structure initializers to be assigned directly to a variable.
|
||||
auto name = UniqueIdentifier("c");
|
||||
{
|
||||
auto decl = line();
|
||||
decl << "const " << StructName(s) << " " << name << " = ";
|
||||
if (!emit_member_values(decl)) {
|
||||
return false;
|
||||
}
|
||||
decl << ";";
|
||||
}
|
||||
out << name;
|
||||
}
|
||||
|
||||
return true;
|
||||
|
|
@ -3891,6 +3917,11 @@ bool GeneratorImpl::EmitTypeAndName(std::ostream& out,
|
|||
}
|
||||
|
||||
bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
|
||||
auto it = emitted_structs_.emplace(str);
|
||||
if (!it.second) {
|
||||
return true;
|
||||
}
|
||||
|
||||
line(b) << "struct " << StructName(str) << " {";
|
||||
{
|
||||
ScopedIndent si(b);
|
||||
|
|
@ -3967,14 +3998,6 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
|
|||
return true;
|
||||
}
|
||||
|
||||
bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
|
||||
auto it = emitted_structs_.emplace(str);
|
||||
if (!it.second) {
|
||||
return true;
|
||||
}
|
||||
return EmitStructType(buffer, str);
|
||||
}
|
||||
|
||||
bool GeneratorImpl::EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr) {
|
||||
switch (expr->op) {
|
||||
case ast::UnaryOp::kIndirection:
|
||||
|
|
|
|||
|
|
@ -348,8 +348,12 @@ class GeneratorImpl : public TextGenerator {
|
|||
/// Handles a constant value
|
||||
/// @param out the output stream
|
||||
/// @param constant the constant value to emit
|
||||
/// @param is_variable_initializer true if the constant is used as the RHS of a variable
|
||||
/// initializer
|
||||
/// @returns true if the constant value was successfully emitted
|
||||
bool EmitConstant(std::ostream& out, const sem::Constant* constant);
|
||||
bool EmitConstant(std::ostream& out,
|
||||
const sem::Constant* constant,
|
||||
bool is_variable_initializer);
|
||||
/// Handles a literal
|
||||
/// @param out the output stream
|
||||
/// @param lit the literal to emit
|
||||
|
|
@ -420,17 +424,12 @@ class GeneratorImpl : public TextGenerator {
|
|||
ast::AddressSpace address_space,
|
||||
ast::Access access,
|
||||
const std::string& name);
|
||||
/// Handles generating a structure declaration
|
||||
/// Handles generating a structure declaration. If the structure has already been emitted, then
|
||||
/// this function will simply return `true` without emitting anything.
|
||||
/// @param buffer the text buffer that the type declaration will be written to
|
||||
/// @param ty the struct to generate
|
||||
/// @returns true if the struct is emitted
|
||||
bool EmitStructType(TextBuffer* buffer, const sem::Struct* ty);
|
||||
/// Handles generating a structure declaration only the first time called. Subsequent calls are
|
||||
/// a no-op and return true.
|
||||
/// @param buffer the text buffer that the type declaration will be written to
|
||||
/// @param ty the struct to generate
|
||||
/// @returns true if the struct is emitted
|
||||
bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty);
|
||||
/// Handles a unary op expression
|
||||
/// @param out the output of the expression stream
|
||||
/// @param expr the expression to emit
|
||||
|
|
|
|||
|
|
@ -381,9 +381,9 @@ TEST_F(HlslGeneratorImplTest_Builtin, Select_Vector) {
|
|||
EXPECT_EQ(out.str(), "(bool2(true, false) ? b : a)");
|
||||
}
|
||||
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, Modf_Scalar_f32) {
|
||||
auto* call = Call("modf", 1_f);
|
||||
WrapInFunction(CallStmt(call));
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, Runtime_Modf_Scalar_f32) {
|
||||
WrapInFunction(Decl(Let("f", Expr(1.5_f))), //
|
||||
Decl(Let("v", Call("modf", "f"))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
|
|
@ -400,17 +400,18 @@ modf_result tint_modf(float param_0) {
|
|||
|
||||
[numthreads(1, 1, 1)]
|
||||
void test_function() {
|
||||
tint_modf(1.0f);
|
||||
const float f = 1.5f;
|
||||
const modf_result v = tint_modf(f);
|
||||
return;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, Modf_Scalar_f16) {
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, Runtime_Modf_Scalar_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
auto* call = Call("modf", 1_h);
|
||||
WrapInFunction(CallStmt(call));
|
||||
WrapInFunction(Decl(Let("f", Expr(1.5_h))), //
|
||||
Decl(Let("v", Call("modf", "f"))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
|
|
@ -427,15 +428,16 @@ modf_result_f16 tint_modf(float16_t param_0) {
|
|||
|
||||
[numthreads(1, 1, 1)]
|
||||
void test_function() {
|
||||
tint_modf(float16_t(1.0h));
|
||||
const float16_t f = float16_t(1.5h);
|
||||
const modf_result_f16 v = tint_modf(f);
|
||||
return;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, Modf_Vector_f32) {
|
||||
auto* call = Call("modf", vec3<f32>());
|
||||
WrapInFunction(CallStmt(call));
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, Runtime_Modf_Vector_f32) {
|
||||
WrapInFunction(Decl(Let("f", vec3<f32>(1.5_f, 2.5_f, 3.5_f))), //
|
||||
Decl(Let("v", Call("modf", "f"))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
|
|
@ -452,17 +454,18 @@ modf_result_vec3 tint_modf(float3 param_0) {
|
|||
|
||||
[numthreads(1, 1, 1)]
|
||||
void test_function() {
|
||||
tint_modf((0.0f).xxx);
|
||||
const float3 f = float3(1.5f, 2.5f, 3.5f);
|
||||
const modf_result_vec3 v = tint_modf(f);
|
||||
return;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, Modf_Vector_f16) {
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, Runtime_Modf_Vector_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
auto* call = Call("modf", vec3<f16>());
|
||||
WrapInFunction(CallStmt(call));
|
||||
WrapInFunction(Decl(Let("f", vec3<f16>(1.5_h, 2.5_h, 3.5_h))), //
|
||||
Decl(Let("v", Call("modf", "f"))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
|
|
@ -479,7 +482,110 @@ modf_result_vec3_f16 tint_modf(vector<float16_t, 3> param_0) {
|
|||
|
||||
[numthreads(1, 1, 1)]
|
||||
void test_function() {
|
||||
tint_modf((float16_t(0.0h)).xxx);
|
||||
const vector<float16_t, 3> f = vector<float16_t, 3>(float16_t(1.5h), float16_t(2.5h), float16_t(3.5h));
|
||||
const modf_result_vec3_f16 v = tint_modf(f);
|
||||
return;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, Const_Modf_Scalar_f32) {
|
||||
WrapInFunction(Decl(Let("v", Call("modf", 1.5_f))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(struct modf_result {
|
||||
float fract;
|
||||
float whole;
|
||||
};
|
||||
[numthreads(1, 1, 1)]
|
||||
void test_function() {
|
||||
const modf_result v = {0.5f, 1.0f};
|
||||
return;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, Const_Modf_Scalar_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
WrapInFunction(Decl(Let("v", Call("modf", 1.5_h))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(struct modf_result_f16 {
|
||||
float16_t fract;
|
||||
float16_t whole;
|
||||
};
|
||||
[numthreads(1, 1, 1)]
|
||||
void test_function() {
|
||||
const modf_result_f16 v = {float16_t(0.5h), float16_t(1.0h)};
|
||||
return;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, Const_Modf_Vector_f32) {
|
||||
WrapInFunction(Decl(Let("v", Call("modf", vec3<f32>(1.5_f, 2.5_f, 3.5_f)))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(struct modf_result_vec3 {
|
||||
float3 fract;
|
||||
float3 whole;
|
||||
};
|
||||
[numthreads(1, 1, 1)]
|
||||
void test_function() {
|
||||
const modf_result_vec3 v = {(0.5f).xxx, float3(1.0f, 2.0f, 3.0f)};
|
||||
return;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, Const_Modf_Vector_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
WrapInFunction(Decl(Let("v", Call("modf", vec3<f16>(1.5_h, 2.5_h, 3.5_h)))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(struct modf_result_vec3_f16 {
|
||||
vector<float16_t, 3> fract;
|
||||
vector<float16_t, 3> whole;
|
||||
};
|
||||
[numthreads(1, 1, 1)]
|
||||
void test_function() {
|
||||
const modf_result_vec3_f16 v = {(float16_t(0.5h)).xxx, vector<float16_t, 3>(float16_t(1.0h), float16_t(2.0h), float16_t(3.0h))};
|
||||
return;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(HlslGeneratorImplTest_Builtin, NonInitializer_Modf_Vector_f32) {
|
||||
WrapInFunction(
|
||||
// Declare a variable with the result of a modf call.
|
||||
// This is required to infer the 'var' type.
|
||||
Decl(Var("v", Call("modf", vec3<f32>(1.5_f, 2.5_f, 3.5_f)))),
|
||||
// Now assign 'v' again with another modf call.
|
||||
// This requires generating a temporary variable for the struct initializer.
|
||||
Assign("v", Call("modf", vec3<f32>(4.5_f, 5.5_f, 6.5_f))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(struct modf_result_vec3 {
|
||||
float3 fract;
|
||||
float3 whole;
|
||||
};
|
||||
[numthreads(1, 1, 1)]
|
||||
void test_function() {
|
||||
modf_result_vec3 v = {(0.5f).xxx, float3(1.0f, 2.0f, 3.0f)};
|
||||
const modf_result_vec3 c = {(0.5f).xxx, float3(4.0f, 5.0f, 6.0f)};
|
||||
v = c;
|
||||
return;
|
||||
}
|
||||
)");
|
||||
|
|
|
|||
|
|
@ -907,9 +907,7 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out,
|
|||
|
||||
auto func = utils::GetOrCreate(
|
||||
atomicCompareExchangeWeak_, ACEWKeyType{{sc, str}}, [&]() -> std::string {
|
||||
// Emit the builtin return type unique to this overload. This does not
|
||||
// exist in the AST, so it will not be generated in Generate().
|
||||
if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
|
||||
if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
|
||||
return "";
|
||||
}
|
||||
|
||||
|
|
@ -1748,7 +1746,11 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan
|
|||
return true;
|
||||
},
|
||||
[&](const sem::Struct* s) {
|
||||
out << program_->Symbols().NameFor(s->Name()) << "{";
|
||||
if (!EmitStructType(&helpers_, s)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
out << StructName(s) << "{";
|
||||
TINT_DEFER(out << "}");
|
||||
|
||||
if (constant->AllZero()) {
|
||||
|
|
@ -2747,6 +2749,11 @@ bool GeneratorImpl::EmitAddressSpace(std::ostream& out, ast::AddressSpace sc) {
|
|||
}
|
||||
|
||||
bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
|
||||
auto it = emitted_structs_.emplace(str);
|
||||
if (!it.second) {
|
||||
return true;
|
||||
}
|
||||
|
||||
line(b) << "struct " << StructName(str) << " {";
|
||||
|
||||
bool is_host_shareable = str->IsHostShareable();
|
||||
|
|
@ -2903,14 +2910,6 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
|
|||
return true;
|
||||
}
|
||||
|
||||
bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
|
||||
auto it = emitted_structs_.emplace(str);
|
||||
if (!it.second) {
|
||||
return true;
|
||||
}
|
||||
return EmitStructType(buffer, str);
|
||||
}
|
||||
|
||||
bool GeneratorImpl::EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr) {
|
||||
// Handle `-e` when `e` is signed, so that we ensure that if `e` is the
|
||||
// largest negative value, it returns `e`.
|
||||
|
|
|
|||
|
|
@ -328,17 +328,12 @@ class GeneratorImpl : public TextGenerator {
|
|||
/// @param sc the address space to generate
|
||||
/// @returns true if the address space is emitted
|
||||
bool EmitAddressSpace(std::ostream& out, ast::AddressSpace sc);
|
||||
/// Handles generating a struct declaration
|
||||
/// Handles generating a struct declaration. If the structure has already been emitted, then
|
||||
/// this function will simply return `true` without emitting anything.
|
||||
/// @param buffer the text buffer that the type declaration will be written to
|
||||
/// @param str the struct to generate
|
||||
/// @returns true if the struct is emitted
|
||||
bool EmitStructType(TextBuffer* buffer, const sem::Struct* str);
|
||||
/// Handles generating a structure declaration only the first time called. Subsequent calls are
|
||||
/// a no-op and return true.
|
||||
/// @param buffer the text buffer that the type declaration will be written to
|
||||
/// @param ty the struct to generate
|
||||
/// @returns true if the struct is emitted
|
||||
bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty);
|
||||
/// Handles a unary op expression
|
||||
/// @param out the output of the expression stream
|
||||
/// @param expr the expression to emit
|
||||
|
|
|
|||
|
|
@ -405,9 +405,9 @@ TEST_F(MslGeneratorImplTest, WorkgroupBarrier) {
|
|||
EXPECT_EQ(out.str(), "threadgroup_barrier(mem_flags::mem_threadgroup)");
|
||||
}
|
||||
|
||||
TEST_F(MslGeneratorImplTest, Modf_Scalar_f32) {
|
||||
auto* call = Call("modf", 1_f);
|
||||
WrapInFunction(CallStmt(call));
|
||||
TEST_F(MslGeneratorImplTest, Runtime_Modf_Scalar_f32) {
|
||||
WrapInFunction(Decl(Let("f", Expr(1.5_f))), //
|
||||
Decl(Let("v", Call("modf", "f"))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
|
|
@ -427,18 +427,19 @@ modf_result tint_modf(float param_0) {
|
|||
}
|
||||
|
||||
kernel void test_function() {
|
||||
tint_modf(1.0f);
|
||||
float const f = 1.5f;
|
||||
modf_result const v = tint_modf(f);
|
||||
return;
|
||||
}
|
||||
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(MslGeneratorImplTest, Modf_Scalar_f16) {
|
||||
TEST_F(MslGeneratorImplTest, Runtime_Modf_Scalar_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
auto* call = Call("modf", 1_h);
|
||||
WrapInFunction(CallStmt(call));
|
||||
WrapInFunction(Decl(Let("f", Expr(1.5_h))), //
|
||||
Decl(Let("v", Call("modf", "f"))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
|
|
@ -458,16 +459,17 @@ modf_result_f16 tint_modf(half param_0) {
|
|||
}
|
||||
|
||||
kernel void test_function() {
|
||||
tint_modf(1.0h);
|
||||
half const f = 1.5h;
|
||||
modf_result_f16 const v = tint_modf(f);
|
||||
return;
|
||||
}
|
||||
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(MslGeneratorImplTest, Modf_Vector_f32) {
|
||||
auto* call = Call("modf", vec3<f32>());
|
||||
WrapInFunction(CallStmt(call));
|
||||
TEST_F(MslGeneratorImplTest, Runtime_Modf_Vector_f32) {
|
||||
WrapInFunction(Decl(Let("f", vec3<f32>(1.5_f, 2.5_f, 3.5_f))), //
|
||||
Decl(Let("v", Call("modf", "f"))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
|
|
@ -487,18 +489,19 @@ modf_result_vec3 tint_modf(float3 param_0) {
|
|||
}
|
||||
|
||||
kernel void test_function() {
|
||||
tint_modf(float3(0.0f));
|
||||
float3 const f = float3(1.5f, 2.5f, 3.5f);
|
||||
modf_result_vec3 const v = tint_modf(f);
|
||||
return;
|
||||
}
|
||||
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(MslGeneratorImplTest, Modf_Vector_f16) {
|
||||
TEST_F(MslGeneratorImplTest, Runtime_Modf_Vector_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
auto* call = Call("modf", vec3<f16>());
|
||||
WrapInFunction(CallStmt(call));
|
||||
WrapInFunction(Decl(Let("f", vec3<f16>(1.5_h, 2.5_h, 3.5_h))), //
|
||||
Decl(Let("v", Call("modf", "f"))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
|
|
@ -518,7 +521,100 @@ modf_result_vec3_f16 tint_modf(half3 param_0) {
|
|||
}
|
||||
|
||||
kernel void test_function() {
|
||||
tint_modf(half3(0.0h));
|
||||
half3 const f = half3(1.5h, 2.5h, 3.5h);
|
||||
modf_result_vec3_f16 const v = tint_modf(f);
|
||||
return;
|
||||
}
|
||||
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(MslGeneratorImplTest, Const_Modf_Scalar_f32) {
|
||||
WrapInFunction(Decl(Let("v", Call("modf", 1.5_f))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct modf_result {
|
||||
float fract;
|
||||
float whole;
|
||||
};
|
||||
kernel void test_function() {
|
||||
modf_result const v = modf_result{.fract=0.5f, .whole=1.0f};
|
||||
return;
|
||||
}
|
||||
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(MslGeneratorImplTest, Const_Modf_Scalar_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
WrapInFunction(Decl(Let("v", Call("modf", 1.5_h))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct modf_result_f16 {
|
||||
half fract;
|
||||
half whole;
|
||||
};
|
||||
kernel void test_function() {
|
||||
modf_result_f16 const v = modf_result_f16{.fract=0.5h, .whole=1.0h};
|
||||
return;
|
||||
}
|
||||
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(MslGeneratorImplTest, Const_Modf_Vector_f32) {
|
||||
WrapInFunction(Decl(Let("v", Call("modf", vec3<f32>(1.5_f, 2.5_f, 3.5_f)))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct modf_result_vec3 {
|
||||
float3 fract;
|
||||
float3 whole;
|
||||
};
|
||||
kernel void test_function() {
|
||||
modf_result_vec3 const v = modf_result_vec3{.fract=float3(0.5f), .whole=float3(1.0f, 2.0f, 3.0f)};
|
||||
return;
|
||||
}
|
||||
|
||||
)");
|
||||
}
|
||||
|
||||
TEST_F(MslGeneratorImplTest, Const_Modf_Vector_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
WrapInFunction(Decl(Let("v", Call("modf", vec3<f16>(1.5_h, 2.5_h, 3.5_h)))));
|
||||
|
||||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct modf_result_vec3_f16 {
|
||||
half3 fract;
|
||||
half3 whole;
|
||||
};
|
||||
kernel void test_function() {
|
||||
modf_result_vec3_f16 const v = modf_result_vec3_f16{.fract=half3(0.5h), .whole=half3(1.0h, 2.0h, 3.0h)};
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -1630,7 +1630,7 @@ OpFunctionEnd
|
|||
EXPECT_EQ(got, expect);
|
||||
}
|
||||
|
||||
TEST_F(BuiltinBuilderTest, Call_Modf_f32) {
|
||||
TEST_F(BuiltinBuilderTest, Runtime_Call_Modf_f32) {
|
||||
auto* vec = Var("vec", vec2<f32>(1_f, 2_f));
|
||||
auto* expr = Call("modf", vec);
|
||||
Func("a_func", utils::Empty, ty.void_(),
|
||||
|
|
@ -1682,7 +1682,7 @@ OpFunctionEnd
|
|||
Validate(b);
|
||||
}
|
||||
|
||||
TEST_F(BuiltinBuilderTest, Call_Modf_f16) {
|
||||
TEST_F(BuiltinBuilderTest, Runtime_Call_Modf_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
auto* vec = Var("vec", vec2<f16>(1_h, 2_h));
|
||||
|
|
@ -1740,6 +1740,100 @@ OpFunctionEnd
|
|||
Validate(b);
|
||||
}
|
||||
|
||||
TEST_F(BuiltinBuilderTest, Const_Call_Modf_f32) {
|
||||
auto* expr = Call("modf", vec2<f32>(1_f, 2_f));
|
||||
Func("a_func", utils::Empty, ty.void_(),
|
||||
utils::Vector{
|
||||
CallStmt(expr),
|
||||
},
|
||||
utils::Vector{
|
||||
Stage(ast::PipelineStage::kFragment),
|
||||
});
|
||||
|
||||
spirv::Builder& b = Build();
|
||||
|
||||
ASSERT_TRUE(b.Build()) << b.error();
|
||||
auto got = DumpBuilder(b);
|
||||
auto* expect = R"(OpCapability Shader
|
||||
%9 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %3 "a_func"
|
||||
OpExecutionMode %3 OriginUpperLeft
|
||||
OpName %3 "a_func"
|
||||
OpName %6 "__modf_result_vec2"
|
||||
OpMemberName %6 0 "fract"
|
||||
OpMemberName %6 1 "whole"
|
||||
OpMemberDecorate %6 0 Offset 0
|
||||
OpMemberDecorate %6 1 Offset 8
|
||||
%2 = OpTypeVoid
|
||||
%1 = OpTypeFunction %2
|
||||
%8 = OpTypeFloat 32
|
||||
%7 = OpTypeVector %8 2
|
||||
%6 = OpTypeStruct %7 %7
|
||||
%10 = OpConstant %8 1
|
||||
%11 = OpConstant %8 2
|
||||
%12 = OpConstantComposite %7 %10 %11
|
||||
%3 = OpFunction %2 None %1
|
||||
%4 = OpLabel
|
||||
%5 = OpExtInst %6 %9 ModfStruct %12
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
)";
|
||||
EXPECT_EQ(expect, got);
|
||||
|
||||
Validate(b);
|
||||
}
|
||||
|
||||
TEST_F(BuiltinBuilderTest, Const_Call_Modf_f16) {
|
||||
Enable(ast::Extension::kF16);
|
||||
|
||||
auto* expr = Call("modf", vec2<f16>(1_h, 2_h));
|
||||
Func("a_func", utils::Empty, ty.void_(),
|
||||
utils::Vector{
|
||||
CallStmt(expr),
|
||||
},
|
||||
utils::Vector{
|
||||
Stage(ast::PipelineStage::kFragment),
|
||||
});
|
||||
|
||||
spirv::Builder& b = Build();
|
||||
|
||||
ASSERT_TRUE(b.Build()) << b.error();
|
||||
auto got = DumpBuilder(b);
|
||||
auto* expect = R"(OpCapability Shader
|
||||
OpCapability Float16
|
||||
OpCapability UniformAndStorageBuffer16BitAccess
|
||||
OpCapability StorageBuffer16BitAccess
|
||||
OpCapability StorageInputOutput16
|
||||
%9 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %3 "a_func"
|
||||
OpExecutionMode %3 OriginUpperLeft
|
||||
OpName %3 "a_func"
|
||||
OpName %6 "__modf_result_vec2_f16"
|
||||
OpMemberName %6 0 "fract"
|
||||
OpMemberName %6 1 "whole"
|
||||
OpMemberDecorate %6 0 Offset 0
|
||||
OpMemberDecorate %6 1 Offset 4
|
||||
%2 = OpTypeVoid
|
||||
%1 = OpTypeFunction %2
|
||||
%8 = OpTypeFloat 16
|
||||
%7 = OpTypeVector %8 2
|
||||
%6 = OpTypeStruct %7 %7
|
||||
%10 = OpConstant %8 0x1p+0
|
||||
%11 = OpConstant %8 0x1p+1
|
||||
%12 = OpConstantComposite %7 %10 %11
|
||||
%3 = OpFunction %2 None %1
|
||||
%4 = OpLabel
|
||||
%5 = OpExtInst %6 %9 ModfStruct %12
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
)";
|
||||
EXPECT_EQ(expect, got);
|
||||
|
||||
Validate(b);
|
||||
}
|
||||
|
||||
TEST_F(BuiltinBuilderTest, Call_Frexp_f32) {
|
||||
auto* vec = Var("vec", vec2<f32>(1_f, 2_f));
|
||||
auto* expr = Call("frexp", vec);
|
||||
|
|
|
|||
|
|
@ -1,18 +1,8 @@
|
|||
struct modf_result {
|
||||
float fract;
|
||||
float whole;
|
||||
};
|
||||
modf_result tint_modf(float param_0) {
|
||||
modf_result result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void unused_entry_point() {
|
||||
return;
|
||||
}
|
||||
|
||||
void i() {
|
||||
const float s = tint_modf(1.0f).whole;
|
||||
const float s = 1.0f;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1,18 +1,8 @@
|
|||
struct modf_result {
|
||||
float fract;
|
||||
float whole;
|
||||
};
|
||||
modf_result tint_modf(float param_0) {
|
||||
modf_result result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void unused_entry_point() {
|
||||
return;
|
||||
}
|
||||
|
||||
void i() {
|
||||
const float s = tint_modf(1.0f).whole;
|
||||
const float s = 1.0f;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1,22 +1,10 @@
|
|||
#version 310 es
|
||||
|
||||
struct modf_result {
|
||||
float fract;
|
||||
float whole;
|
||||
};
|
||||
|
||||
modf_result tint_modf(float param_0) {
|
||||
modf_result result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
void unused_entry_point() {
|
||||
return;
|
||||
}
|
||||
void i() {
|
||||
float s = tint_modf(1.0f).whole;
|
||||
float s = 1.0f;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -1,18 +1,7 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct modf_result {
|
||||
float fract;
|
||||
float whole;
|
||||
};
|
||||
modf_result tint_modf(float param_0) {
|
||||
modf_result result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void i() {
|
||||
float const s = tint_modf(1.0f).whole;
|
||||
float const s = 1.0f;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -1,24 +1,17 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 13
|
||||
; Bound: 9
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%10 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
|
||||
OpExecutionMode %unused_entry_point LocalSize 1 1 1
|
||||
OpName %unused_entry_point "unused_entry_point"
|
||||
OpName %tint_symbol "tint_symbol"
|
||||
OpName %__modf_result "__modf_result"
|
||||
OpMemberName %__modf_result 0 "fract"
|
||||
OpMemberName %__modf_result 1 "whole"
|
||||
OpMemberDecorate %__modf_result 0 Offset 0
|
||||
OpMemberDecorate %__modf_result 1 Offset 4
|
||||
%void = OpTypeVoid
|
||||
%1 = OpTypeFunction %void
|
||||
%float = OpTypeFloat 32
|
||||
%__modf_result = OpTypeStruct %float %float
|
||||
%float_1 = OpConstant %float 1
|
||||
%unused_entry_point = OpFunction %void None %1
|
||||
%4 = OpLabel
|
||||
|
|
@ -26,7 +19,5 @@
|
|||
OpFunctionEnd
|
||||
%tint_symbol = OpFunction %void None %1
|
||||
%6 = OpLabel
|
||||
%7 = OpExtInst %__modf_result %10 ModfStruct %float_1
|
||||
%12 = OpCompositeExtract %float %7 1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
|
|||
|
|
@ -23,7 +23,7 @@
|
|||
|
||||
// fn modf(vec<2, f32>) -> __modf_result_vec<2, f32>
|
||||
fn modf_2d50da() {
|
||||
var res = modf(vec2<f32>(1.f));
|
||||
var res = modf(vec2<f32>(-1.5f));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -2,14 +2,8 @@ struct modf_result_vec2 {
|
|||
float2 fract;
|
||||
float2 whole;
|
||||
};
|
||||
modf_result_vec2 tint_modf(float2 param_0) {
|
||||
modf_result_vec2 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_2d50da() {
|
||||
modf_result_vec2 res = tint_modf((1.0f).xx);
|
||||
modf_result_vec2 res = {(-0.5f).xx, (-1.0f).xx};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -2,14 +2,8 @@ struct modf_result_vec2 {
|
|||
float2 fract;
|
||||
float2 whole;
|
||||
};
|
||||
modf_result_vec2 tint_modf(float2 param_0) {
|
||||
modf_result_vec2 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_2d50da() {
|
||||
modf_result_vec2 res = tint_modf((1.0f).xx);
|
||||
modf_result_vec2 res = {(-0.5f).xx, (-1.0f).xx};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -5,15 +5,9 @@ struct modf_result_vec2 {
|
|||
vec2 whole;
|
||||
};
|
||||
|
||||
modf_result_vec2 tint_modf(vec2 param_0) {
|
||||
modf_result_vec2 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_2d50da() {
|
||||
modf_result_vec2 res = tint_modf(vec2(1.0f));
|
||||
modf_result_vec2 res = modf_result_vec2(vec2(-0.5f), vec2(-1.0f));
|
||||
}
|
||||
|
||||
vec4 vertex_main() {
|
||||
|
|
@ -37,15 +31,9 @@ struct modf_result_vec2 {
|
|||
vec2 whole;
|
||||
};
|
||||
|
||||
modf_result_vec2 tint_modf(vec2 param_0) {
|
||||
modf_result_vec2 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_2d50da() {
|
||||
modf_result_vec2 res = tint_modf(vec2(1.0f));
|
||||
modf_result_vec2 res = modf_result_vec2(vec2(-0.5f), vec2(-1.0f));
|
||||
}
|
||||
|
||||
void fragment_main() {
|
||||
|
|
@ -63,15 +51,9 @@ struct modf_result_vec2 {
|
|||
vec2 whole;
|
||||
};
|
||||
|
||||
modf_result_vec2 tint_modf(vec2 param_0) {
|
||||
modf_result_vec2 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_2d50da() {
|
||||
modf_result_vec2 res = tint_modf(vec2(1.0f));
|
||||
modf_result_vec2 res = modf_result_vec2(vec2(-0.5f), vec2(-1.0f));
|
||||
}
|
||||
|
||||
void compute_main() {
|
||||
|
|
|
|||
|
|
@ -6,14 +6,8 @@ struct modf_result_vec2 {
|
|||
float2 fract;
|
||||
float2 whole;
|
||||
};
|
||||
modf_result_vec2 tint_modf(float2 param_0) {
|
||||
modf_result_vec2 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_2d50da() {
|
||||
modf_result_vec2 res = tint_modf(float2(1.0f));
|
||||
modf_result_vec2 res = modf_result_vec2{.fract=float2(-0.5f), .whole=float2(-1.0f)};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -1,10 +1,9 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 35
|
||||
; Bound: 37
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%16 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
|
||||
OpEntryPoint Fragment %fragment_main "fragment_main"
|
||||
|
|
@ -38,37 +37,40 @@
|
|||
%9 = OpTypeFunction %void
|
||||
%v2float = OpTypeVector %float 2
|
||||
%__modf_result_vec2 = OpTypeStruct %v2float %v2float
|
||||
%float_1 = OpConstant %float 1
|
||||
%18 = OpConstantComposite %v2float %float_1 %float_1
|
||||
%float_n0_5 = OpConstant %float -0.5
|
||||
%16 = OpConstantComposite %v2float %float_n0_5 %float_n0_5
|
||||
%float_n1 = OpConstant %float -1
|
||||
%18 = OpConstantComposite %v2float %float_n1 %float_n1
|
||||
%19 = OpConstantComposite %__modf_result_vec2 %16 %18
|
||||
%_ptr_Function___modf_result_vec2 = OpTypePointer Function %__modf_result_vec2
|
||||
%21 = OpConstantNull %__modf_result_vec2
|
||||
%22 = OpTypeFunction %v4float
|
||||
%22 = OpConstantNull %__modf_result_vec2
|
||||
%23 = OpTypeFunction %v4float
|
||||
%float_1 = OpConstant %float 1
|
||||
%modf_2d50da = OpFunction %void None %9
|
||||
%12 = OpLabel
|
||||
%res = OpVariable %_ptr_Function___modf_result_vec2 Function %21
|
||||
%13 = OpExtInst %__modf_result_vec2 %16 ModfStruct %18
|
||||
OpStore %res %13
|
||||
%res = OpVariable %_ptr_Function___modf_result_vec2 Function %22
|
||||
OpStore %res %19
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%vertex_main_inner = OpFunction %v4float None %22
|
||||
%24 = OpLabel
|
||||
%25 = OpFunctionCall %void %modf_2d50da
|
||||
%vertex_main_inner = OpFunction %v4float None %23
|
||||
%25 = OpLabel
|
||||
%26 = OpFunctionCall %void %modf_2d50da
|
||||
OpReturnValue %5
|
||||
OpFunctionEnd
|
||||
%vertex_main = OpFunction %void None %9
|
||||
%27 = OpLabel
|
||||
%28 = OpFunctionCall %v4float %vertex_main_inner
|
||||
OpStore %value %28
|
||||
%28 = OpLabel
|
||||
%29 = OpFunctionCall %v4float %vertex_main_inner
|
||||
OpStore %value %29
|
||||
OpStore %vertex_point_size %float_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%fragment_main = OpFunction %void None %9
|
||||
%30 = OpLabel
|
||||
%31 = OpFunctionCall %void %modf_2d50da
|
||||
%32 = OpLabel
|
||||
%33 = OpFunctionCall %void %modf_2d50da
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %9
|
||||
%33 = OpLabel
|
||||
%34 = OpFunctionCall %void %modf_2d50da
|
||||
%35 = OpLabel
|
||||
%36 = OpFunctionCall %void %modf_2d50da
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
|
|||
|
|
@ -1,5 +1,5 @@
|
|||
fn modf_2d50da() {
|
||||
var res = modf(vec2<f32>(1.0f));
|
||||
var res = modf(vec2<f32>(-(1.5f)));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -25,7 +25,7 @@ enable f16;
|
|||
|
||||
// fn modf(vec<3, f16>) -> __modf_result_vec<3, f16>
|
||||
fn modf_45005f() {
|
||||
var res = modf(vec3<f16>(1.h));
|
||||
var res = modf(vec3<f16>(-1.5h));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -2,14 +2,8 @@ struct modf_result_vec3_f16 {
|
|||
vector<float16_t, 3> fract;
|
||||
vector<float16_t, 3> whole;
|
||||
};
|
||||
modf_result_vec3_f16 tint_modf(vector<float16_t, 3> param_0) {
|
||||
modf_result_vec3_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_45005f() {
|
||||
modf_result_vec3_f16 res = tint_modf((float16_t(1.0h)).xxx);
|
||||
modf_result_vec3_f16 res = {(float16_t(-0.5h)).xxx, (float16_t(-1.0h)).xxx};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -6,15 +6,9 @@ struct modf_result_vec3_f16 {
|
|||
f16vec3 whole;
|
||||
};
|
||||
|
||||
modf_result_vec3_f16 tint_modf(f16vec3 param_0) {
|
||||
modf_result_vec3_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_45005f() {
|
||||
modf_result_vec3_f16 res = tint_modf(f16vec3(1.0hf));
|
||||
modf_result_vec3_f16 res = modf_result_vec3_f16(f16vec3(-0.5hf), f16vec3(-1.0hf));
|
||||
}
|
||||
|
||||
vec4 vertex_main() {
|
||||
|
|
@ -39,15 +33,9 @@ struct modf_result_vec3_f16 {
|
|||
f16vec3 whole;
|
||||
};
|
||||
|
||||
modf_result_vec3_f16 tint_modf(f16vec3 param_0) {
|
||||
modf_result_vec3_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_45005f() {
|
||||
modf_result_vec3_f16 res = tint_modf(f16vec3(1.0hf));
|
||||
modf_result_vec3_f16 res = modf_result_vec3_f16(f16vec3(-0.5hf), f16vec3(-1.0hf));
|
||||
}
|
||||
|
||||
void fragment_main() {
|
||||
|
|
@ -66,15 +54,9 @@ struct modf_result_vec3_f16 {
|
|||
f16vec3 whole;
|
||||
};
|
||||
|
||||
modf_result_vec3_f16 tint_modf(f16vec3 param_0) {
|
||||
modf_result_vec3_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_45005f() {
|
||||
modf_result_vec3_f16 res = tint_modf(f16vec3(1.0hf));
|
||||
modf_result_vec3_f16 res = modf_result_vec3_f16(f16vec3(-0.5hf), f16vec3(-1.0hf));
|
||||
}
|
||||
|
||||
void compute_main() {
|
||||
|
|
|
|||
|
|
@ -6,14 +6,8 @@ struct modf_result_vec3_f16 {
|
|||
half3 fract;
|
||||
half3 whole;
|
||||
};
|
||||
modf_result_vec3_f16 tint_modf(half3 param_0) {
|
||||
modf_result_vec3_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_45005f() {
|
||||
modf_result_vec3_f16 res = tint_modf(half3(1.0h));
|
||||
modf_result_vec3_f16 res = modf_result_vec3_f16{.fract=half3(-0.5h), .whole=half3(-1.0h)};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -1,14 +1,13 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 37
|
||||
; Bound: 38
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability Float16
|
||||
OpCapability UniformAndStorageBuffer16BitAccess
|
||||
OpCapability StorageBuffer16BitAccess
|
||||
OpCapability StorageInputOutput16
|
||||
%17 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
|
||||
OpEntryPoint Fragment %fragment_main "fragment_main"
|
||||
|
|
@ -43,38 +42,40 @@
|
|||
%half = OpTypeFloat 16
|
||||
%v3half = OpTypeVector %half 3
|
||||
%__modf_result_vec3_f16 = OpTypeStruct %v3half %v3half
|
||||
%half_0x1p_0 = OpConstant %half 0x1p+0
|
||||
%19 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
|
||||
%half_n0x1pn1 = OpConstant %half -0x1p-1
|
||||
%17 = OpConstantComposite %v3half %half_n0x1pn1 %half_n0x1pn1 %half_n0x1pn1
|
||||
%half_n0x1p_0 = OpConstant %half -0x1p+0
|
||||
%19 = OpConstantComposite %v3half %half_n0x1p_0 %half_n0x1p_0 %half_n0x1p_0
|
||||
%20 = OpConstantComposite %__modf_result_vec3_f16 %17 %19
|
||||
%_ptr_Function___modf_result_vec3_f16 = OpTypePointer Function %__modf_result_vec3_f16
|
||||
%22 = OpConstantNull %__modf_result_vec3_f16
|
||||
%23 = OpTypeFunction %v4float
|
||||
%23 = OpConstantNull %__modf_result_vec3_f16
|
||||
%24 = OpTypeFunction %v4float
|
||||
%float_1 = OpConstant %float 1
|
||||
%modf_45005f = OpFunction %void None %9
|
||||
%12 = OpLabel
|
||||
%res = OpVariable %_ptr_Function___modf_result_vec3_f16 Function %22
|
||||
%13 = OpExtInst %__modf_result_vec3_f16 %17 ModfStruct %19
|
||||
OpStore %res %13
|
||||
%res = OpVariable %_ptr_Function___modf_result_vec3_f16 Function %23
|
||||
OpStore %res %20
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%vertex_main_inner = OpFunction %v4float None %23
|
||||
%25 = OpLabel
|
||||
%26 = OpFunctionCall %void %modf_45005f
|
||||
%vertex_main_inner = OpFunction %v4float None %24
|
||||
%26 = OpLabel
|
||||
%27 = OpFunctionCall %void %modf_45005f
|
||||
OpReturnValue %5
|
||||
OpFunctionEnd
|
||||
%vertex_main = OpFunction %void None %9
|
||||
%28 = OpLabel
|
||||
%29 = OpFunctionCall %v4float %vertex_main_inner
|
||||
OpStore %value %29
|
||||
%29 = OpLabel
|
||||
%30 = OpFunctionCall %v4float %vertex_main_inner
|
||||
OpStore %value %30
|
||||
OpStore %vertex_point_size %float_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%fragment_main = OpFunction %void None %9
|
||||
%32 = OpLabel
|
||||
%33 = OpFunctionCall %void %modf_45005f
|
||||
%33 = OpLabel
|
||||
%34 = OpFunctionCall %void %modf_45005f
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %9
|
||||
%35 = OpLabel
|
||||
%36 = OpFunctionCall %void %modf_45005f
|
||||
%36 = OpLabel
|
||||
%37 = OpFunctionCall %void %modf_45005f
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
|
|||
|
|
@ -1,7 +1,7 @@
|
|||
enable f16;
|
||||
|
||||
fn modf_45005f() {
|
||||
var res = modf(vec3<f16>(1.0h));
|
||||
var res = modf(vec3<f16>(-(1.5h)));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -23,7 +23,7 @@
|
|||
|
||||
// fn modf(vec<4, f32>) -> __modf_result_vec<4, f32>
|
||||
fn modf_4bfced() {
|
||||
var res = modf(vec4<f32>(1.f));
|
||||
var res = modf(vec4<f32>(-1.5f));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -2,14 +2,8 @@ struct modf_result_vec4 {
|
|||
float4 fract;
|
||||
float4 whole;
|
||||
};
|
||||
modf_result_vec4 tint_modf(float4 param_0) {
|
||||
modf_result_vec4 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_4bfced() {
|
||||
modf_result_vec4 res = tint_modf((1.0f).xxxx);
|
||||
modf_result_vec4 res = {(-0.5f).xxxx, (-1.0f).xxxx};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -2,14 +2,8 @@ struct modf_result_vec4 {
|
|||
float4 fract;
|
||||
float4 whole;
|
||||
};
|
||||
modf_result_vec4 tint_modf(float4 param_0) {
|
||||
modf_result_vec4 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_4bfced() {
|
||||
modf_result_vec4 res = tint_modf((1.0f).xxxx);
|
||||
modf_result_vec4 res = {(-0.5f).xxxx, (-1.0f).xxxx};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -5,15 +5,9 @@ struct modf_result_vec4 {
|
|||
vec4 whole;
|
||||
};
|
||||
|
||||
modf_result_vec4 tint_modf(vec4 param_0) {
|
||||
modf_result_vec4 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_4bfced() {
|
||||
modf_result_vec4 res = tint_modf(vec4(1.0f));
|
||||
modf_result_vec4 res = modf_result_vec4(vec4(-0.5f), vec4(-1.0f));
|
||||
}
|
||||
|
||||
vec4 vertex_main() {
|
||||
|
|
@ -37,15 +31,9 @@ struct modf_result_vec4 {
|
|||
vec4 whole;
|
||||
};
|
||||
|
||||
modf_result_vec4 tint_modf(vec4 param_0) {
|
||||
modf_result_vec4 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_4bfced() {
|
||||
modf_result_vec4 res = tint_modf(vec4(1.0f));
|
||||
modf_result_vec4 res = modf_result_vec4(vec4(-0.5f), vec4(-1.0f));
|
||||
}
|
||||
|
||||
void fragment_main() {
|
||||
|
|
@ -63,15 +51,9 @@ struct modf_result_vec4 {
|
|||
vec4 whole;
|
||||
};
|
||||
|
||||
modf_result_vec4 tint_modf(vec4 param_0) {
|
||||
modf_result_vec4 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_4bfced() {
|
||||
modf_result_vec4 res = tint_modf(vec4(1.0f));
|
||||
modf_result_vec4 res = modf_result_vec4(vec4(-0.5f), vec4(-1.0f));
|
||||
}
|
||||
|
||||
void compute_main() {
|
||||
|
|
|
|||
|
|
@ -6,14 +6,8 @@ struct modf_result_vec4 {
|
|||
float4 fract;
|
||||
float4 whole;
|
||||
};
|
||||
modf_result_vec4 tint_modf(float4 param_0) {
|
||||
modf_result_vec4 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_4bfced() {
|
||||
modf_result_vec4 res = tint_modf(float4(1.0f));
|
||||
modf_result_vec4 res = modf_result_vec4{.fract=float4(-0.5f), .whole=float4(-1.0f)};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -1,10 +1,9 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 34
|
||||
; Bound: 36
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%15 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
|
||||
OpEntryPoint Fragment %fragment_main "fragment_main"
|
||||
|
|
@ -37,37 +36,40 @@
|
|||
%void = OpTypeVoid
|
||||
%9 = OpTypeFunction %void
|
||||
%__modf_result_vec4 = OpTypeStruct %v4float %v4float
|
||||
%float_1 = OpConstant %float 1
|
||||
%17 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
|
||||
%float_n0_5 = OpConstant %float -0.5
|
||||
%15 = OpConstantComposite %v4float %float_n0_5 %float_n0_5 %float_n0_5 %float_n0_5
|
||||
%float_n1 = OpConstant %float -1
|
||||
%17 = OpConstantComposite %v4float %float_n1 %float_n1 %float_n1 %float_n1
|
||||
%18 = OpConstantComposite %__modf_result_vec4 %15 %17
|
||||
%_ptr_Function___modf_result_vec4 = OpTypePointer Function %__modf_result_vec4
|
||||
%20 = OpConstantNull %__modf_result_vec4
|
||||
%21 = OpTypeFunction %v4float
|
||||
%21 = OpConstantNull %__modf_result_vec4
|
||||
%22 = OpTypeFunction %v4float
|
||||
%float_1 = OpConstant %float 1
|
||||
%modf_4bfced = OpFunction %void None %9
|
||||
%12 = OpLabel
|
||||
%res = OpVariable %_ptr_Function___modf_result_vec4 Function %20
|
||||
%13 = OpExtInst %__modf_result_vec4 %15 ModfStruct %17
|
||||
OpStore %res %13
|
||||
%res = OpVariable %_ptr_Function___modf_result_vec4 Function %21
|
||||
OpStore %res %18
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%vertex_main_inner = OpFunction %v4float None %21
|
||||
%23 = OpLabel
|
||||
%24 = OpFunctionCall %void %modf_4bfced
|
||||
%vertex_main_inner = OpFunction %v4float None %22
|
||||
%24 = OpLabel
|
||||
%25 = OpFunctionCall %void %modf_4bfced
|
||||
OpReturnValue %5
|
||||
OpFunctionEnd
|
||||
%vertex_main = OpFunction %void None %9
|
||||
%26 = OpLabel
|
||||
%27 = OpFunctionCall %v4float %vertex_main_inner
|
||||
OpStore %value %27
|
||||
%27 = OpLabel
|
||||
%28 = OpFunctionCall %v4float %vertex_main_inner
|
||||
OpStore %value %28
|
||||
OpStore %vertex_point_size %float_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%fragment_main = OpFunction %void None %9
|
||||
%29 = OpLabel
|
||||
%30 = OpFunctionCall %void %modf_4bfced
|
||||
%31 = OpLabel
|
||||
%32 = OpFunctionCall %void %modf_4bfced
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %9
|
||||
%32 = OpLabel
|
||||
%33 = OpFunctionCall %void %modf_4bfced
|
||||
%34 = OpLabel
|
||||
%35 = OpFunctionCall %void %modf_4bfced
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
|
|||
|
|
@ -1,5 +1,5 @@
|
|||
fn modf_4bfced() {
|
||||
var res = modf(vec4<f32>(1.0f));
|
||||
var res = modf(vec4<f32>(-(1.5f)));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -23,7 +23,7 @@
|
|||
|
||||
// fn modf(vec<3, f32>) -> __modf_result_vec<3, f32>
|
||||
fn modf_5ea256() {
|
||||
var res = modf(vec3<f32>(1.f));
|
||||
var res = modf(vec3<f32>(-1.5f));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -2,14 +2,8 @@ struct modf_result_vec3 {
|
|||
float3 fract;
|
||||
float3 whole;
|
||||
};
|
||||
modf_result_vec3 tint_modf(float3 param_0) {
|
||||
modf_result_vec3 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_5ea256() {
|
||||
modf_result_vec3 res = tint_modf((1.0f).xxx);
|
||||
modf_result_vec3 res = {(-0.5f).xxx, (-1.0f).xxx};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -2,14 +2,8 @@ struct modf_result_vec3 {
|
|||
float3 fract;
|
||||
float3 whole;
|
||||
};
|
||||
modf_result_vec3 tint_modf(float3 param_0) {
|
||||
modf_result_vec3 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_5ea256() {
|
||||
modf_result_vec3 res = tint_modf((1.0f).xxx);
|
||||
modf_result_vec3 res = {(-0.5f).xxx, (-1.0f).xxx};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -5,15 +5,9 @@ struct modf_result_vec3 {
|
|||
vec3 whole;
|
||||
};
|
||||
|
||||
modf_result_vec3 tint_modf(vec3 param_0) {
|
||||
modf_result_vec3 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_5ea256() {
|
||||
modf_result_vec3 res = tint_modf(vec3(1.0f));
|
||||
modf_result_vec3 res = modf_result_vec3(vec3(-0.5f), vec3(-1.0f));
|
||||
}
|
||||
|
||||
vec4 vertex_main() {
|
||||
|
|
@ -37,15 +31,9 @@ struct modf_result_vec3 {
|
|||
vec3 whole;
|
||||
};
|
||||
|
||||
modf_result_vec3 tint_modf(vec3 param_0) {
|
||||
modf_result_vec3 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_5ea256() {
|
||||
modf_result_vec3 res = tint_modf(vec3(1.0f));
|
||||
modf_result_vec3 res = modf_result_vec3(vec3(-0.5f), vec3(-1.0f));
|
||||
}
|
||||
|
||||
void fragment_main() {
|
||||
|
|
@ -63,15 +51,9 @@ struct modf_result_vec3 {
|
|||
vec3 whole;
|
||||
};
|
||||
|
||||
modf_result_vec3 tint_modf(vec3 param_0) {
|
||||
modf_result_vec3 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_5ea256() {
|
||||
modf_result_vec3 res = tint_modf(vec3(1.0f));
|
||||
modf_result_vec3 res = modf_result_vec3(vec3(-0.5f), vec3(-1.0f));
|
||||
}
|
||||
|
||||
void compute_main() {
|
||||
|
|
|
|||
|
|
@ -6,14 +6,8 @@ struct modf_result_vec3 {
|
|||
float3 fract;
|
||||
float3 whole;
|
||||
};
|
||||
modf_result_vec3 tint_modf(float3 param_0) {
|
||||
modf_result_vec3 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_5ea256() {
|
||||
modf_result_vec3 res = tint_modf(float3(1.0f));
|
||||
modf_result_vec3 res = modf_result_vec3{.fract=float3(-0.5f), .whole=float3(-1.0f)};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -1,10 +1,9 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 35
|
||||
; Bound: 37
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%16 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
|
||||
OpEntryPoint Fragment %fragment_main "fragment_main"
|
||||
|
|
@ -38,37 +37,40 @@
|
|||
%9 = OpTypeFunction %void
|
||||
%v3float = OpTypeVector %float 3
|
||||
%__modf_result_vec3 = OpTypeStruct %v3float %v3float
|
||||
%float_1 = OpConstant %float 1
|
||||
%18 = OpConstantComposite %v3float %float_1 %float_1 %float_1
|
||||
%float_n0_5 = OpConstant %float -0.5
|
||||
%16 = OpConstantComposite %v3float %float_n0_5 %float_n0_5 %float_n0_5
|
||||
%float_n1 = OpConstant %float -1
|
||||
%18 = OpConstantComposite %v3float %float_n1 %float_n1 %float_n1
|
||||
%19 = OpConstantComposite %__modf_result_vec3 %16 %18
|
||||
%_ptr_Function___modf_result_vec3 = OpTypePointer Function %__modf_result_vec3
|
||||
%21 = OpConstantNull %__modf_result_vec3
|
||||
%22 = OpTypeFunction %v4float
|
||||
%22 = OpConstantNull %__modf_result_vec3
|
||||
%23 = OpTypeFunction %v4float
|
||||
%float_1 = OpConstant %float 1
|
||||
%modf_5ea256 = OpFunction %void None %9
|
||||
%12 = OpLabel
|
||||
%res = OpVariable %_ptr_Function___modf_result_vec3 Function %21
|
||||
%13 = OpExtInst %__modf_result_vec3 %16 ModfStruct %18
|
||||
OpStore %res %13
|
||||
%res = OpVariable %_ptr_Function___modf_result_vec3 Function %22
|
||||
OpStore %res %19
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%vertex_main_inner = OpFunction %v4float None %22
|
||||
%24 = OpLabel
|
||||
%25 = OpFunctionCall %void %modf_5ea256
|
||||
%vertex_main_inner = OpFunction %v4float None %23
|
||||
%25 = OpLabel
|
||||
%26 = OpFunctionCall %void %modf_5ea256
|
||||
OpReturnValue %5
|
||||
OpFunctionEnd
|
||||
%vertex_main = OpFunction %void None %9
|
||||
%27 = OpLabel
|
||||
%28 = OpFunctionCall %v4float %vertex_main_inner
|
||||
OpStore %value %28
|
||||
%28 = OpLabel
|
||||
%29 = OpFunctionCall %v4float %vertex_main_inner
|
||||
OpStore %value %29
|
||||
OpStore %vertex_point_size %float_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%fragment_main = OpFunction %void None %9
|
||||
%30 = OpLabel
|
||||
%31 = OpFunctionCall %void %modf_5ea256
|
||||
%32 = OpLabel
|
||||
%33 = OpFunctionCall %void %modf_5ea256
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %9
|
||||
%33 = OpLabel
|
||||
%34 = OpFunctionCall %void %modf_5ea256
|
||||
%35 = OpLabel
|
||||
%36 = OpFunctionCall %void %modf_5ea256
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
|
|||
|
|
@ -1,5 +1,5 @@
|
|||
fn modf_5ea256() {
|
||||
var res = modf(vec3<f32>(1.0f));
|
||||
var res = modf(vec3<f32>(-(1.5f)));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -25,7 +25,7 @@ enable f16;
|
|||
|
||||
// fn modf(f16) -> __modf_result<f16>
|
||||
fn modf_8dbbbf() {
|
||||
var res = modf(1.h);
|
||||
var res = modf(-1.5h);
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -2,14 +2,8 @@ struct modf_result_f16 {
|
|||
float16_t fract;
|
||||
float16_t whole;
|
||||
};
|
||||
modf_result_f16 tint_modf(float16_t param_0) {
|
||||
modf_result_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_8dbbbf() {
|
||||
modf_result_f16 res = tint_modf(float16_t(1.0h));
|
||||
modf_result_f16 res = {float16_t(-0.5h), float16_t(-1.0h)};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -6,15 +6,9 @@ struct modf_result_f16 {
|
|||
float16_t whole;
|
||||
};
|
||||
|
||||
modf_result_f16 tint_modf(float16_t param_0) {
|
||||
modf_result_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_8dbbbf() {
|
||||
modf_result_f16 res = tint_modf(1.0hf);
|
||||
modf_result_f16 res = modf_result_f16(-0.5hf, -1.0hf);
|
||||
}
|
||||
|
||||
vec4 vertex_main() {
|
||||
|
|
@ -39,15 +33,9 @@ struct modf_result_f16 {
|
|||
float16_t whole;
|
||||
};
|
||||
|
||||
modf_result_f16 tint_modf(float16_t param_0) {
|
||||
modf_result_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_8dbbbf() {
|
||||
modf_result_f16 res = tint_modf(1.0hf);
|
||||
modf_result_f16 res = modf_result_f16(-0.5hf, -1.0hf);
|
||||
}
|
||||
|
||||
void fragment_main() {
|
||||
|
|
@ -66,15 +54,9 @@ struct modf_result_f16 {
|
|||
float16_t whole;
|
||||
};
|
||||
|
||||
modf_result_f16 tint_modf(float16_t param_0) {
|
||||
modf_result_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_8dbbbf() {
|
||||
modf_result_f16 res = tint_modf(1.0hf);
|
||||
modf_result_f16 res = modf_result_f16(-0.5hf, -1.0hf);
|
||||
}
|
||||
|
||||
void compute_main() {
|
||||
|
|
|
|||
|
|
@ -6,14 +6,8 @@ struct modf_result_f16 {
|
|||
half fract;
|
||||
half whole;
|
||||
};
|
||||
modf_result_f16 tint_modf(half param_0) {
|
||||
modf_result_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_8dbbbf() {
|
||||
modf_result_f16 res = tint_modf(1.0h);
|
||||
modf_result_f16 res = modf_result_f16{.fract=-0.5h, .whole=-1.0h};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -8,7 +8,6 @@
|
|||
OpCapability UniformAndStorageBuffer16BitAccess
|
||||
OpCapability StorageBuffer16BitAccess
|
||||
OpCapability StorageInputOutput16
|
||||
%16 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
|
||||
OpEntryPoint Fragment %fragment_main "fragment_main"
|
||||
|
|
@ -42,7 +41,9 @@
|
|||
%9 = OpTypeFunction %void
|
||||
%half = OpTypeFloat 16
|
||||
%__modf_result_f16 = OpTypeStruct %half %half
|
||||
%half_0x1p_0 = OpConstant %half 0x1p+0
|
||||
%half_n0x1pn1 = OpConstant %half -0x1p-1
|
||||
%half_n0x1p_0 = OpConstant %half -0x1p+0
|
||||
%17 = OpConstantComposite %__modf_result_f16 %half_n0x1pn1 %half_n0x1p_0
|
||||
%_ptr_Function___modf_result_f16 = OpTypePointer Function %__modf_result_f16
|
||||
%20 = OpConstantNull %__modf_result_f16
|
||||
%21 = OpTypeFunction %v4float
|
||||
|
|
@ -50,8 +51,7 @@
|
|||
%modf_8dbbbf = OpFunction %void None %9
|
||||
%12 = OpLabel
|
||||
%res = OpVariable %_ptr_Function___modf_result_f16 Function %20
|
||||
%13 = OpExtInst %__modf_result_f16 %16 ModfStruct %half_0x1p_0
|
||||
OpStore %res %13
|
||||
OpStore %res %17
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%vertex_main_inner = OpFunction %v4float None %21
|
||||
|
|
|
|||
|
|
@ -1,7 +1,7 @@
|
|||
enable f16;
|
||||
|
||||
fn modf_8dbbbf() {
|
||||
var res = modf(1.0h);
|
||||
var res = modf(-(1.5h));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -25,7 +25,7 @@ enable f16;
|
|||
|
||||
// fn modf(vec<4, f16>) -> __modf_result_vec<4, f16>
|
||||
fn modf_995934() {
|
||||
var res = modf(vec4<f16>(1.h));
|
||||
var res = modf(vec4<f16>(-1.5h));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -2,14 +2,8 @@ struct modf_result_vec4_f16 {
|
|||
vector<float16_t, 4> fract;
|
||||
vector<float16_t, 4> whole;
|
||||
};
|
||||
modf_result_vec4_f16 tint_modf(vector<float16_t, 4> param_0) {
|
||||
modf_result_vec4_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_995934() {
|
||||
modf_result_vec4_f16 res = tint_modf((float16_t(1.0h)).xxxx);
|
||||
modf_result_vec4_f16 res = {(float16_t(-0.5h)).xxxx, (float16_t(-1.0h)).xxxx};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -6,15 +6,9 @@ struct modf_result_vec4_f16 {
|
|||
f16vec4 whole;
|
||||
};
|
||||
|
||||
modf_result_vec4_f16 tint_modf(f16vec4 param_0) {
|
||||
modf_result_vec4_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_995934() {
|
||||
modf_result_vec4_f16 res = tint_modf(f16vec4(1.0hf));
|
||||
modf_result_vec4_f16 res = modf_result_vec4_f16(f16vec4(-0.5hf), f16vec4(-1.0hf));
|
||||
}
|
||||
|
||||
vec4 vertex_main() {
|
||||
|
|
@ -39,15 +33,9 @@ struct modf_result_vec4_f16 {
|
|||
f16vec4 whole;
|
||||
};
|
||||
|
||||
modf_result_vec4_f16 tint_modf(f16vec4 param_0) {
|
||||
modf_result_vec4_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_995934() {
|
||||
modf_result_vec4_f16 res = tint_modf(f16vec4(1.0hf));
|
||||
modf_result_vec4_f16 res = modf_result_vec4_f16(f16vec4(-0.5hf), f16vec4(-1.0hf));
|
||||
}
|
||||
|
||||
void fragment_main() {
|
||||
|
|
@ -66,15 +54,9 @@ struct modf_result_vec4_f16 {
|
|||
f16vec4 whole;
|
||||
};
|
||||
|
||||
modf_result_vec4_f16 tint_modf(f16vec4 param_0) {
|
||||
modf_result_vec4_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_995934() {
|
||||
modf_result_vec4_f16 res = tint_modf(f16vec4(1.0hf));
|
||||
modf_result_vec4_f16 res = modf_result_vec4_f16(f16vec4(-0.5hf), f16vec4(-1.0hf));
|
||||
}
|
||||
|
||||
void compute_main() {
|
||||
|
|
|
|||
|
|
@ -6,14 +6,8 @@ struct modf_result_vec4_f16 {
|
|||
half4 fract;
|
||||
half4 whole;
|
||||
};
|
||||
modf_result_vec4_f16 tint_modf(half4 param_0) {
|
||||
modf_result_vec4_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_995934() {
|
||||
modf_result_vec4_f16 res = tint_modf(half4(1.0h));
|
||||
modf_result_vec4_f16 res = modf_result_vec4_f16{.fract=half4(-0.5h), .whole=half4(-1.0h)};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -1,14 +1,13 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 37
|
||||
; Bound: 38
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability Float16
|
||||
OpCapability UniformAndStorageBuffer16BitAccess
|
||||
OpCapability StorageBuffer16BitAccess
|
||||
OpCapability StorageInputOutput16
|
||||
%17 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
|
||||
OpEntryPoint Fragment %fragment_main "fragment_main"
|
||||
|
|
@ -43,38 +42,40 @@
|
|||
%half = OpTypeFloat 16
|
||||
%v4half = OpTypeVector %half 4
|
||||
%__modf_result_vec4_f16 = OpTypeStruct %v4half %v4half
|
||||
%half_0x1p_0 = OpConstant %half 0x1p+0
|
||||
%19 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
|
||||
%half_n0x1pn1 = OpConstant %half -0x1p-1
|
||||
%17 = OpConstantComposite %v4half %half_n0x1pn1 %half_n0x1pn1 %half_n0x1pn1 %half_n0x1pn1
|
||||
%half_n0x1p_0 = OpConstant %half -0x1p+0
|
||||
%19 = OpConstantComposite %v4half %half_n0x1p_0 %half_n0x1p_0 %half_n0x1p_0 %half_n0x1p_0
|
||||
%20 = OpConstantComposite %__modf_result_vec4_f16 %17 %19
|
||||
%_ptr_Function___modf_result_vec4_f16 = OpTypePointer Function %__modf_result_vec4_f16
|
||||
%22 = OpConstantNull %__modf_result_vec4_f16
|
||||
%23 = OpTypeFunction %v4float
|
||||
%23 = OpConstantNull %__modf_result_vec4_f16
|
||||
%24 = OpTypeFunction %v4float
|
||||
%float_1 = OpConstant %float 1
|
||||
%modf_995934 = OpFunction %void None %9
|
||||
%12 = OpLabel
|
||||
%res = OpVariable %_ptr_Function___modf_result_vec4_f16 Function %22
|
||||
%13 = OpExtInst %__modf_result_vec4_f16 %17 ModfStruct %19
|
||||
OpStore %res %13
|
||||
%res = OpVariable %_ptr_Function___modf_result_vec4_f16 Function %23
|
||||
OpStore %res %20
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%vertex_main_inner = OpFunction %v4float None %23
|
||||
%25 = OpLabel
|
||||
%26 = OpFunctionCall %void %modf_995934
|
||||
%vertex_main_inner = OpFunction %v4float None %24
|
||||
%26 = OpLabel
|
||||
%27 = OpFunctionCall %void %modf_995934
|
||||
OpReturnValue %5
|
||||
OpFunctionEnd
|
||||
%vertex_main = OpFunction %void None %9
|
||||
%28 = OpLabel
|
||||
%29 = OpFunctionCall %v4float %vertex_main_inner
|
||||
OpStore %value %29
|
||||
%29 = OpLabel
|
||||
%30 = OpFunctionCall %v4float %vertex_main_inner
|
||||
OpStore %value %30
|
||||
OpStore %vertex_point_size %float_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%fragment_main = OpFunction %void None %9
|
||||
%32 = OpLabel
|
||||
%33 = OpFunctionCall %void %modf_995934
|
||||
%33 = OpLabel
|
||||
%34 = OpFunctionCall %void %modf_995934
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %9
|
||||
%35 = OpLabel
|
||||
%36 = OpFunctionCall %void %modf_995934
|
||||
%36 = OpLabel
|
||||
%37 = OpFunctionCall %void %modf_995934
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
|
|||
|
|
@ -1,7 +1,7 @@
|
|||
enable f16;
|
||||
|
||||
fn modf_995934() {
|
||||
var res = modf(vec4<f16>(1.0h));
|
||||
var res = modf(vec4<f16>(-(1.5h)));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -25,7 +25,7 @@ enable f16;
|
|||
|
||||
// fn modf(vec<2, f16>) -> __modf_result_vec<2, f16>
|
||||
fn modf_a545b9() {
|
||||
var res = modf(vec2<f16>(1.h));
|
||||
var res = modf(vec2<f16>(-1.5h));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -2,14 +2,8 @@ struct modf_result_vec2_f16 {
|
|||
vector<float16_t, 2> fract;
|
||||
vector<float16_t, 2> whole;
|
||||
};
|
||||
modf_result_vec2_f16 tint_modf(vector<float16_t, 2> param_0) {
|
||||
modf_result_vec2_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_a545b9() {
|
||||
modf_result_vec2_f16 res = tint_modf((float16_t(1.0h)).xx);
|
||||
modf_result_vec2_f16 res = {(float16_t(-0.5h)).xx, (float16_t(-1.0h)).xx};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -6,15 +6,9 @@ struct modf_result_vec2_f16 {
|
|||
f16vec2 whole;
|
||||
};
|
||||
|
||||
modf_result_vec2_f16 tint_modf(f16vec2 param_0) {
|
||||
modf_result_vec2_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_a545b9() {
|
||||
modf_result_vec2_f16 res = tint_modf(f16vec2(1.0hf));
|
||||
modf_result_vec2_f16 res = modf_result_vec2_f16(f16vec2(-0.5hf), f16vec2(-1.0hf));
|
||||
}
|
||||
|
||||
vec4 vertex_main() {
|
||||
|
|
@ -39,15 +33,9 @@ struct modf_result_vec2_f16 {
|
|||
f16vec2 whole;
|
||||
};
|
||||
|
||||
modf_result_vec2_f16 tint_modf(f16vec2 param_0) {
|
||||
modf_result_vec2_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_a545b9() {
|
||||
modf_result_vec2_f16 res = tint_modf(f16vec2(1.0hf));
|
||||
modf_result_vec2_f16 res = modf_result_vec2_f16(f16vec2(-0.5hf), f16vec2(-1.0hf));
|
||||
}
|
||||
|
||||
void fragment_main() {
|
||||
|
|
@ -66,15 +54,9 @@ struct modf_result_vec2_f16 {
|
|||
f16vec2 whole;
|
||||
};
|
||||
|
||||
modf_result_vec2_f16 tint_modf(f16vec2 param_0) {
|
||||
modf_result_vec2_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_a545b9() {
|
||||
modf_result_vec2_f16 res = tint_modf(f16vec2(1.0hf));
|
||||
modf_result_vec2_f16 res = modf_result_vec2_f16(f16vec2(-0.5hf), f16vec2(-1.0hf));
|
||||
}
|
||||
|
||||
void compute_main() {
|
||||
|
|
|
|||
|
|
@ -6,14 +6,8 @@ struct modf_result_vec2_f16 {
|
|||
half2 fract;
|
||||
half2 whole;
|
||||
};
|
||||
modf_result_vec2_f16 tint_modf(half2 param_0) {
|
||||
modf_result_vec2_f16 result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_a545b9() {
|
||||
modf_result_vec2_f16 res = tint_modf(half2(1.0h));
|
||||
modf_result_vec2_f16 res = modf_result_vec2_f16{.fract=half2(-0.5h), .whole=half2(-1.0h)};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -1,14 +1,13 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 37
|
||||
; Bound: 38
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability Float16
|
||||
OpCapability UniformAndStorageBuffer16BitAccess
|
||||
OpCapability StorageBuffer16BitAccess
|
||||
OpCapability StorageInputOutput16
|
||||
%17 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
|
||||
OpEntryPoint Fragment %fragment_main "fragment_main"
|
||||
|
|
@ -43,38 +42,40 @@
|
|||
%half = OpTypeFloat 16
|
||||
%v2half = OpTypeVector %half 2
|
||||
%__modf_result_vec2_f16 = OpTypeStruct %v2half %v2half
|
||||
%half_0x1p_0 = OpConstant %half 0x1p+0
|
||||
%19 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
|
||||
%half_n0x1pn1 = OpConstant %half -0x1p-1
|
||||
%17 = OpConstantComposite %v2half %half_n0x1pn1 %half_n0x1pn1
|
||||
%half_n0x1p_0 = OpConstant %half -0x1p+0
|
||||
%19 = OpConstantComposite %v2half %half_n0x1p_0 %half_n0x1p_0
|
||||
%20 = OpConstantComposite %__modf_result_vec2_f16 %17 %19
|
||||
%_ptr_Function___modf_result_vec2_f16 = OpTypePointer Function %__modf_result_vec2_f16
|
||||
%22 = OpConstantNull %__modf_result_vec2_f16
|
||||
%23 = OpTypeFunction %v4float
|
||||
%23 = OpConstantNull %__modf_result_vec2_f16
|
||||
%24 = OpTypeFunction %v4float
|
||||
%float_1 = OpConstant %float 1
|
||||
%modf_a545b9 = OpFunction %void None %9
|
||||
%12 = OpLabel
|
||||
%res = OpVariable %_ptr_Function___modf_result_vec2_f16 Function %22
|
||||
%13 = OpExtInst %__modf_result_vec2_f16 %17 ModfStruct %19
|
||||
OpStore %res %13
|
||||
%res = OpVariable %_ptr_Function___modf_result_vec2_f16 Function %23
|
||||
OpStore %res %20
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%vertex_main_inner = OpFunction %v4float None %23
|
||||
%25 = OpLabel
|
||||
%26 = OpFunctionCall %void %modf_a545b9
|
||||
%vertex_main_inner = OpFunction %v4float None %24
|
||||
%26 = OpLabel
|
||||
%27 = OpFunctionCall %void %modf_a545b9
|
||||
OpReturnValue %5
|
||||
OpFunctionEnd
|
||||
%vertex_main = OpFunction %void None %9
|
||||
%28 = OpLabel
|
||||
%29 = OpFunctionCall %v4float %vertex_main_inner
|
||||
OpStore %value %29
|
||||
%29 = OpLabel
|
||||
%30 = OpFunctionCall %v4float %vertex_main_inner
|
||||
OpStore %value %30
|
||||
OpStore %vertex_point_size %float_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%fragment_main = OpFunction %void None %9
|
||||
%32 = OpLabel
|
||||
%33 = OpFunctionCall %void %modf_a545b9
|
||||
%33 = OpLabel
|
||||
%34 = OpFunctionCall %void %modf_a545b9
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %9
|
||||
%35 = OpLabel
|
||||
%36 = OpFunctionCall %void %modf_a545b9
|
||||
%36 = OpLabel
|
||||
%37 = OpFunctionCall %void %modf_a545b9
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
|
|||
|
|
@ -1,7 +1,7 @@
|
|||
enable f16;
|
||||
|
||||
fn modf_a545b9() {
|
||||
var res = modf(vec2<f16>(1.0h));
|
||||
var res = modf(vec2<f16>(-(1.5h)));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -23,7 +23,7 @@
|
|||
|
||||
// fn modf(f32) -> __modf_result<f32>
|
||||
fn modf_bbf7f7() {
|
||||
var res = modf(1.f);
|
||||
var res = modf(-1.5f);
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -2,14 +2,8 @@ struct modf_result {
|
|||
float fract;
|
||||
float whole;
|
||||
};
|
||||
modf_result tint_modf(float param_0) {
|
||||
modf_result result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_bbf7f7() {
|
||||
modf_result res = tint_modf(1.0f);
|
||||
modf_result res = {-0.5f, -1.0f};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -2,14 +2,8 @@ struct modf_result {
|
|||
float fract;
|
||||
float whole;
|
||||
};
|
||||
modf_result tint_modf(float param_0) {
|
||||
modf_result result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_bbf7f7() {
|
||||
modf_result res = tint_modf(1.0f);
|
||||
modf_result res = {-0.5f, -1.0f};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -5,15 +5,9 @@ struct modf_result {
|
|||
float whole;
|
||||
};
|
||||
|
||||
modf_result tint_modf(float param_0) {
|
||||
modf_result result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_bbf7f7() {
|
||||
modf_result res = tint_modf(1.0f);
|
||||
modf_result res = modf_result(-0.5f, -1.0f);
|
||||
}
|
||||
|
||||
vec4 vertex_main() {
|
||||
|
|
@ -37,15 +31,9 @@ struct modf_result {
|
|||
float whole;
|
||||
};
|
||||
|
||||
modf_result tint_modf(float param_0) {
|
||||
modf_result result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_bbf7f7() {
|
||||
modf_result res = tint_modf(1.0f);
|
||||
modf_result res = modf_result(-0.5f, -1.0f);
|
||||
}
|
||||
|
||||
void fragment_main() {
|
||||
|
|
@ -63,15 +51,9 @@ struct modf_result {
|
|||
float whole;
|
||||
};
|
||||
|
||||
modf_result tint_modf(float param_0) {
|
||||
modf_result result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
void modf_bbf7f7() {
|
||||
modf_result res = tint_modf(1.0f);
|
||||
modf_result res = modf_result(-0.5f, -1.0f);
|
||||
}
|
||||
|
||||
void compute_main() {
|
||||
|
|
|
|||
|
|
@ -6,14 +6,8 @@ struct modf_result {
|
|||
float fract;
|
||||
float whole;
|
||||
};
|
||||
modf_result tint_modf(float param_0) {
|
||||
modf_result result;
|
||||
result.fract = modf(param_0, result.whole);
|
||||
return result;
|
||||
}
|
||||
|
||||
void modf_bbf7f7() {
|
||||
modf_result res = tint_modf(1.0f);
|
||||
modf_result res = modf_result{.fract=-0.5f, .whole=-1.0f};
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
|
|
|
|||
|
|
@ -1,10 +1,9 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 33
|
||||
; Bound: 34
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%15 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
|
||||
OpEntryPoint Fragment %fragment_main "fragment_main"
|
||||
|
|
@ -37,15 +36,17 @@
|
|||
%void = OpTypeVoid
|
||||
%9 = OpTypeFunction %void
|
||||
%__modf_result = OpTypeStruct %float %float
|
||||
%float_1 = OpConstant %float 1
|
||||
%float_n0_5 = OpConstant %float -0.5
|
||||
%float_n1 = OpConstant %float -1
|
||||
%16 = OpConstantComposite %__modf_result %float_n0_5 %float_n1
|
||||
%_ptr_Function___modf_result = OpTypePointer Function %__modf_result
|
||||
%19 = OpConstantNull %__modf_result
|
||||
%20 = OpTypeFunction %v4float
|
||||
%float_1 = OpConstant %float 1
|
||||
%modf_bbf7f7 = OpFunction %void None %9
|
||||
%12 = OpLabel
|
||||
%res = OpVariable %_ptr_Function___modf_result Function %19
|
||||
%13 = OpExtInst %__modf_result %15 ModfStruct %float_1
|
||||
OpStore %res %13
|
||||
OpStore %res %16
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%vertex_main_inner = OpFunction %v4float None %20
|
||||
|
|
@ -61,12 +62,12 @@
|
|||
OpReturn
|
||||
OpFunctionEnd
|
||||
%fragment_main = OpFunction %void None %9
|
||||
%28 = OpLabel
|
||||
%29 = OpFunctionCall %void %modf_bbf7f7
|
||||
%29 = OpLabel
|
||||
%30 = OpFunctionCall %void %modf_bbf7f7
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %9
|
||||
%31 = OpLabel
|
||||
%32 = OpFunctionCall %void %modf_bbf7f7
|
||||
%32 = OpLabel
|
||||
%33 = OpFunctionCall %void %modf_bbf7f7
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
|
|||
|
|
@ -1,5 +1,5 @@
|
|||
fn modf_bbf7f7() {
|
||||
var res = modf(1.0f);
|
||||
var res = modf(-(1.5f));
|
||||
}
|
||||
|
||||
@vertex
|
||||
|
|
|
|||
|
|
@ -23,7 +23,7 @@
|
|||
|
||||
// fn modf(vec<2, f32>) -> __modf_result_vec<2, f32>
|
||||
fn modf_2d50da() {
|
||||
var arg_0 = vec2<f32>(1.f);
|
||||
var arg_0 = vec2<f32>(-1.5f);
|
||||
var res = modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -9,7 +9,7 @@ modf_result_vec2 tint_modf(float2 param_0) {
|
|||
}
|
||||
|
||||
void modf_2d50da() {
|
||||
float2 arg_0 = (1.0f).xx;
|
||||
float2 arg_0 = (-1.5f).xx;
|
||||
modf_result_vec2 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -9,7 +9,7 @@ modf_result_vec2 tint_modf(float2 param_0) {
|
|||
}
|
||||
|
||||
void modf_2d50da() {
|
||||
float2 arg_0 = (1.0f).xx;
|
||||
float2 arg_0 = (-1.5f).xx;
|
||||
modf_result_vec2 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@ modf_result_vec2 tint_modf(vec2 param_0) {
|
|||
|
||||
|
||||
void modf_2d50da() {
|
||||
vec2 arg_0 = vec2(1.0f);
|
||||
vec2 arg_0 = vec2(-1.5f);
|
||||
modf_result_vec2 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
@ -46,7 +46,7 @@ modf_result_vec2 tint_modf(vec2 param_0) {
|
|||
|
||||
|
||||
void modf_2d50da() {
|
||||
vec2 arg_0 = vec2(1.0f);
|
||||
vec2 arg_0 = vec2(-1.5f);
|
||||
modf_result_vec2 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
@ -73,7 +73,7 @@ modf_result_vec2 tint_modf(vec2 param_0) {
|
|||
|
||||
|
||||
void modf_2d50da() {
|
||||
vec2 arg_0 = vec2(1.0f);
|
||||
vec2 arg_0 = vec2(-1.5f);
|
||||
modf_result_vec2 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@ modf_result_vec2 tint_modf(float2 param_0) {
|
|||
}
|
||||
|
||||
void modf_2d50da() {
|
||||
float2 arg_0 = float2(1.0f);
|
||||
float2 arg_0 = float2(-1.5f);
|
||||
modf_result_vec2 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -1,7 +1,7 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 39
|
||||
; Bound: 40
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%21 = OpExtInstImport "GLSL.std.450"
|
||||
|
|
@ -38,14 +38,15 @@
|
|||
%void = OpTypeVoid
|
||||
%9 = OpTypeFunction %void
|
||||
%v2float = OpTypeVector %float 2
|
||||
%float_1 = OpConstant %float 1
|
||||
%15 = OpConstantComposite %v2float %float_1 %float_1
|
||||
%float_n1_5 = OpConstant %float -1.5
|
||||
%15 = OpConstantComposite %v2float %float_n1_5 %float_n1_5
|
||||
%_ptr_Function_v2float = OpTypePointer Function %v2float
|
||||
%18 = OpConstantNull %v2float
|
||||
%__modf_result_vec2 = OpTypeStruct %v2float %v2float
|
||||
%_ptr_Function___modf_result_vec2 = OpTypePointer Function %__modf_result_vec2
|
||||
%25 = OpConstantNull %__modf_result_vec2
|
||||
%26 = OpTypeFunction %v4float
|
||||
%float_1 = OpConstant %float 1
|
||||
%modf_2d50da = OpFunction %void None %9
|
||||
%12 = OpLabel
|
||||
%arg_0 = OpVariable %_ptr_Function_v2float Function %18
|
||||
|
|
@ -69,12 +70,12 @@
|
|||
OpReturn
|
||||
OpFunctionEnd
|
||||
%fragment_main = OpFunction %void None %9
|
||||
%34 = OpLabel
|
||||
%35 = OpFunctionCall %void %modf_2d50da
|
||||
%35 = OpLabel
|
||||
%36 = OpFunctionCall %void %modf_2d50da
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %9
|
||||
%37 = OpLabel
|
||||
%38 = OpFunctionCall %void %modf_2d50da
|
||||
%38 = OpLabel
|
||||
%39 = OpFunctionCall %void %modf_2d50da
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
|
|||
|
|
@ -1,5 +1,5 @@
|
|||
fn modf_2d50da() {
|
||||
var arg_0 = vec2<f32>(1.0f);
|
||||
var arg_0 = vec2<f32>(-(1.5f));
|
||||
var res = modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -25,7 +25,7 @@ enable f16;
|
|||
|
||||
// fn modf(vec<3, f16>) -> __modf_result_vec<3, f16>
|
||||
fn modf_45005f() {
|
||||
var arg_0 = vec3<f16>(1.h);
|
||||
var arg_0 = vec3<f16>(-1.5h);
|
||||
var res = modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -9,7 +9,7 @@ modf_result_vec3_f16 tint_modf(vector<float16_t, 3> param_0) {
|
|||
}
|
||||
|
||||
void modf_45005f() {
|
||||
vector<float16_t, 3> arg_0 = (float16_t(1.0h)).xxx;
|
||||
vector<float16_t, 3> arg_0 = (float16_t(-1.5h)).xxx;
|
||||
modf_result_vec3_f16 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -14,7 +14,7 @@ modf_result_vec3_f16 tint_modf(f16vec3 param_0) {
|
|||
|
||||
|
||||
void modf_45005f() {
|
||||
f16vec3 arg_0 = f16vec3(1.0hf);
|
||||
f16vec3 arg_0 = f16vec3(-1.5hf);
|
||||
modf_result_vec3_f16 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
@ -48,7 +48,7 @@ modf_result_vec3_f16 tint_modf(f16vec3 param_0) {
|
|||
|
||||
|
||||
void modf_45005f() {
|
||||
f16vec3 arg_0 = f16vec3(1.0hf);
|
||||
f16vec3 arg_0 = f16vec3(-1.5hf);
|
||||
modf_result_vec3_f16 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
@ -76,7 +76,7 @@ modf_result_vec3_f16 tint_modf(f16vec3 param_0) {
|
|||
|
||||
|
||||
void modf_45005f() {
|
||||
f16vec3 arg_0 = f16vec3(1.0hf);
|
||||
f16vec3 arg_0 = f16vec3(-1.5hf);
|
||||
modf_result_vec3_f16 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@ modf_result_vec3_f16 tint_modf(half3 param_0) {
|
|||
}
|
||||
|
||||
void modf_45005f() {
|
||||
half3 arg_0 = half3(1.0h);
|
||||
half3 arg_0 = half3(-1.5h);
|
||||
modf_result_vec3_f16 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -43,8 +43,8 @@
|
|||
%9 = OpTypeFunction %void
|
||||
%half = OpTypeFloat 16
|
||||
%v3half = OpTypeVector %half 3
|
||||
%half_0x1p_0 = OpConstant %half 0x1p+0
|
||||
%16 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
|
||||
%half_n0x1_8p_0 = OpConstant %half -0x1.8p+0
|
||||
%16 = OpConstantComposite %v3half %half_n0x1_8p_0 %half_n0x1_8p_0 %half_n0x1_8p_0
|
||||
%_ptr_Function_v3half = OpTypePointer Function %v3half
|
||||
%19 = OpConstantNull %v3half
|
||||
%__modf_result_vec3_f16 = OpTypeStruct %v3half %v3half
|
||||
|
|
|
|||
|
|
@ -1,7 +1,7 @@
|
|||
enable f16;
|
||||
|
||||
fn modf_45005f() {
|
||||
var arg_0 = vec3<f16>(1.0h);
|
||||
var arg_0 = vec3<f16>(-(1.5h));
|
||||
var res = modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -23,7 +23,7 @@
|
|||
|
||||
// fn modf(vec<4, f32>) -> __modf_result_vec<4, f32>
|
||||
fn modf_4bfced() {
|
||||
var arg_0 = vec4<f32>(1.f);
|
||||
var arg_0 = vec4<f32>(-1.5f);
|
||||
var res = modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -9,7 +9,7 @@ modf_result_vec4 tint_modf(float4 param_0) {
|
|||
}
|
||||
|
||||
void modf_4bfced() {
|
||||
float4 arg_0 = (1.0f).xxxx;
|
||||
float4 arg_0 = (-1.5f).xxxx;
|
||||
modf_result_vec4 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -9,7 +9,7 @@ modf_result_vec4 tint_modf(float4 param_0) {
|
|||
}
|
||||
|
||||
void modf_4bfced() {
|
||||
float4 arg_0 = (1.0f).xxxx;
|
||||
float4 arg_0 = (-1.5f).xxxx;
|
||||
modf_result_vec4 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@ modf_result_vec4 tint_modf(vec4 param_0) {
|
|||
|
||||
|
||||
void modf_4bfced() {
|
||||
vec4 arg_0 = vec4(1.0f);
|
||||
vec4 arg_0 = vec4(-1.5f);
|
||||
modf_result_vec4 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
@ -46,7 +46,7 @@ modf_result_vec4 tint_modf(vec4 param_0) {
|
|||
|
||||
|
||||
void modf_4bfced() {
|
||||
vec4 arg_0 = vec4(1.0f);
|
||||
vec4 arg_0 = vec4(-1.5f);
|
||||
modf_result_vec4 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
@ -73,7 +73,7 @@ modf_result_vec4 tint_modf(vec4 param_0) {
|
|||
|
||||
|
||||
void modf_4bfced() {
|
||||
vec4 arg_0 = vec4(1.0f);
|
||||
vec4 arg_0 = vec4(-1.5f);
|
||||
modf_result_vec4 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@ modf_result_vec4 tint_modf(float4 param_0) {
|
|||
}
|
||||
|
||||
void modf_4bfced() {
|
||||
float4 arg_0 = float4(1.0f);
|
||||
float4 arg_0 = float4(-1.5f);
|
||||
modf_result_vec4 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -1,7 +1,7 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 37
|
||||
; Bound: 38
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%19 = OpExtInstImport "GLSL.std.450"
|
||||
|
|
@ -37,13 +37,14 @@
|
|||
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
|
||||
%void = OpTypeVoid
|
||||
%9 = OpTypeFunction %void
|
||||
%float_1 = OpConstant %float 1
|
||||
%14 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
|
||||
%float_n1_5 = OpConstant %float -1.5
|
||||
%14 = OpConstantComposite %v4float %float_n1_5 %float_n1_5 %float_n1_5 %float_n1_5
|
||||
%_ptr_Function_v4float = OpTypePointer Function %v4float
|
||||
%__modf_result_vec4 = OpTypeStruct %v4float %v4float
|
||||
%_ptr_Function___modf_result_vec4 = OpTypePointer Function %__modf_result_vec4
|
||||
%23 = OpConstantNull %__modf_result_vec4
|
||||
%24 = OpTypeFunction %v4float
|
||||
%float_1 = OpConstant %float 1
|
||||
%modf_4bfced = OpFunction %void None %9
|
||||
%12 = OpLabel
|
||||
%arg_0 = OpVariable %_ptr_Function_v4float Function %5
|
||||
|
|
@ -67,12 +68,12 @@
|
|||
OpReturn
|
||||
OpFunctionEnd
|
||||
%fragment_main = OpFunction %void None %9
|
||||
%32 = OpLabel
|
||||
%33 = OpFunctionCall %void %modf_4bfced
|
||||
%33 = OpLabel
|
||||
%34 = OpFunctionCall %void %modf_4bfced
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %9
|
||||
%35 = OpLabel
|
||||
%36 = OpFunctionCall %void %modf_4bfced
|
||||
%36 = OpLabel
|
||||
%37 = OpFunctionCall %void %modf_4bfced
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
|
|||
|
|
@ -1,5 +1,5 @@
|
|||
fn modf_4bfced() {
|
||||
var arg_0 = vec4<f32>(1.0f);
|
||||
var arg_0 = vec4<f32>(-(1.5f));
|
||||
var res = modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -23,7 +23,7 @@
|
|||
|
||||
// fn modf(vec<3, f32>) -> __modf_result_vec<3, f32>
|
||||
fn modf_5ea256() {
|
||||
var arg_0 = vec3<f32>(1.f);
|
||||
var arg_0 = vec3<f32>(-1.5f);
|
||||
var res = modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -9,7 +9,7 @@ modf_result_vec3 tint_modf(float3 param_0) {
|
|||
}
|
||||
|
||||
void modf_5ea256() {
|
||||
float3 arg_0 = (1.0f).xxx;
|
||||
float3 arg_0 = (-1.5f).xxx;
|
||||
modf_result_vec3 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -9,7 +9,7 @@ modf_result_vec3 tint_modf(float3 param_0) {
|
|||
}
|
||||
|
||||
void modf_5ea256() {
|
||||
float3 arg_0 = (1.0f).xxx;
|
||||
float3 arg_0 = (-1.5f).xxx;
|
||||
modf_result_vec3 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@ modf_result_vec3 tint_modf(vec3 param_0) {
|
|||
|
||||
|
||||
void modf_5ea256() {
|
||||
vec3 arg_0 = vec3(1.0f);
|
||||
vec3 arg_0 = vec3(-1.5f);
|
||||
modf_result_vec3 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
@ -46,7 +46,7 @@ modf_result_vec3 tint_modf(vec3 param_0) {
|
|||
|
||||
|
||||
void modf_5ea256() {
|
||||
vec3 arg_0 = vec3(1.0f);
|
||||
vec3 arg_0 = vec3(-1.5f);
|
||||
modf_result_vec3 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
@ -73,7 +73,7 @@ modf_result_vec3 tint_modf(vec3 param_0) {
|
|||
|
||||
|
||||
void modf_5ea256() {
|
||||
vec3 arg_0 = vec3(1.0f);
|
||||
vec3 arg_0 = vec3(-1.5f);
|
||||
modf_result_vec3 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@ modf_result_vec3 tint_modf(float3 param_0) {
|
|||
}
|
||||
|
||||
void modf_5ea256() {
|
||||
float3 arg_0 = float3(1.0f);
|
||||
float3 arg_0 = float3(-1.5f);
|
||||
modf_result_vec3 res = tint_modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -1,7 +1,7 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 39
|
||||
; Bound: 40
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%21 = OpExtInstImport "GLSL.std.450"
|
||||
|
|
@ -38,14 +38,15 @@
|
|||
%void = OpTypeVoid
|
||||
%9 = OpTypeFunction %void
|
||||
%v3float = OpTypeVector %float 3
|
||||
%float_1 = OpConstant %float 1
|
||||
%15 = OpConstantComposite %v3float %float_1 %float_1 %float_1
|
||||
%float_n1_5 = OpConstant %float -1.5
|
||||
%15 = OpConstantComposite %v3float %float_n1_5 %float_n1_5 %float_n1_5
|
||||
%_ptr_Function_v3float = OpTypePointer Function %v3float
|
||||
%18 = OpConstantNull %v3float
|
||||
%__modf_result_vec3 = OpTypeStruct %v3float %v3float
|
||||
%_ptr_Function___modf_result_vec3 = OpTypePointer Function %__modf_result_vec3
|
||||
%25 = OpConstantNull %__modf_result_vec3
|
||||
%26 = OpTypeFunction %v4float
|
||||
%float_1 = OpConstant %float 1
|
||||
%modf_5ea256 = OpFunction %void None %9
|
||||
%12 = OpLabel
|
||||
%arg_0 = OpVariable %_ptr_Function_v3float Function %18
|
||||
|
|
@ -69,12 +70,12 @@
|
|||
OpReturn
|
||||
OpFunctionEnd
|
||||
%fragment_main = OpFunction %void None %9
|
||||
%34 = OpLabel
|
||||
%35 = OpFunctionCall %void %modf_5ea256
|
||||
%35 = OpLabel
|
||||
%36 = OpFunctionCall %void %modf_5ea256
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %9
|
||||
%37 = OpLabel
|
||||
%38 = OpFunctionCall %void %modf_5ea256
|
||||
%38 = OpLabel
|
||||
%39 = OpFunctionCall %void %modf_5ea256
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
|
|||
|
|
@ -1,5 +1,5 @@
|
|||
fn modf_5ea256() {
|
||||
var arg_0 = vec3<f32>(1.0f);
|
||||
var arg_0 = vec3<f32>(-(1.5f));
|
||||
var res = modf(arg_0);
|
||||
}
|
||||
|
||||
|
|
|
|||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in New Issue