encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

[spirv-reader] Use builtin vector aliases 

This makes the generated WGSL more concise and readable.

Change-Id: Ia486a74796a4029aaac1c4d051d304d05f1d5ea2
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/132140
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: James Price <jrprice@google.com>
This commit is contained in:
James Price 2023-05-09 10:36:41 +00:00 committed by Dawn LUCI CQ
parent 5696997251
commit 114bae24b0
39 changed files with 707 additions and 744 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

22
src/tint/reader/spirv/function.cc
View File

@ -4047,8 +4047,8 @@ TypedExpression FunctionEmitter::EmitGlslStd450ExtInst(const spvtools::opt::Inst
builder_.MemberAccessor(
builder_.Call("refract",
utils::Vector{
builder_.vec2<tint::f32>(incident.expr, 0_f),
builder_.vec2<tint::f32>(normal.expr, 0_f),
builder_.Call("vec2f", incident.expr, 0_f),
builder_.Call("vec2f", normal.expr, 0_f),
eta.expr,
}),
"x"),
@ -4137,8 +4137,8 @@ TypedExpression FunctionEmitter::EmitGlslStd450MatrixInverse(
// s * d, -s * b, -s * c, s * a
auto* r = pb.mat2x2<f32>( //
pb.vec2<f32>(pb.Mul(s, d), pb.Mul(neg(s), b)),
pb.vec2<f32>(pb.Mul(neg(s), c), pb.Mul(s, a)));
pb.Call("vec2f", pb.Mul(s, d), pb.Mul(neg(s), b)),
pb.Call("vec2f", pb.Mul(neg(s), c), pb.Mul(s, a)));
return {mat.type, r};
}
@ -4158,21 +4158,21 @@ TypedExpression FunctionEmitter::EmitGlslStd450MatrixInverse(
auto r = pb.Mul(s, //
pb.mat3x3<f32>( //
pb.vec3<f32>(
pb.Call("vec3f",
// e * i - f * h
sub_mul2(e, i, f, h),
// c * h - b * i
sub_mul2(c, h, b, i),
// b * f - c * e
sub_mul2(b, f, c, e)),
pb.vec3<f32>(
pb.Call("vec3f",
// f * g - d * i
sub_mul2(f, g, d, i),
// a * i - c * g
sub_mul2(a, i, c, g),
// c * d - a * f
sub_mul2(c, d, a, f)),
pb.vec3<f32>(
pb.Call("vec3f",
// d * h - e * g
sub_mul2(d, h, e, g),
// b * g - a * h
@ -4236,7 +4236,7 @@ TypedExpression FunctionEmitter::EmitGlslStd450MatrixInverse(
auto r = pb.Mul(s, //
pb.mat4x4<f32>( //
pb.vec4<f32>(
pb.Call("vec4f",
// f * kplo - g * jpln + h * jokn
sub_add_mul3(f, kplo, g, jpln, h, jokn),
// -b * kplo + c * jpln - d * jokn
@ -4245,7 +4245,7 @@ TypedExpression FunctionEmitter::EmitGlslStd450MatrixInverse(
sub_add_mul3(b, gpho, c, fphn, d, fogn),
// -b * glhk + c * flhj - d * fkgj
add_sub_mul3(neg(b), glhk, c, flhj, d, fkgj)),
pb.vec4<f32>(
pb.Call("vec4f",
// -e * kplo + g * iplm - h * iokm
add_sub_mul3(neg(e), kplo, g, iplm, h, iokm),
// a * kplo - c * iplm + d * iokm
@ -4254,7 +4254,7 @@ TypedExpression FunctionEmitter::EmitGlslStd450MatrixInverse(
add_sub_mul3(neg(a), gpho, c, ephm, d, eogm),
// a * glhk - c * elhi + d * ekgi
sub_add_mul3(a, glhk, c, elhi, d, ekgi)),
pb.vec4<f32>(
pb.Call("vec4f",
// e * jpln - f * iplm + h * injm
sub_add_mul3(e, jpln, f, iplm, h, injm),
// -a * jpln + b * iplm - d * injm
@ -4263,7 +4263,7 @@ TypedExpression FunctionEmitter::EmitGlslStd450MatrixInverse(
sub_add_mul3(a, fphn, b, ephm, d, enfm),
// -a * flhj + b * elhi - d * ejfi
add_sub_mul3(neg(a), flhj, b, elhi, d, ejfi)),
pb.vec4<f32>(
pb.Call("vec4f",
// -e * jokn + f * iokm - g * injm
add_sub_mul3(neg(e), jokn, f, iokm, g, injm),
// a * jokn - b * iokm + c * injm

139
src/tint/reader/spirv/function_arithmetic_test.cc
View File

@ -81,22 +81,22 @@ std::string Preamble() {
// Returns the AST dump for a given SPIR-V assembly constant.
std::string AstFor(std::string assembly) {
if (assembly == "v2uint_10_20") {
return "vec2<u32>(10u, 20u)";
return "vec2u(10u, 20u)";
}
if (assembly == "v2uint_20_10") {
return "vec2<u32>(20u, 10u)";
return "vec2u(20u, 10u)";
}
if (assembly == "v2int_30_40") {
return "vec2<i32>(30i, 40i)";
return "vec2i(30i, 40i)";
}
if (assembly == "v2int_40_30") {
return "vec2<i32>(40i, 30i)";
return "vec2i(40i, 30i)";
}
if (assembly == "cast_int_v2uint_10_20") {
return "bitcast<vec2<i32>>(vec2<u32>(10u, 20u))";
return "bitcast<vec2i>(vec2u(10u, 20u))";
}
if (assembly == "cast_uint_v2int_40_30") {
return "bitcast<vec2<u32>>(vec2<i32>(40i, 30i))";
return "bitcast<vec2u>(vec2i(40i, 30i))";
}
if (assembly == "v2float_50_60") {
return "v2float_50_60";
@ -190,7 +190,7 @@ TEST_F(SpvUnaryArithTest, SNegate_SignedVec_SignedVec) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<i32> = -(vec2<i32>(30i, 40i));"));
HasSubstr("let x_1 : vec2i = -(vec2i(30i, 40i));"));
}
TEST_F(SpvUnaryArithTest, SNegate_SignedVec_UnsignedVec) {
@ -207,7 +207,7 @@ TEST_F(SpvUnaryArithTest, SNegate_SignedVec_UnsignedVec) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<i32> = -(bitcast<vec2<i32>>(vec2<u32>(10u, 20u)));"));
HasSubstr("let x_1 : vec2i = -(bitcast<vec2i>(vec2u(10u, 20u)));"));
}
TEST_F(SpvUnaryArithTest, SNegate_UnsignedVec_SignedVec) {
@ -224,7 +224,7 @@ TEST_F(SpvUnaryArithTest, SNegate_UnsignedVec_SignedVec) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<u32> = bitcast<vec2<u32>>(-(vec2<i32>(30i, 40i)));"));
HasSubstr("let x_1 : vec2u = bitcast<vec2u>(-(vec2i(30i, 40i)));"));
}
TEST_F(SpvUnaryArithTest, SNegate_UnsignedVec_UnsignedVec) {
@ -242,8 +242,7 @@ TEST_F(SpvUnaryArithTest, SNegate_UnsignedVec_UnsignedVec) {
auto ast_body = fe.ast_body();
EXPECT_THAT(
test::ToString(p->program(), ast_body),
HasSubstr(
R"(let x_1 : vec2<u32> = bitcast<vec2<u32>>(-(bitcast<vec2<i32>>(vec2<u32>(10u, 20u))));)"));
HasSubstr(R"(let x_1 : vec2u = bitcast<vec2u>(-(bitcast<vec2i>(vec2u(10u, 20u))));)"));
}
TEST_F(SpvUnaryArithTest, FNegate_Scalar) {
@ -276,7 +275,7 @@ TEST_F(SpvUnaryArithTest, FNegate_Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<f32> = -(v2float_50_60);"));
HasSubstr("let x_1 : vec2f = -(v2float_50_60);"));
}
struct BinaryData {
@ -366,11 +365,11 @@ INSTANTIATE_TEST_SUITE_P(
// Both uint
BinaryData{"uint", "uint_10", "OpIAdd", "uint_20", "u32", "10u", "+", "20u"}, // Both int
BinaryData{"int", "int_30", "OpIAdd", "int_40", "i32", "30i", "+", "40i"}, // Both v2uint
BinaryData{"v2uint", "v2uint_10_20", "OpIAdd", "v2uint_20_10", "vec2<u32>",
BinaryData{"v2uint", "v2uint_10_20", "OpIAdd", "v2uint_20_10", "vec2u",
AstFor("v2uint_10_20"), "+", AstFor("v2uint_20_10")},
// Both v2int
BinaryData{"v2int", "v2int_30_40", "OpIAdd", "v2int_40_30", "vec2<i32>",
AstFor("v2int_30_40"), "+", AstFor("v2int_40_30")}));
BinaryData{"v2int", "v2int_30_40", "OpIAdd", "v2int_40_30", "vec2i", AstFor("v2int_30_40"),
"+", AstFor("v2int_40_30")}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_IAdd_MixedSignedness,
@ -388,13 +387,12 @@ INSTANTIATE_TEST_SUITE_P(
BinaryDataGeneral{"int", "uint_20", "OpIAdd", "uint_10", "i32",
"bitcast<i32>((20u + 10u))"},
// Mixed, returning v2uint
BinaryDataGeneral{
"v2uint", "v2int_30_40", "OpIAdd", "v2uint_10_20", "vec2<u32>",
R"(bitcast<vec2<u32>>((vec2<i32>(30i, 40i) + bitcast<vec2<i32>>(vec2<u32>(10u, 20u)))))"},
BinaryDataGeneral{"v2uint", "v2int_30_40", "OpIAdd", "v2uint_10_20", "vec2u",
R"(bitcast<vec2u>((vec2i(30i, 40i) + bitcast<vec2i>(vec2u(10u, 20u)))))"},
// Mixed, returning v2int
BinaryDataGeneral{
"v2int", "v2uint_10_20", "OpIAdd", "v2int_40_30", "vec2<i32>",
R"(bitcast<vec2<i32>>((vec2<u32>(10u, 20u) + bitcast<vec2<u32>>(vec2<i32>(40i, 30i)))))"}));
"v2int", "v2uint_10_20", "OpIAdd", "v2int_40_30", "vec2i",
R"(bitcast<vec2i>((vec2u(10u, 20u) + bitcast<vec2u>(vec2i(40i, 30i)))))"}));
INSTANTIATE_TEST_SUITE_P(SpvParserTest_FAdd,
SpvBinaryArithTest,
@ -403,7 +401,7 @@ INSTANTIATE_TEST_SUITE_P(SpvParserTest_FAdd,
BinaryData{"float", "float_50", "OpFAdd", "float_60", "f32", "50.0f",
"+", "60.0f"}, // Vector float
BinaryData{"v2float", "v2float_50_60", "OpFAdd", "v2float_60_50",
"vec2<f32>", AstFor("v2float_50_60"), "+",
"vec2f", AstFor("v2float_50_60"), "+",
AstFor("v2float_60_50")}));
INSTANTIATE_TEST_SUITE_P(
@ -413,11 +411,11 @@ INSTANTIATE_TEST_SUITE_P(
// Both uint
BinaryData{"uint", "uint_10", "OpISub", "uint_20", "u32", "10u", "-", "20u"}, // Both int
BinaryData{"int", "int_30", "OpISub", "int_40", "i32", "30i", "-", "40i"}, // Both v2uint
BinaryData{"v2uint", "v2uint_10_20", "OpISub", "v2uint_20_10", "vec2<u32>",
BinaryData{"v2uint", "v2uint_10_20", "OpISub", "v2uint_20_10", "vec2u",
AstFor("v2uint_10_20"), "-", AstFor("v2uint_20_10")},
// Both v2int
BinaryData{"v2int", "v2int_30_40", "OpISub", "v2int_40_30", "vec2<i32>",
AstFor("v2int_30_40"), "-", AstFor("v2int_40_30")}));
BinaryData{"v2int", "v2int_30_40", "OpISub", "v2int_40_30", "vec2i", AstFor("v2int_30_40"),
"-", AstFor("v2int_40_30")}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_ISub_MixedSignedness,
@ -435,13 +433,12 @@ INSTANTIATE_TEST_SUITE_P(
BinaryDataGeneral{"int", "uint_20", "OpISub", "uint_10", "i32",
"bitcast<i32>((20u - 10u))"},
// Mixed, returning v2uint
BinaryDataGeneral{
"v2uint", "v2int_30_40", "OpISub", "v2uint_10_20", "vec2<u32>",
R"(bitcast<vec2<u32>>((vec2<i32>(30i, 40i) - bitcast<vec2<i32>>(vec2<u32>(10u, 20u)))))"},
BinaryDataGeneral{"v2uint", "v2int_30_40", "OpISub", "v2uint_10_20", "vec2u",
R"(bitcast<vec2u>((vec2i(30i, 40i) - bitcast<vec2i>(vec2u(10u, 20u)))))"},
// Mixed, returning v2int
BinaryDataGeneral{
"v2int", "v2uint_10_20", "OpISub", "v2int_40_30", "vec2<i32>",
R"(bitcast<vec2<i32>>((vec2<u32>(10u, 20u) - bitcast<vec2<u32>>(vec2<i32>(40i, 30i)))))"}));
"v2int", "v2uint_10_20", "OpISub", "v2int_40_30", "vec2i",
R"(bitcast<vec2i>((vec2u(10u, 20u) - bitcast<vec2u>(vec2i(40i, 30i)))))"}));
INSTANTIATE_TEST_SUITE_P(SpvParserTest_FSub,
SpvBinaryArithTest,
@ -450,7 +447,7 @@ INSTANTIATE_TEST_SUITE_P(SpvParserTest_FSub,
BinaryData{"float", "float_50", "OpFSub", "float_60", "f32", "50.0f",
"-", "60.0f"}, // Vector float
BinaryData{"v2float", "v2float_50_60", "OpFSub", "v2float_60_50",
"vec2<f32>", AstFor("v2float_50_60"), "-",
"vec2f", AstFor("v2float_50_60"), "-",
AstFor("v2float_60_50")}));
INSTANTIATE_TEST_SUITE_P(
@ -460,11 +457,11 @@ INSTANTIATE_TEST_SUITE_P(
// Both uint
BinaryData{"uint", "uint_10", "OpIMul", "uint_20", "u32", "10u", "*", "20u"}, // Both int
BinaryData{"int", "int_30", "OpIMul", "int_40", "i32", "30i", "*", "40i"}, // Both v2uint
BinaryData{"v2uint", "v2uint_10_20", "OpIMul", "v2uint_20_10", "vec2<u32>",
BinaryData{"v2uint", "v2uint_10_20", "OpIMul", "v2uint_20_10", "vec2u",
AstFor("v2uint_10_20"), "*", AstFor("v2uint_20_10")},
// Both v2int
BinaryData{"v2int", "v2int_30_40", "OpIMul", "v2int_40_30", "vec2<i32>",
AstFor("v2int_30_40"), "*", AstFor("v2int_40_30")}));
BinaryData{"v2int", "v2int_30_40", "OpIMul", "v2int_40_30", "vec2i", AstFor("v2int_30_40"),
"*", AstFor("v2int_40_30")}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_IMul_MixedSignedness,
@ -482,13 +479,12 @@ INSTANTIATE_TEST_SUITE_P(
BinaryDataGeneral{"int", "uint_20", "OpIMul", "uint_10", "i32",
"bitcast<i32>((20u * 10u))"},
// Mixed, returning v2uint
BinaryDataGeneral{
"v2uint", "v2int_30_40", "OpIMul", "v2uint_10_20", "vec2<u32>",
R"(bitcast<vec2<u32>>((vec2<i32>(30i, 40i) * bitcast<vec2<i32>>(vec2<u32>(10u, 20u)))))"},
BinaryDataGeneral{"v2uint", "v2int_30_40", "OpIMul", "v2uint_10_20", "vec2u",
R"(bitcast<vec2u>((vec2i(30i, 40i) * bitcast<vec2i>(vec2u(10u, 20u)))))"},
// Mixed, returning v2int
BinaryDataGeneral{
"v2int", "v2uint_10_20", "OpIMul", "v2int_40_30", "vec2<i32>",
R"(bitcast<vec2<i32>>((vec2<u32>(10u, 20u) * bitcast<vec2<u32>>(vec2<i32>(40i, 30i)))))"}));
"v2int", "v2uint_10_20", "OpIMul", "v2int_40_30", "vec2i",
R"(bitcast<vec2i>((vec2u(10u, 20u) * bitcast<vec2u>(vec2i(40i, 30i)))))"}));
INSTANTIATE_TEST_SUITE_P(SpvParserTest_FMul,
SpvBinaryArithTest,
@ -497,7 +493,7 @@ INSTANTIATE_TEST_SUITE_P(SpvParserTest_FMul,
BinaryData{"float", "float_50", "OpFMul", "float_60", "f32", "50.0f",
"*", "60.0f"}, // Vector float
BinaryData{"v2float", "v2float_50_60", "OpFMul", "v2float_60_50",
"vec2<f32>", AstFor("v2float_50_60"), "*",
"vec2f", AstFor("v2float_50_60"), "*",
AstFor("v2float_60_50")}));
INSTANTIATE_TEST_SUITE_P(SpvParserTest_UDiv,
@ -506,17 +502,16 @@ INSTANTIATE_TEST_SUITE_P(SpvParserTest_UDiv,
// Both uint
BinaryData{"uint", "uint_10", "OpUDiv", "uint_20", "u32", "10u", "/",
"20u"}, // Both v2uint
BinaryData{"v2uint", "v2uint_10_20", "OpUDiv", "v2uint_20_10",
"vec2<u32>", AstFor("v2uint_10_20"), "/",
AstFor("v2uint_20_10")}));
BinaryData{"v2uint", "v2uint_10_20", "OpUDiv", "v2uint_20_10", "vec2u",
AstFor("v2uint_10_20"), "/", AstFor("v2uint_20_10")}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_SDiv,
INSTANTIATE_TEST_SUITE_P(SpvParserTest_SDiv,
SpvBinaryArithTest,
::testing::Values(
// Both int
BinaryData{"int", "int_30", "OpSDiv", "int_40", "i32", "30i", "/", "40i"}, // Both v2int
BinaryData{"v2int", "v2int_30_40", "OpSDiv", "v2int_40_30", "vec2<i32>",
BinaryData{"int", "int_30", "OpSDiv", "int_40", "i32", "30i", "/",
"40i"}, // Both v2int
BinaryData{"v2int", "v2int_30_40", "OpSDiv", "v2int_40_30", "vec2i",
AstFor("v2int_30_40"), "/", AstFor("v2int_40_30")}));
INSTANTIATE_TEST_SUITE_P(
@ -528,11 +523,11 @@ INSTANTIATE_TEST_SUITE_P(
// Mixed, returning int, first arg uint
BinaryData{"int", "uint_10", "OpSDiv", "int_30", "i32", "bitcast<i32>(10u)", "/",
"30i"}, // Mixed, returning v2int, first arg v2uint
BinaryData{"v2int", "v2uint_10_20", "OpSDiv", "v2int_30_40", "vec2<i32>",
BinaryData{"v2int", "v2uint_10_20", "OpSDiv", "v2int_30_40", "vec2i",
AstFor("cast_int_v2uint_10_20"), "/", AstFor("v2int_30_40")},
// Mixed, returning v2int, second arg v2uint
BinaryData{"v2int", "v2int_30_40", "OpSDiv", "v2uint_10_20", "vec2<i32>",
AstFor("v2int_30_40"), "/", AstFor("cast_int_v2uint_10_20")}));
BinaryData{"v2int", "v2int_30_40", "OpSDiv", "v2uint_10_20", "vec2i", AstFor("v2int_30_40"),
"/", AstFor("cast_int_v2uint_10_20")}));
TEST_F(SpvBinaryArithTestBasic, SDiv_Scalar_UnsignedResult) {
// The WGSL signed division operator expects both operands to be signed
@ -574,8 +569,7 @@ TEST_F(SpvBinaryArithTestBasic, SDiv_Vector_UnsignedResult) {
auto ast_body = fe.ast_body();
EXPECT_THAT(
test::ToString(p->program(), ast_body),
HasSubstr(
R"(let x_1 : vec2<u32> = bitcast<vec2<u32>>((vec2<i32>(30i, 40i) / vec2<i32>(40i, 30i)));)"));
HasSubstr(R"(let x_1 : vec2u = bitcast<vec2u>((vec2i(30i, 40i) / vec2i(40i, 30i)));)"));
}
INSTANTIATE_TEST_SUITE_P(SpvParserTest_FDiv,
@ -585,7 +579,7 @@ INSTANTIATE_TEST_SUITE_P(SpvParserTest_FDiv,
BinaryData{"float", "float_50", "OpFDiv", "float_60", "f32", "50.0f",
"/", "60.0f"}, // Vector float
BinaryData{"v2float", "v2float_50_60", "OpFDiv", "v2float_60_50",
"vec2<f32>", AstFor("v2float_50_60"), "/",
"vec2f", AstFor("v2float_50_60"), "/",
AstFor("v2float_60_50")}));
INSTANTIATE_TEST_SUITE_P(SpvParserTest_UMod,
@ -594,20 +588,19 @@ INSTANTIATE_TEST_SUITE_P(SpvParserTest_UMod,
// Both uint
BinaryData{"uint", "uint_10", "OpUMod", "uint_20", "u32", "10u", "%",
"20u"}, // Both v2uint
BinaryData{"v2uint", "v2uint_10_20", "OpUMod", "v2uint_20_10",
"vec2<u32>", AstFor("v2uint_10_20"), "%",
AstFor("v2uint_20_10")}));
BinaryData{"v2uint", "v2uint_10_20", "OpUMod", "v2uint_20_10", "vec2u",
AstFor("v2uint_10_20"), "%", AstFor("v2uint_20_10")}));
// Currently WGSL is missing a mapping for OpSRem
// https://github.com/gpuweb/gpuweb/issues/702
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_SMod,
INSTANTIATE_TEST_SUITE_P(SpvParserTest_SMod,
SpvBinaryArithTest,
::testing::Values(
// Both int
BinaryData{"int", "int_30", "OpSMod", "int_40", "i32", "30i", "%", "40i"}, // Both v2int
BinaryData{"v2int", "v2int_30_40", "OpSMod", "v2int_40_30", "vec2<i32>",
BinaryData{"int", "int_30", "OpSMod", "int_40", "i32", "30i", "%",
"40i"}, // Both v2int
BinaryData{"v2int", "v2int_30_40", "OpSMod", "v2int_40_30", "vec2i",
AstFor("v2int_30_40"), "%", AstFor("v2int_40_30")}));
INSTANTIATE_TEST_SUITE_P(
@ -619,11 +612,11 @@ INSTANTIATE_TEST_SUITE_P(
// Mixed, returning int, first arg uint
BinaryData{"int", "uint_10", "OpSMod", "int_30", "i32", "bitcast<i32>(10u)", "%",
"30i"}, // Mixed, returning v2int, first arg v2uint
BinaryData{"v2int", "v2uint_10_20", "OpSMod", "v2int_30_40", "vec2<i32>",
BinaryData{"v2int", "v2uint_10_20", "OpSMod", "v2int_30_40", "vec2i",
AstFor("cast_int_v2uint_10_20"), "%", AstFor("v2int_30_40")},
// Mixed, returning v2int, second arg v2uint
BinaryData{"v2int", "v2int_30_40", "OpSMod", "v2uint_10_20", "vec2<i32>",
AstFor("v2int_30_40"), "%", AstFor("cast_int_v2uint_10_20")}));
BinaryData{"v2int", "v2int_30_40", "OpSMod", "v2uint_10_20", "vec2i", AstFor("v2int_30_40"),
"%", AstFor("cast_int_v2uint_10_20")}));
TEST_F(SpvBinaryArithTestBasic, SMod_Scalar_UnsignedResult) {
// The WGSL signed modulus operator expects both operands to be signed
@ -665,8 +658,7 @@ TEST_F(SpvBinaryArithTestBasic, SMod_Vector_UnsignedResult) {
auto ast_body = fe.ast_body();
EXPECT_THAT(
test::ToString(p->program(), ast_body),
HasSubstr(
R"(let x_1 : vec2<u32> = bitcast<vec2<u32>>((vec2<i32>(30i, 40i) % vec2<i32>(40i, 30i)));)"));
HasSubstr(R"(let x_1 : vec2u = bitcast<vec2u>((vec2i(30i, 40i) % vec2i(40i, 30i)));)"));
}
INSTANTIATE_TEST_SUITE_P(SpvParserTest_FRem,
@ -676,7 +668,7 @@ INSTANTIATE_TEST_SUITE_P(SpvParserTest_FRem,
BinaryData{"float", "float_50", "OpFRem", "float_60", "f32", "50.0f",
"%", "60.0f"}, // Vector float
BinaryData{"v2float", "v2float_50_60", "OpFRem", "v2float_60_50",
"vec2<f32>", AstFor("v2float_50_60"), "%",
"vec2f", AstFor("v2float_50_60"), "%",
AstFor("v2float_60_50")}));
TEST_F(SpvBinaryArithTestBasic, FMod_Scalar) {
@ -710,7 +702,7 @@ TEST_F(SpvBinaryArithTestBasic, FMod_Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<f32> = (v2float_50_60 - (v2float_60_50 * "
HasSubstr("let x_1 : vec2f = (v2float_50_60 - (v2float_60_50 * "
"floor((v2float_50_60 / v2float_60_50))));"));
}
@ -730,7 +722,7 @@ TEST_F(SpvBinaryArithTestBasic, VectorTimesScalar) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_10 : vec2<f32> = (x_1 * x_2);"));
HasSubstr("let x_10 : vec2f = (x_1 * x_2);"));
}
TEST_F(SpvBinaryArithTestBasic, MatrixTimesScalar) {
@ -768,7 +760,7 @@ TEST_F(SpvBinaryArithTestBasic, VectorTimesMatrix) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_10 : vec2<f32> = (x_1 * x_2);"));
HasSubstr("let x_10 : vec2f = (x_1 * x_2);"));
}
TEST_F(SpvBinaryArithTestBasic, MatrixTimesVector) {
@ -787,7 +779,7 @@ TEST_F(SpvBinaryArithTestBasic, MatrixTimesVector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_10 : vec2<f32> = (x_1 * x_2);"));
HasSubstr("let x_10 : vec2f = (x_1 * x_2);"));
}
TEST_F(SpvBinaryArithTestBasic, MatrixTimesMatrix) {
@ -847,8 +839,8 @@ TEST_F(SpvBinaryArithTestBasic, OuterProduct) {
auto ast_body = fe.ast_body();
auto got = test::ToString(p->program(), ast_body);
EXPECT_THAT(got, HasSubstr("let x_3 : mat2x3<f32> = mat2x3<f32>("
"vec3<f32>((x_2.x * x_1.x), (x_2.x * x_1.y), (x_2.x * x_1.z)), "
"vec3<f32>((x_2.y * x_1.x), (x_2.y * x_1.y), (x_2.y * x_1.z)));"))
"vec3f((x_2.x * x_1.x), (x_2.x * x_1.y), (x_2.x * x_1.z)), "
"vec3f((x_2.y * x_1.x), (x_2.y * x_1.y), (x_2.y * x_1.z)));"))
<< got;
}
@ -912,9 +904,8 @@ INSTANTIATE_TEST_SUITE_P(
BuiltinData{"OpDPdyCoarse", "dpdyCoarse"},
BuiltinData{"OpFwidthCoarse", "fwidthCoarse"}),
::testing::Values(ArgAndTypeData{"float", "float_50", "f32"},
ArgAndTypeData{"v2float", "v2float_50_60", "vec2<f32>"},
ArgAndTypeData{"v3float", "v3float_50_60_70",
"vec3<f32>"})));
ArgAndTypeData{"v2float", "v2float_50_60", "vec2f"},
ArgAndTypeData{"v3float", "v3float_50_60_70", "vec3f"})));
TEST_F(SpvUnaryArithTest, Transpose_2x2) {
const auto assembly = Preamble() + R"(

178
src/tint/reader/spirv/function_bit_test.cc
View File

@ -69,25 +69,25 @@ std::string SimplePreamble() {
// Returns the AST dump for a given SPIR-V assembly constant.
std::string AstFor(std::string assembly) {
if (assembly == "v2uint_10_20") {
return "vec2<u32>(10u, 20u)";
return "vec2u(10u, 20u)";
}
if (assembly == "v2uint_20_10") {
return "vec2<u32>(20u, 10u)";
return "vec2u(20u, 10u)";
}
if (assembly == "v2int_30_40") {
return "vec2<i32>(30i, 40i)";
return "vec2i(30i, 40i)";
}
if (assembly == "v2int_40_30") {
return "vec2<i32>(40i, 30i)";
return "vec2i(40i, 30i)";
}
if (assembly == "cast_int_v2uint_10_20") {
return "bitcast<vec2<i32>(vec2<u32>(10u, 20u))";
return "bitcast<vec2i(vec2u(10u, 20u))";
}
if (assembly == "v2float_50_60") {
return "vec2<f32>(50.0, 60.0))";
return "vec2f(50.0, 60.0))";
}
if (assembly == "v2float_60_50") {
return "vec2<f32>(60.0, 50.0))";
return "vec2f(60.0, 50.0))";
}
return "bad case";
}
@ -182,10 +182,10 @@ INSTANTIATE_TEST_SUITE_P(
// int, uint -> int
BinaryData{"int", "int_30", "OpShiftLeftLogical", "uint_20", "i32", "30i", "<<", "20u"},
// v2uint v2uint -> v2uint
BinaryData{"v2uint", "v2uint_10_20", "OpShiftLeftLogical", "v2uint_20_10", "vec2<u32>",
BinaryData{"v2uint", "v2uint_10_20", "OpShiftLeftLogical", "v2uint_20_10", "vec2u",
AstFor("v2uint_10_20"), "<<", AstFor("v2uint_20_10")},
// v2int, v2uint -> v2int
BinaryData{"v2int", "v2int_30_40", "OpShiftLeftLogical", "v2uint_20_10", "vec2<i32>",
BinaryData{"v2int", "v2int_30_40", "OpShiftLeftLogical", "v2uint_20_10", "vec2i",
AstFor("v2int_30_40"), "<<", AstFor("v2uint_20_10")}));
INSTANTIATE_TEST_SUITE_P(
@ -200,23 +200,22 @@ INSTANTIATE_TEST_SUITE_P(
BinaryDataGeneral{"uint", "uint_10", "OpShiftLeftLogical", "int_40", "u32",
"(10u << bitcast<u32>(40i))"},
// v2uint, v2int -> v2uint
BinaryDataGeneral{"v2uint", "v2uint_10_20", "OpShiftLeftLogical", "v2uint_20_10",
"vec2<u32>", "(vec2<u32>(10u, 20u) << vec2<u32>(20u, 10u))"},
BinaryDataGeneral{"v2uint", "v2uint_10_20", "OpShiftLeftLogical", "v2uint_20_10", "vec2u",
"(vec2u(10u, 20u) << vec2u(20u, 10u))"},
// v2int, v2int -> v2int
BinaryDataGeneral{"v2int", "v2int_30_40", "OpShiftLeftLogical", "v2int_40_30", "vec2<i32>",
"(vec2<i32>(30i, 40i) << bitcast<vec2<u32>>(vec2<i32>(40i, 30i)))"}));
BinaryDataGeneral{"v2int", "v2int_30_40", "OpShiftLeftLogical", "v2int_40_30", "vec2i",
"(vec2i(30i, 40i) << bitcast<vec2u>(vec2i(40i, 30i)))"}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_ShiftLeftLogical_BitcastResult,
INSTANTIATE_TEST_SUITE_P(SpvParserTest_ShiftLeftLogical_BitcastResult,
SpvBinaryBitGeneralTest,
::testing::Values(
// int, int -> uint
BinaryDataGeneral{"uint", "int_30", "OpShiftLeftLogical", "uint_10", "u32",
"bitcast<u32>((30i << 10u))"},
BinaryDataGeneral{"uint", "int_30", "OpShiftLeftLogical", "uint_10",
"u32", "bitcast<u32>((30i << 10u))"},
// v2uint, v2int -> v2uint
BinaryDataGeneral{"v2uint", "v2int_30_40", "OpShiftLeftLogical", "v2uint_20_10",
"vec2<u32>",
"bitcast<vec2<u32>>((vec2<i32>(30i, 40i) << vec2<u32>(20u, 10u)))"}));
BinaryDataGeneral{
"v2uint", "v2int_30_40", "OpShiftLeftLogical", "v2uint_20_10",
"vec2u", "bitcast<vec2u>((vec2i(30i, 40i) << vec2u(20u, 10u)))"}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_ShiftRightLogical_Arg2Unsigned,
@ -229,12 +228,12 @@ INSTANTIATE_TEST_SUITE_P(
BinaryDataGeneral{"int", "int_30", "OpShiftRightLogical", "uint_20", "i32",
"bitcast<i32>((bitcast<u32>(30i) >> 20u))"},
// v2uint, v2uint -> v2uint
BinaryDataGeneral{"v2uint", "v2uint_10_20", "OpShiftRightLogical", "v2uint_20_10",
"vec2<u32>", "(vec2<u32>(10u, 20u) >> vec2<u32>(20u, 10u))"},
BinaryDataGeneral{"v2uint", "v2uint_10_20", "OpShiftRightLogical", "v2uint_20_10", "vec2u",
"(vec2u(10u, 20u) >> vec2u(20u, 10u))"},
// v2int, v2uint -> v2int
BinaryDataGeneral{
"v2int", "v2int_30_40", "OpShiftRightLogical", "v2uint_10_20", "vec2<i32>",
R"(bitcast<vec2<i32>>((bitcast<vec2<u32>>(vec2<i32>(30i, 40i)) >> vec2<u32>(10u, 20u))))"}));
"v2int", "v2int_30_40", "OpShiftRightLogical", "v2uint_10_20", "vec2i",
R"(bitcast<vec2i>((bitcast<vec2u>(vec2i(30i, 40i)) >> vec2u(10u, 20u))))"}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_ShiftRightLogical_Arg2Signed,
@ -247,13 +246,12 @@ INSTANTIATE_TEST_SUITE_P(
BinaryDataGeneral{"int", "int_30", "OpShiftRightLogical", "int_40", "i32",
"bitcast<i32>((bitcast<u32>(30i) >> bitcast<u32>(40i)))"},
// v2uint, v2int -> v2uint
BinaryDataGeneral{"v2uint", "v2uint_10_20", "OpShiftRightLogical", "v2int_30_40",
"vec2<u32>",
"(vec2<u32>(10u, 20u) >> bitcast<vec2<u32>>(vec2<i32>(30i, 40i)))"},
BinaryDataGeneral{"v2uint", "v2uint_10_20", "OpShiftRightLogical", "v2int_30_40", "vec2u",
"(vec2u(10u, 20u) >> bitcast<vec2u>(vec2i(30i, 40i)))"},
// v2int, v2int -> v2int
BinaryDataGeneral{
"v2int", "v2int_40_30", "OpShiftRightLogical", "v2int_30_40", "vec2<i32>",
R"(bitcast<vec2<i32>>((bitcast<vec2<u32>>(vec2<i32>(40i, 30i)) >> bitcast<vec2<u32>>(vec2<i32>(30i, 40i)))))"}));
"v2int", "v2int_40_30", "OpShiftRightLogical", "v2int_30_40", "vec2i",
R"(bitcast<vec2i>((bitcast<vec2u>(vec2i(40i, 30i)) >> bitcast<vec2u>(vec2i(30i, 40i)))))"}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_ShiftRightLogical_BitcastResult,
@ -263,9 +261,8 @@ INSTANTIATE_TEST_SUITE_P(
BinaryDataGeneral{"int", "uint_20", "OpShiftRightLogical", "uint_10", "i32",
"bitcast<i32>((20u >> 10u))"},
// v2uint, v2uint -> v2int
BinaryDataGeneral{"v2int", "v2uint_10_20", "OpShiftRightLogical", "v2uint_20_10",
"vec2<i32>",
R"(bitcast<vec2<i32>>((vec2<u32>(10u, 20u) >> vec2<u32>(20u, 10u))))"}));
BinaryDataGeneral{"v2int", "v2uint_10_20", "OpShiftRightLogical", "v2uint_20_10", "vec2i",
R"(bitcast<vec2i>((vec2u(10u, 20u) >> vec2u(20u, 10u))))"}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_ShiftRightArithmetic_Arg2Unsigned,
@ -279,11 +276,11 @@ INSTANTIATE_TEST_SUITE_P(
"(30i >> 10u)"},
// v2uint, v2uint -> v2uint
BinaryDataGeneral{
"v2uint", "v2uint_10_20", "OpShiftRightArithmetic", "v2uint_20_10", "vec2<u32>",
R"(bitcast<vec2<u32>>((bitcast<vec2<i32>>(vec2<u32>(10u, 20u)) >> vec2<u32>(20u, 10u))))"},
"v2uint", "v2uint_10_20", "OpShiftRightArithmetic", "v2uint_20_10", "vec2u",
R"(bitcast<vec2u>((bitcast<vec2i>(vec2u(10u, 20u)) >> vec2u(20u, 10u))))"},
// v2int, v2uint -> v2int
BinaryDataGeneral{"v2int", "v2int_40_30", "OpShiftRightArithmetic", "v2uint_20_10",
"vec2<i32>", "(vec2<i32>(40i, 30i) >> vec2<u32>(20u, 10u))"}));
BinaryDataGeneral{"v2int", "v2int_40_30", "OpShiftRightArithmetic", "v2uint_20_10", "vec2i",
"(vec2i(40i, 30i) >> vec2u(20u, 10u))"}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_ShiftRightArithmetic_Arg2Signed,
@ -297,24 +294,22 @@ INSTANTIATE_TEST_SUITE_P(
"(30i >> bitcast<u32>(40i))"},
// v2uint, v2int -> v2uint
BinaryDataGeneral{
"v2uint", "v2uint_10_20", "OpShiftRightArithmetic", "v2int_30_40", "vec2<u32>",
R"(bitcast<vec2<u32>>((bitcast<vec2<i32>>(vec2<u32>(10u, 20u)) >> bitcast<vec2<u32>>(vec2<i32>(30i, 40i)))))"},
"v2uint", "v2uint_10_20", "OpShiftRightArithmetic", "v2int_30_40", "vec2u",
R"(bitcast<vec2u>((bitcast<vec2i>(vec2u(10u, 20u)) >> bitcast<vec2u>(vec2i(30i, 40i)))))"},
// v2int, v2int -> v2int
BinaryDataGeneral{"v2int", "v2int_40_30", "OpShiftRightArithmetic", "v2int_30_40",
"vec2<i32>",
"(vec2<i32>(40i, 30i) >> bitcast<vec2<u32>>(vec2<i32>(30i, 40i)))"}));
BinaryDataGeneral{"v2int", "v2int_40_30", "OpShiftRightArithmetic", "v2int_30_40", "vec2i",
"(vec2i(40i, 30i) >> bitcast<vec2u>(vec2i(30i, 40i)))"}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_ShiftRightArithmetic_BitcastResult,
INSTANTIATE_TEST_SUITE_P(SpvParserTest_ShiftRightArithmetic_BitcastResult,
SpvBinaryBitGeneralTest,
::testing::Values(
// int, uint -> uint
BinaryDataGeneral{"uint", "int_30", "OpShiftRightArithmetic", "uint_10", "u32",
"bitcast<u32>((30i >> 10u))"},
BinaryDataGeneral{"uint", "int_30", "OpShiftRightArithmetic",
"uint_10", "u32", "bitcast<u32>((30i >> 10u))"},
// v2int, v2uint -> v2uint
BinaryDataGeneral{"v2uint", "v2int_30_40", "OpShiftRightArithmetic", "v2uint_20_10",
"vec2<u32>",
"bitcast<vec2<u32>>((vec2<i32>(30i, 40i) >> vec2<u32>(20u, 10u)))"}));
BinaryDataGeneral{
"v2uint", "v2int_30_40", "OpShiftRightArithmetic", "v2uint_20_10",
"vec2u", "bitcast<vec2u>((vec2i(30i, 40i) >> vec2u(20u, 10u)))"}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_BitwiseAnd,
@ -326,10 +321,10 @@ INSTANTIATE_TEST_SUITE_P(
BinaryData{"int", "int_30", "OpBitwiseAnd", "int_40", "i32", "30i", "&", "40i"},
// TODO(crbug.com/tint/678): Resolver fails on vector bitwise operations
// Both v2uint
BinaryData{"v2uint", "v2uint_10_20", "OpBitwiseAnd", "v2uint_20_10", "vec2<u32>",
BinaryData{"v2uint", "v2uint_10_20", "OpBitwiseAnd", "v2uint_20_10", "vec2u",
AstFor("v2uint_10_20"), "&", AstFor("v2uint_20_10")},
// Both v2int
BinaryData{"v2int", "v2int_30_40", "OpBitwiseAnd", "v2int_40_30", "vec2<i32>",
BinaryData{"v2int", "v2int_30_40", "OpBitwiseAnd", "v2int_40_30", "vec2i",
AstFor("v2int_30_40"), "&", AstFor("v2int_40_30")}));
INSTANTIATE_TEST_SUITE_P(
@ -349,13 +344,12 @@ INSTANTIATE_TEST_SUITE_P(
BinaryDataGeneral{"int", "uint_20", "OpBitwiseAnd", "uint_10", "i32",
"bitcast<i32>((20u & 10u))"},
// Mixed, returning v2uint
BinaryDataGeneral{
"v2uint", "v2int_30_40", "OpBitwiseAnd", "v2uint_10_20", "vec2<u32>",
R"(bitcast<vec2<u32>>((vec2<i32>(30i, 40i) & bitcast<vec2<i32>>(vec2<u32>(10u, 20u)))))"},
BinaryDataGeneral{"v2uint", "v2int_30_40", "OpBitwiseAnd", "v2uint_10_20", "vec2u",
R"(bitcast<vec2u>((vec2i(30i, 40i) & bitcast<vec2i>(vec2u(10u, 20u)))))"},
// Mixed, returning v2int
BinaryDataGeneral{
"v2int", "v2uint_10_20", "OpBitwiseAnd", "v2int_40_30", "vec2<i32>",
R"(bitcast<vec2<i32>>((vec2<u32>(10u, 20u) & bitcast<vec2<u32>>(vec2<i32>(40i, 30i)))))"}));
"v2int", "v2uint_10_20", "OpBitwiseAnd", "v2int_40_30", "vec2i",
R"(bitcast<vec2i>((vec2u(10u, 20u) & bitcast<vec2u>(vec2i(40i, 30i)))))"}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_BitwiseOr,
@ -367,10 +361,10 @@ INSTANTIATE_TEST_SUITE_P(
BinaryData{"int", "int_30", "OpBitwiseOr", "int_40", "i32", "30i", "|", "40i"},
// TODO(crbug.com/tint/678): Resolver fails on vector bitwise operations
// Both v2uint
BinaryData{"v2uint", "v2uint_10_20", "OpBitwiseOr", "v2uint_20_10", "vec2<u32>",
BinaryData{"v2uint", "v2uint_10_20", "OpBitwiseOr", "v2uint_20_10", "vec2u",
AstFor("v2uint_10_20"), "|", AstFor("v2uint_20_10")},
// Both v2int
BinaryData{"v2int", "v2int_30_40", "OpBitwiseOr", "v2int_40_30", "vec2<i32>",
BinaryData{"v2int", "v2int_30_40", "OpBitwiseOr", "v2int_40_30", "vec2i",
AstFor("v2int_30_40"), "|", AstFor("v2int_40_30")}));
INSTANTIATE_TEST_SUITE_P(
@ -390,13 +384,12 @@ INSTANTIATE_TEST_SUITE_P(
BinaryDataGeneral{"int", "uint_20", "OpBitwiseOr", "uint_10", "i32",
"bitcast<i32>((20u | 10u))"},
// Mixed, returning v2uint
BinaryDataGeneral{
"v2uint", "v2int_30_40", "OpBitwiseOr", "v2uint_10_20", "vec2<u32>",
R"(bitcast<vec2<u32>>((vec2<i32>(30i, 40i) | bitcast<vec2<i32>>(vec2<u32>(10u, 20u)))))"},
BinaryDataGeneral{"v2uint", "v2int_30_40", "OpBitwiseOr", "v2uint_10_20", "vec2u",
R"(bitcast<vec2u>((vec2i(30i, 40i) | bitcast<vec2i>(vec2u(10u, 20u)))))"},
// Mixed, returning v2int
BinaryDataGeneral{
"v2int", "v2uint_10_20", "OpBitwiseOr", "v2int_40_30", "vec2<i32>",
R"(bitcast<vec2<i32>>((vec2<u32>(10u, 20u) | bitcast<vec2<u32>>(vec2<i32>(40i, 30i)))))"}));
"v2int", "v2uint_10_20", "OpBitwiseOr", "v2int_40_30", "vec2i",
R"(bitcast<vec2i>((vec2u(10u, 20u) | bitcast<vec2u>(vec2i(40i, 30i)))))"}));
INSTANTIATE_TEST_SUITE_P(
SpvParserTest_BitwiseXor,
@ -408,10 +401,10 @@ INSTANTIATE_TEST_SUITE_P(
BinaryData{"int", "int_30", "OpBitwiseXor", "int_40", "i32", "30i", "^", "40i"},
// TODO(crbug.com/tint/678): Resolver fails on vector bitwise operations
// Both v2uint
BinaryData{"v2uint", "v2uint_10_20", "OpBitwiseXor", "v2uint_20_10", "vec2<u32>",
BinaryData{"v2uint", "v2uint_10_20", "OpBitwiseXor", "v2uint_20_10", "vec2u",
AstFor("v2uint_10_20"), "^", AstFor("v2uint_20_10")},
// Both v2int
BinaryData{"v2int", "v2int_30_40", "OpBitwiseXor", "v2int_40_30", "vec2<i32>",
BinaryData{"v2int", "v2int_30_40", "OpBitwiseXor", "v2int_40_30", "vec2i",
AstFor("v2int_30_40"), "^", AstFor("v2int_40_30")}));
INSTANTIATE_TEST_SUITE_P(
@ -431,13 +424,12 @@ INSTANTIATE_TEST_SUITE_P(
BinaryDataGeneral{"int", "uint_20", "OpBitwiseXor", "uint_10", "i32",
"bitcast<i32>((20u ^ 10u))"},
// Mixed, returning v2uint
BinaryDataGeneral{
"v2uint", "v2int_30_40", "OpBitwiseXor", "v2uint_10_20", "vec2<u32>",
R"(bitcast<vec2<u32>>((vec2<i32>(30i, 40i) ^ bitcast<vec2<i32>>(vec2<u32>(10u, 20u)))))"},
BinaryDataGeneral{"v2uint", "v2int_30_40", "OpBitwiseXor", "v2uint_10_20", "vec2u",
R"(bitcast<vec2u>((vec2i(30i, 40i) ^ bitcast<vec2i>(vec2u(10u, 20u)))))"},
// Mixed, returning v2int
BinaryDataGeneral{
"v2int", "v2uint_10_20", "OpBitwiseXor", "v2int_40_30", "vec2<i32>",
R"(bitcast<vec2<i32>>((vec2<u32>(10u, 20u) ^ bitcast<vec2<u32>>(vec2<i32>(40i, 30i)))))"}));
"v2int", "v2uint_10_20", "OpBitwiseXor", "v2int_40_30", "vec2i",
R"(bitcast<vec2i>((vec2u(10u, 20u) ^ bitcast<vec2u>(vec2i(40i, 30i)))))"}));
TEST_F(SpvUnaryBitTest, Not_Int_Int) {
const auto assembly = SimplePreamble() + R"(
@ -521,7 +513,7 @@ TEST_F(SpvUnaryBitTest, Not_SignedVec_SignedVec) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<i32> = ~(vec2<i32>(30i, 40i));"));
EXPECT_THAT(body, HasSubstr("let x_1 : vec2i = ~(vec2i(30i, 40i));"));
}
TEST_F(SpvUnaryBitTest, Not_SignedVec_UnsignedVec) {
@ -538,8 +530,7 @@ TEST_F(SpvUnaryBitTest, Not_SignedVec_UnsignedVec) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body,
HasSubstr("let x_1 : vec2<i32> = bitcast<vec2<i32>>(~(vec2<u32>(10u, 20u)));"));
EXPECT_THAT(body, HasSubstr("let x_1 : vec2i = bitcast<vec2i>(~(vec2u(10u, 20u)));"));
}
TEST_F(SpvUnaryBitTest, Not_UnsignedVec_SignedVec) {
@ -556,8 +547,7 @@ TEST_F(SpvUnaryBitTest, Not_UnsignedVec_SignedVec) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body,
HasSubstr("let x_1 : vec2<u32> = bitcast<vec2<u32>>(~(vec2<i32>(30i, 40i)));"));
EXPECT_THAT(body, HasSubstr("let x_1 : vec2u = bitcast<vec2u>(~(vec2i(30i, 40i)));"));
}
TEST_F(SpvUnaryBitTest, Not_UnsignedVec_UnsignedVec) {
const auto assembly = SimplePreamble() + R"(
@ -573,7 +563,7 @@ TEST_F(SpvUnaryBitTest, Not_UnsignedVec_UnsignedVec) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<u32> = ~(vec2<u32>(10u, 20u));"));
EXPECT_THAT(body, HasSubstr("let x_1 : vec2u = ~(vec2u(10u, 20u));"));
}
std::string BitTestPreamble() {
@ -674,7 +664,7 @@ TEST_F(SpvUnaryBitTest, BitCount_UintVector_UintVector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<u32> = countOneBits(v2u1);")) << body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec2u = countOneBits(v2u1);")) << body;
}
TEST_F(SpvUnaryBitTest, BitCount_UintVector_IntVector) {
@ -689,8 +679,7 @@ TEST_F(SpvUnaryBitTest, BitCount_UintVector_IntVector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<u32> = bitcast<vec2<u32>>(countOneBits(v2i1));"))
<< body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec2u = bitcast<vec2u>(countOneBits(v2i1));")) << body;
}
TEST_F(SpvUnaryBitTest, BitCount_IntVector_UintVector) {
@ -705,8 +694,7 @@ TEST_F(SpvUnaryBitTest, BitCount_IntVector_UintVector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<i32> = bitcast<vec2<i32>>(countOneBits(v2u1));"))
<< body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec2i = bitcast<vec2i>(countOneBits(v2u1));")) << body;
}
TEST_F(SpvUnaryBitTest, BitCount_IntVector_IntVector) {
@ -721,7 +709,7 @@ TEST_F(SpvUnaryBitTest, BitCount_IntVector_IntVector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<i32> = countOneBits(v2i1);")) << body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec2i = countOneBits(v2i1);")) << body;
}
TEST_F(SpvUnaryBitTest, BitReverse_Uint_Uint) {
@ -790,7 +778,7 @@ TEST_F(SpvUnaryBitTest, BitReverse_UintVector_UintVector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<u32> = reverseBits(v2u1);")) << body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec2u = reverseBits(v2u1);")) << body;
}
TEST_F(SpvUnaryBitTest, BitReverse_UintVector_IntVector) {
@ -829,7 +817,7 @@ TEST_F(SpvUnaryBitTest, BitReverse_IntVector_IntVector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<i32> = reverseBits(v2i1);")) << body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec2i = reverseBits(v2i1);")) << body;
}
TEST_F(SpvUnaryBitTest, InsertBits_Int) {
@ -875,9 +863,8 @@ TEST_F(SpvUnaryBitTest, InsertBits_IntVector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(
body,
HasSubstr(R"(let x_1 : vec2<i32> = insertBits(x_28, vec2<i32>(40i, 30i), 10u, 20u);)"))
EXPECT_THAT(body,
HasSubstr(R"(let x_1 : vec2i = insertBits(x_28, vec2i(40i, 30i), 10u, 20u);)"))
<< body;
}
@ -895,8 +882,7 @@ TEST_F(SpvUnaryBitTest, InsertBits_IntVector_SignedOffsetAndCount) {
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(
body,
HasSubstr(
R"(let x_1 : vec2<i32> = insertBits(x_28, vec2<i32>(40i, 30i), u32(10i), u32(20i));)"))
HasSubstr(R"(let x_1 : vec2i = insertBits(x_28, vec2i(40i, 30i), u32(10i), u32(20i));)"))
<< body;
}
@ -943,9 +929,8 @@ TEST_F(SpvUnaryBitTest, InsertBits_UintVector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(
body,
HasSubstr(R"(let x_1 : vec2<u32> = insertBits(x_26, vec2<u32>(20u, 10u), 10u, 20u);)"))
EXPECT_THAT(body,
HasSubstr(R"(let x_1 : vec2u = insertBits(x_26, vec2u(20u, 10u), 10u, 20u);)"))
<< body;
}
@ -963,8 +948,7 @@ TEST_F(SpvUnaryBitTest, InsertBits_UintVector_SignedOffsetAndCount) {
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(
body,
HasSubstr(
R"(let x_1 : vec2<u32> = insertBits(x_26, vec2<u32>(20u, 10u), u32(10i), u32(20i));)"))
HasSubstr(R"(let x_1 : vec2u = insertBits(x_26, vec2u(20u, 10u), u32(10i), u32(20i));)"))
<< body;
}
@ -1010,7 +994,7 @@ TEST_F(SpvUnaryBitTest, ExtractBits_IntVector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<i32> = extractBits(x_28, 10u, 20u);")) << body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec2i = extractBits(x_28, 10u, 20u);")) << body;
}
TEST_F(SpvUnaryBitTest, ExtractBits_IntVector_SignedOffsetAndCount) {
@ -1025,7 +1009,7 @@ TEST_F(SpvUnaryBitTest, ExtractBits_IntVector_SignedOffsetAndCount) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<i32> = extractBits(x_28, u32(10i), u32(20i));"))
EXPECT_THAT(body, HasSubstr("let x_1 : vec2i = extractBits(x_28, u32(10i), u32(20i));"))
<< body;
}
@ -1071,7 +1055,7 @@ TEST_F(SpvUnaryBitTest, ExtractBits_UintVector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<u32> = extractBits(x_26, 10u, 20u);")) << body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec2u = extractBits(x_26, 10u, 20u);")) << body;
}
TEST_F(SpvUnaryBitTest, ExtractBits_UintVector_SignedOffsetAndCount) {
@ -1086,7 +1070,7 @@ TEST_F(SpvUnaryBitTest, ExtractBits_UintVector_SignedOffsetAndCount) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<u32> = extractBits(x_26, u32(10i), u32(20i));"))
EXPECT_THAT(body, HasSubstr("let x_1 : vec2u = extractBits(x_26, u32(10i), u32(20i));"))
<< body;
}

48
src/tint/reader/spirv/function_composite_test.cc
View File

@ -96,9 +96,9 @@ TEST_F(SpvParserTest_Composite_Construct, Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr(R"(let x_1 : vec2<u32> = vec2<u32>(10u, 20u);
let x_2 : vec2<i32> = vec2<i32>(30i, 40i);
let x_3 : vec2<f32> = vec2<f32>(50.0f, 60.0f);
HasSubstr(R"(let x_1 : vec2u = vec2u(10u, 20u);
let x_2 : vec2i = vec2i(30i, 40i);
let x_3 : vec2f = vec2f(50.0f, 60.0f);
)"));
}
@ -117,9 +117,9 @@ TEST_F(SpvParserTest_Composite_Construct, Matrix) {
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : mat3x2<f32> = mat3x2<f32>("
"vec2<f32>(50.0f, 60.0f), "
"vec2<f32>(60.0f, 50.0f), "
"vec2<f32>(70.0f, 70.0f));"));
"vec2f(50.0f, 60.0f), "
"vec2f(60.0f, 50.0f), "
"vec2f(70.0f, 70.0f));"));
}
TEST_F(SpvParserTest_Composite_Construct, Array) {
@ -153,7 +153,7 @@ TEST_F(SpvParserTest_Composite_Construct, Struct) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : S = S(vec2<f32>(50.0f, 60.0f), 5u, 30i);"));
HasSubstr("let x_1 : S = S(vec2f(50.0f, 60.0f), 5u, 30i);"));
}
TEST_F(SpvParserTest_Composite_Construct, ConstantComposite_Struct_NoDeduplication) {
@ -201,7 +201,7 @@ TEST_F(SpvParserTest_CompositeExtract, Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : f32 = vec2<f32>(50.0f, 60.0f).y;"));
HasSubstr("let x_1 : f32 = vec2f(50.0f, 60.0f).y;"));
}
TEST_F(SpvParserTest_CompositeExtract, Vector_IndexTooBigError) {
@ -238,8 +238,7 @@ TEST_F(SpvParserTest_CompositeExtract, Matrix) {
auto fe = p->function_emitter(100);
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_2 : vec2<f32> = x_1[2u];"));
EXPECT_THAT(test::ToString(p->program(), ast_body), HasSubstr("let x_2 : vec2f = x_1[2u];"));
}
TEST_F(SpvParserTest_CompositeExtract, Matrix_IndexTooBigError) {
@ -448,9 +447,9 @@ TEST_F(SpvParserTest_CompositeInsert, Vector) {
auto ast_body = fe.ast_body();
auto got = test::ToString(p->program(), ast_body);
const auto* expected =
R"(var x_1_1 : vec2<f32> = vec2<f32>(50.0f, 60.0f);
R"(var x_1_1 : vec2f = vec2f(50.0f, 60.0f);
x_1_1.y = 70.0f;
let x_1 : vec2<f32> = x_1_1;
let x_1 : vec2f = x_1_1;
return;
)";
EXPECT_EQ(got, expected);
@ -492,7 +491,7 @@ TEST_F(SpvParserTest_CompositeInsert, Matrix) {
auto ast_body = fe.ast_body();
auto body_str = test::ToString(p->program(), ast_body);
EXPECT_THAT(body_str, HasSubstr(R"(var x_2_1 : mat3x2<f32> = x_1;
x_2_1[2u] = vec2<f32>(50.0f, 60.0f);
x_2_1[2u] = vec2f(50.0f, 60.0f);
let x_2 : mat3x2<f32> = x_2_1;
)")) << body_str;
}
@ -537,7 +536,7 @@ TEST_F(SpvParserTest_CompositeInsert, Matrix_Vector) {
auto ast_body = fe.ast_body();
auto body_str = test::ToString(p->program(), ast_body);
EXPECT_THAT(body_str, HasSubstr(R"(var x_2_1 : mat3x2<f32> = x_1;
x_2_1[2u] = vec2<f32>(50.0f, 60.0f);
x_2_1[2u] = vec2f(50.0f, 60.0f);
let x_2 : mat3x2<f32> = x_2_1;
return;
)")) << body_str;
@ -782,7 +781,7 @@ TEST_F(SpvParserTest_VectorShuffle, FunctionScopeOperands_UseBoth) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_10 : vec4<u32> = vec4<u32>(x_2.y, x_2.x, x_1.y, x_1.x);"));
HasSubstr("let x_10 : vec4u = vec4u(x_2.y, x_2.x, x_1.y, x_1.x);"));
}
TEST_F(SpvParserTest_VectorShuffle, ConstantOperands_UseBoth) {
@ -799,12 +798,11 @@ TEST_F(SpvParserTest_VectorShuffle, ConstantOperands_UseBoth) {
auto fe = p->function_emitter(100);
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_10 : vec4<u32> = vec4<u32>("
"vec2<u32>(4u, 3u).y, "
"vec2<u32>(4u, 3u).x, "
"vec2<u32>(3u, 4u).y, "
"vec2<u32>(3u, 4u).x);"));
EXPECT_THAT(test::ToString(p->program(), ast_body), HasSubstr("let x_10 : vec4u = vec4u("
"vec2u(4u, 3u).y, "
"vec2u(4u, 3u).x, "
"vec2u(3u, 4u).y, "
"vec2u(3u, 4u).x);"));
}
TEST_F(SpvParserTest_VectorShuffle, ConstantOperands_AllOnesMapToNull) {
@ -823,7 +821,7 @@ TEST_F(SpvParserTest_VectorShuffle, ConstantOperands_AllOnesMapToNull) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_10 : vec2<u32> = vec2<u32>(0u, x_1.y);"));
HasSubstr("let x_10 : vec2u = vec2u(0u, x_1.y);"));
}
TEST_F(SpvParserTest_VectorShuffle, FunctionScopeOperands_MixedInputOperandSizes) {
@ -845,7 +843,7 @@ TEST_F(SpvParserTest_VectorShuffle, FunctionScopeOperands_MixedInputOperandSizes
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_10 : vec2<u32> = vec2<u32>(x_1.y, x_3.z);"));
HasSubstr("let x_10 : vec2u = vec2u(x_1.y, x_3.z);"));
}
TEST_F(SpvParserTest_VectorShuffle, IndexTooBig_IsError) {
@ -926,9 +924,9 @@ TEST_F(SpvParserTest_VectorInsertDynamic, Sample) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto got = test::ToString(p->program(), ast_body);
EXPECT_THAT(got, HasSubstr(R"(var x_10_1 : vec2<u32> = x_1;
EXPECT_THAT(got, HasSubstr(R"(var x_10_1 : vec2u = x_1;
x_10_1[x_3] = x_2;
let x_10 : vec2<u32> = x_10_1;
let x_10 : vec2u = x_10_1;
)")) << got
<< assembly;
}

16
src/tint/reader/spirv/function_conversion_test.cc
View File

@ -100,7 +100,7 @@ TEST_F(SpvUnaryConversionTest, Bitcast_Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<f32> = bitcast<vec2<f32>>(vec2<u32>(10u, 20u));"));
HasSubstr("let x_1 : vec2f = bitcast<vec2f>(vec2u(10u, 20u));"));
}
TEST_F(SpvUnaryConversionTest, ConvertSToF_BadArg) {
@ -245,7 +245,7 @@ TEST_F(SpvUnaryConversionTest, ConvertSToF_Vector_FromSigned) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<f32> = vec2<f32>(x_30);"));
HasSubstr("let x_1 : vec2f = vec2f(x_30);"));
}
TEST_F(SpvUnaryConversionTest, ConvertSToF_Vector_FromUnsigned) {
@ -263,7 +263,7 @@ TEST_F(SpvUnaryConversionTest, ConvertSToF_Vector_FromUnsigned) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<f32> = vec2<f32>(bitcast<vec2<i32>>(x_30));"));
HasSubstr("let x_1 : vec2f = vec2f(bitcast<vec2i>(x_30));"));
}
TEST_F(SpvUnaryConversionTest, ConvertUToF_Scalar_BadArgType) {
@ -348,7 +348,7 @@ TEST_F(SpvUnaryConversionTest, ConvertUToF_Vector_FromSigned) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<f32> = vec2<f32>(bitcast<vec2<u32>>(x_30));"));
HasSubstr("let x_1 : vec2f = vec2f(bitcast<vec2u>(x_30));"));
}
TEST_F(SpvUnaryConversionTest, ConvertUToF_Vector_FromUnsigned) {
@ -366,7 +366,7 @@ TEST_F(SpvUnaryConversionTest, ConvertUToF_Vector_FromUnsigned) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<f32> = vec2<f32>(x_30);"));
HasSubstr("let x_1 : vec2f = vec2f(x_30);"));
}
TEST_F(SpvUnaryConversionTest, ConvertFToS_Scalar_BadArgType) {
@ -451,7 +451,7 @@ TEST_F(SpvUnaryConversionTest, ConvertFToS_Vector_ToSigned) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<i32> = vec2<i32>(x_30);"));
HasSubstr("let x_1 : vec2i = vec2i(x_30);"));
}
TEST_F(SpvUnaryConversionTest, ConvertFToS_Vector_ToUnsigned) {
@ -469,7 +469,7 @@ TEST_F(SpvUnaryConversionTest, ConvertFToS_Vector_ToUnsigned) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<u32> = bitcast<vec2<u32>>(vec2<i32>(x_30));"));
HasSubstr("let x_1 : vec2u = bitcast<vec2u>(vec2i(x_30));"));
}
TEST_F(SpvUnaryConversionTest, ConvertFToU_Scalar_BadArgType) {
@ -568,7 +568,7 @@ TEST_F(SpvUnaryConversionTest, ConvertFToU_Vector_ToUnsigned) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<u32> = vec2<u32>(x_30);"));
HasSubstr("let x_1 : vec2u = vec2u(x_30);"));
}
TEST_F(SpvUnaryConversionTest, ConvertFToU_HoistedValue) {

4
src/tint/reader/spirv/function_decl_test.cc
View File

@ -186,7 +186,7 @@ TEST_F(SpvParserTest, Emit_FunctionDecl_ParamPtrTexture_ParamPtrSampler) {
auto got = test::ToString(p->program());
std::string expect = R"(fn x_200(x_14 : texture_2d<f32>, x_15 : sampler) {
let x_20 : vec4<f32> = textureSample(x_14, x_15, vec2<f32>());
let x_20 : vec4f = textureSample(x_14, x_15, vec2f());
return;
}
)";
@ -216,7 +216,7 @@ TEST_F(SpvParserTest, Emit_FunctionDecl_ParamTexture_ParamSampler) {
auto got = test::ToString(p->program());
std::string expect = R"(fn x_200(x_14 : texture_2d<f32>, x_15 : sampler) {
let x_20 : vec4<f32> = textureSample(x_14, x_15, vec2<f32>());
let x_20 : vec4f = textureSample(x_14, x_15, vec2f());
return;
}
)";

120
src/tint/reader/spirv/function_glsl_std_450_test.cc
View File

@ -282,8 +282,7 @@ TEST_P(SpvParserTest_GlslStd450_Floating_Floating, Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<f32> = " + GetParam().wgsl_func + "(v2f1);"))
<< body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec2f = " + GetParam().wgsl_func + "(v2f1);")) << body;
}
TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloating, Scalar) {
@ -315,7 +314,7 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloating, Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<f32> = " + GetParam().wgsl_func + "(v2f1, v2f2);"))
EXPECT_THAT(body, HasSubstr("let x_1 : vec2f = " + GetParam().wgsl_func + "(v2f1, v2f2);"))
<< body;
}
@ -351,7 +350,7 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloatingFloating, Vector) {
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body,
HasSubstr("let x_1 : vec2<f32> = " + GetParam().wgsl_func + "(v2f1, v2f2, v2f3);"))
HasSubstr("let x_1 : vec2f = " + GetParam().wgsl_func + "(v2f1, v2f2, v2f3);"))
<< body;
}
@ -385,7 +384,7 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingInting, Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<f32> = " + GetParam().wgsl_func + "(v2f1, v2i1);"))
EXPECT_THAT(body, HasSubstr("let x_1 : vec2f = " + GetParam().wgsl_func + "(v2f1, v2i1);"))
<< body;
}
@ -403,7 +402,7 @@ TEST_P(SpvParserTest_GlslStd450_Float3_Float3Float3, Samples) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec3<f32> = " + GetParam().wgsl_func + "(v3f1, v3f2);"))
EXPECT_THAT(body, HasSubstr("let x_1 : vec3f = " + GetParam().wgsl_func + "(v3f1, v3f2);"))
<< body;
}
@ -542,8 +541,7 @@ TEST_P(SpvParserTest_GlslStd450_Inting_Inting, Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<i32> = " + GetParam().wgsl_func + "(v2i1);"))
<< body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec2i = " + GetParam().wgsl_func + "(v2i1);")) << body;
}
TEST_P(SpvParserTest_GlslStd450_Inting_Inting_SignednessCoercing, Vector_UnsignedArg) {
@ -560,8 +558,8 @@ TEST_P(SpvParserTest_GlslStd450_Inting_Inting_SignednessCoercing, Vector_Unsigne
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<i32> = " + GetParam().wgsl_func +
"(bitcast<vec2<i32>>(v2u1));"))
EXPECT_THAT(body,
HasSubstr("let x_1 : vec2i = " + GetParam().wgsl_func + "(bitcast<vec2i>(v2u1));"))
<< body;
}
@ -579,8 +577,8 @@ TEST_P(SpvParserTest_GlslStd450_Inting_Inting_SignednessCoercing, Vector_Unsigne
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<u32> = bitcast<vec2<u32>>(" + GetParam().wgsl_func +
"(v2i1));"))
EXPECT_THAT(body,
HasSubstr("let x_1 : vec2u = bitcast<vec2u>(" + GetParam().wgsl_func + "(v2i1));"))
<< body;
}
@ -615,7 +613,7 @@ TEST_P(SpvParserTest_GlslStd450_Inting_IntingInting, Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<i32> = " + GetParam().wgsl_func + "(v2i1, v2i2);"))
EXPECT_THAT(body, HasSubstr("let x_1 : vec2i = " + GetParam().wgsl_func + "(v2i1, v2i2);"))
<< body;
}
@ -652,7 +650,7 @@ TEST_P(SpvParserTest_GlslStd450_Inting_IntingIntingInting, Vector) {
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body,
HasSubstr("let x_1 : vec2<i32> = " + GetParam().wgsl_func + "(v2i1, v2i2, v2i3);"))
HasSubstr("let x_1 : vec2i = " + GetParam().wgsl_func + "(v2i1, v2i2, v2i3);"))
<< body;
}
@ -707,8 +705,7 @@ TEST_P(SpvParserTest_GlslStd450_Uinting_Uinting, Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<u32> = " + GetParam().wgsl_func + "(v2u1);"))
<< body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec2u = " + GetParam().wgsl_func + "(v2u1);")) << body;
}
TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUinting, Scalar) {
@ -741,7 +738,7 @@ TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUinting, Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<u32> = " + GetParam().wgsl_func + "(v2u1, v2u2);"))
EXPECT_THAT(body, HasSubstr("let x_1 : vec2u = " + GetParam().wgsl_func + "(v2u1, v2u2);"))
<< body;
}
@ -777,7 +774,7 @@ TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUintingUinting, Vector) {
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body,
HasSubstr("let x_1 : vec2<u32> = " + GetParam().wgsl_func + "(v2u1, v2u2, v2u3);"))
HasSubstr("let x_1 : vec2u = " + GetParam().wgsl_func + "(v2u1, v2u2, v2u3);"))
<< body;
}
@ -828,7 +825,7 @@ TEST_F(SpvParserTest, Normalize_Vector2) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec2<f32> = normalize(v2f1);")) << body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec2f = normalize(v2f1);")) << body;
}
TEST_F(SpvParserTest, Normalize_Vector3) {
@ -844,7 +841,7 @@ TEST_F(SpvParserTest, Normalize_Vector3) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec3<f32> = normalize(v3f1);")) << body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec3f = normalize(v3f1);")) << body;
}
TEST_F(SpvParserTest, Normalize_Vector4) {
@ -860,7 +857,7 @@ TEST_F(SpvParserTest, Normalize_Vector4) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : vec4<f32> = normalize(v4f1);")) << body;
EXPECT_THAT(body, HasSubstr("let x_1 : vec4f = normalize(v4f1);")) << body;
}
// Check that we convert signedness of operands and result type.
@ -880,9 +877,7 @@ TEST_F(SpvParserTest, RectifyOperandsAndResult_SAbs) {
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr(R"(let x_1 : u32 = bitcast<u32>(abs(bitcast<i32>(u1)));)")) << body;
EXPECT_THAT(
body,
HasSubstr(R"(let x_2 : vec2<u32> = bitcast<vec2<u32>>(abs(bitcast<vec2<i32>>(v2u1)));)"))
EXPECT_THAT(body, HasSubstr(R"(let x_2 : vec2u = bitcast<vec2u>(abs(bitcast<vec2i>(v2u1)));)"))
<< body;
}
@ -906,7 +901,7 @@ TEST_F(SpvParserTest, RectifyOperandsAndResult_SMax) {
EXPECT_THAT(
body,
HasSubstr(
R"(let x_2 : vec2<u32> = bitcast<vec2<u32>>(max(bitcast<vec2<i32>>(v2u1), bitcast<vec2<i32>>(v2u2)));)"))
R"(let x_2 : vec2u = bitcast<vec2u>(max(bitcast<vec2i>(v2u1), bitcast<vec2i>(v2u2)));)"))
<< body;
}
@ -930,7 +925,7 @@ TEST_F(SpvParserTest, RectifyOperandsAndResult_SMin) {
EXPECT_THAT(
body,
HasSubstr(
R"(let x_2 : vec2<u32> = bitcast<vec2<u32>>(min(bitcast<vec2<i32>>(v2u1), bitcast<vec2<i32>>(v2u2)));)"))
R"(let x_2 : vec2u = bitcast<vec2u>(min(bitcast<vec2i>(v2u1), bitcast<vec2i>(v2u2)));)"))
<< body;
}
@ -955,7 +950,7 @@ TEST_F(SpvParserTest, RectifyOperandsAndResult_SClamp) {
EXPECT_THAT(
body,
HasSubstr(
R"(let x_2 : vec2<u32> = bitcast<vec2<u32>>(clamp(bitcast<vec2<i32>>(v2u1), v2i2, bitcast<vec2<i32>>(v2u3)));)"))
R"(let x_2 : vec2u = bitcast<vec2u>(clamp(bitcast<vec2i>(v2u1), v2i2, bitcast<vec2i>(v2u3)));)"))
<< body;
}
@ -979,7 +974,7 @@ TEST_F(SpvParserTest, RectifyOperandsAndResult_UMax) {
EXPECT_THAT(
body,
HasSubstr(
R"(let x_2 : vec2<i32> = bitcast<vec2<i32>>(max(bitcast<vec2<u32>>(v2i1), bitcast<vec2<u32>>(v2i2)));)"))
R"(let x_2 : vec2i = bitcast<vec2i>(max(bitcast<vec2u>(v2i1), bitcast<vec2u>(v2i2)));)"))
<< body;
}
@ -1003,7 +998,7 @@ TEST_F(SpvParserTest, RectifyOperandsAndResult_UMin) {
EXPECT_THAT(
body,
HasSubstr(
R"(let x_2 : vec2<i32> = bitcast<vec2<i32>>(min(bitcast<vec2<u32>>(v2i1), bitcast<vec2<u32>>(v2i2)));)"))
R"(let x_2 : vec2i = bitcast<vec2i>(min(bitcast<vec2u>(v2i1), bitcast<vec2u>(v2i2)));)"))
<< body;
}
@ -1028,7 +1023,7 @@ TEST_F(SpvParserTest, RectifyOperandsAndResult_UClamp) {
EXPECT_THAT(
body,
HasSubstr(
R"(let x_2 : vec2<i32> = bitcast<vec2<i32>>(clamp(bitcast<vec2<u32>>(v2i1), v2u2, bitcast<vec2<u32>>(v2i3)));)"))
R"(let x_2 : vec2i = bitcast<vec2i>(clamp(bitcast<vec2u>(v2i1), v2u2, bitcast<vec2u>(v2i3)));)"))
<< body;
}
@ -1054,9 +1049,9 @@ TEST_F(SpvParserTest, RectifyOperandsAndResult_FindILsb) {
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr(R"(
let x_1 : u32 = bitcast<u32>(firstTrailingBit(i1));
let x_2 : vec2<u32> = bitcast<vec2<u32>>(firstTrailingBit(v2i1));
let x_2 : vec2u = bitcast<vec2u>(firstTrailingBit(v2i1));
let x_3 : i32 = bitcast<i32>(firstTrailingBit(u1));
let x_4 : vec2<i32> = bitcast<vec2<i32>>(firstTrailingBit(v2u1));)"))
let x_4 : vec2i = bitcast<vec2i>(firstTrailingBit(v2u1));)"))
<< body;
}
@ -1100,13 +1095,13 @@ TEST_F(SpvParserTest, RectifyOperandsAndResult_FindSMsb) {
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr(R"(
let x_1 : i32 = firstLeadingBit(i1);
let x_2 : vec2<i32> = firstLeadingBit(v2i1);
let x_2 : vec2i = firstLeadingBit(v2i1);
let x_3 : u32 = bitcast<u32>(firstLeadingBit(i1));
let x_4 : vec2<u32> = bitcast<vec2<u32>>(firstLeadingBit(v2i1));
let x_4 : vec2u = bitcast<vec2u>(firstLeadingBit(v2i1));
let x_5 : i32 = firstLeadingBit(bitcast<i32>(u1));
let x_6 : vec2<i32> = firstLeadingBit(bitcast<vec2<i32>>(v2u1));
let x_6 : vec2i = firstLeadingBit(bitcast<vec2i>(v2u1));
let x_7 : u32 = bitcast<u32>(firstLeadingBit(bitcast<i32>(u1)));
let x_8 : vec2<u32> = bitcast<vec2<u32>>(firstLeadingBit(bitcast<vec2<i32>>(v2u1)));
let x_8 : vec2u = bitcast<vec2u>(firstLeadingBit(bitcast<vec2i>(v2u1)));
)")) << body;
}
@ -1150,13 +1145,13 @@ TEST_F(SpvParserTest, RectifyOperandsAndResult_FindUMsb) {
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr(R"(
let x_1 : i32 = bitcast<i32>(firstLeadingBit(bitcast<u32>(i1)));
let x_2 : vec2<i32> = bitcast<vec2<i32>>(firstLeadingBit(bitcast<vec2<u32>>(v2i1)));
let x_2 : vec2i = bitcast<vec2i>(firstLeadingBit(bitcast<vec2u>(v2i1)));
let x_3 : u32 = firstLeadingBit(bitcast<u32>(i1));
let x_4 : vec2<u32> = firstLeadingBit(bitcast<vec2<u32>>(v2i1));
let x_4 : vec2u = firstLeadingBit(bitcast<vec2u>(v2i1));
let x_5 : i32 = bitcast<i32>(firstLeadingBit(u1));
let x_6 : vec2<i32> = bitcast<vec2<i32>>(firstLeadingBit(v2u1));
let x_6 : vec2i = bitcast<vec2i>(firstLeadingBit(v2u1));
let x_7 : u32 = firstLeadingBit(u1);
let x_8 : vec2<u32> = firstLeadingBit(v2u1);
let x_8 : vec2u = firstLeadingBit(v2u1);
)")) << body;
}
@ -1219,8 +1214,8 @@ TEST_P(SpvParserTest_GlslStd450_DataUnpacking, Valid) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
EXPECT_THAT(body, HasSubstr("let x_1 : " +
std::string(param.vec_size == 2 ? "vec2<f32>" : "vec4<f32>") +
EXPECT_THAT(body,
HasSubstr("let x_1 : " + std::string(param.vec_size == 2 ? "vec2f" : "vec4f") +
+" = " + param.wgsl_func + "(u1);"))
<< body;
@ -1247,8 +1242,7 @@ TEST_F(SpvParserTest, GlslStd450_Refract_Scalar) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
const auto* expected =
R"(let x_1 : f32 = refract(vec2<f32>(f1, 0.0f), vec2<f32>(f2, 0.0f), f3).x;)";
const auto* expected = R"(let x_1 : f32 = refract(vec2f(f1, 0.0f), vec2f(f2, 0.0f), f3).x;)";
EXPECT_THAT(body, HasSubstr(expected)) << body;
}
@ -1265,7 +1259,7 @@ TEST_F(SpvParserTest, GlslStd450_Refract_Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
const auto* expected = R"(let x_1 : vec2<f32> = refract(v2f1, v2f2, f3);)";
const auto* expected = R"(let x_1 : vec2f = refract(v2f1, v2f2, f3);)";
EXPECT_THAT(body, HasSubstr(expected)) << body;
}
@ -1304,7 +1298,7 @@ TEST_F(SpvParserTest, GlslStd450_FaceForward_Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
const auto* expected = R"(let x_1 : vec2<f32> = faceForward(v2f1, v2f2, v2f3);)";
const auto* expected = R"(let x_1 : vec2f = faceForward(v2f1, v2f2, v2f3);)";
EXPECT_THAT(body, HasSubstr(expected)) << body;
}
@ -1346,9 +1340,9 @@ TEST_F(SpvParserTest, GlslStd450_Reflect_Vector) {
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
const auto* expected = R"(
let x_98 : vec2<f32> = (v2f1 + v2f1);
let x_99 : vec2<f32> = (v2f2 + v2f2);
let x_1 : vec2<f32> = reflect(x_98, x_99);
let x_98 : vec2f = (v2f1 + v2f1);
let x_99 : vec2f = (v2f2 + v2f2);
let x_1 : vec2f = reflect(x_98, x_99);
)";
EXPECT_THAT(body, HasSubstr(expected)) << body;
@ -1384,7 +1378,7 @@ TEST_F(SpvParserTest, GlslStd450_Ldexp_Vector_Floatvec_Uintvec) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
const auto body = test::ToString(p->program(), ast_body);
const auto* expected = "let x_1 : vec2<f32> = ldexp(v2f1, vec2<i32>(v2u1));";
const auto* expected = "let x_1 : vec2f = ldexp(v2f1, vec2i(v2u1));";
EXPECT_THAT(body, HasSubstr(expected)) << body;
}
@ -1426,8 +1420,8 @@ TEST_F(SpvParserTest, GlslStd450_MatrixInverse_mat2x2) {
std::string expected =
"let s = (1.0f / determinant(m2x2f1));\n"
"let x_1 : mat2x2<f32> = mat2x2<f32>(vec2<f32>((s * m2x2f1[1u][1u]), (-(s) * "
"m2x2f1[0u][1u])), vec2<f32>((-(s) * m2x2f1[1u][0u]), (s * m2x2f1[0u][0u])));";
"let x_1 : mat2x2<f32> = mat2x2<f32>(vec2f((s * m2x2f1[1u][1u]), (-(s) * "
"m2x2f1[0u][1u])), vec2f((-(s) * m2x2f1[1u][0u]), (s * m2x2f1[0u][0u])));";
EXPECT_THAT(body, HasSubstr(expected)) << body;
}
@ -1447,13 +1441,13 @@ TEST_F(SpvParserTest, GlslStd450_MatrixInverse_mat3x3) {
std::string expected =
"let s = (1.0f / determinant(m3x3f1));\n"
"let x_1 : mat3x3<f32> = (s * mat3x3<f32>(vec3<f32>(((m3x3f1[1u][1u] * m3x3f1[2u][2u]) - "
"let x_1 : mat3x3<f32> = (s * mat3x3<f32>(vec3f(((m3x3f1[1u][1u] * m3x3f1[2u][2u]) - "
"(m3x3f1[1u][2u] * m3x3f1[2u][1u])), ((m3x3f1[0u][2u] * m3x3f1[2u][1u]) - (m3x3f1[0u][1u] "
"* m3x3f1[2u][2u])), ((m3x3f1[0u][1u] * m3x3f1[1u][2u]) - (m3x3f1[0u][2u] * "
"m3x3f1[1u][1u]))), vec3<f32>(((m3x3f1[1u][2u] * m3x3f1[2u][0u]) - (m3x3f1[1u][0u] * "
"m3x3f1[1u][1u]))), vec3f(((m3x3f1[1u][2u] * m3x3f1[2u][0u]) - (m3x3f1[1u][0u] * "
"m3x3f1[2u][2u])), ((m3x3f1[0u][0u] * m3x3f1[2u][2u]) - (m3x3f1[0u][2u] * "
"m3x3f1[2u][0u])), ((m3x3f1[0u][2u] * m3x3f1[1u][0u]) - (m3x3f1[0u][0u] * "
"m3x3f1[1u][2u]))), vec3<f32>(((m3x3f1[1u][0u] * m3x3f1[2u][1u]) - (m3x3f1[1u][1u] * "
"m3x3f1[1u][2u]))), vec3f(((m3x3f1[1u][0u] * m3x3f1[2u][1u]) - (m3x3f1[1u][1u] * "
"m3x3f1[2u][0u])), ((m3x3f1[0u][1u] * m3x3f1[2u][0u]) - (m3x3f1[0u][0u] * "
"m3x3f1[2u][1u])), ((m3x3f1[0u][0u] * m3x3f1[1u][1u]) - (m3x3f1[0u][1u] * "
"m3x3f1[1u][0u])))));";
@ -1476,7 +1470,7 @@ TEST_F(SpvParserTest, GlslStd450_MatrixInverse_mat4x4) {
std::string expected =
"let s = (1.0f / determinant(m4x4f1));\n"
"let x_1 : mat4x4<f32> = (s * mat4x4<f32>(vec4<f32>((((m4x4f1[1u][1u] * ((m4x4f1[2u][2u] * "
"let x_1 : mat4x4<f32> = (s * mat4x4<f32>(vec4f((((m4x4f1[1u][1u] * ((m4x4f1[2u][2u] * "
"m4x4f1[3u][3u]) - (m4x4f1[2u][3u] * m4x4f1[3u][2u]))) - (m4x4f1[1u][2u] * "
"((m4x4f1[2u][1u] * m4x4f1[3u][3u]) - (m4x4f1[2u][3u] * m4x4f1[3u][1u])))) + "
"(m4x4f1[1u][3u] * ((m4x4f1[2u][1u] * m4x4f1[3u][2u]) - (m4x4f1[2u][2u] * "
@ -1491,7 +1485,7 @@ TEST_F(SpvParserTest, GlslStd450_MatrixInverse_mat4x4) {
"((m4x4f1[1u][2u] * m4x4f1[2u][3u]) - (m4x4f1[1u][3u] * m4x4f1[2u][2u]))) + "
"(m4x4f1[0u][2u] * ((m4x4f1[1u][1u] * m4x4f1[2u][3u]) - (m4x4f1[1u][3u] * "
"m4x4f1[2u][1u])))) - (m4x4f1[0u][3u] * ((m4x4f1[1u][1u] * m4x4f1[2u][2u]) - "
"(m4x4f1[1u][2u] * m4x4f1[2u][1u]))))), vec4<f32>((((-(m4x4f1[1u][0u]) * ((m4x4f1[2u][2u] "
"(m4x4f1[1u][2u] * m4x4f1[2u][1u]))))), vec4f((((-(m4x4f1[1u][0u]) * ((m4x4f1[2u][2u] "
"* m4x4f1[3u][3u]) - (m4x4f1[2u][3u] * m4x4f1[3u][2u]))) + (m4x4f1[1u][2u] * "
"((m4x4f1[2u][0u] * m4x4f1[3u][3u]) - (m4x4f1[2u][3u] * m4x4f1[3u][0u])))) - "
"(m4x4f1[1u][3u] * ((m4x4f1[2u][0u] * m4x4f1[3u][2u]) - (m4x4f1[2u][2u] * "
@ -1506,7 +1500,7 @@ TEST_F(SpvParserTest, GlslStd450_MatrixInverse_mat4x4) {
"((m4x4f1[1u][2u] * m4x4f1[2u][3u]) - (m4x4f1[1u][3u] * m4x4f1[2u][2u]))) - "
"(m4x4f1[0u][2u] * ((m4x4f1[1u][0u] * m4x4f1[2u][3u]) - (m4x4f1[1u][3u] * "
"m4x4f1[2u][0u])))) + (m4x4f1[0u][3u] * ((m4x4f1[1u][0u] * m4x4f1[2u][2u]) - "
"(m4x4f1[1u][2u] * m4x4f1[2u][0u]))))), vec4<f32>((((m4x4f1[1u][0u] * ((m4x4f1[2u][1u] * "
"(m4x4f1[1u][2u] * m4x4f1[2u][0u]))))), vec4f((((m4x4f1[1u][0u] * ((m4x4f1[2u][1u] * "
"m4x4f1[3u][3u]) - (m4x4f1[2u][3u] * m4x4f1[3u][1u]))) - (m4x4f1[1u][1u] * "
"((m4x4f1[2u][0u] * m4x4f1[3u][3u]) - (m4x4f1[2u][3u] * m4x4f1[3u][0u])))) + "
"(m4x4f1[1u][3u] * ((m4x4f1[2u][0u] * m4x4f1[3u][1u]) - (m4x4f1[2u][1u] * "
@ -1521,7 +1515,7 @@ TEST_F(SpvParserTest, GlslStd450_MatrixInverse_mat4x4) {
"((m4x4f1[1u][1u] * m4x4f1[2u][3u]) - (m4x4f1[1u][3u] * m4x4f1[2u][1u]))) + "
"(m4x4f1[0u][1u] * ((m4x4f1[1u][0u] * m4x4f1[2u][3u]) - (m4x4f1[1u][3u] * "
"m4x4f1[2u][0u])))) - (m4x4f1[0u][3u] * ((m4x4f1[1u][0u] * m4x4f1[2u][1u]) - "
"(m4x4f1[1u][1u] * m4x4f1[2u][0u]))))), vec4<f32>((((-(m4x4f1[1u][0u]) * ((m4x4f1[2u][1u] "
"(m4x4f1[1u][1u] * m4x4f1[2u][0u]))))), vec4f((((-(m4x4f1[1u][0u]) * ((m4x4f1[2u][1u] "
"* m4x4f1[3u][2u]) - (m4x4f1[2u][2u] * m4x4f1[3u][1u]))) + (m4x4f1[1u][1u] * "
"((m4x4f1[2u][0u] * m4x4f1[3u][2u]) - (m4x4f1[2u][2u] * m4x4f1[3u][0u])))) - "
"(m4x4f1[1u][2u] * ((m4x4f1[2u][0u] * m4x4f1[3u][1u]) - (m4x4f1[2u][1u] * "
@ -1557,11 +1551,11 @@ TEST_F(SpvParserTest, GlslStd450_MatrixInverse_MultipleInScope) {
std::string expected =
"let s = (1.0f / determinant(m2x2f1));\n"
"let x_1 : mat2x2<f32> = mat2x2<f32>(vec2<f32>((s * m2x2f1[1u][1u]), (-(s) * "
"m2x2f1[0u][1u])), vec2<f32>((-(s) * m2x2f1[1u][0u]), (s * m2x2f1[0u][0u])));\n"
"let x_1 : mat2x2<f32> = mat2x2<f32>(vec2f((s * m2x2f1[1u][1u]), (-(s) * "
"m2x2f1[0u][1u])), vec2f((-(s) * m2x2f1[1u][0u]), (s * m2x2f1[0u][0u])));\n"
"let s_1 = (1.0f / determinant(m2x2f1));\n"
"let x_2 : mat2x2<f32> = mat2x2<f32>(vec2<f32>((s_1 * m2x2f1[1u][1u]), (-(s_1) * "
"m2x2f1[0u][1u])), vec2<f32>((-(s_1) * m2x2f1[1u][0u]), (s_1 * m2x2f1[0u][0u])));";
"let x_2 : mat2x2<f32> = mat2x2<f32>(vec2f((s_1 * m2x2f1[1u][1u]), (-(s_1) * "
"m2x2f1[0u][1u])), vec2f((-(s_1) * m2x2f1[1u][0u]), (s_1 * m2x2f1[0u][0u])));";
EXPECT_THAT(body, HasSubstr(expected)) << body;
}

54
src/tint/reader/spirv/function_logical_test.cc
View File

@ -77,7 +77,7 @@ std::string AstFor(std::string assembly) {
return "vec2<bool>(false, true)";
}
if (assembly == "v2uint_10_20") {
return "vec2<u32>(10u, 20u)";
return "vec2u(10u, 20u)";
}
if (assembly == "cast_uint_10") {
return "bitcast<i32>(10u)";
@ -86,13 +86,13 @@ std::string AstFor(std::string assembly) {
return "bitcast<i32>(20u)";
}
if (assembly == "cast_v2uint_10_20") {
return "bitcast<vec2<i32>>(vec2<u32>(10u, 20u))";
return "bitcast<vec2i>(vec2u(10u, 20u))";
}
if (assembly == "v2uint_20_10") {
return "vec2<u32>(20u, 10u)";
return "vec2u(20u, 10u)";
}
if (assembly == "cast_v2uint_20_10") {
return "bitcast<vec2<i32>>(vec2<u32>(20u, 10u))";
return "bitcast<vec2i>(vec2u(20u, 10u))";
}
if (assembly == "cast_int_30") {
return "bitcast<u32>(30i)";
@ -101,22 +101,22 @@ std::string AstFor(std::string assembly) {
return "bitcast<u32>(40i)";
}
if (assembly == "v2int_30_40") {
return "vec2<i32>(30i, 40i)";
return "vec2i(30i, 40i)";
}
if (assembly == "cast_v2int_30_40") {
return "bitcast<vec2<u32>>(vec2<i32>(30i, 40i))";
return "bitcast<vec2u>(vec2i(30i, 40i))";
}
if (assembly == "v2int_40_30") {
return "vec2<i32>(40i, 30i)";
return "vec2i(40i, 30i)";
}
if (assembly == "cast_v2int_40_30") {
return "bitcast<vec2<u32>>(vec2<i32>(40i, 30i))";
return "bitcast<vec2u>(vec2i(40i, 30i))";
}
if (assembly == "v2float_50_60") {
return "vec2<f32>(50.0f, 60.0f)";
return "vec2f(50.0f, 60.0f)";
}
if (assembly == "v2float_60_50") {
return "vec2<f32>(60.0f, 50.0f)";
return "vec2f(60.0f, 50.0f)";
}
return "bad case";
}
@ -534,8 +534,7 @@ TEST_F(SpvFUnordTest, FUnordEqual_Vector) {
auto ast_body = fe.ast_body();
EXPECT_THAT(
test::ToString(p->program(), ast_body),
HasSubstr(
"let x_1 : vec2<bool> = !((vec2<f32>(50.0f, 60.0f) != vec2<f32>(60.0f, 50.0f)));"));
HasSubstr("let x_1 : vec2<bool> = !((vec2f(50.0f, 60.0f) != vec2f(60.0f, 50.0f)));"));
}
TEST_F(SpvFUnordTest, FUnordNotEqual_Scalar) {
@ -570,8 +569,7 @@ TEST_F(SpvFUnordTest, FUnordNotEqual_Vector) {
auto ast_body = fe.ast_body();
EXPECT_THAT(
test::ToString(p->program(), ast_body),
HasSubstr(
"let x_1 : vec2<bool> = !((vec2<f32>(50.0f, 60.0f) == vec2<f32>(60.0f, 50.0f)));"));
HasSubstr("let x_1 : vec2<bool> = !((vec2f(50.0f, 60.0f) == vec2f(60.0f, 50.0f)));"));
}
TEST_F(SpvFUnordTest, FUnordLessThan_Scalar) {
@ -606,8 +604,7 @@ TEST_F(SpvFUnordTest, FUnordLessThan_Vector) {
auto ast_body = fe.ast_body();
EXPECT_THAT(
test::ToString(p->program(), ast_body),
HasSubstr(
"let x_1 : vec2<bool> = !((vec2<f32>(50.0f, 60.0f) >= vec2<f32>(60.0f, 50.0f)));"));
HasSubstr("let x_1 : vec2<bool> = !((vec2f(50.0f, 60.0f) >= vec2f(60.0f, 50.0f)));"));
}
TEST_F(SpvFUnordTest, FUnordLessThanEqual_Scalar) {
@ -642,8 +639,7 @@ TEST_F(SpvFUnordTest, FUnordLessThanEqual_Vector) {
auto ast_body = fe.ast_body();
EXPECT_THAT(
test::ToString(p->program(), ast_body),
HasSubstr(
"let x_1 : vec2<bool> = !((vec2<f32>(50.0f, 60.0f) > vec2<f32>(60.0f, 50.0f)));"));
HasSubstr("let x_1 : vec2<bool> = !((vec2f(50.0f, 60.0f) > vec2f(60.0f, 50.0f)));"));
}
TEST_F(SpvFUnordTest, FUnordGreaterThan_Scalar) {
@ -678,8 +674,7 @@ TEST_F(SpvFUnordTest, FUnordGreaterThan_Vector) {
auto ast_body = fe.ast_body();
EXPECT_THAT(
test::ToString(p->program(), ast_body),
HasSubstr(
"let x_1 : vec2<bool> = !((vec2<f32>(50.0f, 60.0f) <= vec2<f32>(60.0f, 50.0f)));"));
HasSubstr("let x_1 : vec2<bool> = !((vec2f(50.0f, 60.0f) <= vec2f(60.0f, 50.0f)));"));
}
TEST_F(SpvFUnordTest, FUnordGreaterThanEqual_Scalar) {
@ -714,8 +709,7 @@ TEST_F(SpvFUnordTest, FUnordGreaterThanEqual_Vector) {
auto ast_body = fe.ast_body();
EXPECT_THAT(
test::ToString(p->program(), ast_body),
HasSubstr(
"let x_1 : vec2<bool> = !((vec2<f32>(50.0f, 60.0f) < vec2<f32>(60.0f, 50.0f)));"));
HasSubstr("let x_1 : vec2<bool> = !((vec2f(50.0f, 60.0f) < vec2f(60.0f, 50.0f)));"));
}
using SpvLogicalTest = SpvParserTestBase<::testing::Test>;
@ -787,9 +781,9 @@ TEST_F(SpvLogicalTest, Select_BoolCond_VectorParams) {
auto fe = p->function_emitter(100);
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body), HasSubstr("let x_1 : vec2<u32> = select("
"vec2<u32>(20u, 10u), "
"vec2<u32>(10u, 20u), "
EXPECT_THAT(test::ToString(p->program(), ast_body), HasSubstr("let x_1 : vec2u = select("
"vec2u(20u, 10u), "
"vec2u(10u, 20u), "
"true);"));
// Fails validation prior to SPIR-V 1.4: If the value operands are vectors,
@ -812,9 +806,9 @@ TEST_F(SpvLogicalTest, Select_VecBoolCond_VectorParams) {
auto fe = p->function_emitter(100);
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body), HasSubstr("let x_1 : vec2<u32> = select("
"vec2<u32>(20u, 10u), "
"vec2<u32>(10u, 20u), "
EXPECT_THAT(test::ToString(p->program(), ast_body), HasSubstr("let x_1 : vec2u = select("
"vec2u(20u, 10u), "
"vec2u(10u, 20u), "
"vec2<bool>(true, false));"));
}
@ -883,7 +877,7 @@ TEST_F(SpvLogicalTest, IsNan_Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<bool> = isNan(vec2<f32>(50.0f, 60.0f));"));
HasSubstr("let x_1 : vec2<bool> = isNan(vec2f(50.0f, 60.0f));"));
}
TEST_F(SpvLogicalTest, IsInf_Scalar) {
@ -917,7 +911,7 @@ TEST_F(SpvLogicalTest, IsInf_Vector) {
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("let x_1 : vec2<bool> = isInf(vec2<f32>(50.0f, 60.0f));"));
HasSubstr("let x_1 : vec2<bool> = isInf(vec2f(50.0f, 60.0f));"));
}
// TODO(dneto): Kernel-guarded instructions.

6
src/tint/reader/spirv/function_memory_test.cc
View File

@ -498,7 +498,7 @@ TEST_F(SpvParserMemoryTest, EmitStatement_AccessChain_Matrix) {
EXPECT_TRUE(fe.EmitBody());
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("myvar[2u] = vec4<f32>(42.0f, 42.0f, 42.0f, 42.0f);"));
HasSubstr("myvar[2u] = vec4f(42.0f, 42.0f, 42.0f, 42.0f);"));
}
TEST_F(SpvParserMemoryTest, EmitStatement_AccessChain_Array) {
@ -530,7 +530,7 @@ TEST_F(SpvParserMemoryTest, EmitStatement_AccessChain_Array) {
EXPECT_TRUE(fe.EmitBody());
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("myvar[2u] = vec4<f32>(42.0f, 42.0f, 42.0f, 42.0f);"));
HasSubstr("myvar[2u] = vec4f(42.0f, 42.0f, 42.0f, 42.0f);"));
}
TEST_F(SpvParserMemoryTest, EmitStatement_AccessChain_Struct) {
@ -801,7 +801,7 @@ TEST_F(SpvParserMemoryTest, EmitStatement_AccessChain_DereferenceBase) {
auto p = parser(test::Assemble(assembly));
ASSERT_TRUE(p->BuildAndParseInternalModule());
const auto got = test::ToString(p->program());
const std::string expected = R"(fn x_200(x_1 : ptr<private, vec2<u32>>) {
const std::string expected = R"(fn x_200(x_1 : ptr<private, vec2u>) {
let x_3 : u32 = (*(x_1)).x;
return;
}

13
src/tint/reader/spirv/function_misc_test.cc
View File

@ -104,9 +104,9 @@ TEST_F(SpvParserTestMiscInstruction, OpUndef_BeforeFunction_Vector) {
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr(R"(let x_14 : vec2<bool> = vec2<bool>();
let x_11 : vec2<u32> = vec2<u32>();
let x_12 : vec2<i32> = vec2<i32>();
let x_13 : vec2<f32> = vec2<f32>();
let x_11 : vec2u = vec2u();
let x_12 : vec2i = vec2i();
let x_13 : vec2f = vec2f();
)"));
}
@ -158,10 +158,9 @@ TEST_F(SpvParserTestMiscInstruction, OpUndef_InFunction_Vector) {
auto fe = p->function_emitter(100);
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr(R"(let x_11 : vec2<u32> = vec2<u32>();
let x_12 : vec2<i32> = vec2<i32>();
let x_13 : vec2<f32> = vec2<f32>();
EXPECT_THAT(test::ToString(p->program(), ast_body), HasSubstr(R"(let x_11 : vec2u = vec2u();
let x_12 : vec2i = vec2i();
let x_13 : vec2f = vec2f();
)"));
}

20
src/tint/reader/spirv/function_var_test.cc
View File

@ -235,7 +235,7 @@ TEST_F(SpvParserFunctionVarTest, EmitFunctionVariables_VectorInitializer) {
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("var x_200 : vec2<f32> = vec2<f32>(1.5f, 2.0f);"));
HasSubstr("var x_200 : vec2f = vec2f(1.5f, 2.0f);"));
}
TEST_F(SpvParserFunctionVarTest, EmitFunctionVariables_MatrixInitializer) {
@ -262,9 +262,9 @@ TEST_F(SpvParserFunctionVarTest, EmitFunctionVariables_MatrixInitializer) {
auto ast_body = fe.ast_body();
EXPECT_THAT(test::ToString(p->program(), ast_body),
HasSubstr("var x_200 : mat3x2<f32> = mat3x2<f32>("
"vec2<f32>(1.5f, 2.0f), "
"vec2<f32>(2.0f, 3.0f), "
"vec2<f32>(3.0f, 4.0f));"));
"vec2f(1.5f, 2.0f), "
"vec2f(2.0f, 3.0f), "
"vec2f(3.0f, 4.0f));"));
}
TEST_F(SpvParserFunctionVarTest, EmitFunctionVariables_ArrayInitializer) {
@ -1731,14 +1731,14 @@ TEST_F(SpvParserFunctionVarTest, EmitStatement_Hoist_CompositeInsert) {
auto fe = p->function_emitter(100);
EXPECT_TRUE(fe.EmitBody()) << p->error();
const auto* expected = R"(var x_200 : vec2<i32>;
const auto* expected = R"(var x_200 : vec2i;
if (true) {
x_200 = vec2<i32>();
x_200 = vec2i();
x_200.x = 0i;
} else {
return;
}
let x_201 : vec2<i32> = x_200;
let x_201 : vec2i = x_200;
return;
)";
auto ast_body = fe.ast_body();
@ -1774,14 +1774,14 @@ TEST_F(SpvParserFunctionVarTest, EmitStatement_Hoist_VectorInsertDynamic) {
auto ast_body = fe.ast_body();
const auto got = test::ToString(p->program(), ast_body);
const auto* expected = R"(var x_200 : vec2<i32>;
const auto* expected = R"(var x_200 : vec2i;
if (true) {
x_200 = vec2<i32>();
x_200 = vec2i();
x_200[1i] = 3i;
} else {
return;
}
let x_201 : vec2<i32> = x_200;
let x_201 : vec2i = x_200;
return;
)";
EXPECT_EQ(got, expected) << got;

4
src/tint/reader/spirv/parser_impl.h
View File

@ -586,7 +586,7 @@ class ParserImpl : Reader {
uint32_t position_member_index = 0;
/// The member index for the PointSize builtin within the struct.
uint32_t pointsize_member_index = 0;
/// The ID for the member type, which should map to vec4<f32>.
/// The ID for the member type, which should map to vec4f.
uint32_t position_member_type_id = 0;
/// The ID of the type of a pointer to the struct in the Output storage
/// class class.
@ -863,7 +863,7 @@ class ParserImpl : Reader {
// Bookkeeping for the gl_Position builtin.
// In Vulkan SPIR-V, it's the 0 member of the gl_PerVertex structure.
// But in WGSL we make a module-scope variable:
// [[position]] var<in> gl_Position : vec4<f32>;
// [[position]] var<in> gl_Position : vec4f;
// The builtin variable was detected if and only if the struct_id is non-zero.
BuiltInPositionInfo builtin_position_;

2
src/tint/reader/spirv/parser_impl_function_decl_test.cc
View File

@ -122,7 +122,7 @@ OpFunctionEnd)";
EXPECT_THAT(program_ast, HasSubstr(R"(
struct main_out {
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
}
)")) << program_ast;

163
src/tint/reader/spirv/parser_impl_handle_test.cc
View File

@ -1514,7 +1514,7 @@ INSTANTIATE_TEST_SUITE_P(
@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
"textureGather(1i, x_20, x_10, coords12, "
"vec2<i32>(u_offsets2d))"},
"vec2i(u_offsets2d))"},
// OpImageGather 2D Array
ImageAccessCase{"%float 2D 0 1 0 1 Unknown",
"%result = OpImageGather "
@ -1543,7 +1543,7 @@ INSTANTIATE_TEST_SUITE_P(
@group(2) @binding(1) var x_20 : texture_2d_array<f32>;)",
"textureGather(1i, x_20, x_10, coords123.xy, "
"i32(round(coords123.z)), "
"vec2<i32>(u_offsets2d))"},
"vec2i(u_offsets2d))"},
// OpImageGather Cube
ImageAccessCase{"%float Cube 0 0 0 1 Unknown",
"%result = OpImageGather "
@ -1586,7 +1586,7 @@ INSTANTIATE_TEST_SUITE_P(
@group(2) @binding(1) var x_20 : texture_depth_2d;)",
"textureGather(x_20, x_10, coords12, "
"vec2<i32>(u_offsets2d))"},
"vec2i(u_offsets2d))"},
// OpImageGather 2DDepth Array
ImageAccessCase{"%float 2D 1 1 0 1 Unknown",
"%result = OpImageGather "
@ -1615,7 +1615,7 @@ INSTANTIATE_TEST_SUITE_P(
@group(2) @binding(1) var x_20 : texture_depth_2d_array;)",
"textureGather(x_20, x_10, coords123.xy, "
"i32(round(coords123.z)), "
"vec2<i32>(u_offsets2d))"},
"vec2i(u_offsets2d))"},
// OpImageGather DepthCube
ImageAccessCase{"%float Cube 1 0 0 1 Unknown",
"%result = OpImageGather "
@ -1664,7 +1664,7 @@ INSTANTIATE_TEST_SUITE_P(
@group(2) @binding(1) var x_20 : texture_depth_2d;)",
"textureGatherCompare(x_20, x_10, coords12, 0.20000000298023223877f, "
"vec2<i32>(u_offsets2d))"},
"vec2i(u_offsets2d))"},
// OpImageDrefGather 2DDepth Array
ImageAccessCase{"%float 2D 1 1 0 1 Unknown",
"%result = OpImageDrefGather "
@ -1693,7 +1693,7 @@ INSTANTIATE_TEST_SUITE_P(
@group(2) @binding(1) var x_20 : texture_depth_2d_array;)",
"textureGatherCompare(x_20, x_10, coords123.xy, "
"i32(round(coords123.z)), 0.20000000298023223877f, "
"vec2<i32>(u_offsets2d))"},
"vec2i(u_offsets2d))"},
// OpImageDrefGather DepthCube
ImageAccessCase{"%float Cube 1 0 0 1 Unknown",
"%result = OpImageDrefGather "
@ -1792,7 +1792,7 @@ INSTANTIATE_TEST_SUITE_P(
R"(@group(0) @binding(0) var x_10 : sampler;
@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(textureSampleBias(x_20, x_10, coords12, 7.0f, vec2<i32>(u_offsets2d))"},
R"(textureSampleBias(x_20, x_10, coords12, 7.0f, vec2i(u_offsets2d))"},
// OpImageSampleImplicitLod arrayed with Bias
ImageAccessCase{
"%float 2D 0 1 0 1 Unknown",
@ -1827,7 +1827,7 @@ INSTANTIATE_TEST_SUITE_P(
@group(0) @binding(1) var x_30 : sampler_comparison;
)",
R"(
let x_200 : vec4<f32> = vec4<f32>(textureSample(x_20, x_10, coords12), 0.0f, 0.0f, 0.0f);
let x_200 : vec4f = vec4f(textureSample(x_20, x_10, coords12), 0.0f, 0.0f, 0.0f);
let x_210 : f32 = textureSampleCompare(x_20, x_30, coords12, 0.20000000298023223877f);
)"}));
@ -1981,7 +1981,7 @@ INSTANTIATE_TEST_SUITE_P(
R"(@group(0) @binding(0) var x_10 : sampler;
@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(textureSampleLevel(x_20, x_10, coords12, 0.0f, vec2<i32>(u_offsets2d))"},
R"(textureSampleLevel(x_20, x_10, coords12, 0.0f, vec2i(u_offsets2d))"},
// OpImageSampleExplicitLod arrayed - using Lod and ConstOffset
ImageAccessCase{
@ -2037,7 +2037,7 @@ INSTANTIATE_TEST_SUITE_P(
R"(@group(0) @binding(0) var x_10 : sampler;
@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(textureSampleGrad(x_20, x_10, coords12, vf12, vf21, vec2<i32>(u_offsets2d))"},
R"(textureSampleGrad(x_20, x_10, coords12, vf12, vf21, vec2i(u_offsets2d))"},
// OpImageSampleExplicitLod arrayed - using Grad and ConstOffset
ImageAccessCase{
@ -2060,7 +2060,7 @@ INSTANTIATE_TEST_SUITE_P(
R"(@group(0) @binding(0) var x_10 : sampler;
@group(2) @binding(1) var x_20 : texture_2d_array<f32>;)",
R"(textureSampleGrad(x_20, x_10, coords123.xy, i32(round(coords123.z)), vf12, vf21, vec2<i32>(u_offsets2d)))"}));
R"(textureSampleGrad(x_20, x_10, coords123.xy, i32(round(coords123.z)), vf12, vf21, vec2i(u_offsets2d)))"}));
// Test crbug.com/378:
// In WGSL, sampling from depth texture with explicit level of detail
@ -2088,7 +2088,7 @@ INSTANTIATE_TEST_SUITE_P(
@group(2) @binding(1) var x_20 : texture_depth_2d;
)",
R"(vec4<f32>(textureSampleLevel(x_20, x_10, vf12, i32(f1)), 0.0f, 0.0f, 0.0f))"}}));
R"(vec4f(textureSampleLevel(x_20, x_10, vf12, i32(f1)), 0.0f, 0.0f, 0.0f))"}}));
/////
// Projection sampling
@ -2194,7 +2194,7 @@ INSTANTIATE_TEST_SUITE_P(
R"(@group(0) @binding(0) var x_10 : sampler;
@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(textureSampleBias(x_20, x_10, (coords123.xy / coords123.z), 7.0f, vec2<i32>(u_offsets2d)))"}));
R"(textureSampleBias(x_20, x_10, (coords123.xy / coords123.z), 7.0f, vec2i(u_offsets2d)))"}));
INSTANTIATE_TEST_SUITE_P(
ImageSampleProjExplicitLod_Lod,
@ -2259,9 +2259,9 @@ INSTANTIATE_TEST_SUITE_P(
@group(2) @binding(1) var x_20 : texture_depth_2d;
)",
// Sampling the depth texture yields an f32, but the
// SPIR-V operation yiedls vec4<f32>, so fill out the
// SPIR-V operation yiedls vec4f, so fill out the
// remaining components with 0.
R"(vec4<f32>(textureSample(x_20, x_10, (coords123.xy / coords123.z)), 0.0f, 0.0f, 0.0f))"}));
R"(vec4f(textureSample(x_20, x_10, (coords123.xy / coords123.z)), 0.0f, 0.0f, 0.0f))"}));
INSTANTIATE_TEST_SUITE_P(
ImageSampleProjDrefImplicitLod,
@ -2318,7 +2318,7 @@ INSTANTIATE_TEST_SUITE_P(
@group(2) @binding(1) var x_20 : texture_depth_2d;
)",
R"(textureSampleCompareLevel(x_20, x_10, (coords123.xy / coords123.z), 0.20000000298023223877f, 0.0f, vec2<i32>(3i, 4i)))"}));
R"(textureSampleCompareLevel(x_20, x_10, (coords123.xy / coords123.z), 0.20000000298023223877f, 0.0f, vec2i(3i, 4i)))"}));
/////
// End projection sampling
@ -2411,15 +2411,15 @@ INSTANTIATE_TEST_SUITE_P(
// Source 1 component
{"%float 2D 0 0 0 2 R32f", "OpImageWrite %im %vi12 %f1",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32float, write>;)",
"textureStore(x_20, vi12, vec4<f32>(f1, 0.0f, 0.0f, 0.0f));"},
"textureStore(x_20, vi12, vec4f(f1, 0.0f, 0.0f, 0.0f));"},
// Source 2 component, dest 1 component
{"%float 2D 0 0 0 2 R32f", "OpImageWrite %im %vi12 %vf12",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32float, write>;)",
"textureStore(x_20, vi12, vec4<f32>(vf12, 0.0f, 0.0f));"},
"textureStore(x_20, vi12, vec4f(vf12, 0.0f, 0.0f));"},
// Source 3 component, dest 1 component
{"%float 2D 0 0 0 2 R32f", "OpImageWrite %im %vi12 %vf123",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32float, write>;)",
"textureStore(x_20, vi12, vec4<f32>(vf123, 0.0f));"},
"textureStore(x_20, vi12, vec4f(vf123, 0.0f));"},
// Source 4 component, dest 1 component
{"%float 2D 0 0 0 2 R32f", "OpImageWrite %im %vi12 %vf1234",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32float, write>;)",
@ -2427,11 +2427,11 @@ INSTANTIATE_TEST_SUITE_P(
// Source 2 component, dest 2 component
{"%float 2D 0 0 0 2 Rg32f", "OpImageWrite %im %vi12 %vf12",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<rg32float, write>;)",
"textureStore(x_20, vi12, vec4<f32>(vf12, 0.0f, 0.0f));"},
"textureStore(x_20, vi12, vec4f(vf12, 0.0f, 0.0f));"},
// Source 3 component, dest 2 component
{"%float 2D 0 0 0 2 Rg32f", "OpImageWrite %im %vi12 %vf123",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<rg32float, write>;)",
"textureStore(x_20, vi12, vec4<f32>(vf123, 0.0f));"},
"textureStore(x_20, vi12, vec4f(vf123, 0.0f));"},
// Source 4 component, dest 2 component
{"%float 2D 0 0 0 2 Rg32f", "OpImageWrite %im %vi12 %vf1234",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<rg32float, write>;)",
@ -2451,15 +2451,15 @@ INSTANTIATE_TEST_SUITE_P(
// Source 1 component
{"%uint 2D 0 0 0 2 R32ui", "OpImageWrite %im %vi12 %u1",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32uint, write>;)",
"textureStore(x_20, vi12, vec4<u32>(u1, 0u, 0u, 0u));"},
"textureStore(x_20, vi12, vec4u(u1, 0u, 0u, 0u));"},
// Source 2 component, dest 1 component
{"%uint 2D 0 0 0 2 R32ui", "OpImageWrite %im %vi12 %vu12",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32uint, write>;)",
"textureStore(x_20, vi12, vec4<u32>(vu12, 0u, 0u));"},
"textureStore(x_20, vi12, vec4u(vu12, 0u, 0u));"},
// Source 3 component, dest 1 component
{"%uint 2D 0 0 0 2 R32ui", "OpImageWrite %im %vi12 %vu123",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32uint, write>;)",
"textureStore(x_20, vi12, vec4<u32>(vu123, 0u));"},
"textureStore(x_20, vi12, vec4u(vu123, 0u));"},
// Source 4 component, dest 1 component
{"%uint 2D 0 0 0 2 R32ui", "OpImageWrite %im %vi12 %vu1234",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32uint, write>;)",
@ -2467,11 +2467,11 @@ INSTANTIATE_TEST_SUITE_P(
// Source 2 component, dest 2 component
{"%uint 2D 0 0 0 2 Rg32ui", "OpImageWrite %im %vi12 %vu12",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<rg32uint, write>;)",
"textureStore(x_20, vi12, vec4<u32>(vu12, 0u, 0u));"},
"textureStore(x_20, vi12, vec4u(vu12, 0u, 0u));"},
// Source 3 component, dest 2 component
{"%uint 2D 0 0 0 2 Rg32ui", "OpImageWrite %im %vi12 %vu123",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<rg32uint, write>;)",
"textureStore(x_20, vi12, vec4<u32>(vu123, 0u));"},
"textureStore(x_20, vi12, vec4u(vu123, 0u));"},
// Source 4 component, dest 2 component
{"%uint 2D 0 0 0 2 Rg32ui", "OpImageWrite %im %vi12 %vu1234",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<rg32uint, write>;)",
@ -2491,15 +2491,15 @@ INSTANTIATE_TEST_SUITE_P(
// Source 1 component
{"%int 2D 0 0 0 2 R32i", "OpImageWrite %im %vi12 %i1",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32sint, write>;)",
"textureStore(x_20, vi12, vec4<i32>(i1, 0i, 0i, 0i));"},
"textureStore(x_20, vi12, vec4i(i1, 0i, 0i, 0i));"},
// Source 2 component, dest 1 component
{"%int 2D 0 0 0 2 R32i", "OpImageWrite %im %vi12 %vi12",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32sint, write>;)",
"textureStore(x_20, vi12, vec4<i32>(vi12, 0i, 0i));"},
"textureStore(x_20, vi12, vec4i(vi12, 0i, 0i));"},
// Source 3 component, dest 1 component
{"%int 2D 0 0 0 2 R32i", "OpImageWrite %im %vi12 %vi123",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32sint, write>;)",
"textureStore(x_20, vi12, vec4<i32>(vi123, 0i));"},
"textureStore(x_20, vi12, vec4i(vi123, 0i));"},
// Source 4 component, dest 1 component
{"%int 2D 0 0 0 2 R32i", "OpImageWrite %im %vi12 %vi1234",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32sint, write>;)",
@ -2507,11 +2507,11 @@ INSTANTIATE_TEST_SUITE_P(
// Source 2 component, dest 2 component
{"%int 2D 0 0 0 2 Rg32i", "OpImageWrite %im %vi12 %vi12",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<rg32sint, write>;)",
"textureStore(x_20, vi12, vec4<i32>(vi12, 0i, 0i));"},
"textureStore(x_20, vi12, vec4i(vi12, 0i, 0i));"},
// Source 3 component, dest 2 component
{"%int 2D 0 0 0 2 Rg32i", "OpImageWrite %im %vi12 %vi123",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<rg32sint, write>;)",
"textureStore(x_20, vi12, vec4<i32>(vi123, 0i));"},
"textureStore(x_20, vi12, vec4i(vi123, 0i));"},
// Source 4 component, dest 2 component
{"%int 2D 0 0 0 2 Rg32i", "OpImageWrite %im %vi12 %vi1234",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<rg32sint, write>;)",
@ -2635,11 +2635,11 @@ INSTANTIATE_TEST_SUITE_P(
// Source unsigned, dest unsigned
{"%uint 2D 0 0 0 2 R32ui", "OpImageWrite %im %vi12 %vu12",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32uint, write>;)",
R"(textureStore(x_20, vi12, vec4<u32>(vu12, 0u, 0u)))"},
R"(textureStore(x_20, vi12, vec4u(vu12, 0u, 0u)))"},
// Source signed, dest signed
{"%int 2D 0 0 0 2 R32i", "OpImageWrite %im %vi12 %vi12",
R"(@group(2) @binding(1) var x_20 : texture_storage_2d<r32sint, write>;)",
R"(textureStore(x_20, vi12, vec4<i32>(vi12, 0i, 0i)))"}}));
R"(textureStore(x_20, vi12, vec4i(vi12, 0i, 0i)))"}}));
INSTANTIATE_TEST_SUITE_P(
ImageFetch_OptionalParams,
@ -2649,20 +2649,20 @@ INSTANTIATE_TEST_SUITE_P(
// Level of detail is injected for sampled texture
{"%float 2D 0 0 0 1 Unknown", "%99 = OpImageFetch %v4float %im %vi12",
R"(@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(let x_99 : vec4<f32> = textureLoad(x_20, vi12, 0i);)"},
R"(let x_99 : vec4f = textureLoad(x_20, vi12, 0i);)"},
// OpImageFetch with explicit level, on sampled texture
{"%float 2D 0 0 0 1 Unknown", "%99 = OpImageFetch %v4float %im %vi12 Lod %int_3",
R"(@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(let x_99 : vec4<f32> = textureLoad(x_20, vi12, 3i);)"},
R"(let x_99 : vec4f = textureLoad(x_20, vi12, 3i);)"},
// OpImageFetch with no extra params, on depth texture
// Level of detail is injected for depth texture
{"%float 2D 1 0 0 1 Unknown", "%99 = OpImageFetch %v4float %im %vi12",
R"(@group(2) @binding(1) var x_20 : texture_depth_2d;)",
R"(let x_99 : vec4<f32> = vec4<f32>(textureLoad(x_20, vi12, 0i), 0.0f, 0.0f, 0.0f);)"},
R"(let x_99 : vec4f = vec4f(textureLoad(x_20, vi12, 0i), 0.0f, 0.0f, 0.0f);)"},
// OpImageFetch with extra params, on depth texture
{"%float 2D 1 0 0 1 Unknown", "%99 = OpImageFetch %v4float %im %vi12 Lod %int_3",
R"(@group(2) @binding(1) var x_20 : texture_depth_2d;)",
R"(let x_99 : vec4<f32> = vec4<f32>(textureLoad(x_20, vi12, 3i), 0.0f, 0.0f, 0.0f);)"}}));
R"(let x_99 : vec4f = vec4f(textureLoad(x_20, vi12, 3i), 0.0f, 0.0f, 0.0f);)"}}));
INSTANTIATE_TEST_SUITE_P(
ImageFetch_Depth,
@ -2675,7 +2675,7 @@ INSTANTIATE_TEST_SUITE_P(
// ImageFetch on depth image.
{"%float 2D 1 0 0 1 Unknown", "%99 = OpImageFetch %v4float %im %vi12 ",
R"(@group(2) @binding(1) var x_20 : texture_depth_2d;)",
R"(let x_99 : vec4<f32> = vec4<f32>(textureLoad(x_20, vi12, 0i), 0.0f, 0.0f, 0.0f);)"}}));
R"(let x_99 : vec4f = vec4f(textureLoad(x_20, vi12, 0i), 0.0f, 0.0f, 0.0f);)"}}));
INSTANTIATE_TEST_SUITE_P(
ImageFetch_DepthMultisampled,
@ -2688,7 +2688,7 @@ INSTANTIATE_TEST_SUITE_P(
// ImageFetch on multisampled depth image.
{"%float 2D 1 0 1 1 Unknown", "%99 = OpImageFetch %v4float %im %vi12 Sample %i1",
R"(@group(2) @binding(1) var x_20 : texture_depth_multisampled_2d;)",
R"(let x_99 : vec4<f32> = vec4<f32>(textureLoad(x_20, vi12, i1), 0.0f, 0.0f, 0.0f);)"}}));
R"(let x_99 : vec4f = vec4f(textureLoad(x_20, vi12, i1), 0.0f, 0.0f, 0.0f);)"}}));
INSTANTIATE_TEST_SUITE_P(ImageFetch_Multisampled,
SpvParserHandleTest_ImageAccessTest,
@ -2703,7 +2703,7 @@ INSTANTIATE_TEST_SUITE_P(ImageFetch_Multisampled,
{"%float 2D 0 0 1 1 Unknown",
"%99 = OpImageFetch %v4float %im %vi12 Sample %i1",
R"(@group(2) @binding(1) var x_20 : texture_multisampled_2d<f32>;)",
R"(let x_99 : vec4<f32> = textureLoad(x_20, vi12, i1);)"}}));
R"(let x_99 : vec4f = textureLoad(x_20, vi12, i1);)"}}));
INSTANTIATE_TEST_SUITE_P(ImageFetch_Multisampled_ConvertSampleOperand,
SpvParserHandleTest_ImageAccessTest,
@ -2711,7 +2711,7 @@ INSTANTIATE_TEST_SUITE_P(ImageFetch_Multisampled_ConvertSampleOperand,
{"%float 2D 0 0 1 1 Unknown",
"%99 = OpImageFetch %v4float %im %vi12 Sample %u1",
R"(@group(2) @binding(1) var x_20 : texture_multisampled_2d<f32>;)",
R"(let x_99 : vec4<f32> = textureLoad(x_20, vi12, i32(u1));)"}}));
R"(let x_99 : vec4f = textureLoad(x_20, vi12, i32(u1));)"}}));
INSTANTIATE_TEST_SUITE_P(ConvertResultSignedness,
SpvParserHandleTest_SampledImageAccessTest,
@ -2733,11 +2733,11 @@ INSTANTIATE_TEST_SUITE_P(ConvertResultSignedness,
// OpImageFetch requires no conversion, float -> v4float
{"%float 2D 0 0 0 1 Unknown", "%99 = OpImageFetch %v4float %im %vi12",
R"(@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(let x_99 : vec4<f32> = textureLoad(x_20, vi12, 0i);)"},
R"(let x_99 : vec4f = textureLoad(x_20, vi12, 0i);)"},
// OpImageFetch requires no conversion, uint -> v4uint
{"%uint 2D 0 0 0 1 Unknown", "%99 = OpImageFetch %v4uint %im %vi12",
R"(@group(2) @binding(1) var x_20 : texture_2d<u32>;)",
R"(let x_99 : vec4<u32> = textureLoad(x_20, vi12, 0i);)"},
R"(let x_99 : vec4u = textureLoad(x_20, vi12, 0i);)"},
// OpImageFetch requires conversion, uint -> v4int
// is invalid SPIR-V:
// "Expected Image 'Sampled Type' to be the same as Result Type
@ -2746,7 +2746,7 @@ INSTANTIATE_TEST_SUITE_P(ConvertResultSignedness,
// OpImageFetch requires no conversion, int -> v4int
{"%int 2D 0 0 0 1 Unknown", "%99 = OpImageFetch %v4int %im %vi12",
R"(@group(2) @binding(1) var x_20 : texture_2d<i32>;)",
R"(let x_99 : vec4<i32> = textureLoad(x_20, vi12, 0i);)"},
R"(let x_99 : vec4i = textureLoad(x_20, vi12, 0i);)"},
// OpImageFetch requires conversion, int -> v4uint
// is invalid SPIR-V:
// "Expected Image 'Sampled Type' to be the same as Result Type
@ -2759,11 +2759,11 @@ INSTANTIATE_TEST_SUITE_P(ConvertResultSignedness,
// OpImageRead requires no conversion, float -> v4float
{"%float 2D 0 0 0 2 Rgba32f", "%99 = OpImageRead %v4float %im %vi12",
R"(@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(let x_99 : vec4<f32> = textureLoad(x_20, vi12, 0i);)"},
R"(let x_99 : vec4f = textureLoad(x_20, vi12, 0i);)"},
// OpImageRead requires no conversion, uint -> v4uint
{"%uint 2D 0 0 0 2 Rgba32ui", "%99 = OpImageRead %v4uint %im %vi12",
R"(@group(2) @binding(1) var x_20 : texture_2d<u32>;)",
R"(let x_99 : vec4<u32> = textureLoad(x_20, vi12, 0i);)"},
R"(let x_99 : vec4u = textureLoad(x_20, vi12, 0i);)"},
// OpImageRead requires conversion, uint -> v4int
// is invalid SPIR-V:
@ -2773,7 +2773,7 @@ INSTANTIATE_TEST_SUITE_P(ConvertResultSignedness,
// OpImageRead requires no conversion, int -> v4int
{"%int 2D 0 0 0 2 Rgba32i", "%99 = OpImageRead %v4int %im %vi12",
R"(@group(2) @binding(1) var x_20 : texture_2d<i32>;)",
R"(let x_99 : vec4<i32> = textureLoad(x_20, vi12, 0i);)"},
R"(let x_99 : vec4i = textureLoad(x_20, vi12, 0i);)"},
// OpImageRead requires conversion, int -> v4uint
// is invalid SPIR-V:
@ -2792,7 +2792,7 @@ INSTANTIATE_TEST_SUITE_P(ConvertResultSignedness,
R"(@group(0) @binding(0) var x_10 : sampler;
@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(let x_99 : vec4<f32> = textureSample(x_20, x_10, vf12);)"}}));
R"(let x_99 : vec4f = textureSample(x_20, x_10, vf12);)"}}));
INSTANTIATE_TEST_SUITE_P(ImageQuerySize_NonArrayed_SignedResult,
// ImageQuerySize requires storage image or multisampled
@ -2813,19 +2813,19 @@ INSTANTIATE_TEST_SUITE_P(ImageQuerySize_NonArrayed_SignedResult,
"%98 = OpImageRead %v4float %im %vi12\n", // Implicitly mark as
// NonWritable
R"(@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(let x_99 : vec2<i32> = vec2<i32>(textureDimensions(x_20))"},
R"(let x_99 : vec2i = vec2i(textureDimensions(x_20))"},
// 3D storage image
{"%float 3D 0 0 0 2 Rgba32f",
"%99 = OpImageQuerySize %v3int %im \n"
"%98 = OpImageRead %v4float %im %vi123\n", // Implicitly mark as
// NonWritable
R"(@group(2) @binding(1) var x_20 : texture_3d<f32>;)",
R"(let x_99 : vec3<i32> = vec3<i32>(textureDimensions(x_20));)"},
R"(let x_99 : vec3i = vec3i(textureDimensions(x_20));)"},
// Multisampled
{"%float 2D 0 0 1 1 Unknown", "%99 = OpImageQuerySize %v2int %im \n",
R"(@group(2) @binding(1) var x_20 : texture_multisampled_2d<f32>;)",
R"(let x_99 : vec2<i32> = vec2<i32>(textureDimensions(x_20));)"}}));
R"(let x_99 : vec2i = vec2i(textureDimensions(x_20));)"}}));
INSTANTIATE_TEST_SUITE_P(
ImageQuerySize_Arrayed_SignedResult,
@ -2841,7 +2841,7 @@ INSTANTIATE_TEST_SUITE_P(
"%99 = OpImageQuerySize %v3int %im \n"
"%98 = OpImageRead %v4float %im %vi123\n",
R"(@group(2) @binding(1) var x_20 : texture_2d_array<f32>;)",
R"(let x_99 : vec3<i32> = vec3<i32>(vec3<u32>(textureDimensions(x_20), textureNumLayers(x_20)));)"}
R"(let x_99 : vec3i = vec3i(vec3u(textureDimensions(x_20), textureNumLayers(x_20)));)"}
// 3D array storage image doesn't exist.
// Multisampled array
@ -2862,27 +2862,27 @@ INSTANTIATE_TEST_SUITE_P(
// 2D
{"%float 2D 0 0 0 1 Unknown", "%99 = OpImageQuerySizeLod %v2int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(let x_99 : vec2<i32> = vec2<i32>(textureDimensions(x_20, i1));)"},
R"(let x_99 : vec2i = vec2i(textureDimensions(x_20, i1));)"},
// 3D
{"%float 3D 0 0 0 1 Unknown", "%99 = OpImageQuerySizeLod %v3int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_3d<f32>;)",
R"(let x_99 : vec3<i32> = vec3<i32>(textureDimensions(x_20, i1));)"},
R"(let x_99 : vec3i = vec3i(textureDimensions(x_20, i1));)"},
// Cube
{"%float Cube 0 0 0 1 Unknown", "%99 = OpImageQuerySizeLod %v2int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_cube<f32>;)",
R"(let x_99 : vec2<i32> = vec2<i32>(textureDimensions(x_20, i1).xy);)"},
R"(let x_99 : vec2i = vec2i(textureDimensions(x_20, i1).xy);)"},
// Depth 2D
{"%float 2D 1 0 0 1 Unknown", "%99 = OpImageQuerySizeLod %v2int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_depth_2d;)",
R"(let x_99 : vec2<i32> = vec2<i32>(textureDimensions(x_20, i1));)"},
R"(let x_99 : vec2i = vec2i(textureDimensions(x_20, i1));)"},
// Depth Cube
{"%float Cube 1 0 0 1 Unknown", "%99 = OpImageQuerySizeLod %v2int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_depth_cube;)",
R"(let x_99 : vec2<i32> = vec2<i32>(textureDimensions(x_20, i1).xy);)"}}));
R"(let x_99 : vec2i = vec2i(textureDimensions(x_20, i1).xy);)"}}));
INSTANTIATE_TEST_SUITE_P(
ImageQuerySizeLod_Arrayed_SignedResult_SignedLevel,
@ -2897,7 +2897,7 @@ INSTANTIATE_TEST_SUITE_P(
// 2D array
{"%float 2D 0 1 0 1 Unknown", "%99 = OpImageQuerySizeLod %v3int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_2d_array<f32>;)",
R"(let x_99 : vec3<i32> = vec3<i32>(vec3<u32>(textureDimensions(x_20, i1), textureNumLayers(x_20)));)"},
R"(let x_99 : vec3i = vec3i(vec3u(textureDimensions(x_20, i1), textureNumLayers(x_20)));)"},
// There is no 3D array
@ -2908,12 +2908,12 @@ INSTANTIATE_TEST_SUITE_P(
// https://github.com/gpuweb/gpuweb/issues/1345
{"%float Cube 0 1 0 1 Unknown", "%99 = OpImageQuerySizeLod %v3int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_cube_array<f32>;)",
R"(let x_99 : vec3<i32> = vec3<i32>(vec3<u32>(textureDimensions(x_20, i1).xy, textureNumLayers(x_20)));)"},
R"(let x_99 : vec3i = vec3i(vec3u(textureDimensions(x_20, i1).xy, textureNumLayers(x_20)));)"},
// Depth 2D array
{"%float 2D 1 1 0 1 Unknown", "%99 = OpImageQuerySizeLod %v3int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_depth_2d_array;)",
R"(let x_99 : vec3<i32> = vec3<i32>(vec3<u32>(textureDimensions(x_20, i1), textureNumLayers(x_20)));)"},
R"(let x_99 : vec3i = vec3i(vec3u(textureDimensions(x_20, i1), textureNumLayers(x_20)));)"},
// Depth Cube Array
//
@ -2922,7 +2922,7 @@ INSTANTIATE_TEST_SUITE_P(
// https://github.com/gpuweb/gpuweb/issues/1345
{"%float Cube 1 1 0 1 Unknown", "%99 = OpImageQuerySizeLod %v3int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_depth_cube_array;)",
R"(let x_99 : vec3<i32> = vec3<i32>(vec3<u32>(textureDimensions(x_20, i1).xy, textureNumLayers(x_20)));)"}}));
R"(let x_99 : vec3i = vec3i(vec3u(textureDimensions(x_20, i1).xy, textureNumLayers(x_20)));)"}}));
INSTANTIATE_TEST_SUITE_P(
// textureDimensions accepts both signed and unsigned the level-of-detail values.
@ -3414,12 +3414,9 @@ INSTANTIATE_TEST_SUITE_P(
{"%float 2D 0 0 0 1 Unknown",
"%result = OpImageFetch %v4float %im %vu12",
"",
{"vec2<i32>(vu12)"}},
{"%float 2D 0 0 0 2 R32f",
"%result = OpImageRead %v4float %im %vu12",
"",
{"vec2<i32>(vu12)"}},
{"%float 2D 0 0 0 2 R32f", "OpImageWrite %im %vu12 %vf1234", "", {"vec2<i32>(vu12)"}}}));
{"vec2i(vu12)"}},
{"%float 2D 0 0 0 2 R32f", "%result = OpImageRead %v4float %im %vu12", "", {"vec2i(vu12)"}},
{"%float 2D 0 0 0 2 R32f", "OpImageWrite %im %vu12 %vf1234", "", {"vec2i(vu12)"}}}));
INSTANTIATE_TEST_SUITE_P(ConvertUintCoords_Arrayed,
// In SPIR-V, image read, fetch, and write use integer coordinates.
@ -3429,15 +3426,15 @@ INSTANTIATE_TEST_SUITE_P(ConvertUintCoords_Arrayed,
{"%float 2D 0 1 0 1 Unknown",
"%result = OpImageFetch %v4float %im %vu123",
"",
{"vec2<i32>(vu123.xy)", "i32(vu123.z)"}},
{"vec2i(vu123.xy)", "i32(vu123.z)"}},
{"%float 2D 0 1 0 2 R32f",
"%result = OpImageRead %v4float %im %vu123",
"",
{"vec2<i32>(vu123.xy)", "i32(vu123.z)"}},
{"vec2i(vu123.xy)", "i32(vu123.z)"}},
{"%float 2D 0 1 0 2 R32f",
"OpImageWrite %im %vu123 %vf1234",
"",
{"vec2<i32>(vu123.xy)", "i32(vu123.z)"}}}));
{"vec2i(vu123.xy)", "i32(vu123.z)"}}}));
INSTANTIATE_TEST_SUITE_P(
BadInstructions,
@ -3817,9 +3814,9 @@ TEST_F(SpvParserHandleTest, NeverGenerateConstDeclForHandle_UseVariableDirectly)
EXPECT_TRUE(p->error().empty()) << p->error();
auto ast_body = fe.ast_body();
const auto got = test::ToString(p->program(), ast_body);
auto* expect = R"(var var_1 : vec4<f32>;
let x_22 : vec4<f32> = textureSample(x_2, x_3, vec2<f32>());
let x_26 : vec4<f32> = textureSample(x_2, x_3, vec2<f32>());
auto* expect = R"(var var_1 : vec4f;
let x_22 : vec4f = textureSample(x_2, x_3, vec2f());
let x_26 : vec4f = textureSample(x_2, x_3, vec2f());
var_1 = (x_22 + x_26);
return;
)";
@ -3893,7 +3890,7 @@ TEST_F(SpvParserHandleTest, SamplerLoadedInEnclosingConstruct_DontGenerateVar) {
auto* expect = R"(switch(0i) {
default: {
if (true) {
let x_24 : vec4<f32> = textureSample(var_im, var_s, vec2<f32>());
let x_24 : vec4f = textureSample(var_im, var_s, vec2f());
}
}
}
@ -3980,7 +3977,7 @@ OpFunctionEnd
x_900 = 0.0f;
if (true) {
if (true) {
let x_18 : vec4<f32> = textureSample(x_20, x_10, x_900);
let x_18 : vec4f = textureSample(x_20, x_10, x_900);
}
}
return;
@ -4063,11 +4060,11 @@ OpFunctionEnd
EXPECT_TRUE(p->error().empty()) << p->error();
auto ast_body = fe.ast_body();
const auto got = test::ToString(p->program(), ast_body);
auto* expect = R"(var x_900 : vec2<f32>;
x_900 = vec2<f32>();
auto* expect = R"(var x_900 : vec2f;
x_900 = vec2f();
if (true) {
if (true) {
let x_19 : vec4<f32> = textureSample(x_20, x_10, x_900.x);
let x_19 : vec4f = textureSample(x_20, x_10, x_900.x);
}
}
return;
@ -4138,9 +4135,9 @@ TEST_F(SpvParserHandleTest, TexelTypeWhenLoop) {
EXPECT_TRUE(p->error().empty()) << p->error();
auto ast_body = fe.ast_body();
const auto got = test::ToString(p->program(), ast_body);
auto* expect = R"(var x_24 : vec2<f32>;
auto* expect = R"(var x_24 : vec2f;
var x_26 : i32;
x_24 = vec2<f32>(0.0f, 0.0f);
x_24 = vec2f(0.0f, 0.0f);
x_26 = 0i;
loop {
var x_27 : i32;
@ -4151,11 +4148,11 @@ loop {
continuing {
x_27 = (x_26 + 1i);
x_24 = vec2<f32>(1.0f, 1.0f);
x_24 = vec2f(1.0f, 1.0f);
x_26 = x_27;
}
}
textureStore(Output2Texture2D, vec2<i32>(vec2<u32>(1u, 1u)), vec4<f32>(x_24, 0.0f, 0.0f));
textureStore(Output2Texture2D, vec2i(vec2u(1u, 1u)), vec4f(x_24, 0.0f, 0.0f));
return;
)";
ASSERT_EQ(expect, got);
@ -4230,7 +4227,7 @@ loop {
}
}
let x_21 : f32 = select(0.0f, x_14, (x_14 > 1.0f));
x_1 = vec4<f32>(x_21, x_21, x_21, x_21);
x_1 = vec4f(x_21, x_21, x_21, x_21);
return;
)";
ASSERT_EQ(expect, got);

246
src/tint/reader/spirv/parser_impl_module_var_test.cc
View File

@ -347,7 +347,7 @@ TEST_F(SpvModuleScopeVarParserTest, BuiltinPosition_StorePosition) {
EXPECT_TRUE(p->BuildAndParseInternalModule());
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_THAT(module_str, HasSubstr("gl_Position = vec4<f32>();")) << module_str;
EXPECT_THAT(module_str, HasSubstr("gl_Position = vec4f();")) << module_str;
}
TEST_F(SpvModuleScopeVarParserTest, BuiltinPosition_StorePosition_PerVertexStructOutOfOrderDecl) {
@ -389,7 +389,7 @@ TEST_F(SpvModuleScopeVarParserTest, BuiltinPosition_StorePosition_PerVertexStruc
EXPECT_TRUE(p->BuildAndParseInternalModule());
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_THAT(module_str, HasSubstr("gl_Position = vec4<f32>();")) << module_str;
EXPECT_THAT(module_str, HasSubstr("gl_Position = vec4f();")) << module_str;
}
TEST_F(SpvModuleScopeVarParserTest, BuiltinPosition_StorePositionMember_OneAccessChain) {
@ -449,7 +449,7 @@ TEST_F(SpvModuleScopeVarParserTest, BuiltinPointSize_Write1_IsErased) {
EXPECT_TRUE(p->BuildAndParseInternalModule());
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_EQ(module_str, R"(var<private> gl_Position : vec4<f32>;
EXPECT_EQ(module_str, R"(var<private> gl_Position : vec4f;
fn main_1() {
return;
@ -457,7 +457,7 @@ fn main_1() {
struct main_out {
@builtin(position)
gl_Position : vec4<f32>,
gl_Position : vec4f,
}
@vertex
@ -507,7 +507,7 @@ TEST_F(SpvModuleScopeVarParserTest, BuiltinPointSize_ReadReplaced) {
const auto module_str = test::ToString(p->program());
EXPECT_EQ(module_str, R"(var<private> x_900 : f32;
var<private> gl_Position : vec4<f32>;
var<private> gl_Position : vec4f;
fn main_1() {
x_900 = 1.0f;
@ -516,7 +516,7 @@ fn main_1() {
struct main_out {
@builtin(position)
gl_Position : vec4<f32>,
gl_Position : vec4f,
}
@vertex
@ -564,7 +564,7 @@ TEST_F(SpvModuleScopeVarParserTest, BuiltinPointSize_WriteViaCopyObjectPostAcces
EXPECT_TRUE(p->BuildAndParseInternalModule()) << p->error();
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_EQ(module_str, R"(var<private> gl_Position : vec4<f32>;
EXPECT_EQ(module_str, R"(var<private> gl_Position : vec4f;
fn main_1() {
return;
@ -572,7 +572,7 @@ fn main_1() {
struct main_out {
@builtin(position)
gl_Position : vec4<f32>,
gl_Position : vec4f,
}
@vertex
@ -623,7 +623,7 @@ TEST_F(SpvModuleScopeVarParserTest, BuiltinPointSize_Loose_Write1_IsErased) {
EXPECT_TRUE(p->BuildAndParseInternalModule());
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_EQ(module_str, R"(var<private> x_2 : vec4<f32>;
EXPECT_EQ(module_str, R"(var<private> x_2 : vec4f;
fn main_1() {
return;
@ -631,7 +631,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
}
@vertex
@ -676,7 +676,7 @@ TEST_F(SpvModuleScopeVarParserTest, BuiltinPointSize_Loose_ReadReplaced_Vertex)
EXPECT_TRUE(p->BuildAndParseInternalModule());
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_EQ(module_str, R"(var<private> x_2 : vec4<f32>;
EXPECT_EQ(module_str, R"(var<private> x_2 : vec4f;
var<private> x_900 : f32;
@ -687,7 +687,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
}
@vertex
@ -735,7 +735,7 @@ TEST_F(SpvModuleScopeVarParserTest, BuiltinPointSize_Loose_WriteViaCopyObjectPri
EXPECT_TRUE(p->BuildAndParseInternalModule()) << p->error();
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_EQ(module_str, R"(var<private> x_2 : vec4<f32>;
EXPECT_EQ(module_str, R"(var<private> x_2 : vec4f;
fn main_1() {
return;
@ -743,7 +743,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
}
@vertex
@ -771,7 +771,7 @@ TEST_F(SpvModuleScopeVarParserTest,
EXPECT_TRUE(p->BuildAndParseInternalModule()) << p->error();
EXPECT_TRUE(p->error().empty()) << p->error();
const auto module_str = test::ToString(p->program());
EXPECT_EQ(module_str, R"(var<private> x_2 : vec4<f32>;
EXPECT_EQ(module_str, R"(var<private> x_2 : vec4f;
fn main_1() {
return;
@ -779,7 +779,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
}
@vertex
@ -956,7 +956,7 @@ TEST_F(SpvModuleScopeVarParserTest, VectorInitializer) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2<f32> = vec2<f32>(1.5f, 2.0f);"));
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2f = vec2f(1.5f, 2.0f);"));
}
TEST_F(SpvModuleScopeVarParserTest, VectorBoolNullInitializer) {
@ -995,7 +995,7 @@ TEST_F(SpvModuleScopeVarParserTest, VectorUintNullInitializer) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2<u32> = vec2<u32>();"));
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2u = vec2u();"));
}
TEST_F(SpvModuleScopeVarParserTest, VectorUintUndefInitializer) {
@ -1007,7 +1007,7 @@ TEST_F(SpvModuleScopeVarParserTest, VectorUintUndefInitializer) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2<u32> = vec2<u32>();"));
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2u = vec2u();"));
// This example module emits ok, but is not valid SPIR-V in the first place.
p->DeliberatelyInvalidSpirv();
@ -1022,7 +1022,7 @@ TEST_F(SpvModuleScopeVarParserTest, VectorIntNullInitializer) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2<i32> = vec2<i32>();"));
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2i = vec2i();"));
}
TEST_F(SpvModuleScopeVarParserTest, VectorIntUndefInitializer) {
@ -1034,7 +1034,7 @@ TEST_F(SpvModuleScopeVarParserTest, VectorIntUndefInitializer) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2<i32> = vec2<i32>();"));
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2i = vec2i();"));
// This example module emits ok, but is not valid SPIR-V in the first place.
p->DeliberatelyInvalidSpirv();
@ -1049,7 +1049,7 @@ TEST_F(SpvModuleScopeVarParserTest, VectorFloatNullInitializer) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2<f32> = vec2<f32>();"));
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2f = vec2f();"));
}
TEST_F(SpvModuleScopeVarParserTest, VectorFloatUndefInitializer) {
@ -1061,7 +1061,7 @@ TEST_F(SpvModuleScopeVarParserTest, VectorFloatUndefInitializer) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2<f32> = vec2<f32>();"));
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : vec2f = vec2f();"));
// This example module emits ok, but is not valid SPIR-V in the first place.
p->DeliberatelyInvalidSpirv();
@ -1083,9 +1083,9 @@ TEST_F(SpvModuleScopeVarParserTest, MatrixInitializer) {
EXPECT_TRUE(p->error().empty());
const auto module_str = test::ToString(p->program());
EXPECT_THAT(module_str, HasSubstr("var<private> x_200 : mat3x2<f32> = mat3x2<f32>("
"vec2<f32>(1.5f, 2.0f), "
"vec2<f32>(2.0f, 3.0f), "
"vec2<f32>(3.0f, 4.0f));"));
"vec2f(1.5f, 2.0f), "
"vec2f(2.0f, 3.0f), "
"vec2f(3.0f, 4.0f));"));
}
TEST_F(SpvModuleScopeVarParserTest, MatrixNullInitializer) {
@ -2529,7 +2529,7 @@ TEST_F(SpvModuleScopeVarParserTest, VertexIndex_I32_Load_Direct) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : i32;
var<private> x_4 : vec4<f32>;
var<private> x_4 : vec4f;
fn main_1() {
let x_2 : i32 = x_1;
@ -2538,7 +2538,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_4_1 : vec4<f32>,
x_4_1 : vec4f,
}
@vertex
@ -2586,7 +2586,7 @@ TEST_F(SpvModuleScopeVarParserTest, VertexIndex_UsedTwice_DifferentConstructs) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : u32;
var<private> x_5 : vec4<f32>;
var<private> x_5 : vec4f;
fn main_1() {
let x_2 : u32 = x_1;
@ -2600,7 +2600,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_5_1 : vec4<f32>,
x_5_1 : vec4f,
}
@vertex
@ -2628,7 +2628,7 @@ TEST_F(SpvModuleScopeVarParserTest, VertexIndex_I32_Load_CopyObject) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : i32;
var<private> x_4 : vec4<f32>;
var<private> x_4 : vec4f;
fn main_1() {
let x_14 : ptr<private, i32> = &(x_1);
@ -2638,7 +2638,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_4_1 : vec4<f32>,
x_4_1 : vec4f,
}
@vertex
@ -2666,7 +2666,7 @@ TEST_F(SpvModuleScopeVarParserTest, VertexIndex_I32_Load_AccessChain) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : i32;
var<private> x_4 : vec4<f32>;
var<private> x_4 : vec4f;
fn main_1() {
let x_2 : i32 = x_1;
@ -2675,7 +2675,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_4_1 : vec4<f32>,
x_4_1 : vec4f,
}
@vertex
@ -2702,7 +2702,7 @@ TEST_F(SpvModuleScopeVarParserTest, VertexIndex_U32_Load_Direct) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : u32;
var<private> x_4 : vec4<f32>;
var<private> x_4 : vec4f;
fn main_1() {
let x_2 : u32 = x_1;
@ -2711,7 +2711,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_4_1 : vec4<f32>,
x_4_1 : vec4f,
}
@vertex
@ -2739,7 +2739,7 @@ TEST_F(SpvModuleScopeVarParserTest, VertexIndex_U32_Load_CopyObject) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : u32;
var<private> x_4 : vec4<f32>;
var<private> x_4 : vec4f;
fn main_1() {
let x_14 : ptr<private, u32> = &(x_1);
@ -2749,7 +2749,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_4_1 : vec4<f32>,
x_4_1 : vec4f,
}
@vertex
@ -2777,7 +2777,7 @@ TEST_F(SpvModuleScopeVarParserTest, VertexIndex_U32_Load_AccessChain) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : u32;
var<private> x_4 : vec4<f32>;
var<private> x_4 : vec4f;
fn main_1() {
let x_2 : u32 = x_1;
@ -2786,7 +2786,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_4_1 : vec4<f32>,
x_4_1 : vec4f,
}
@vertex
@ -2861,7 +2861,7 @@ TEST_F(SpvModuleScopeVarParserTest, InstanceIndex_I32_Load_Direct) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : i32;
var<private> position_1 : vec4<f32>;
var<private> position_1 : vec4f;
fn main_1() {
let x_2 : i32 = x_1;
@ -2870,7 +2870,7 @@ fn main_1() {
struct main_out {
@builtin(position)
position_1_1 : vec4<f32>,
position_1_1 : vec4f,
}
@vertex
@ -2898,7 +2898,7 @@ TEST_F(SpvModuleScopeVarParserTest, InstanceIndex_I32_Load_CopyObject) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : i32;
var<private> position_1 : vec4<f32>;
var<private> position_1 : vec4f;
fn main_1() {
let x_14 : ptr<private, i32> = &(x_1);
@ -2908,7 +2908,7 @@ fn main_1() {
struct main_out {
@builtin(position)
position_1_1 : vec4<f32>,
position_1_1 : vec4f,
}
@vertex
@ -2936,7 +2936,7 @@ TEST_F(SpvModuleScopeVarParserTest, InstanceIndex_I32_Load_AccessChain) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : i32;
var<private> position_1 : vec4<f32>;
var<private> position_1 : vec4f;
fn main_1() {
let x_2 : i32 = x_1;
@ -2945,7 +2945,7 @@ fn main_1() {
struct main_out {
@builtin(position)
position_1_1 : vec4<f32>,
position_1_1 : vec4f,
}
@vertex
@ -2995,7 +2995,7 @@ TEST_F(SpvModuleScopeVarParserTest, InstanceIndex_U32_Load_Direct) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : u32;
var<private> position_1 : vec4<f32>;
var<private> position_1 : vec4f;
fn main_1() {
let x_2 : u32 = x_1;
@ -3004,7 +3004,7 @@ fn main_1() {
struct main_out {
@builtin(position)
position_1_1 : vec4<f32>,
position_1_1 : vec4f,
}
@vertex
@ -3032,7 +3032,7 @@ TEST_F(SpvModuleScopeVarParserTest, InstanceIndex_U32_Load_CopyObject) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : u32;
var<private> position_1 : vec4<f32>;
var<private> position_1 : vec4f;
fn main_1() {
let x_14 : ptr<private, u32> = &(x_1);
@ -3042,7 +3042,7 @@ fn main_1() {
struct main_out {
@builtin(position)
position_1_1 : vec4<f32>,
position_1_1 : vec4f,
}
@vertex
@ -3070,7 +3070,7 @@ TEST_F(SpvModuleScopeVarParserTest, InstanceIndex_U32_Load_AccessChain) {
const auto module_str = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : u32;
var<private> position_1 : vec4<f32>;
var<private> position_1 : vec4f;
fn main_1() {
let x_2 : u32 = x_1;
@ -3079,7 +3079,7 @@ fn main_1() {
struct main_out {
@builtin(position)
position_1_1 : vec4<f32>,
position_1_1 : vec4f,
}
@vertex
@ -3165,10 +3165,10 @@ std::string WgslType(std::string spirv_type) {
return "i32";
}
if (spirv_type == "%v3uint") {
return "vec3<u32>";
return "vec3u";
}
if (spirv_type == "%v3int") {
return "vec3<i32>";
return "vec3i";
}
return "error";
}
@ -3180,11 +3180,11 @@ std::string UnsignedWgslType(std::string wgsl_type) {
if (wgsl_type == "i32") {
return "u32";
}
if (wgsl_type == "vec3<u32>") {
return "vec3<u32>";
if (wgsl_type == "vec3u") {
return "vec3u";
}
if (wgsl_type == "vec3<i32>") {
return "vec3<u32>";
if (wgsl_type == "vec3i") {
return "vec3u";
}
return "error";
}
@ -3196,11 +3196,11 @@ std::string SignedWgslType(std::string wgsl_type) {
if (wgsl_type == "i32") {
return "i32";
}
if (wgsl_type == "vec3<u32>") {
return "vec3<i32>";
if (wgsl_type == "vec3u") {
return "vec3i";
}
if (wgsl_type == "vec3<i32>") {
return "vec3<i32>";
if (wgsl_type == "vec3i") {
return "vec3i";
}
return "error";
}
@ -3746,7 +3746,7 @@ TEST_F(SpvModuleScopeVarParserTest, EntryPointWrapping_BuiltinVar_Input_SameSign
const auto got = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : u32;
var<private> x_4 : vec4<f32>;
var<private> x_4 : vec4f;
fn main_1() {
let x_2 : u32 = x_1;
@ -3755,7 +3755,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_4_1 : vec4<f32>,
x_4_1 : vec4f,
}
@vertex
@ -3795,7 +3795,7 @@ TEST_F(SpvModuleScopeVarParserTest, EntryPointWrapping_BuiltinVar_Input_Opposite
const auto got = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : i32;
var<private> x_4 : vec4<f32>;
var<private> x_4 : vec4f;
fn main_1() {
let x_2 : i32 = x_1;
@ -3804,7 +3804,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_4_1 : vec4<f32>,
x_4_1 : vec4f,
}
@vertex
@ -4043,7 +4043,7 @@ TEST_F(SpvModuleScopeVarParserTest, BuiltinPosition_BuiltIn_Position) {
EXPECT_TRUE(p->error().empty());
const auto got = test::ToString(p->program());
const std::string expected = R"(var<private> gl_Position : vec4<f32>;
const std::string expected = R"(var<private> gl_Position : vec4f;
fn main_1() {
return;
@ -4051,7 +4051,7 @@ fn main_1() {
struct main_out {
@builtin(position)
gl_Position : vec4<f32>,
gl_Position : vec4f,
}
@vertex
@ -4112,7 +4112,7 @@ TEST_F(SpvModuleScopeVarParserTest, BuiltinPosition_BuiltIn_Position_Initializer
const auto got = test::ToString(p->program());
const std::string expected =
R"(var<private> gl_Position : vec4<f32> = vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f);
R"(var<private> gl_Position : vec4f = vec4f(1.0f, 2.0f, 3.0f, 4.0f);
fn main_1() {
return;
@ -4120,7 +4120,7 @@ fn main_1() {
struct main_out {
@builtin(position)
gl_Position : vec4<f32>,
gl_Position : vec4f,
}
@vertex
@ -4198,7 +4198,7 @@ TEST_F(SpvModuleScopeVarParserTest, Input_FlattenArray_OneLevel) {
const auto got = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : array<f32, 3u>;
var<private> x_2 : vec4<f32>;
var<private> x_2 : vec4f;
fn main_1() {
return;
@ -4206,7 +4206,7 @@ fn main_1() {
struct main_out {
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
}
@vertex
@ -4256,7 +4256,7 @@ TEST_F(SpvModuleScopeVarParserTest, Input_FlattenMatrix) {
const auto got = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : mat2x4<f32>;
var<private> x_2 : vec4<f32>;
var<private> x_2 : vec4f;
fn main_1() {
return;
@ -4264,11 +4264,11 @@ fn main_1() {
struct main_out {
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
}
@vertex
fn main(@location(9) x_1_param : vec4<f32>, @location(10) x_1_param_1 : vec4<f32>) -> main_out {
fn main(@location(9) x_1_param : vec4f, @location(10) x_1_param_1 : vec4f) -> main_out {
x_1[0i] = x_1_param;
x_1[1i] = x_1_param_1;
main_1();
@ -4318,12 +4318,12 @@ TEST_F(SpvModuleScopeVarParserTest, Input_FlattenStruct_LocOnVariable) {
const auto got = test::ToString(p->program());
const std::string expected = R"(struct Communicators {
alice : f32,
bob : vec4<f32>,
bob : vec4f,
}
var<private> x_1 : Communicators;
var<private> x_2 : vec4<f32>;
var<private> x_2 : vec4f;
fn main_1() {
return;
@ -4331,11 +4331,11 @@ fn main_1() {
struct main_out {
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
}
@vertex
fn main(@location(9) x_1_param : f32, @location(10) x_1_param_1 : vec4<f32>) -> main_out {
fn main(@location(9) x_1_param : f32, @location(10) x_1_param_1 : vec4f) -> main_out {
x_1.alice = x_1_param;
x_1.bob = x_1_param_1;
main_1();
@ -4382,7 +4382,7 @@ TEST_F(SpvModuleScopeVarParserTest, Input_FlattenNested) {
const auto got = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : array<mat2x4<f32>, 2u>;
var<private> x_2 : vec4<f32>;
var<private> x_2 : vec4f;
fn main_1() {
return;
@ -4390,11 +4390,11 @@ fn main_1() {
struct main_out {
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
}
@vertex
fn main(@location(7) x_1_param : vec4<f32>, @location(8) x_1_param_1 : vec4<f32>, @location(9) x_1_param_2 : vec4<f32>, @location(10) x_1_param_3 : vec4<f32>) -> main_out {
fn main(@location(7) x_1_param : vec4f, @location(8) x_1_param_1 : vec4f, @location(9) x_1_param_2 : vec4f, @location(10) x_1_param_3 : vec4f) -> main_out {
x_1[0i][0i] = x_1_param;
x_1[0i][1i] = x_1_param_1;
x_1[1i][0i] = x_1_param_2;
@ -4443,7 +4443,7 @@ TEST_F(SpvModuleScopeVarParserTest, Output_FlattenArray_OneLevel) {
const auto got = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : array<f32, 3u>;
var<private> x_2 : vec4<f32>;
var<private> x_2 : vec4f;
fn main_1() {
return;
@ -4457,7 +4457,7 @@ struct main_out {
@location(6)
x_1_3 : f32,
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
}
@vertex
@ -4504,7 +4504,7 @@ TEST_F(SpvModuleScopeVarParserTest, Output_FlattenMatrix) {
const auto got = test::ToString(p->program());
const std::string expected = R"(var<private> x_1 : mat2x4<f32>;
var<private> x_2 : vec4<f32>;
var<private> x_2 : vec4f;
fn main_1() {
return;
@ -4512,11 +4512,11 @@ fn main_1() {
struct main_out {
@location(9)
x_1_1 : vec4<f32>,
x_1_1 : vec4f,
@location(10)
x_1_2 : vec4<f32>,
x_1_2 : vec4f,
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
}
@vertex
@ -4568,12 +4568,12 @@ TEST_F(SpvModuleScopeVarParserTest, Output_FlattenStruct_LocOnVariable) {
const auto got = test::ToString(p->program());
const std::string expected = R"(struct Communicators {
alice : f32,
bob : vec4<f32>,
bob : vec4f,
}
var<private> x_1 : Communicators;
var<private> x_2 : vec4<f32>;
var<private> x_2 : vec4f;
fn main_1() {
return;
@ -4583,9 +4583,9 @@ struct main_out {
@location(9)
x_1_1 : f32,
@location(10)
x_1_2 : vec4<f32>,
x_1_2 : vec4f,
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
}
@vertex
@ -4641,14 +4641,14 @@ TEST_F(SpvModuleScopeVarParserTest, FlattenStruct_LocOnMembers) {
const auto got = test::ToString(p->program());
const std::string expected = R"(struct Communicators {
alice : f32,
bob : vec4<f32>,
bob : vec4f,
}
var<private> x_1 : Communicators;
var<private> x_3 : Communicators;
var<private> x_2 : vec4<f32>;
var<private> x_2 : vec4f;
fn main_1() {
return;
@ -4656,15 +4656,15 @@ fn main_1() {
struct main_out {
@builtin(position)
x_2_1 : vec4<f32>,
x_2_1 : vec4f,
@location(9)
x_3_1 : f32,
@location(11)
x_3_2 : vec4<f32>,
x_3_2 : vec4f,
}
@vertex
fn main(@location(9) x_1_param : f32, @location(11) x_1_param_1 : vec4<f32>) -> main_out {
fn main(@location(9) x_1_param : f32, @location(11) x_1_param_1 : vec4f) -> main_out {
x_1.alice = x_1_param;
x_1.bob = x_1_param_1;
main_1();
@ -4721,17 +4721,17 @@ TEST_F(SpvModuleScopeVarParserTest, EntryPointWrapping_Interpolation_Flat_Vertex
const std::string expected =
R"(var<private> x_1 : u32;
var<private> x_2 : vec2<u32>;
var<private> x_2 : vec2u;
var<private> x_3 : i32;
var<private> x_4 : vec2<i32>;
var<private> x_4 : vec2i;
var<private> x_5 : f32;
var<private> x_6 : vec2<f32>;
var<private> x_6 : vec2f;
var<private> x_10 : vec4<f32>;
var<private> x_10 : vec4f;
fn main_1() {
return;
@ -4739,11 +4739,11 @@ fn main_1() {
struct main_out {
@builtin(position)
x_10_1 : vec4<f32>,
x_10_1 : vec4f,
}
@vertex
fn main(@location(1) @interpolate(flat) x_1_param : u32, @location(2) @interpolate(flat) x_2_param : vec2<u32>, @location(3) @interpolate(flat) x_3_param : i32, @location(4) @interpolate(flat) x_4_param : vec2<i32>, @location(5) @interpolate(flat) x_5_param : f32, @location(6) @interpolate(flat) x_6_param : vec2<f32>) -> main_out {
fn main(@location(1) @interpolate(flat) x_1_param : u32, @location(2) @interpolate(flat) x_2_param : vec2u, @location(3) @interpolate(flat) x_3_param : i32, @location(4) @interpolate(flat) x_4_param : vec2i, @location(5) @interpolate(flat) x_5_param : f32, @location(6) @interpolate(flat) x_6_param : vec2f) -> main_out {
x_1 = x_1_param;
x_2 = x_2_param;
x_3 = x_3_param;
@ -4804,17 +4804,17 @@ TEST_F(SpvModuleScopeVarParserTest, EntryPointWrapping_Interpolation_Flat_Vertex
const std::string expected =
R"(var<private> x_1 : u32;
var<private> x_2 : vec2<u32>;
var<private> x_2 : vec2u;
var<private> x_3 : i32;
var<private> x_4 : vec2<i32>;
var<private> x_4 : vec2i;
var<private> x_5 : f32;
var<private> x_6 : vec2<f32>;
var<private> x_6 : vec2f;
var<private> x_10 : vec4<f32>;
var<private> x_10 : vec4f;
fn main_1() {
return;
@ -4824,17 +4824,17 @@ struct main_out {
@location(1) @interpolate(flat)
x_1_1 : u32,
@location(2) @interpolate(flat)
x_2_1 : vec2<u32>,
x_2_1 : vec2u,
@location(3) @interpolate(flat)
x_3_1 : i32,
@location(4) @interpolate(flat)
x_4_1 : vec2<i32>,
x_4_1 : vec2i,
@location(5) @interpolate(flat)
x_5_1 : f32,
@location(6) @interpolate(flat)
x_6_1 : vec2<f32>,
x_6_1 : vec2f,
@builtin(position)
x_10_1 : vec4<f32>,
x_10_1 : vec4f,
}
@vertex
@ -5251,17 +5251,17 @@ TEST_F(SpvModuleScopeVarParserTest, EntryPointWrapping_Interpolation_Default_Ver
const std::string expected =
R"(var<private> x_1 : u32;
var<private> x_2 : vec2<u32>;
var<private> x_2 : vec2u;
var<private> x_3 : i32;
var<private> x_4 : vec2<i32>;
var<private> x_4 : vec2i;
var<private> x_5 : f32;
var<private> x_6 : vec2<f32>;
var<private> x_6 : vec2f;
var<private> x_10 : vec4<f32>;
var<private> x_10 : vec4f;
fn main_1() {
return;
@ -5271,17 +5271,17 @@ struct main_out {
@location(1) @interpolate(flat)
x_1_1 : u32,
@location(2) @interpolate(flat)
x_2_1 : vec2<u32>,
x_2_1 : vec2u,
@location(3) @interpolate(flat)
x_3_1 : i32,
@location(4) @interpolate(flat)
x_4_1 : vec2<i32>,
x_4_1 : vec2i,
@location(5)
x_5_1 : f32,
@location(6)
x_6_1 : vec2<f32>,
x_6_1 : vec2f,
@builtin(position)
x_10_1 : vec4<f32>,
x_10_1 : vec4f,
}
@vertex
@ -5333,22 +5333,22 @@ TEST_F(SpvModuleScopeVarParserTest, EntryPointWrapping_Interpolation_Default_Fra
const std::string expected =
R"(var<private> x_1 : u32;
var<private> x_2 : vec2<u32>;
var<private> x_2 : vec2u;
var<private> x_3 : i32;
var<private> x_4 : vec2<i32>;
var<private> x_4 : vec2i;
var<private> x_5 : f32;
var<private> x_6 : vec2<f32>;
var<private> x_6 : vec2f;
fn main_1() {
return;
}
@fragment
fn main(@location(1) @interpolate(flat) x_1_param : u32, @location(2) @interpolate(flat) x_2_param : vec2<u32>, @location(3) @interpolate(flat) x_3_param : i32, @location(4) @interpolate(flat) x_4_param : vec2<i32>, @location(5) x_5_param : f32, @location(6) x_6_param : vec2<f32>) {
fn main(@location(1) @interpolate(flat) x_1_param : u32, @location(2) @interpolate(flat) x_2_param : vec2u, @location(3) @interpolate(flat) x_3_param : i32, @location(4) @interpolate(flat) x_4_param : vec2i, @location(5) x_5_param : f32, @location(6) x_6_param : vec2f) {
x_1 = x_1_param;
x_2 = x_2_param;
x_3 = x_3_param;

8
src/tint/reader/spirv/parser_type.cc
View File

@ -203,7 +203,13 @@ Vector::Vector(const Type* t, uint32_t s) : type(t), size(s) {}
Vector::Vector(const Vector&) = default;
ast::Type Vector::Build(ProgramBuilder& b) const {
return b.ty.vec(type->Build(b), size);
auto prefix = "vec" + std::to_string(size);
return Switch(
type, //
[&](const I32*) { return b.ty(prefix + "i"); },
[&](const U32*) { return b.ty(prefix + "u"); },
[&](const F32*) { return b.ty(prefix + "f"); },
[&](Default) { return b.ty.vec(type->Build(b), size); });
}
Matrix::Matrix(const Type* t, uint32_t c, uint32_t r) : type(t), columns(c), rows(r) {}

2
test/tint/access/let/matrix.spvasm.expected.wgsl
View File

@ -1,5 +1,5 @@
fn main_1() {
let x_24 : f32 = mat3x3<f32>(vec3<f32>(1.0f, 2.0f, 3.0f), vec3<f32>(4.0f, 5.0f, 6.0f), vec3<f32>(7.0f, 8.0f, 9.0f))[1u].y;
let x_24 : f32 = mat3x3<f32>(vec3f(1.0f, 2.0f, 3.0f), vec3f(4.0f, 5.0f, 6.0f), vec3f(7.0f, 8.0f, 9.0f))[1u].y;
return;
}

6
test/tint/access/let/vector.spvasm.expected.wgsl
View File

@ -1,9 +1,9 @@
const x_10 = vec3<f32>(1.0f, 2.0f, 3.0f);
const x_10 = vec3f(1.0f, 2.0f, 3.0f);
fn main_1() {
let x_11 : f32 = x_10.y;
let x_13 : vec2<f32> = vec2<f32>(x_10.x, x_10.z);
let x_14 : vec3<f32> = vec3<f32>(x_10.x, x_10.z, x_10.y);
let x_13 : vec2f = vec2f(x_10.x, x_10.z);
let x_14 : vec3f = vec3f(x_10.x, x_10.z, x_10.y);
return;
}

2
test/tint/access/var/matrix.spvasm.expected.wgsl
View File

@ -1,6 +1,6 @@
fn main_1() {
var m : mat3x3<f32> = mat3x3<f32>();
let x_15 : vec3<f32> = m[1i];
let x_15 : vec3f = m[1i];
let x_16 : f32 = x_15.y;
return;
}

10
test/tint/access/var/vector.spvasm.expected.wgsl
View File

@ -1,10 +1,10 @@
fn main_1() {
var v : vec3<f32> = vec3<f32>();
var v : vec3f = vec3f();
let x_14 : f32 = v.y;
let x_16 : vec3<f32> = v;
let x_17 : vec2<f32> = vec2<f32>(x_16.x, x_16.z);
let x_18 : vec3<f32> = v;
let x_19 : vec3<f32> = vec3<f32>(x_18.x, x_18.z, x_18.y);
let x_16 : vec3f = v;
let x_17 : vec2f = vec2f(x_16.x, x_16.z);
let x_18 : vec3f = v;
let x_19 : vec3f = vec3f(x_18.x, x_18.z, x_18.y);
return;
}

18
test/tint/bug/oss-fuzz/57795.spv.expected.wgsl
View File

@ -7,7 +7,7 @@ struct S_1 {
}
struct S_2 {
field0 : array<array<vec3<f32>, 37u>, 95u>,
field0 : array<array<vec3f, 37u>, 95u>,
}
struct S_3 {
@ -15,7 +15,7 @@ struct S_3 {
}
struct S_4 {
field0 : array<vec2<i32>, 56u>,
field0 : array<vec2i, 56u>,
}
struct S_5 {
@ -23,28 +23,28 @@ struct S_5 {
}
struct S_6 {
field0 : array<array<vec3<f32>, 18u>, 13u>,
field0 : array<array<vec3f, 18u>, 13u>,
}
struct S_7 {
field0 : array<vec2<i32>, 88u>,
field0 : array<vec2i, 88u>,
}
const x_72 = vec4<f32>(0.0f, 0.0f, 0.0f, 0.0f);
const x_72 = vec4f(0.0f, 0.0f, 0.0f, 0.0f);
const x_73 = mat4x4<f32>(x_72, x_72, x_72, x_72);
var<private> x_75 : array<mat4x4<f32>, 58u> = array<mat4x4<f32>, 58u>(x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73, x_73);
const x_77 = vec3<f32>(0.0f, 0.0f, 0.0f);
const x_77 = vec3f(0.0f, 0.0f, 0.0f);
const x_78 = array<vec3<f32>, 18u>(x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77);
const x_78 = array<vec3f, 18u>(x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77);
const x_80 = S_6(array<array<vec3<f32>, 18u>, 13u>(x_78, x_78, x_78, x_78, x_78, x_78, x_78, x_78, x_78, x_78, x_78, x_78, x_78));
const x_80 = S_6(array<array<vec3f, 18u>, 13u>(x_78, x_78, x_78, x_78, x_78, x_78, x_78, x_78, x_78, x_78, x_78, x_78, x_78));
var<private> x_82 : array<S_6, 46u> = array<S_6, 46u>(x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80, x_80);
var<private> x_85 : array<vec3<f32>, 37u> = array<vec3<f32>, 37u>(x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77);
var<private> x_85 : array<vec3f, 37u> = array<vec3f, 37u>(x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77, x_77);
fn main_1() {
let x_88 : u32 = 58u;

20
test/tint/bug/tint/1061.spvasm.expected.wgsl
View File

@ -1,34 +1,34 @@
struct buf0 {
/* @offset(0) */
r : vec4<f32>,
r : vec4f,
}
@group(0) @binding(0) var<uniform> x_7 : buf0;
var<private> x_GLF_color : vec4<f32>;
var<private> x_GLF_color : vec4f;
fn main_1() {
var f : f32;
var v : vec4<f32>;
f = determinant(mat3x3<f32>(vec3<f32>(1.0f, 0.0f, 0.0f), vec3<f32>(0.0f, 1.0f, 0.0f), vec3<f32>(0.0f, 0.0f, 1.0f)));
var v : vec4f;
f = determinant(mat3x3<f32>(vec3f(1.0f, 0.0f, 0.0f), vec3f(0.0f, 1.0f, 0.0f), vec3f(0.0f, 0.0f, 1.0f)));
let x_33 : f32 = f;
let x_35 : f32 = f;
let x_37 : f32 = f;
let x_39 : f32 = f;
v = vec4<f32>(sin(x_33), cos(x_35), exp2(x_37), log(x_39));
let x_42 : vec4<f32> = v;
let x_44 : vec4<f32> = x_7.r;
v = vec4f(sin(x_33), cos(x_35), exp2(x_37), log(x_39));
let x_42 : vec4f = v;
let x_44 : vec4f = x_7.r;
if ((distance(x_42, x_44) < 0.10000000149011611938f)) {
x_GLF_color = vec4<f32>(1.0f, 0.0f, 0.0f, 1.0f);
x_GLF_color = vec4f(1.0f, 0.0f, 0.0f, 1.0f);
} else {
x_GLF_color = vec4<f32>(0.0f, 0.0f, 0.0f, 0.0f);
x_GLF_color = vec4f(0.0f, 0.0f, 0.0f, 0.0f);
}
return;
}
struct main_out {
@location(0)
x_GLF_color_1 : vec4<f32>,
x_GLF_color_1 : vec4f,
}
@fragment

34
test/tint/bug/tint/1088.spvasm.expected.wgsl
View File

@ -18,25 +18,25 @@ struct LeftOver {
test : Arr_1,
}
var<private> position_1 : vec3<f32>;
var<private> position_1 : vec3f;
@group(2) @binding(2) var<uniform> x_14 : LeftOver;
var<private> vUV : vec2<f32>;
var<private> vUV : vec2f;
var<private> uv : vec2<f32>;
var<private> uv : vec2f;
var<private> normal : vec3<f32>;
var<private> normal : vec3f;
var<private> gl_Position : vec4<f32>;
var<private> gl_Position : vec4f;
fn main_1() {
var q : vec4<f32>;
var p : vec3<f32>;
let x_13 : vec3<f32> = position_1;
q = vec4<f32>(x_13.x, x_13.y, x_13.z, 1.0f);
let x_21 : vec4<f32> = q;
p = vec3<f32>(x_21.x, x_21.y, x_21.z);
var q : vec4f;
var p : vec3f;
let x_13 : vec3f = position_1;
q = vec4f(x_13.x, x_13.y, x_13.z, 1.0f);
let x_21 : vec4f = q;
p = vec3f(x_21.x, x_21.y, x_21.z);
let x_27 : f32 = p.x;
let x_41 : f32 = x_14.test[0i].el;
let x_45 : f32 = position_1.y;
@ -46,9 +46,9 @@ fn main_1() {
let x_57 : f32 = x_14.time;
p.y = (x_55 + sin((x_57 + 4.0f)));
let x_69 : mat4x4<f32> = x_14.worldViewProjection;
let x_70 : vec3<f32> = p;
gl_Position = (x_69 * vec4<f32>(x_70.x, x_70.y, x_70.z, 1.0f));
let x_83 : vec2<f32> = uv;
let x_70 : vec3f = p;
gl_Position = (x_69 * vec4f(x_70.x, x_70.y, x_70.z, 1.0f));
let x_83 : vec2f = uv;
vUV = x_83;
let x_87 : f32 = gl_Position.y;
gl_Position.y = (x_87 * -1.0f);
@ -57,13 +57,13 @@ fn main_1() {
struct main_out {
@builtin(position)
gl_Position : vec4<f32>,
gl_Position : vec4f,
@location(0)
vUV_1 : vec2<f32>,
vUV_1 : vec2f,
}
@vertex
fn main(@location(0) position_1_param : vec3<f32>, @location(2) uv_param : vec2<f32>, @location(1) normal_param : vec3<f32>) -> main_out {
fn main(@location(0) position_1_param : vec3f, @location(2) uv_param : vec2f, @location(1) normal_param : vec3f) -> main_out {
position_1 = position_1_param;
uv = uv_param;
normal = normal_param;

82
test/tint/bug/tint/1520.spvasm.expected.wgsl
View File

@ -4,31 +4,31 @@ struct UniformBuffer {
/* @offset(16) */
unknownInput_S1_c0 : f32,
/* @offset(32) */
ucolorRed_S1_c0 : vec4<f32>,
ucolorRed_S1_c0 : vec4f,
/* @offset(48) */
ucolorGreen_S1_c0 : vec4<f32>,
ucolorGreen_S1_c0 : vec4f,
/* @offset(64) */
umatrix_S1 : mat3x3<f32>,
}
@binding(0) @group(0) var<uniform> x_4 : UniformBuffer;
var<private> sk_FragColor : vec4<f32>;
var<private> sk_FragColor : vec4f;
var<private> sk_Clockwise : bool;
var<private> vcolor_S0 : vec4<f32>;
var<private> vcolor_S0 : vec4f;
const x_35 = vec4<i32>(0i, 0i, 0i, 0i);
const x_35 = vec4i(0i, 0i, 0i, 0i);
const x_46 = vec4<i32>(1i, 1i, 1i, 1i);
const x_46 = vec4i(1i, 1i, 1i, 1i);
const x_57 = vec4<i32>(2i, 2i, 2i, 2i);
const x_57 = vec4i(2i, 2i, 2i, 2i);
fn test_int_S1_c0_b() -> bool {
var unknown : i32;
var ok : bool;
var val : vec4<i32>;
var val : vec4i;
var x_40 : bool;
var x_41 : bool;
var x_54 : bool;
@ -41,19 +41,19 @@ fn test_int_S1_c0_b() -> bool {
ok = true;
x_41 = false;
if (true) {
x_40 = all(((x_35 / vec4<i32>(x_27, x_27, x_27, x_27)) == x_35));
x_40 = all(((x_35 / vec4i(x_27, x_27, x_27, x_27)) == x_35));
x_41 = x_40;
}
ok = x_41;
let x_44 : vec4<i32> = vec4<i32>(x_27, x_27, x_27, x_27);
let x_44 : vec4i = vec4i(x_27, x_27, x_27, x_27);
val = x_44;
let x_47 : vec4<i32> = (x_44 + x_46);
let x_47 : vec4i = (x_44 + x_46);
val = x_47;
let x_48 : vec4<i32> = (x_47 - x_46);
let x_48 : vec4i = (x_47 - x_46);
val = x_48;
let x_49 : vec4<i32> = (x_48 + x_46);
let x_49 : vec4i = (x_48 + x_46);
val = x_49;
let x_50 : vec4<i32> = (x_49 - x_46);
let x_50 : vec4i = (x_49 - x_46);
val = x_50;
x_55 = false;
if (x_41) {
@ -61,13 +61,13 @@ fn test_int_S1_c0_b() -> bool {
x_55 = x_54;
}
ok = x_55;
let x_58 : vec4<i32> = (x_50 * x_57);
let x_58 : vec4i = (x_50 * x_57);
val = x_58;
let x_59 : vec4<i32> = (x_58 / x_57);
let x_59 : vec4i = (x_58 / x_57);
val = x_59;
let x_60 : vec4<i32> = (x_59 * x_57);
let x_60 : vec4i = (x_59 * x_57);
val = x_60;
let x_61 : vec4<i32> = (x_60 / x_57);
let x_61 : vec4i = (x_60 / x_57);
val = x_61;
x_66 = false;
if (x_55) {
@ -78,19 +78,19 @@ fn test_int_S1_c0_b() -> bool {
return x_66;
}
const x_82 = vec4<f32>(0.0f, 0.0f, 0.0f, 0.0f);
const x_82 = vec4f(0.0f, 0.0f, 0.0f, 0.0f);
const x_91 = vec4<f32>(1.0f, 1.0f, 1.0f, 1.0f);
const x_91 = vec4f(1.0f, 1.0f, 1.0f, 1.0f);
const x_102 = vec4<f32>(2.0f, 2.0f, 2.0f, 2.0f);
const x_102 = vec4f(2.0f, 2.0f, 2.0f, 2.0f);
fn main_1() {
var outputColor_S0 : vec4<f32>;
var output_S1 : vec4<f32>;
var outputColor_S0 : vec4f;
var output_S1 : vec4f;
var x_8_unknown : f32;
var x_9_ok : bool;
var x_10_val : vec4<f32>;
var x_116 : vec4<f32>;
var x_10_val : vec4f;
var x_116 : vec4f;
var x_86 : bool;
var x_87 : bool;
var x_99 : bool;
@ -99,26 +99,26 @@ fn main_1() {
var x_111 : bool;
var x_114 : bool;
var x_115 : bool;
let x_72 : vec4<f32> = vcolor_S0;
let x_72 : vec4f = vcolor_S0;
outputColor_S0 = x_72;
let x_77 : f32 = x_4.unknownInput_S1_c0;
x_8_unknown = x_77;
x_9_ok = true;
x_87 = false;
if (true) {
x_86 = all(((x_82 / vec4<f32>(x_77, x_77, x_77, x_77)) == x_82));
x_86 = all(((x_82 / vec4f(x_77, x_77, x_77, x_77)) == x_82));
x_87 = x_86;
}
x_9_ok = x_87;
let x_89 : vec4<f32> = vec4<f32>(x_77, x_77, x_77, x_77);
let x_89 : vec4f = vec4f(x_77, x_77, x_77, x_77);
x_10_val = x_89;
let x_92 : vec4<f32> = (x_89 + x_91);
let x_92 : vec4f = (x_89 + x_91);
x_10_val = x_92;
let x_93 : vec4<f32> = (x_92 - x_91);
let x_93 : vec4f = (x_92 - x_91);
x_10_val = x_93;
let x_94 : vec4<f32> = (x_93 + x_91);
let x_94 : vec4f = (x_93 + x_91);
x_10_val = x_94;
let x_95 : vec4<f32> = (x_94 - x_91);
let x_95 : vec4f = (x_94 - x_91);
x_10_val = x_95;
x_100 = false;
if (x_87) {
@ -126,13 +126,13 @@ fn main_1() {
x_100 = x_99;
}
x_9_ok = x_100;
let x_103 : vec4<f32> = (x_95 * x_102);
let x_103 : vec4f = (x_95 * x_102);
x_10_val = x_103;
let x_104 : vec4<f32> = (x_103 / x_102);
let x_104 : vec4f = (x_103 / x_102);
x_10_val = x_104;
let x_105 : vec4<f32> = (x_104 * x_102);
let x_105 : vec4f = (x_104 * x_102);
x_10_val = x_105;
let x_106 : vec4<f32> = (x_105 / x_102);
let x_106 : vec4f = (x_105 / x_102);
x_10_val = x_106;
x_111 = false;
if (x_100) {
@ -146,13 +146,13 @@ fn main_1() {
x_115 = x_114;
}
if (x_115) {
let x_122 : vec4<f32> = x_4.ucolorGreen_S1_c0;
let x_122 : vec4f = x_4.ucolorGreen_S1_c0;
x_116 = x_122;
} else {
let x_124 : vec4<f32> = x_4.ucolorRed_S1_c0;
let x_124 : vec4f = x_4.ucolorRed_S1_c0;
x_116 = x_124;
}
let x_125 : vec4<f32> = x_116;
let x_125 : vec4f = x_116;
output_S1 = x_125;
sk_FragColor = x_125;
return;
@ -160,11 +160,11 @@ fn main_1() {
struct main_out {
@location(0)
sk_FragColor_1 : vec4<f32>,
sk_FragColor_1 : vec4f,
}
@fragment
fn main(@builtin(front_facing) sk_Clockwise_param : bool, @location(0) vcolor_S0_param : vec4<f32>) -> main_out {
fn main(@builtin(front_facing) sk_Clockwise_param : bool, @location(0) vcolor_S0_param : vec4f) -> main_out {
sk_Clockwise = sk_Clockwise_param;
vcolor_S0 = vcolor_S0_param;
main_1();

6
test/tint/bug/tint/1670.spvasm.expected.wgsl
View File

@ -7,13 +7,13 @@ fn main_1() {
var m4 : mat4x4<f32>;
let x_12 : mat2x2<f32> = m2;
let s = (1.0f / determinant(x_12));
m2i = mat2x2<f32>(vec2<f32>((s * x_12[1u][1u]), (-(s) * x_12[0u][1u])), vec2<f32>((-(s) * x_12[1u][0u]), (s * x_12[0u][0u])));
m2i = mat2x2<f32>(vec2f((s * x_12[1u][1u]), (-(s) * x_12[0u][1u])), vec2f((-(s) * x_12[1u][0u]), (s * x_12[0u][0u])));
let x_19 : mat3x3<f32> = m3;
let s_1 = (1.0f / determinant(x_19));
m3i = (s_1 * mat3x3<f32>(vec3<f32>(((x_19[1u][1u] * x_19[2u][2u]) - (x_19[1u][2u] * x_19[2u][1u])), ((x_19[0u][2u] * x_19[2u][1u]) - (x_19[0u][1u] * x_19[2u][2u])), ((x_19[0u][1u] * x_19[1u][2u]) - (x_19[0u][2u] * x_19[1u][1u]))), vec3<f32>(((x_19[1u][2u] * x_19[2u][0u]) - (x_19[1u][0u] * x_19[2u][2u])), ((x_19[0u][0u] * x_19[2u][2u]) - (x_19[0u][2u] * x_19[2u][0u])), ((x_19[0u][2u] * x_19[1u][0u]) - (x_19[0u][0u] * x_19[1u][2u]))), vec3<f32>(((x_19[1u][0u] * x_19[2u][1u]) - (x_19[1u][1u] * x_19[2u][0u])), ((x_19[0u][1u] * x_19[2u][0u]) - (x_19[0u][0u] * x_19[2u][1u])), ((x_19[0u][0u] * x_19[1u][1u]) - (x_19[0u][1u] * x_19[1u][0u])))));
m3i = (s_1 * mat3x3<f32>(vec3f(((x_19[1u][1u] * x_19[2u][2u]) - (x_19[1u][2u] * x_19[2u][1u])), ((x_19[0u][2u] * x_19[2u][1u]) - (x_19[0u][1u] * x_19[2u][2u])), ((x_19[0u][1u] * x_19[1u][2u]) - (x_19[0u][2u] * x_19[1u][1u]))), vec3f(((x_19[1u][2u] * x_19[2u][0u]) - (x_19[1u][0u] * x_19[2u][2u])), ((x_19[0u][0u] * x_19[2u][2u]) - (x_19[0u][2u] * x_19[2u][0u])), ((x_19[0u][2u] * x_19[1u][0u]) - (x_19[0u][0u] * x_19[1u][2u]))), vec3f(((x_19[1u][0u] * x_19[2u][1u]) - (x_19[1u][1u] * x_19[2u][0u])), ((x_19[0u][1u] * x_19[2u][0u]) - (x_19[0u][0u] * x_19[2u][1u])), ((x_19[0u][0u] * x_19[1u][1u]) - (x_19[0u][1u] * x_19[1u][0u])))));
let x_26 : mat4x4<f32> = m4;
let s_2 = (1.0f / determinant(x_26));
m4i = (s_2 * mat4x4<f32>(vec4<f32>((((x_26[1u][1u] * ((x_26[2u][2u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][2u]))) - (x_26[1u][2u] * ((x_26[2u][1u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][1u])))) + (x_26[1u][3u] * ((x_26[2u][1u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][1u])))), (((-(x_26[0u][1u]) * ((x_26[2u][2u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][2u]))) + (x_26[0u][2u] * ((x_26[2u][1u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][1u])))) - (x_26[0u][3u] * ((x_26[2u][1u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][1u])))), (((x_26[0u][1u] * ((x_26[1u][2u] * x_26[3u][3u]) - (x_26[1u][3u] * x_26[3u][2u]))) - (x_26[0u][2u] * ((x_26[1u][1u] * x_26[3u][3u]) - (x_26[1u][3u] * x_26[3u][1u])))) + (x_26[0u][3u] * ((x_26[1u][1u] * x_26[3u][2u]) - (x_26[1u][2u] * x_26[3u][1u])))), (((-(x_26[0u][1u]) * ((x_26[1u][2u] * x_26[2u][3u]) - (x_26[1u][3u] * x_26[2u][2u]))) + (x_26[0u][2u] * ((x_26[1u][1u] * x_26[2u][3u]) - (x_26[1u][3u] * x_26[2u][1u])))) - (x_26[0u][3u] * ((x_26[1u][1u] * x_26[2u][2u]) - (x_26[1u][2u] * x_26[2u][1u]))))), vec4<f32>((((-(x_26[1u][0u]) * ((x_26[2u][2u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][2u]))) + (x_26[1u][2u] * ((x_26[2u][0u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][0u])))) - (x_26[1u][3u] * ((x_26[2u][0u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][0u])))), (((x_26[0u][0u] * ((x_26[2u][2u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][2u]))) - (x_26[0u][2u] * ((x_26[2u][0u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][0u])))) + (x_26[0u][3u] * ((x_26[2u][0u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][0u])))), (((-(x_26[0u][0u]) * ((x_26[1u][2u] * x_26[3u][3u]) - (x_26[1u][3u] * x_26[3u][2u]))) + (x_26[0u][2u] * ((x_26[1u][0u] * x_26[3u][3u]) - (x_26[1u][3u] * x_26[3u][0u])))) - (x_26[0u][3u] * ((x_26[1u][0u] * x_26[3u][2u]) - (x_26[1u][2u] * x_26[3u][0u])))), (((x_26[0u][0u] * ((x_26[1u][2u] * x_26[2u][3u]) - (x_26[1u][3u] * x_26[2u][2u]))) - (x_26[0u][2u] * ((x_26[1u][0u] * x_26[2u][3u]) - (x_26[1u][3u] * x_26[2u][0u])))) + (x_26[0u][3u] * ((x_26[1u][0u] * x_26[2u][2u]) - (x_26[1u][2u] * x_26[2u][0u]))))), vec4<f32>((((x_26[1u][0u] * ((x_26[2u][1u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][1u]))) - (x_26[1u][1u] * ((x_26[2u][0u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][0u])))) + (x_26[1u][3u] * ((x_26[2u][0u] * x_26[3u][1u]) - (x_26[2u][1u] * x_26[3u][0u])))), (((-(x_26[0u][0u]) * ((x_26[2u][1u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][1u]))) + (x_26[0u][1u] * ((x_26[2u][0u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][0u])))) - (x_26[0u][3u] * ((x_26[2u][0u] * x_26[3u][1u]) - (x_26[2u][1u] * x_26[3u][0u])))), (((x_26[0u][0u] * ((x_26[1u][1u] * x_26[3u][3u]) - (x_26[1u][3u] * x_26[3u][1u]))) - (x_26[0u][1u] * ((x_26[1u][0u] * x_26[3u][3u]) - (x_26[1u][3u] * x_26[3u][0u])))) + (x_26[0u][3u] * ((x_26[1u][0u] * x_26[3u][1u]) - (x_26[1u][1u] * x_26[3u][0u])))), (((-(x_26[0u][0u]) * ((x_26[1u][1u] * x_26[2u][3u]) - (x_26[1u][3u] * x_26[2u][1u]))) + (x_26[0u][1u] * ((x_26[1u][0u] * x_26[2u][3u]) - (x_26[1u][3u] * x_26[2u][0u])))) - (x_26[0u][3u] * ((x_26[1u][0u] * x_26[2u][1u]) - (x_26[1u][1u] * x_26[2u][0u]))))), vec4<f32>((((-(x_26[1u][0u]) * ((x_26[2u][1u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][1u]))) + (x_26[1u][1u] * ((x_26[2u][0u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][0u])))) - (x_26[1u][2u] * ((x_26[2u][0u] * x_26[3u][1u]) - (x_26[2u][1u] * x_26[3u][0u])))), (((x_26[0u][0u] * ((x_26[2u][1u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][1u]))) - (x_26[0u][1u] * ((x_26[2u][0u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][0u])))) + (x_26[0u][2u] * ((x_26[2u][0u] * x_26[3u][1u]) - (x_26[2u][1u] * x_26[3u][0u])))), (((-(x_26[0u][0u]) * ((x_26[1u][1u] * x_26[3u][2u]) - (x_26[1u][2u] * x_26[3u][1u]))) + (x_26[0u][1u] * ((x_26[1u][0u] * x_26[3u][2u]) - (x_26[1u][2u] * x_26[3u][0u])))) - (x_26[0u][2u] * ((x_26[1u][0u] * x_26[3u][1u]) - (x_26[1u][1u] * x_26[3u][0u])))), (((x_26[0u][0u] * ((x_26[1u][1u] * x_26[2u][2u]) - (x_26[1u][2u] * x_26[2u][1u]))) - (x_26[0u][1u] * ((x_26[1u][0u] * x_26[2u][2u]) - (x_26[1u][2u] * x_26[2u][0u])))) + (x_26[0u][2u] * ((x_26[1u][0u] * x_26[2u][1u]) - (x_26[1u][1u] * x_26[2u][0u])))))));
m4i = (s_2 * mat4x4<f32>(vec4f((((x_26[1u][1u] * ((x_26[2u][2u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][2u]))) - (x_26[1u][2u] * ((x_26[2u][1u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][1u])))) + (x_26[1u][3u] * ((x_26[2u][1u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][1u])))), (((-(x_26[0u][1u]) * ((x_26[2u][2u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][2u]))) + (x_26[0u][2u] * ((x_26[2u][1u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][1u])))) - (x_26[0u][3u] * ((x_26[2u][1u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][1u])))), (((x_26[0u][1u] * ((x_26[1u][2u] * x_26[3u][3u]) - (x_26[1u][3u] * x_26[3u][2u]))) - (x_26[0u][2u] * ((x_26[1u][1u] * x_26[3u][3u]) - (x_26[1u][3u] * x_26[3u][1u])))) + (x_26[0u][3u] * ((x_26[1u][1u] * x_26[3u][2u]) - (x_26[1u][2u] * x_26[3u][1u])))), (((-(x_26[0u][1u]) * ((x_26[1u][2u] * x_26[2u][3u]) - (x_26[1u][3u] * x_26[2u][2u]))) + (x_26[0u][2u] * ((x_26[1u][1u] * x_26[2u][3u]) - (x_26[1u][3u] * x_26[2u][1u])))) - (x_26[0u][3u] * ((x_26[1u][1u] * x_26[2u][2u]) - (x_26[1u][2u] * x_26[2u][1u]))))), vec4f((((-(x_26[1u][0u]) * ((x_26[2u][2u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][2u]))) + (x_26[1u][2u] * ((x_26[2u][0u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][0u])))) - (x_26[1u][3u] * ((x_26[2u][0u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][0u])))), (((x_26[0u][0u] * ((x_26[2u][2u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][2u]))) - (x_26[0u][2u] * ((x_26[2u][0u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][0u])))) + (x_26[0u][3u] * ((x_26[2u][0u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][0u])))), (((-(x_26[0u][0u]) * ((x_26[1u][2u] * x_26[3u][3u]) - (x_26[1u][3u] * x_26[3u][2u]))) + (x_26[0u][2u] * ((x_26[1u][0u] * x_26[3u][3u]) - (x_26[1u][3u] * x_26[3u][0u])))) - (x_26[0u][3u] * ((x_26[1u][0u] * x_26[3u][2u]) - (x_26[1u][2u] * x_26[3u][0u])))), (((x_26[0u][0u] * ((x_26[1u][2u] * x_26[2u][3u]) - (x_26[1u][3u] * x_26[2u][2u]))) - (x_26[0u][2u] * ((x_26[1u][0u] * x_26[2u][3u]) - (x_26[1u][3u] * x_26[2u][0u])))) + (x_26[0u][3u] * ((x_26[1u][0u] * x_26[2u][2u]) - (x_26[1u][2u] * x_26[2u][0u]))))), vec4f((((x_26[1u][0u] * ((x_26[2u][1u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][1u]))) - (x_26[1u][1u] * ((x_26[2u][0u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][0u])))) + (x_26[1u][3u] * ((x_26[2u][0u] * x_26[3u][1u]) - (x_26[2u][1u] * x_26[3u][0u])))), (((-(x_26[0u][0u]) * ((x_26[2u][1u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][1u]))) + (x_26[0u][1u] * ((x_26[2u][0u] * x_26[3u][3u]) - (x_26[2u][3u] * x_26[3u][0u])))) - (x_26[0u][3u] * ((x_26[2u][0u] * x_26[3u][1u]) - (x_26[2u][1u] * x_26[3u][0u])))), (((x_26[0u][0u] * ((x_26[1u][1u] * x_26[3u][3u]) - (x_26[1u][3u] * x_26[3u][1u]))) - (x_26[0u][1u] * ((x_26[1u][0u] * x_26[3u][3u]) - (x_26[1u][3u] * x_26[3u][0u])))) + (x_26[0u][3u] * ((x_26[1u][0u] * x_26[3u][1u]) - (x_26[1u][1u] * x_26[3u][0u])))), (((-(x_26[0u][0u]) * ((x_26[1u][1u] * x_26[2u][3u]) - (x_26[1u][3u] * x_26[2u][1u]))) + (x_26[0u][1u] * ((x_26[1u][0u] * x_26[2u][3u]) - (x_26[1u][3u] * x_26[2u][0u])))) - (x_26[0u][3u] * ((x_26[1u][0u] * x_26[2u][1u]) - (x_26[1u][1u] * x_26[2u][0u]))))), vec4f((((-(x_26[1u][0u]) * ((x_26[2u][1u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][1u]))) + (x_26[1u][1u] * ((x_26[2u][0u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][0u])))) - (x_26[1u][2u] * ((x_26[2u][0u] * x_26[3u][1u]) - (x_26[2u][1u] * x_26[3u][0u])))), (((x_26[0u][0u] * ((x_26[2u][1u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][1u]))) - (x_26[0u][1u] * ((x_26[2u][0u] * x_26[3u][2u]) - (x_26[2u][2u] * x_26[3u][0u])))) + (x_26[0u][2u] * ((x_26[2u][0u] * x_26[3u][1u]) - (x_26[2u][1u] * x_26[3u][0u])))), (((-(x_26[0u][0u]) * ((x_26[1u][1u] * x_26[3u][2u]) - (x_26[1u][2u] * x_26[3u][1u]))) + (x_26[0u][1u] * ((x_26[1u][0u] * x_26[3u][2u]) - (x_26[1u][2u] * x_26[3u][0u])))) - (x_26[0u][2u] * ((x_26[1u][0u] * x_26[3u][1u]) - (x_26[1u][1u] * x_26[3u][0u])))), (((x_26[0u][0u] * ((x_26[1u][1u] * x_26[2u][2u]) - (x_26[1u][2u] * x_26[2u][1u]))) - (x_26[0u][1u] * ((x_26[1u][0u] * x_26[2u][2u]) - (x_26[1u][2u] * x_26[2u][0u])))) + (x_26[0u][2u] * ((x_26[1u][0u] * x_26[2u][1u]) - (x_26[1u][1u] * x_26[2u][0u])))))));
return;
}

2
test/tint/bug/tint/1776.spvasm.expected.wgsl
View File

@ -1,6 +1,6 @@
struct S {
/* @offset(0) */
a : vec4<f32>,
a : vec4f,
/* @offset(16) */
b : i32,
}

4
test/tint/bug/tint/219.spvasm.expected.wgsl
View File

@ -1,10 +1,10 @@
fn x_200(x_201 : ptr<function, vec2<f32>>) -> f32 {
fn x_200(x_201 : ptr<function, vec2f>) -> f32 {
let x_212 : f32 = (*(x_201)).x;
return x_212;
}
fn main_1() {
var x_11 : vec2<f32>;
var x_11 : vec2f;
let x_12 : f32 = x_200(&(x_11));
return;
}

8
test/tint/bug/tint/413.spvasm.expected.wgsl
View File

@ -2,15 +2,15 @@
@group(0) @binding(1) var Dst : texture_storage_2d<r32uint, write>;
const x_17 = vec2<i32>(0i, 0i);
const x_17 = vec2i(0i, 0i);
fn main_1() {
var srcValue : vec4<u32>;
let x_18 : vec4<u32> = textureLoad(Src, x_17, 0i);
var srcValue : vec4u;
let x_18 : vec4u = textureLoad(Src, x_17, 0i);
srcValue = x_18;
let x_22 : u32 = srcValue.x;
srcValue.x = (x_22 + bitcast<u32>(1i));
let x_27 : vec4<u32> = srcValue;
let x_27 : vec4u = srcValue;
textureStore(Dst, x_17, x_27);
return;
}

4
test/tint/bug/tint/870.spvasm.expected.wgsl
View File

@ -2,9 +2,9 @@ alias Arr = array<i32, 6u>;
struct sspp962805860buildInformationS {
/* @offset(0) */
footprint : vec4<f32>,
footprint : vec4f,
/* @offset(16) */
offset : vec4<f32>,
offset : vec4f,
/* @offset(32) */
essence : i32,
/* @offset(36) */

4
test/tint/bug/tint/977.spvasm.expected.wgsl
View File

@ -28,7 +28,7 @@ struct Uniforms {
sizeB : i32,
}
var<private> gl_GlobalInvocationID : vec3<u32>;
var<private> gl_GlobalInvocationID : vec3u;
@group(0) @binding(2) var<storage, read_write> resultMatrix : ResultMatrix;
@ -76,7 +76,7 @@ fn main_1() {
}
@compute @workgroup_size(1i, 1i, 1i)
fn main(@builtin(global_invocation_id) gl_GlobalInvocationID_param : vec3<u32>) {
fn main(@builtin(global_invocation_id) gl_GlobalInvocationID_param : vec3u) {
gl_GlobalInvocationID = gl_GlobalInvocationID_param;
main_1();
}

6
test/tint/builtins/extractBits/vec3/i32.spvasm.expected.wgsl
View File

@ -1,11 +1,11 @@
fn f_1() {
var v : vec3<i32> = vec3<i32>();
var v : vec3i = vec3i();
var offset_1 : u32 = 0u;
var count : u32 = 0u;
let x_17 : vec3<i32> = v;
let x_17 : vec3i = v;
let x_18 : u32 = offset_1;
let x_19 : u32 = count;
let x_15 : vec3<i32> = extractBits(x_17, x_18, x_19);
let x_15 : vec3i = extractBits(x_17, x_18, x_19);
return;
}

6
test/tint/builtins/extractBits/vec3/u32.spvasm.expected.wgsl
View File

@ -1,11 +1,11 @@
fn f_1() {
var v : vec3<u32> = vec3<u32>();
var v : vec3u = vec3u();
var offset_1 : u32 = 0u;
var count : u32 = 0u;
let x_16 : vec3<u32> = v;
let x_16 : vec3u = v;
let x_17 : u32 = offset_1;
let x_18 : u32 = count;
let x_14 : vec3<u32> = extractBits(x_16, x_17, x_18);
let x_14 : vec3u = extractBits(x_16, x_17, x_18);
return;
}

10
test/tint/builtins/insertBits/vec3/i32.spvasm.expected.wgsl
View File

@ -1,13 +1,13 @@
fn f_1() {
var v : vec3<i32> = vec3<i32>();
var n : vec3<i32> = vec3<i32>();
var v : vec3i = vec3i();
var n : vec3i = vec3i();
var offset_1 : u32 = 0u;
var count : u32 = 0u;
let x_18 : vec3<i32> = v;
let x_19 : vec3<i32> = n;
let x_18 : vec3i = v;
let x_19 : vec3i = n;
let x_20 : u32 = offset_1;
let x_21 : u32 = count;
let x_16 : vec3<i32> = insertBits(x_18, x_19, x_20, x_21);
let x_16 : vec3i = insertBits(x_18, x_19, x_20, x_21);
return;
}

10
test/tint/builtins/insertBits/vec3/u32.spvasm.expected.wgsl
View File

@ -1,13 +1,13 @@
fn f_1() {
var v : vec3<u32> = vec3<u32>();
var n : vec3<u32> = vec3<u32>();
var v : vec3u = vec3u();
var n : vec3u = vec3u();
var offset_1 : u32 = 0u;
var count : u32 = 0u;
let x_17 : vec3<u32> = v;
let x_18 : vec3<u32> = n;
let x_17 : vec3u = v;
let x_18 : vec3u = n;
let x_19 : u32 = offset_1;
let x_20 : u32 = count;
let x_15 : vec3<u32> = insertBits(x_17, x_18, x_19, x_20);
let x_15 : vec3u = insertBits(x_17, x_18, x_19, x_20);
return;
}

12
test/tint/builtins/textureDimensions/depth_ms.spvasm.expected.wgsl
View File

@ -1,28 +1,28 @@
@group(1) @binding(0) var arg_0 : texture_depth_multisampled_2d;
var<private> tint_symbol_1 : vec4<f32> = vec4<f32>();
var<private> tint_symbol_1 : vec4f = vec4f();
fn textureDimensions_f60bdb() {
var res : vec2<i32> = vec2<i32>();
let x_16 : vec2<i32> = vec2<i32>(textureDimensions(arg_0));
var res : vec2i = vec2i();
let x_16 : vec2i = vec2i(textureDimensions(arg_0));
res = x_16;
return;
}
fn tint_symbol_2(tint_symbol : vec4<f32>) {
fn tint_symbol_2(tint_symbol : vec4f) {
tint_symbol_1 = tint_symbol;
return;
}
fn vertex_main_1() {
textureDimensions_f60bdb();
tint_symbol_2(vec4<f32>());
tint_symbol_2(vec4f());
return;
}
struct vertex_main_out {
@builtin(position)
tint_symbol_1_1 : vec4<f32>,
tint_symbol_1_1 : vec4f,
}
@vertex

10
test/tint/builtins/textureLoad/depth_ms.spvasm.expected.wgsl
View File

@ -1,28 +1,28 @@
@group(1) @binding(0) var arg_0 : texture_depth_multisampled_2d;
var<private> tint_symbol_1 : vec4<f32> = vec4<f32>();
var<private> tint_symbol_1 : vec4f = vec4f();
fn textureLoad_6273b1() {
var res : f32 = 0.0f;
let x_17 : vec4<f32> = vec4<f32>(textureLoad(arg_0, vec2<i32>(), 1i), 0.0f, 0.0f, 0.0f);
let x_17 : vec4f = vec4f(textureLoad(arg_0, vec2i(), 1i), 0.0f, 0.0f, 0.0f);
res = x_17.x;
return;
}
fn tint_symbol_2(tint_symbol : vec4<f32>) {
fn tint_symbol_2(tint_symbol : vec4f) {
tint_symbol_1 = tint_symbol;
return;
}
fn vertex_main_1() {
textureLoad_6273b1();
tint_symbol_2(vec4<f32>());
tint_symbol_2(vec4f());
return;
}
struct vertex_main_out {
@builtin(position)
tint_symbol_1_1 : vec4<f32>,
tint_symbol_1_1 : vec4f,
}
@vertex

8
test/tint/builtins/textureNumSamples/depth_ms.spvasm.expected.wgsl
View File

@ -1,6 +1,6 @@
@group(1) @binding(0) var arg_0 : texture_depth_multisampled_2d;
var<private> tint_symbol_1 : vec4<f32> = vec4<f32>();
var<private> tint_symbol_1 : vec4f = vec4f();
fn textureNumSamples_a3c8a0() {
var res : i32 = 0i;
@ -9,20 +9,20 @@ fn textureNumSamples_a3c8a0() {
return;
}
fn tint_symbol_2(tint_symbol : vec4<f32>) {
fn tint_symbol_2(tint_symbol : vec4f) {
tint_symbol_1 = tint_symbol;
return;
}
fn vertex_main_1() {
textureNumSamples_a3c8a0();
tint_symbol_2(vec4<f32>());
tint_symbol_2(vec4f());
return;
}
struct vertex_main_out {
@builtin(position)
tint_symbol_1_1 : vec4<f32>,
tint_symbol_1_1 : vec4f,
}
@vertex

4
test/tint/ptr_ref/access/matrix.spvasm.expected.wgsl
View File

@ -1,7 +1,7 @@
fn main_1() {
var m : mat3x3<f32> = mat3x3<f32>();
m = mat3x3<f32>(vec3<f32>(1.0f, 2.0f, 3.0f), vec3<f32>(4.0f, 5.0f, 6.0f), vec3<f32>(7.0f, 8.0f, 9.0f));
m[1i] = vec3<f32>(5.0f, 5.0f, 5.0f);
m = mat3x3<f32>(vec3f(1.0f, 2.0f, 3.0f), vec3f(4.0f, 5.0f, 6.0f), vec3f(7.0f, 8.0f, 9.0f));
m[1i] = vec3f(5.0f, 5.0f, 5.0f);
return;
}

6
test/tint/samples/simple_vertex.spvasm.expected.wgsl
View File

@ -1,13 +1,13 @@
var<private> gl_Position : vec4<f32>;
var<private> gl_Position : vec4f;
fn main_1() {
gl_Position = vec4<f32>(0.0f, 0.0f, 0.0f, 0.0f);
gl_Position = vec4f(0.0f, 0.0f, 0.0f, 0.0f);
return;
}
struct main_out {
@builtin(position)
gl_Position : vec4<f32>,
gl_Position : vec4f,
}
@vertex