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

tint/resolver: Add f16 types, constructor, and conversions 

This patch add f16 types and their constructors and conversions in
resolver and intrinsic table. Also implement relating unit tests.

Bug: tint:1473, tint:1502
Change-Id: Ida1336193a72a73959e50e6a3eb12be44c0396b5
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/94642
Commit-Queue: Zhaoming Jiang <zhaoming.jiang@intel.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
Zhaoming Jiang 2022-06-28 14:03:36 +00:00 committed by Dawn LUCI CQ
parent 53af158366
commit 6058882d4b
45 changed files with 6086 additions and 4399 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

152
src/tint/intrinsics.def
View File

@ -71,11 +71,12 @@ enum texel_format {
// https://gpuweb.github.io/gpuweb/wgsl/#plain-types-section
type bool
@precedence(4) @display("abstract-float") type af
@precedence(3) @display("abstract-int") type ai
@precedence(2) type i32
@precedence(1) type u32
@precedence(0) type f32
@precedence(5) @display("abstract-float") type af
@precedence(4) @display("abstract-int") type ai
@precedence(3) type i32
@precedence(2) type u32
@precedence(1) type f32
@precedence(0) type f16
type vec2<T>
type vec3<T>
type vec4<T>
@ -126,16 +127,19 @@ type __atomic_compare_exchange_result<T>
// A type matcher that can match one or more types. //
////////////////////////////////////////////////////////////////////////////////
match f32f16: f32 | f16
match fiu32: f32 | i32 | u32
match fi32: f32 | i32
match iu32: i32 | u32
match scalar: f32 | i32 | u32 | bool
match abstract_or_scalar: ai | af | f32 | i32 | u32 | bool
match scalar: f32 | f16 | i32 | u32 | bool
match abstract_or_scalar: ai | af | f32 | f16 | i32 | u32 | bool
match af_f32: af | f32
match scalar_no_f32: i32 | u32 | bool
match scalar_no_i32: f32 | u32 | bool
match scalar_no_u32: f32 | i32 | bool
match scalar_no_bool: f32 | i32 | u32
match af_f32f16: af | f32 | f16
match scalar_no_f32: i32 | f16 | u32 | bool
match scalar_no_f16: f32 | i32 | u32 | bool
match scalar_no_i32: f32 | f16 | u32 | bool
match scalar_no_u32: f32 | f16 | i32 | bool
match scalar_no_bool: f32 | f16 | i32 | u32
////////////////////////////////////////////////////////////////////////////////
// Enum matchers //
@ -437,9 +441,9 @@ fn reverseBits<T: iu32>(T) -> T
fn reverseBits<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
fn round(f32) -> f32
fn round<N: num>(vec<N, f32>) -> vec<N, f32>
fn select<T: scalar>(T, T, bool) -> T
fn select<T: scalar, N: num>(vec<N, T>, vec<N, T>, bool) -> vec<N, T>
fn select<N: num, T: scalar>(vec<N, T>, vec<N, T>, vec<N, bool>) -> vec<N, T>
fn select<T: scalar_no_f16>(T, T, bool) -> T
fn select<T: scalar_no_f16, N: num>(vec<N, T>, vec<N, T>, bool) -> vec<N, T>
fn select<N: num, T: scalar_no_f16>(vec<N, T>, vec<N, T>, vec<N, bool>) -> vec<N, T>
fn sign(f32) -> f32
fn sign<N: num>(vec<N, f32>) -> vec<N, f32>
fn sin(f32) -> f32
@ -629,37 +633,39 @@ fn textureLoad(texture: texture_external, coords: vec2<i32>) -> vec4<f32>
ctor i32() -> i32
ctor u32() -> u32
ctor f32() -> f32
ctor f16() -> f16
ctor bool() -> bool
ctor vec2<T: scalar>() -> vec2<T>
ctor vec3<T: scalar>() -> vec3<T>
ctor vec4<T: scalar>() -> vec4<T>
ctor mat2x2() -> mat2x2<f32>
ctor mat2x3() -> mat2x3<f32>
ctor mat2x4() -> mat2x4<f32>
ctor mat3x2() -> mat3x2<f32>
ctor mat3x3() -> mat3x3<f32>
ctor mat3x4() -> mat3x4<f32>
ctor mat4x2() -> mat4x2<f32>
ctor mat4x3() -> mat4x3<f32>
ctor mat4x4() -> mat4x4<f32>
ctor mat2x2<T: f32f16>() -> mat2x2<T>
ctor mat2x3<T: f32f16>() -> mat2x3<T>
ctor mat2x4<T: f32f16>() -> mat2x4<T>
ctor mat3x2<T: f32f16>() -> mat3x2<T>
ctor mat3x3<T: f32f16>() -> mat3x3<T>
ctor mat3x4<T: f32f16>() -> mat3x4<T>
ctor mat4x2<T: f32f16>() -> mat4x2<T>
ctor mat4x3<T: f32f16>() -> mat4x3<T>
ctor mat4x4<T: f32f16>() -> mat4x4<T>
// Identity constructors
ctor i32(i32) -> i32
ctor u32(u32) -> u32
ctor f32(f32) -> f32
ctor f16(f16) -> f16
ctor bool(bool) -> bool
ctor vec2<T: scalar>(vec2<T>) -> vec2<T>
ctor vec3<T: scalar>(vec3<T>) -> vec3<T>
ctor vec4<T: scalar>(vec4<T>) -> vec4<T>
ctor mat2x2<f32>(mat2x2<f32>) -> mat2x2<f32>
ctor mat2x3<f32>(mat2x3<f32>) -> mat2x3<f32>
ctor mat2x4<f32>(mat2x4<f32>) -> mat2x4<f32>
ctor mat3x2<f32>(mat3x2<f32>) -> mat3x2<f32>
ctor mat3x3<f32>(mat3x3<f32>) -> mat3x3<f32>
ctor mat3x4<f32>(mat3x4<f32>) -> mat3x4<f32>
ctor mat4x2<f32>(mat4x2<f32>) -> mat4x2<f32>
ctor mat4x3<f32>(mat4x3<f32>) -> mat4x3<f32>
ctor mat4x4<f32>(mat4x4<f32>) -> mat4x4<f32>
ctor mat2x2<T: f32f16>(mat2x2<T>) -> mat2x2<T>
ctor mat2x3<T: f32f16>(mat2x3<T>) -> mat2x3<T>
ctor mat2x4<T: f32f16>(mat2x4<T>) -> mat2x4<T>
ctor mat3x2<T: f32f16>(mat3x2<T>) -> mat3x2<T>
ctor mat3x3<T: f32f16>(mat3x3<T>) -> mat3x3<T>
ctor mat3x4<T: f32f16>(mat3x4<T>) -> mat3x4<T>
ctor mat4x2<T: f32f16>(mat4x2<T>) -> mat4x2<T>
ctor mat4x3<T: f32f16>(mat4x3<T>) -> mat4x3<T>
ctor mat4x4<T: f32f16>(mat4x4<T>) -> mat4x4<T>
// Vector constructors
ctor vec2<T: abstract_or_scalar>(T) -> vec2<T>
@ -679,82 +685,106 @@ ctor vec4<T: abstract_or_scalar>(x: T, zyw: vec3<T>) -> vec4<T>
// Matrix constructors
ctor mat2x2<T: af_f32>(T) -> mat2x2<T>
ctor mat2x2<T: af_f32>(T, T,
T, T) -> mat2x2<T>
ctor mat2x2<T: af_f32>(vec2<T>, vec2<T>) -> mat2x2<T>
ctor mat2x3<T: af_f32>(T) -> mat2x3<T>
ctor mat2x3<T: af_f32>(T, T, T,
T, T, T) -> mat2x3<T>
ctor mat2x3<T: af_f32>(vec3<T>, vec3<T>) -> mat2x3<T>
ctor mat2x4<T: af_f32>(T) -> mat2x4<T>
ctor mat2x4<T: af_f32>(T, T, T, T,
T, T, T, T) -> mat2x4<T>
ctor mat2x4<T: af_f32>(vec4<T>, vec4<T>) -> mat2x4<T>
ctor mat3x2<T: af_f32>(T) -> mat3x2<T>
ctor mat3x2<T: af_f32>(T, T,
ctor mat3x3<T: af_f32>(T) -> mat3x3<T>
ctor mat3x4<T: af_f32>(T) -> mat3x4<T>
ctor mat4x2<T: af_f32>(T) -> mat4x2<T>
ctor mat4x3<T: af_f32>(T) -> mat4x3<T>
ctor mat4x4<T: af_f32>(T) -> mat4x4<T>
ctor mat2x2<T: af_f32f16>(T, T,
T, T) -> mat2x2<T>
ctor mat2x2<T: af_f32f16>(vec2<T>, vec2<T>) -> mat2x2<T>
ctor mat2x3<T: af_f32f16>(T, T, T,
T, T, T) -> mat2x3<T>
ctor mat2x3<T: af_f32f16>(vec3<T>, vec3<T>) -> mat2x3<T>
ctor mat2x4<T: af_f32f16>(T, T, T, T,
T, T, T, T) -> mat2x4<T>
ctor mat2x4<T: af_f32f16>(vec4<T>, vec4<T>) -> mat2x4<T>
ctor mat3x2<T: af_f32f16>(T, T,
T, T,
T, T) -> mat3x2<T>
ctor mat3x2<T: af_f32>(vec2<T>, vec2<T>, vec2<T>) -> mat3x2<T>
ctor mat3x2<T: af_f32f16>(vec2<T>, vec2<T>, vec2<T>) -> mat3x2<T>
ctor mat3x3<T: af_f32>(T) -> mat3x3<T>
ctor mat3x3<T: af_f32>(T, T, T,
ctor mat3x3<T: af_f32f16>(T, T, T,
T, T, T,
T, T, T) -> mat3x3<T>
ctor mat3x3<T: af_f32>(vec3<T>, vec3<T>, vec3<T>) -> mat3x3<T>
ctor mat3x3<T: af_f32f16>(vec3<T>, vec3<T>, vec3<T>) -> mat3x3<T>
ctor mat3x4<T: af_f32>(T) -> mat3x4<T>
ctor mat3x4<T: af_f32>(T, T, T, T,
ctor mat3x4<T: af_f32f16>(T, T, T, T,
T, T, T, T,
T, T, T, T) -> mat3x4<T>
ctor mat3x4<T: af_f32>(vec4<T>, vec4<T>, vec4<T>) -> mat3x4<T>
ctor mat3x4<T: af_f32f16>(vec4<T>, vec4<T>, vec4<T>) -> mat3x4<T>
ctor mat4x2<T: af_f32>(T) -> mat4x2<T>
ctor mat4x2<T: af_f32>(T, T,
ctor mat4x2<T: af_f32f16>(T, T,
T, T,
T, T,
T, T) -> mat4x2<T>
ctor mat4x2<T: af_f32>(vec2<T>, vec2<T>, vec2<T>, vec2<T>) -> mat4x2<T>
ctor mat4x2<T: af_f32f16>(vec2<T>, vec2<T>, vec2<T>, vec2<T>) -> mat4x2<T>
ctor mat4x3<T: af_f32>(T) -> mat4x3<T>
ctor mat4x3<T: af_f32>(T, T, T,
ctor mat4x3<T: af_f32f16>(T, T, T,
T, T, T,
T, T, T,
T, T, T) -> mat4x3<T>
ctor mat4x3<T: af_f32>(vec3<T>, vec3<T>, vec3<T>, vec3<T>) -> mat4x3<T>
ctor mat4x3<T: af_f32f16>(vec3<T>, vec3<T>, vec3<T>, vec3<T>) -> mat4x3<T>
ctor mat4x4<T: af_f32>(T) -> mat4x4<T>
ctor mat4x4<T: af_f32>(T, T, T, T,
ctor mat4x4<T: af_f32f16>(T, T, T, T,
T, T, T, T,
T, T, T, T,
T, T, T, T) -> mat4x4<T>
ctor mat4x4<T: af_f32>(vec4<T>, vec4<T>, vec4<T>, vec4<T>) -> mat4x4<T>
ctor mat4x4<T: af_f32f16>(vec4<T>, vec4<T>, vec4<T>, vec4<T>) -> mat4x4<T>
////////////////////////////////////////////////////////////////////////////////
// Type conversions //
////////////////////////////////////////////////////////////////////////////////
conv f32<T: scalar_no_f32>(T) -> f32
conv f16<T: scalar_no_f16>(T) -> f16
conv i32<T: scalar_no_i32>(T) -> i32
conv u32<T: scalar_no_u32>(T) -> u32
conv bool<T: scalar_no_bool>(T) -> bool
conv vec2<T: f32, U: scalar_no_f32>(vec2<U>) -> vec2<f32>
conv vec2<T: f16, U: scalar_no_f16>(vec2<U>) -> vec2<f16>
conv vec2<T: i32, U: scalar_no_i32>(vec2<U>) -> vec2<i32>
conv vec2<T: u32, U: scalar_no_u32>(vec2<U>) -> vec2<u32>
conv vec2<T: bool, U: scalar_no_bool>(vec2<U>) -> vec2<bool>
conv vec3<T: f32, U: scalar_no_f32>(vec3<U>) -> vec3<f32>
conv vec3<T: f16, U: scalar_no_f16>(vec3<U>) -> vec3<f16>
conv vec3<T: i32, U: scalar_no_i32>(vec3<U>) -> vec3<i32>
conv vec3<T: u32, U: scalar_no_u32>(vec3<U>) -> vec3<u32>
conv vec3<T: bool, U: scalar_no_bool>(vec3<U>) -> vec3<bool>
conv vec4<T: f32, U: scalar_no_f32>(vec4<U>) -> vec4<f32>
conv vec4<T: f16, U: scalar_no_f16>(vec4<U>) -> vec4<f16>
conv vec4<T: i32, U: scalar_no_i32>(vec4<U>) -> vec4<i32>
conv vec4<T: u32, U: scalar_no_u32>(vec4<U>) -> vec4<u32>
conv vec4<T: bool, U: scalar_no_bool>(vec4<U>) -> vec4<bool>
conv mat2x2<T: f16>(mat2x2<f32>) -> mat2x2<f16>
conv mat2x2<T: f32>(mat2x2<f16>) -> mat2x2<f32>
conv mat2x3<T: f16>(mat2x3<f32>) -> mat2x3<f16>
conv mat2x3<T: f32>(mat2x3<f16>) -> mat2x3<f32>
conv mat2x4<T: f16>(mat2x4<f32>) -> mat2x4<f16>
conv mat2x4<T: f32>(mat2x4<f16>) -> mat2x4<f32>
conv mat3x2<T: f16>(mat3x2<f32>) -> mat3x2<f16>
conv mat3x2<T: f32>(mat3x2<f16>) -> mat3x2<f32>
conv mat3x3<T: f16>(mat3x3<f32>) -> mat3x3<f16>
conv mat3x3<T: f32>(mat3x3<f16>) -> mat3x3<f32>
conv mat3x4<T: f16>(mat3x4<f32>) -> mat3x4<f16>
conv mat3x4<T: f32>(mat3x4<f16>) -> mat3x4<f32>
conv mat4x2<T: f16>(mat4x2<f32>) -> mat4x2<f16>
conv mat4x2<T: f32>(mat4x2<f16>) -> mat4x2<f32>
conv mat4x3<T: f16>(mat4x3<f32>) -> mat4x3<f16>
conv mat4x3<T: f32>(mat4x3<f16>) -> mat4x3<f32>
conv mat4x4<T: f16>(mat4x4<f32>) -> mat4x4<f16>
conv mat4x4<T: f32>(mat4x4<f16>) -> mat4x4<f32>
////////////////////////////////////////////////////////////////////////////////
// Operators //
// //

3
src/tint/reader/wgsl/token.h
View File

@ -351,7 +351,8 @@ class Token {
bool IsLiteral() const {
return type_ == Type::kIntLiteral || type_ == Type::kIntLiteral_I ||
type_ == Type::kIntLiteral_U || type_ == Type::kFalse || type_ == Type::kTrue ||
type_ == Type::kFloatLiteral || type_ == Type::kFloatLiteral_F;
type_ == Type::kFloatLiteral || type_ == Type::kFloatLiteral_F ||
type_ == Type::kFloatLiteral_H;
}
/// @returns true if token is a 'matNxM'
bool IsMatrix() const {

2
src/tint/reader/wgsl/token_test.cc
View File

@ -85,6 +85,8 @@ TEST_F(TokenTest, ToStr) {
EXPECT_THAT(Token(Token::Type::kFloatLiteral, Source{}, d).to_str(), Not(EndsWith("f")));
EXPECT_THAT(Token(Token::Type::kFloatLiteral_F, Source{}, d).to_str(), StartsWith("123"));
EXPECT_THAT(Token(Token::Type::kFloatLiteral_F, Source{}, d).to_str(), EndsWith("f"));
EXPECT_THAT(Token(Token::Type::kFloatLiteral_H, Source{}, d).to_str(), StartsWith("123"));
EXPECT_THAT(Token(Token::Type::kFloatLiteral_H, Source{}, d).to_str(), EndsWith("h"));
EXPECT_EQ(Token(Token::Type::kIntLiteral, Source{}, i).to_str(), "123");
EXPECT_EQ(Token(Token::Type::kIntLiteral_I, Source{}, i).to_str(), "123i");
EXPECT_EQ(Token(Token::Type::kIntLiteral_U, Source{}, i).to_str(), "123u");

3
src/tint/resolver/call_test.cc
View File

@ -71,6 +71,7 @@ static constexpr Params all_param_types[] = {
ParamsFor<u32>(), //
ParamsFor<i32>(), //
ParamsFor<f32>(), //
ParamsFor<f16>(), //
ParamsFor<vec3<bool>>(), //
ParamsFor<vec3<i32>>(), //
ParamsFor<vec3<u32>>(), //
@ -81,6 +82,8 @@ static constexpr Params all_param_types[] = {
};
TEST_F(ResolverCallTest, Valid) {
Enable(ast::Extension::kF16);
ast::ParameterList params;
ast::ExpressionList args;
for (auto& p : all_param_types) {

2
src/tint/resolver/ctor_conv_intrinsic.cc
View File

@ -36,6 +36,8 @@ const char* str(CtorConvIntrinsic i) {
return "u32";
case CtorConvIntrinsic::kF32:
return "f32";
case CtorConvIntrinsic::kF16:
return "f16";
case CtorConvIntrinsic::kBool:
return "bool";
case CtorConvIntrinsic::kVec2:

1
src/tint/resolver/ctor_conv_intrinsic.h
View File

@ -36,6 +36,7 @@ enum class CtorConvIntrinsic {
kI32,
kU32,
kF32,
kF16,
kBool,
kVec2,
kVec3,

14
src/tint/resolver/host_shareable_validation_test.cc
View File

@ -104,15 +104,19 @@ TEST_F(ResolverHostShareableValidationTest, NestedStructures) {
}
TEST_F(ResolverHostShareableValidationTest, NoError) {
Enable(ast::Extension::kF16);
auto* i1 = Structure("I1", {
Member(Source{{1, 1}}, "x1", ty.f32()),
Member(Source{{2, 1}}, "y1", ty.vec3<f32>()),
Member(Source{{3, 1}}, "z1", ty.array<i32, 4>()),
Member(Source{{1, 1}}, "w1", ty.f16()),
Member(Source{{2, 1}}, "x1", ty.f32()),
Member(Source{{3, 1}}, "y1", ty.vec3<f32>()),
Member(Source{{4, 1}}, "z1", ty.array<i32, 4>()),
});
auto* a1 = Alias("a1", ty.Of(i1));
auto* i2 = Structure("I2", {
Member(Source{{4, 1}}, "x2", ty.mat2x2<f32>()),
Member(Source{{5, 1}}, "y2", ty.Of(i1)),
Member(Source{{5, 1}}, "x2", ty.mat2x2<f32>()),
Member(Source{{6, 1}}, "w2", ty.mat3x4<f16>()),
Member(Source{{7, 1}}, "z2", ty.Of(i1)),
});
auto* a2 = Alias("a2", ty.Of(i2));
auto* i3 = Structure("I3", {

8
src/tint/resolver/intrinsic_table.cc
View File

@ -341,6 +341,14 @@ const sem::Bool* build_bool(MatchState& state) {
return state.builder.create<sem::Bool>();
}
const sem::F16* build_f16(MatchState& state) {
return state.builder.create<sem::F16>();
}
bool match_f16(const sem::Type* ty) {
return ty->IsAnyOf<Any, sem::F16, sem::AbstractNumeric>();
}
const sem::F32* build_f32(MatchState& state) {
return state.builder.create<sem::F32>();
}

8732
src/tint/resolver/intrinsic_table.inl
View File

File diff suppressed because it is too large Load Diff

69
src/tint/resolver/intrinsic_table_test.cc
View File

@ -709,18 +709,19 @@ TEST_F(IntrinsicTableTest, MismatchTypeConstructorImplicit) {
EXPECT_EQ(Diagnostics().str(), R"(12:34 error: no matching constructor for vec3(i32, f32, i32)
6 candidate constructors:
vec3(x: T, y: T, z: T) -> vec3<T> where: T is abstract-int, abstract-float, f32, i32, u32 or bool
vec3(xy: vec2<T>, z: T) -> vec3<T> where: T is abstract-int, abstract-float, f32, i32, u32 or bool
vec3(x: T, yz: vec2<T>) -> vec3<T> where: T is abstract-int, abstract-float, f32, i32, u32 or bool
vec3(T) -> vec3<T> where: T is abstract-int, abstract-float, f32, i32, u32 or bool
vec3(vec3<T>) -> vec3<T> where: T is f32, i32, u32 or bool
vec3() -> vec3<T> where: T is f32, i32, u32 or bool
vec3(x: T, y: T, z: T) -> vec3<T> where: T is abstract-int, abstract-float, f32, f16, i32, u32 or bool
vec3(xy: vec2<T>, z: T) -> vec3<T> where: T is abstract-int, abstract-float, f32, f16, i32, u32 or bool
vec3(x: T, yz: vec2<T>) -> vec3<T> where: T is abstract-int, abstract-float, f32, f16, i32, u32 or bool
vec3(T) -> vec3<T> where: T is abstract-int, abstract-float, f32, f16, i32, u32 or bool
vec3(vec3<T>) -> vec3<T> where: T is f32, f16, i32, u32 or bool
vec3() -> vec3<T> where: T is f32, f16, i32, u32 or bool
4 candidate conversions:
vec3(vec3<U>) -> vec3<f32> where: T is f32, U is i32, u32 or bool
vec3(vec3<U>) -> vec3<i32> where: T is i32, U is f32, u32 or bool
vec3(vec3<U>) -> vec3<u32> where: T is u32, U is f32, i32 or bool
vec3(vec3<U>) -> vec3<bool> where: T is bool, U is f32, i32 or u32
5 candidate conversions:
vec3(vec3<U>) -> vec3<f32> where: T is f32, U is i32, f16, u32 or bool
vec3(vec3<U>) -> vec3<f16> where: T is f16, U is f32, i32, u32 or bool
vec3(vec3<U>) -> vec3<i32> where: T is i32, U is f32, f16, u32 or bool
vec3(vec3<U>) -> vec3<u32> where: T is u32, U is f32, f16, i32 or bool
vec3(vec3<U>) -> vec3<bool> where: T is bool, U is f32, f16, i32 or u32
)");
}
@ -733,18 +734,19 @@ TEST_F(IntrinsicTableTest, MismatchTypeConstructorExplicit) {
R"(12:34 error: no matching constructor for vec3<i32>(i32, f32, i32)
6 candidate constructors:
vec3(x: T, y: T, z: T) -> vec3<T> where: T is abstract-int, abstract-float, f32, i32, u32 or bool
vec3(x: T, yz: vec2<T>) -> vec3<T> where: T is abstract-int, abstract-float, f32, i32, u32 or bool
vec3(T) -> vec3<T> where: T is abstract-int, abstract-float, f32, i32, u32 or bool
vec3(xy: vec2<T>, z: T) -> vec3<T> where: T is abstract-int, abstract-float, f32, i32, u32 or bool
vec3(vec3<T>) -> vec3<T> where: T is f32, i32, u32 or bool
vec3() -> vec3<T> where: T is f32, i32, u32 or bool
vec3(x: T, y: T, z: T) -> vec3<T> where: T is abstract-int, abstract-float, f32, f16, i32, u32 or bool
vec3(x: T, yz: vec2<T>) -> vec3<T> where: T is abstract-int, abstract-float, f32, f16, i32, u32 or bool
vec3(T) -> vec3<T> where: T is abstract-int, abstract-float, f32, f16, i32, u32 or bool
vec3(xy: vec2<T>, z: T) -> vec3<T> where: T is abstract-int, abstract-float, f32, f16, i32, u32 or bool
vec3(vec3<T>) -> vec3<T> where: T is f32, f16, i32, u32 or bool
vec3() -> vec3<T> where: T is f32, f16, i32, u32 or bool
4 candidate conversions:
vec3(vec3<U>) -> vec3<f32> where: T is f32, U is i32, u32 or bool
vec3(vec3<U>) -> vec3<i32> where: T is i32, U is f32, u32 or bool
vec3(vec3<U>) -> vec3<u32> where: T is u32, U is f32, i32 or bool
vec3(vec3<U>) -> vec3<bool> where: T is bool, U is f32, i32 or u32
5 candidate conversions:
vec3(vec3<U>) -> vec3<f32> where: T is f32, U is i32, f16, u32 or bool
vec3(vec3<U>) -> vec3<f16> where: T is f16, U is f32, i32, u32 or bool
vec3(vec3<U>) -> vec3<i32> where: T is i32, U is f32, f16, u32 or bool
vec3(vec3<U>) -> vec3<u32> where: T is u32, U is f32, f16, i32 or bool
vec3(vec3<U>) -> vec3<bool> where: T is bool, U is f32, f16, i32 or u32
)");
}
@ -770,18 +772,19 @@ TEST_F(IntrinsicTableTest, MismatchTypeConversion) {
R"(12:34 error: no matching constructor for vec3<f32>(array<u32>)
6 candidate constructors:
vec3(vec3<T>) -> vec3<T> where: T is f32, i32, u32 or bool
vec3(T) -> vec3<T> where: T is abstract-int, abstract-float, f32, i32, u32 or bool
vec3() -> vec3<T> where: T is f32, i32, u32 or bool
vec3(xy: vec2<T>, z: T) -> vec3<T> where: T is abstract-int, abstract-float, f32, i32, u32 or bool
vec3(x: T, yz: vec2<T>) -> vec3<T> where: T is abstract-int, abstract-float, f32, i32, u32 or bool
vec3(x: T, y: T, z: T) -> vec3<T> where: T is abstract-int, abstract-float, f32, i32, u32 or bool
vec3(vec3<T>) -> vec3<T> where: T is f32, f16, i32, u32 or bool
vec3(T) -> vec3<T> where: T is abstract-int, abstract-float, f32, f16, i32, u32 or bool
vec3() -> vec3<T> where: T is f32, f16, i32, u32 or bool
vec3(xy: vec2<T>, z: T) -> vec3<T> where: T is abstract-int, abstract-float, f32, f16, i32, u32 or bool
vec3(x: T, yz: vec2<T>) -> vec3<T> where: T is abstract-int, abstract-float, f32, f16, i32, u32 or bool
vec3(x: T, y: T, z: T) -> vec3<T> where: T is abstract-int, abstract-float, f32, f16, i32, u32 or bool
4 candidate conversions:
vec3(vec3<U>) -> vec3<f32> where: T is f32, U is i32, u32 or bool
vec3(vec3<U>) -> vec3<i32> where: T is i32, U is f32, u32 or bool
vec3(vec3<U>) -> vec3<u32> where: T is u32, U is f32, i32 or bool
vec3(vec3<U>) -> vec3<bool> where: T is bool, U is f32, i32 or u32
5 candidate conversions:
vec3(vec3<U>) -> vec3<f32> where: T is f32, U is i32, f16, u32 or bool
vec3(vec3<U>) -> vec3<f16> where: T is f16, U is f32, i32, u32 or bool
vec3(vec3<U>) -> vec3<i32> where: T is i32, U is f32, f16, u32 or bool
vec3(vec3<U>) -> vec3<u32> where: T is u32, U is f32, f16, i32 or bool
vec3(vec3<U>) -> vec3<bool> where: T is bool, U is f32, f16, i32 or u32
)");
}

41
src/tint/resolver/is_host_shareable_test.cc
View File

@ -35,6 +35,7 @@ TEST_F(ResolverIsHostShareable, NumericScalar) {
EXPECT_TRUE(r()->IsHostShareable(create<sem::I32>()));
EXPECT_TRUE(r()->IsHostShareable(create<sem::U32>()));
EXPECT_TRUE(r()->IsHostShareable(create<sem::F32>()));
EXPECT_TRUE(r()->IsHostShareable(create<sem::F16>()));
}
TEST_F(ResolverIsHostShareable, NumericVector) {
@ -47,6 +48,9 @@ TEST_F(ResolverIsHostShareable, NumericVector) {
EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::F32>(), 2u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::F32>(), 3u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::F32>(), 4u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::F16>(), 2u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::F16>(), 3u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::F16>(), 4u)));
}
TEST_F(ResolverIsHostShareable, BoolVector) {
@ -62,19 +66,32 @@ TEST_F(ResolverIsHostShareable, BoolVector) {
}
TEST_F(ResolverIsHostShareable, Matrix) {
auto* vec2 = create<sem::Vector>(create<sem::F32>(), 2u);
auto* vec3 = create<sem::Vector>(create<sem::F32>(), 3u);
auto* vec4 = create<sem::Vector>(create<sem::F32>(), 4u);
auto* vec2_f32 = create<sem::Vector>(create<sem::F32>(), 2u);
auto* vec3_f32 = create<sem::Vector>(create<sem::F32>(), 3u);
auto* vec4_f32 = create<sem::Vector>(create<sem::F32>(), 4u);
auto* vec2_f16 = create<sem::Vector>(create<sem::F16>(), 2u);
auto* vec3_f16 = create<sem::Vector>(create<sem::F16>(), 3u);
auto* vec4_f16 = create<sem::Vector>(create<sem::F16>(), 4u);
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2, 2u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2, 3u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2, 4u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3, 2u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3, 3u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3, 4u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4, 2u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4, 3u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4, 4u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2_f32, 2u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2_f32, 3u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2_f32, 4u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3_f32, 2u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3_f32, 3u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3_f32, 4u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4_f32, 2u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4_f32, 3u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4_f32, 4u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2_f16, 2u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2_f16, 3u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2_f16, 4u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3_f16, 2u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3_f16, 3u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3_f16, 4u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4_f16, 2u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4_f16, 3u)));
EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4_f16, 4u)));
}
TEST_F(ResolverIsHostShareable, Pointer) {

40
src/tint/resolver/is_storeable_test.cc
View File

@ -32,6 +32,7 @@ TEST_F(ResolverIsStorableTest, Scalar) {
EXPECT_TRUE(r()->IsStorable(create<sem::I32>()));
EXPECT_TRUE(r()->IsStorable(create<sem::U32>()));
EXPECT_TRUE(r()->IsStorable(create<sem::F32>()));
EXPECT_TRUE(r()->IsStorable(create<sem::F16>()));
}
TEST_F(ResolverIsStorableTest, Vector) {
@ -44,21 +45,36 @@ TEST_F(ResolverIsStorableTest, Vector) {
EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::F32>(), 2u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::F32>(), 3u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::F32>(), 4u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::F16>(), 2u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::F16>(), 3u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::F16>(), 4u)));
}
TEST_F(ResolverIsStorableTest, Matrix) {
auto* vec2 = create<sem::Vector>(create<sem::F32>(), 2u);
auto* vec3 = create<sem::Vector>(create<sem::F32>(), 3u);
auto* vec4 = create<sem::Vector>(create<sem::F32>(), 4u);
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2, 2u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2, 3u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2, 4u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3, 2u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3, 3u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3, 4u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4, 2u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4, 3u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4, 4u)));
auto* vec2_f32 = create<sem::Vector>(create<sem::F32>(), 2u);
auto* vec3_f32 = create<sem::Vector>(create<sem::F32>(), 3u);
auto* vec4_f32 = create<sem::Vector>(create<sem::F32>(), 4u);
auto* vec2_f16 = create<sem::Vector>(create<sem::F16>(), 2u);
auto* vec3_f16 = create<sem::Vector>(create<sem::F16>(), 3u);
auto* vec4_f16 = create<sem::Vector>(create<sem::F16>(), 4u);
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2_f32, 2u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2_f32, 3u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2_f32, 4u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3_f32, 2u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3_f32, 3u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3_f32, 4u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4_f32, 2u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4_f32, 3u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4_f32, 4u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2_f16, 2u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2_f16, 3u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2_f16, 4u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3_f16, 2u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3_f16, 3u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3_f16, 4u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4_f16, 2u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4_f16, 3u)));
EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4_f16, 4u)));
}
TEST_F(ResolverIsStorableTest, Pointer) {

50
src/tint/resolver/materialize_test.cc
View File

@ -41,11 +41,16 @@ constexpr double kHighestI32 = static_cast<double>(i32::kHighest);
constexpr double kLowestI32 = static_cast<double>(i32::kLowest);
constexpr double kHighestF32 = static_cast<double>(f32::kHighest);
constexpr double kLowestF32 = static_cast<double>(f32::kLowest);
// constexpr double kHighestF16 = static_cast<double>(f16::kHighest);
// constexpr double kLowestF16 = static_cast<double>(f16::kLowest);
constexpr double kTooBigF32 = static_cast<double>(3.5e+38);
// constexpr double kTooBigF16 = static_cast<double>(6.6e+4);
constexpr double kPiF64 = 3.141592653589793;
constexpr double kPiF32 = 3.1415927410125732; // kPiF64 quantized to f32
// constexpr double kPiF16 = 3.140625; // kPiF64 quantized to f16
constexpr double kSubnormalF32 = 0x1.0p-128;
// constexpr double kSubnormalF16 = 0x1.0p-16;
enum class Expectation {
kMaterialize,
@ -245,7 +250,7 @@ using MaterializeAbstractNumericToConcreteType =
resolver::ResolverTestWithParam<std::tuple<Expectation, Method, Data>>;
TEST_P(MaterializeAbstractNumericToConcreteType, Test) {
// Once F16 is properly supported, we'll need to enable this:
// Once built-in and ops using f16 is properly supported, we'll need to enable this:
// Enable(ast::Extension::kF16);
const auto& param = GetParam();
@ -433,8 +438,13 @@ INSTANTIATE_TEST_SUITE_P(
Types<f32, AFloat>(AFloat(kPiF32), kPiF64), //
Types<f32, AFloat>(AFloat(kSubnormalF32), kSubnormalF32), //
Types<f32, AFloat>(AFloat(-kSubnormalF32), -kSubnormalF32), //
/* Types<f16, AFloat>(1.0_a), */ //
/* Types<f16, AFloat>(1.0_a), */ //
/* Types<f16, AFloat>(0.0_a, 0.0), */ //
/* Types<f16, AFloat>(1.0_a, 1.0), */ //
/* Types<f16, AFloat>(AFloat(kHighestF16), kHighestF16), */ //
/* Types<f16, AFloat>(AFloat(kLowestF16), kLowestF16), */ //
/* Types<f16, AFloat>(AFloat(kPiF16), kPiF64), */ //
/* Types<f16, AFloat>(AFloat(kSubnormalF16), kSubnormalF16), */ //
/* Types<f16, AFloat>(AFloat(-kSubnormalF16), -kSubnormalF16), */ //
})));
INSTANTIATE_TEST_SUITE_P(
@ -460,8 +470,14 @@ INSTANTIATE_TEST_SUITE_P(
Types<f32V, AFloatV>(AFloat(kPiF32), kPiF64), //
Types<f32V, AFloatV>(AFloat(kSubnormalF32), kSubnormalF32), //
Types<f32V, AFloatV>(AFloat(-kSubnormalF32), -kSubnormalF32), //
/* Types<f16V, AFloatV>(1.0_a), */ //
/* Types<f16V, AFloatV>(1.0_a), */ //
/* Types<f16V, AFloatV>(0.0_a, 0.0), */ //
/* Types<f16V, AFloatV>(1.0_a, 1.0), */ //
/* Types<f16V, AFloatV>(-1.0_a, -1.0), */ //
/* Types<f16V, AFloatV>(AFloat(kHighestF16), kHighestF16), */ //
/* Types<f16V, AFloatV>(AFloat(kLowestF16), kLowestF16), */ //
/* Types<f16V, AFloatV>(AFloat(kPiF16), kPiF64), */ //
/* Types<f16V, AFloatV>(AFloat(kSubnormalF16), kSubnormalF16), */ //
/* Types<f16V, AFloatV>(AFloat(-kSubnormalF16), -kSubnormalF16), */ //
})));
INSTANTIATE_TEST_SUITE_P(
@ -478,7 +494,14 @@ INSTANTIATE_TEST_SUITE_P(
Types<f32M, AFloatM>(AFloat(kPiF32), kPiF64), //
Types<f32M, AFloatM>(AFloat(kSubnormalF32), kSubnormalF32), //
Types<f32M, AFloatM>(AFloat(-kSubnormalF32), -kSubnormalF32), //
/* Types<f16V, AFloatM>(1.0_a), */ //
/* Types<f16M, AFloatM>(0.0_a, 0.0), */ //
/* Types<f16M, AFloatM>(1.0_a, 1.0), */ //
/* Types<f16M, AFloatM>(-1.0_a, -1.0), */ //
/* Types<f16M, AFloatM>(AFloat(kHighestF16), kHighestF16), */ //
/* Types<f16M, AFloatM>(AFloat(kLowestF16), kLowestF16), */ //
/* Types<f16M, AFloatM>(AFloat(kPiF16), kPiF64), */ //
/* Types<f16M, AFloatM>(AFloat(kSubnormalF16), kSubnormalF16), */ //
/* Types<f16M, AFloatM>(AFloat(-kSubnormalF16), -kSubnormalF16), */ //
})));
INSTANTIATE_TEST_SUITE_P(MaterializeSwitch,
@ -548,8 +571,8 @@ INSTANTIATE_TEST_SUITE_P(ScalarValueCannotBeRepresented,
Types<u32, AInt>(0_a, kLowestU32 - 1), //
Types<f32, AFloat>(0.0_a, kTooBigF32), //
Types<f32, AFloat>(0.0_a, -kTooBigF32), //
/* Types<f16, AFloat>(), */ //
/* Types<f16, AFloat>(), */ //
/* Types<f16, AFloat>(0.0_a, kTooBigF16), */ //
/* Types<f16, AFloat>(0.0_a, -kTooBigF16), */ //
})));
INSTANTIATE_TEST_SUITE_P(VectorValueCannotBeRepresented,
@ -563,8 +586,8 @@ INSTANTIATE_TEST_SUITE_P(VectorValueCannotBeRepresented,
Types<u32V, AIntV>(0_a, kLowestU32 - 1), //
Types<f32V, AFloatV>(0.0_a, kTooBigF32), //
Types<f32V, AFloatV>(0.0_a, -kTooBigF32), //
/* Types<f16V, AFloatV>(), */ //
/* Types<f16V, AFloatV>(), */ //
/* Types<f16V, AFloatV>(0.0_a, kTooBigF16), */ //
/* Types<f16V, AFloatV>(0.0_a, -kTooBigF16), */ //
})));
INSTANTIATE_TEST_SUITE_P(MatrixValueCannotBeRepresented,
@ -574,8 +597,8 @@ INSTANTIATE_TEST_SUITE_P(MatrixValueCannotBeRepresented,
testing::ValuesIn(std::vector<Data>{
Types<f32M, AFloatM>(0.0_a, kTooBigF32), //
Types<f32M, AFloatM>(0.0_a, -kTooBigF32), //
/* Types<f16M, AFloatM>(), */ //
/* Types<f16M, AFloatM>(), */ //
/* Types<f16M, AFloatM>(0.0_a, kTooBigF16), */ //
/* Types<f16M, AFloatM>(0.0_a, -kTooBigF16), */ //
})));
} // namespace materialize_abstract_numeric_to_concrete_type
@ -683,9 +706,6 @@ using MaterializeAbstractNumericToDefaultType =
resolver::ResolverTestWithParam<std::tuple<Expectation, Method, Data>>;
TEST_P(MaterializeAbstractNumericToDefaultType, Test) {
// Once F16 is properly supported, we'll need to enable this:
// Enable(ast::Extension::kF16);
const auto& param = GetParam();
const auto& expectation = std::get<0>(param);
const auto& method = std::get<1>(param);

1
src/tint/resolver/resolver.cc
View File

@ -1506,6 +1506,7 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
},
[&](const sem::I32*) { return ct_ctor_or_conv(CtorConvIntrinsic::kI32, nullptr); },
[&](const sem::U32*) { return ct_ctor_or_conv(CtorConvIntrinsic::kU32, nullptr); },
[&](const sem::F16*) { return ct_ctor_or_conv(CtorConvIntrinsic::kF16, nullptr); },
[&](const sem::F32*) { return ct_ctor_or_conv(CtorConvIntrinsic::kF32, nullptr); },
[&](const sem::Bool*) { return ct_ctor_or_conv(CtorConvIntrinsic::kBool, nullptr); },
[&](const sem::Array* arr) -> sem::Call* {

180
src/tint/resolver/resolver_constants_test.cc
View File

@ -75,6 +75,23 @@ TEST_F(ResolverConstantsTest, Scalar_f32) {
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 9.9f);
}
TEST_F(ResolverConstantsTest, Scalar_f16) {
Enable(ast::Extension::kF16);
auto* expr = Expr(9.9_h);
WrapInFunction(expr);
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get(expr);
EXPECT_NE(sem, nullptr);
EXPECT_TRUE(sem->Type()->Is<sem::F16>());
EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
EXPECT_EQ(sem->ConstantValue().ElementType(), sem->Type());
ASSERT_EQ(sem->ConstantValue().ElementCount(), 1u);
// 9.9 is not exactly representable by f16, and should be quantized to 9.8984375
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 9.8984375f);
}
TEST_F(ResolverConstantsTest, Scalar_bool) {
auto* expr = Expr(true);
WrapInFunction(expr);
@ -150,6 +167,27 @@ TEST_F(ResolverConstantsTest, Vec3_ZeroInit_f32) {
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 0.0);
}
TEST_F(ResolverConstantsTest, Vec3_ZeroInit_f16) {
Enable(ast::Extension::kF16);
auto* expr = vec3<f16>();
WrapInFunction(expr);
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get(expr);
EXPECT_NE(sem, nullptr);
auto* vec = sem->Type()->As<sem::Vector>();
ASSERT_NE(vec, nullptr);
EXPECT_TRUE(vec->type()->Is<sem::F16>());
EXPECT_EQ(vec->Width(), 3u);
EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 0.0);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 0.0);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 0.0);
}
TEST_F(ResolverConstantsTest, Vec3_ZeroInit_bool) {
auto* expr = vec3<bool>();
WrapInFunction(expr);
@ -230,6 +268,28 @@ TEST_F(ResolverConstantsTest, Vec3_Splat_f32) {
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 9.9f);
}
TEST_F(ResolverConstantsTest, Vec3_Splat_f16) {
Enable(ast::Extension::kF16);
auto* expr = vec3<f16>(9.9_h);
WrapInFunction(expr);
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get(expr);
EXPECT_NE(sem, nullptr);
auto* vec = sem->Type()->As<sem::Vector>();
ASSERT_NE(vec, nullptr);
EXPECT_TRUE(vec->type()->Is<sem::F16>());
EXPECT_EQ(vec->Width(), 3u);
EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
// 9.9 is not exactly representable by f16, and should be quantized to 9.8984375
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 9.8984375f);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 9.8984375f);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 9.8984375f);
}
TEST_F(ResolverConstantsTest, Vec3_Splat_bool) {
auto* expr = vec3<bool>(true);
WrapInFunction(expr);
@ -310,6 +370,27 @@ TEST_F(ResolverConstantsTest, Vec3_FullConstruct_f32) {
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 3.f);
}
TEST_F(ResolverConstantsTest, Vec3_FullConstruct_f16) {
Enable(ast::Extension::kF16);
auto* expr = vec3<f16>(1_h, 2_h, 3_h);
WrapInFunction(expr);
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get(expr);
EXPECT_NE(sem, nullptr);
auto* vec = sem->Type()->As<sem::Vector>();
ASSERT_NE(vec, nullptr);
EXPECT_TRUE(vec->type()->Is<sem::F16>());
EXPECT_EQ(vec->Width(), 3u);
EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 1.f);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 2.f);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 3.f);
}
TEST_F(ResolverConstantsTest, Vec3_FullConstruct_bool) {
auto* expr = vec3<bool>(true, false, true);
WrapInFunction(expr);
@ -390,6 +471,27 @@ TEST_F(ResolverConstantsTest, Vec3_MixConstruct_f32) {
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 3.f);
}
TEST_F(ResolverConstantsTest, Vec3_MixConstruct_f16) {
Enable(ast::Extension::kF16);
auto* expr = vec3<f16>(1_h, vec2<f16>(2_h, 3_h));
WrapInFunction(expr);
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get(expr);
EXPECT_NE(sem, nullptr);
auto* vec = sem->Type()->As<sem::Vector>();
ASSERT_NE(vec, nullptr);
EXPECT_TRUE(vec->type()->Is<sem::F16>());
EXPECT_EQ(vec->Width(), 3u);
EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 1.f);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 2.f);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 3.f);
}
TEST_F(ResolverConstantsTest, Vec3_MixConstruct_bool) {
auto* expr = vec3<bool>(vec2<bool>(true, false), true);
WrapInFunction(expr);
@ -450,6 +552,48 @@ TEST_F(ResolverConstantsTest, Vec3_Convert_u32_to_f32) {
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 30.f);
}
TEST_F(ResolverConstantsTest, Vec3_Convert_f16_to_i32) {
Enable(ast::Extension::kF16);
auto* expr = vec3<i32>(vec3<f16>(1.1_h, 2.2_h, 3.3_h));
WrapInFunction(expr);
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get(expr);
EXPECT_NE(sem, nullptr);
auto* vec = sem->Type()->As<sem::Vector>();
ASSERT_NE(vec, nullptr);
EXPECT_TRUE(vec->type()->Is<sem::I32>());
EXPECT_EQ(vec->Width(), 3u);
EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 1);
EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, 2);
EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, 3);
}
TEST_F(ResolverConstantsTest, Vec3_Convert_u32_to_f16) {
Enable(ast::Extension::kF16);
auto* expr = vec3<f16>(vec3<u32>(10_u, 20_u, 30_u));
WrapInFunction(expr);
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get(expr);
EXPECT_NE(sem, nullptr);
auto* vec = sem->Type()->As<sem::Vector>();
ASSERT_NE(vec, nullptr);
EXPECT_TRUE(vec->type()->Is<sem::F16>());
EXPECT_EQ(vec->Width(), 3u);
EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 10.f);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 20.f);
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 30.f);
}
TEST_F(ResolverConstantsTest, Vec3_Convert_Large_f32_to_i32) {
auto* expr = vec3<i32>(vec3<f32>(1e10_f, -1e20_f, 1e30_f));
WrapInFunction(expr);
@ -490,11 +634,10 @@ TEST_F(ResolverConstantsTest, Vec3_Convert_Large_f32_to_u32) {
EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, u32::kHighest);
}
// TODO(crbug.com/tint/1502): Enable when f16 overloads are implemented
TEST_F(ResolverConstantsTest, DISABLED_Vec3_Convert_Large_f32_to_f16) {
TEST_F(ResolverConstantsTest, Vec3_Convert_Large_f32_to_f16) {
Enable(ast::Extension::kF16);
auto* expr = vec3<f16>(vec3<f32>(0.00001_f, -0.00002_f, 0.00003_f));
auto* expr = vec3<f16>(vec3<f32>(1e10_f, -1e20_f, 1e30_f));
WrapInFunction(expr);
EXPECT_TRUE(r()->Resolve()) << r()->error();
@ -515,11 +658,10 @@ TEST_F(ResolverConstantsTest, DISABLED_Vec3_Convert_Large_f32_to_f16) {
EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, kInf);
}
// TODO(crbug.com/tint/1502): Enable when f16 overloads are implemented
TEST_F(ResolverConstantsTest, DISABLED_Vec3_Convert_Small_f32_to_f16) {
TEST_F(ResolverConstantsTest, Vec3_Convert_Small_f32_to_f16) {
Enable(ast::Extension::kF16);
auto* expr = vec3<f16>(vec3<f32>(1e-10_f, -1e20_f, 1e30_f));
auto* expr = vec3<f16>(vec3<f32>(1e-20_f, -2e-30_f, 3e-40_f));
WrapInFunction(expr);
EXPECT_TRUE(r()->Resolve()) << r()->error();
@ -562,6 +704,32 @@ TEST_F(ResolverConstantsTest, Mat2x3_ZeroInit_f32) {
EXPECT_EQ(sem->ConstantValue().Element<f32>(5).value, 0._f);
}
TEST_F(ResolverConstantsTest, Mat2x3_ZeroInit_f16) {
Enable(ast::Extension::kF16);
auto* expr = mat2x3<f16>();
WrapInFunction(expr);
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get(expr);
EXPECT_NE(sem, nullptr);
auto* mat = sem->Type()->As<sem::Matrix>();
ASSERT_NE(mat, nullptr);
EXPECT_TRUE(mat->type()->Is<sem::F16>());
EXPECT_EQ(mat->columns(), 2u);
EXPECT_EQ(mat->rows(), 3u);
EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
ASSERT_EQ(sem->ConstantValue().ElementCount(), 6u);
EXPECT_EQ(sem->ConstantValue().Element<f16>(0).value, 0._h);
EXPECT_EQ(sem->ConstantValue().Element<f16>(1).value, 0._h);
EXPECT_EQ(sem->ConstantValue().Element<f16>(2).value, 0._h);
EXPECT_EQ(sem->ConstantValue().Element<f16>(3).value, 0._h);
EXPECT_EQ(sem->ConstantValue().Element<f16>(4).value, 0._h);
EXPECT_EQ(sem->ConstantValue().Element<f16>(5).value, 0._h);
}
TEST_F(ResolverConstantsTest, Mat3x2_Construct_Scalars_af) {
auto* expr = Construct(ty.mat(nullptr, 3, 2), 1.0_a, 2.0_a, 3.0_a, 4.0_a, 5.0_a, 6.0_a);
WrapInFunction(expr);

198
src/tint/resolver/struct_layout_test.cc
View File

@ -54,6 +54,57 @@ TEST_F(ResolverStructLayoutTest, Scalars) {
}
}
TEST_F(ResolverStructLayoutTest, ScalarsWithF16) {
Enable(ast::Extension::kF16);
auto* s = Structure("S", {
Member("a", ty.f32()),
Member("b", ty.f16()),
Member("c", ty.u32()),
Member("d", ty.f16()),
Member("e", ty.f16()),
Member("f", ty.i32()),
Member("g", ty.f16()),
});
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = TypeOf(s)->As<sem::Struct>();
ASSERT_NE(sem, nullptr);
EXPECT_EQ(sem->Size(), 24u);
EXPECT_EQ(sem->SizeNoPadding(), 22u);
EXPECT_EQ(sem->Align(), 4u);
ASSERT_EQ(sem->Members().size(), 7u);
// f32
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
EXPECT_EQ(sem->Members()[0]->Size(), 4u);
// f16
EXPECT_EQ(sem->Members()[1]->Offset(), 4u);
EXPECT_EQ(sem->Members()[1]->Align(), 2u);
EXPECT_EQ(sem->Members()[1]->Size(), 2u);
// u32
EXPECT_EQ(sem->Members()[2]->Offset(), 8u);
EXPECT_EQ(sem->Members()[2]->Align(), 4u);
EXPECT_EQ(sem->Members()[2]->Size(), 4u);
// f16
EXPECT_EQ(sem->Members()[3]->Offset(), 12u);
EXPECT_EQ(sem->Members()[3]->Align(), 2u);
EXPECT_EQ(sem->Members()[3]->Size(), 2u);
// f16
EXPECT_EQ(sem->Members()[4]->Offset(), 14u);
EXPECT_EQ(sem->Members()[4]->Align(), 2u);
EXPECT_EQ(sem->Members()[4]->Size(), 2u);
// i32
EXPECT_EQ(sem->Members()[5]->Offset(), 16u);
EXPECT_EQ(sem->Members()[5]->Align(), 4u);
EXPECT_EQ(sem->Members()[5]->Size(), 4u);
// f16
EXPECT_EQ(sem->Members()[6]->Offset(), 20u);
EXPECT_EQ(sem->Members()[6]->Align(), 2u);
EXPECT_EQ(sem->Members()[6]->Size(), 2u);
}
TEST_F(ResolverStructLayoutTest, Alias) {
auto* alias_a = Alias("a", ty.f32());
auto* alias_b = Alias("b", ty.f32());
@ -83,58 +134,80 @@ TEST_F(ResolverStructLayoutTest, Alias) {
}
TEST_F(ResolverStructLayoutTest, ImplicitStrideArrayStaticSize) {
Enable(ast::Extension::kF16);
auto* s = Structure("S", {
Member("a", ty.array<i32, 3>()),
Member("b", ty.array<f32, 5>()),
Member("c", ty.array<f32, 1>()),
Member("c", ty.array<f16, 7>()),
Member("d", ty.array<f32, 1>()),
});
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = TypeOf(s)->As<sem::Struct>();
ASSERT_NE(sem, nullptr);
EXPECT_EQ(sem->Size(), 36u);
EXPECT_EQ(sem->SizeNoPadding(), 36u);
EXPECT_EQ(sem->Size(), 52u);
EXPECT_EQ(sem->SizeNoPadding(), 52u);
EXPECT_EQ(sem->Align(), 4u);
ASSERT_EQ(sem->Members().size(), 3u);
ASSERT_EQ(sem->Members().size(), 4u);
// array<i32, 3>
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
EXPECT_EQ(sem->Members()[0]->Size(), 12u);
// array<f32, 5>
EXPECT_EQ(sem->Members()[1]->Offset(), 12u);
EXPECT_EQ(sem->Members()[1]->Align(), 4u);
EXPECT_EQ(sem->Members()[1]->Size(), 20u);
// array<f16, 7>
EXPECT_EQ(sem->Members()[2]->Offset(), 32u);
EXPECT_EQ(sem->Members()[2]->Align(), 4u);
EXPECT_EQ(sem->Members()[2]->Size(), 4u);
EXPECT_EQ(sem->Members()[2]->Align(), 2u);
EXPECT_EQ(sem->Members()[2]->Size(), 14u);
// array<f32, 1>
EXPECT_EQ(sem->Members()[3]->Offset(), 48u);
EXPECT_EQ(sem->Members()[3]->Align(), 4u);
EXPECT_EQ(sem->Members()[3]->Size(), 4u);
for (auto& m : sem->Members()) {
EXPECT_EQ(m->Struct()->Declaration(), s);
}
}
TEST_F(ResolverStructLayoutTest, ExplicitStrideArrayStaticSize) {
Enable(ast::Extension::kF16);
auto* s = Structure("S", {
Member("a", ty.array<i32, 3>(/*stride*/ 8)),
Member("b", ty.array<f32, 5>(/*stride*/ 16)),
Member("c", ty.array<f32, 1>(/*stride*/ 32)),
Member("c", ty.array<f16, 7>(/*stride*/ 4)),
Member("d", ty.array<f32, 1>(/*stride*/ 32)),
});
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = TypeOf(s)->As<sem::Struct>();
ASSERT_NE(sem, nullptr);
EXPECT_EQ(sem->Size(), 136u);
EXPECT_EQ(sem->SizeNoPadding(), 136u);
EXPECT_EQ(sem->Size(), 164u);
EXPECT_EQ(sem->SizeNoPadding(), 164u);
EXPECT_EQ(sem->Align(), 4u);
ASSERT_EQ(sem->Members().size(), 3u);
ASSERT_EQ(sem->Members().size(), 4u);
// array<i32, 3>, stride = 8
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
EXPECT_EQ(sem->Members()[0]->Size(), 24u);
// array<f32, 5>, stride = 16
EXPECT_EQ(sem->Members()[1]->Offset(), 24u);
EXPECT_EQ(sem->Members()[1]->Align(), 4u);
EXPECT_EQ(sem->Members()[1]->Size(), 80u);
// array<f16, 7>, stride = 4
EXPECT_EQ(sem->Members()[2]->Offset(), 104u);
EXPECT_EQ(sem->Members()[2]->Align(), 4u);
EXPECT_EQ(sem->Members()[2]->Size(), 32u);
EXPECT_EQ(sem->Members()[2]->Align(), 2u);
EXPECT_EQ(sem->Members()[2]->Size(), 28u);
// array<f32, 1>, stride = 32
EXPECT_EQ(sem->Members()[3]->Offset(), 132u);
EXPECT_EQ(sem->Members()[3]->Align(), 4u);
EXPECT_EQ(sem->Members()[3]->Size(), 32u);
for (auto& m : sem->Members()) {
EXPECT_EQ(m->Struct()->Declaration(), s);
}
@ -262,53 +335,92 @@ TEST_F(ResolverStructLayoutTest, Vector) {
}
TEST_F(ResolverStructLayoutTest, Matrix) {
Enable(ast::Extension::kF16);
auto* s = Structure("S", {
Member("a", ty.mat2x2<f32>()),
Member("b", ty.mat2x3<f32>()),
Member("c", ty.mat2x4<f32>()),
Member("d", ty.mat3x2<f32>()),
Member("e", ty.mat3x3<f32>()),
Member("f", ty.mat3x4<f32>()),
Member("g", ty.mat4x2<f32>()),
Member("h", ty.mat4x3<f32>()),
Member("i", ty.mat4x4<f32>()),
Member("a_1", ty.mat2x2<f32>()),
Member("a_2", ty.mat2x2<f16>()),
Member("b_1", ty.mat2x3<f32>()),
Member("b_2", ty.mat2x3<f16>()),
Member("c_1", ty.mat2x4<f32>()),
Member("c_2", ty.mat2x4<f16>()),
Member("d_1", ty.mat3x2<f32>()),
Member("d_2", ty.mat3x2<f16>()),
Member("e_1", ty.mat3x3<f32>()),
Member("e_2", ty.mat3x3<f16>()),
Member("f_1", ty.mat3x4<f32>()),
Member("f_2", ty.mat3x4<f16>()),
Member("g_1", ty.mat4x2<f32>()),
Member("g_2", ty.mat4x2<f16>()),
Member("h_1", ty.mat4x3<f32>()),
Member("h_2", ty.mat4x3<f16>()),
Member("i_1", ty.mat4x4<f32>()),
Member("i_2", ty.mat4x4<f16>()),
});
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = TypeOf(s)->As<sem::Struct>();
ASSERT_NE(sem, nullptr);
EXPECT_EQ(sem->Size(), 368u);
EXPECT_EQ(sem->SizeNoPadding(), 368u);
EXPECT_EQ(sem->Size(), 576u);
EXPECT_EQ(sem->SizeNoPadding(), 576u);
EXPECT_EQ(sem->Align(), 16u);
ASSERT_EQ(sem->Members().size(), 9u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u); // mat2x2
ASSERT_EQ(sem->Members().size(), 18u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u); // mat2x2<f32>
EXPECT_EQ(sem->Members()[0]->Align(), 8u);
EXPECT_EQ(sem->Members()[0]->Size(), 16u);
EXPECT_EQ(sem->Members()[1]->Offset(), 16u); // mat2x3
EXPECT_EQ(sem->Members()[1]->Align(), 16u);
EXPECT_EQ(sem->Members()[1]->Size(), 32u);
EXPECT_EQ(sem->Members()[2]->Offset(), 48u); // mat2x4
EXPECT_EQ(sem->Members()[1]->Offset(), 16u); // mat2x2<f16>
EXPECT_EQ(sem->Members()[1]->Align(), 4u);
EXPECT_EQ(sem->Members()[1]->Size(), 8u);
EXPECT_EQ(sem->Members()[2]->Offset(), 32u); // mat2x3<f32>
EXPECT_EQ(sem->Members()[2]->Align(), 16u);
EXPECT_EQ(sem->Members()[2]->Size(), 32u);
EXPECT_EQ(sem->Members()[3]->Offset(), 80u); // mat3x2
EXPECT_EQ(sem->Members()[3]->Offset(), 64u); // mat2x3<f16>
EXPECT_EQ(sem->Members()[3]->Align(), 8u);
EXPECT_EQ(sem->Members()[3]->Size(), 24u);
EXPECT_EQ(sem->Members()[4]->Offset(), 112u); // mat3x3
EXPECT_EQ(sem->Members()[3]->Size(), 16u);
EXPECT_EQ(sem->Members()[4]->Offset(), 80u); // mat2x4<f32>
EXPECT_EQ(sem->Members()[4]->Align(), 16u);
EXPECT_EQ(sem->Members()[4]->Size(), 48u);
EXPECT_EQ(sem->Members()[5]->Offset(), 160u); // mat3x4
EXPECT_EQ(sem->Members()[5]->Align(), 16u);
EXPECT_EQ(sem->Members()[5]->Size(), 48u);
EXPECT_EQ(sem->Members()[6]->Offset(), 208u); // mat4x2
EXPECT_EQ(sem->Members()[4]->Size(), 32u);
EXPECT_EQ(sem->Members()[5]->Offset(), 112u); // mat2x4<f16>
EXPECT_EQ(sem->Members()[5]->Align(), 8u);
EXPECT_EQ(sem->Members()[5]->Size(), 16u);
EXPECT_EQ(sem->Members()[6]->Offset(), 128u); // mat3x2<f32>
EXPECT_EQ(sem->Members()[6]->Align(), 8u);
EXPECT_EQ(sem->Members()[6]->Size(), 32u);
EXPECT_EQ(sem->Members()[7]->Offset(), 240u); // mat4x3
EXPECT_EQ(sem->Members()[7]->Align(), 16u);
EXPECT_EQ(sem->Members()[7]->Size(), 64u);
EXPECT_EQ(sem->Members()[8]->Offset(), 304u); // mat4x4
EXPECT_EQ(sem->Members()[6]->Size(), 24u);
EXPECT_EQ(sem->Members()[7]->Offset(), 152u); // mat3x2<f16>
EXPECT_EQ(sem->Members()[7]->Align(), 4u);
EXPECT_EQ(sem->Members()[7]->Size(), 12u);
EXPECT_EQ(sem->Members()[8]->Offset(), 176u); // mat3x3<f32>
EXPECT_EQ(sem->Members()[8]->Align(), 16u);
EXPECT_EQ(sem->Members()[8]->Size(), 64u);
EXPECT_EQ(sem->Members()[8]->Size(), 48u);
EXPECT_EQ(sem->Members()[9]->Offset(), 224u); // mat3x3<f16>
EXPECT_EQ(sem->Members()[9]->Align(), 8u);
EXPECT_EQ(sem->Members()[9]->Size(), 24u);
EXPECT_EQ(sem->Members()[10]->Offset(), 256u); // mat3x4<f32>
EXPECT_EQ(sem->Members()[10]->Align(), 16u);
EXPECT_EQ(sem->Members()[10]->Size(), 48u);
EXPECT_EQ(sem->Members()[11]->Offset(), 304u); // mat3x4<f16>
EXPECT_EQ(sem->Members()[11]->Align(), 8u);
EXPECT_EQ(sem->Members()[11]->Size(), 24u);
EXPECT_EQ(sem->Members()[12]->Offset(), 328u); // mat4x2<f32>
EXPECT_EQ(sem->Members()[12]->Align(), 8u);
EXPECT_EQ(sem->Members()[12]->Size(), 32u);
EXPECT_EQ(sem->Members()[13]->Offset(), 360u); // mat4x2<f16>
EXPECT_EQ(sem->Members()[13]->Align(), 4u);
EXPECT_EQ(sem->Members()[13]->Size(), 16u);
EXPECT_EQ(sem->Members()[14]->Offset(), 384u); // mat4x3<f32>
EXPECT_EQ(sem->Members()[14]->Align(), 16u);
EXPECT_EQ(sem->Members()[14]->Size(), 64u);
EXPECT_EQ(sem->Members()[15]->Offset(), 448u); // mat4x3<f16>
EXPECT_EQ(sem->Members()[15]->Align(), 8u);
EXPECT_EQ(sem->Members()[15]->Size(), 32u);
EXPECT_EQ(sem->Members()[16]->Offset(), 480u); // mat4x4<f32>
EXPECT_EQ(sem->Members()[16]->Align(), 16u);
EXPECT_EQ(sem->Members()[16]->Size(), 64u);
EXPECT_EQ(sem->Members()[17]->Offset(), 544u); // mat4x4<f16>
EXPECT_EQ(sem->Members()[17]->Align(), 8u);
EXPECT_EQ(sem->Members()[17]->Size(), 32u);
for (auto& m : sem->Members()) {
EXPECT_EQ(m->Struct()->Declaration(), s);
}

520
src/tint/resolver/type_constructor_validation_test.cc
View File

File diff suppressed because it is too large Load Diff

44
src/tint/resolver/type_validation_test.cc
View File

@ -1120,6 +1120,9 @@ using ValidMatrixTypes = ResolverTestWithParam<Params>;
TEST_P(ValidMatrixTypes, Okay) {
// var a : matNxM<EL_TY>;
auto& params = GetParam();
Enable(ast::Extension::kF16);
GlobalVar("a", ty.mat(params.elem_ty(*this), params.columns, params.rows),
ast::StorageClass::kPrivate);
EXPECT_TRUE(r()->Resolve()) << r()->error();
@ -1137,16 +1140,31 @@ INSTANTIATE_TEST_SUITE_P(ResolverTypeValidationTest,
ParamsFor<f32>(4, 4),
ParamsFor<alias<f32>>(4, 2),
ParamsFor<alias<f32>>(4, 3),
ParamsFor<alias<f32>>(4, 4)));
ParamsFor<alias<f32>>(4, 4),
ParamsFor<f16>(2, 2),
ParamsFor<f16>(2, 3),
ParamsFor<f16>(2, 4),
ParamsFor<f16>(3, 2),
ParamsFor<f16>(3, 3),
ParamsFor<f16>(3, 4),
ParamsFor<f16>(4, 2),
ParamsFor<f16>(4, 3),
ParamsFor<f16>(4, 4),
ParamsFor<alias<f16>>(4, 2),
ParamsFor<alias<f16>>(4, 3),
ParamsFor<alias<f16>>(4, 4)));
using InvalidMatrixElementTypes = ResolverTestWithParam<Params>;
TEST_P(InvalidMatrixElementTypes, InvalidElementType) {
// var a : matNxM<EL_TY>;
auto& params = GetParam();
Enable(ast::Extension::kF16);
GlobalVar("a", ty.mat(Source{{12, 34}}, params.elem_ty(*this), params.columns, params.rows),
ast::StorageClass::kPrivate);
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), "12:34 error: matrix element type must be 'f32'");
EXPECT_EQ(r()->error(), "12:34 error: matrix element type must be 'f32' or 'f16'");
}
INSTANTIATE_TEST_SUITE_P(ResolverTypeValidationTest,
InvalidMatrixElementTypes,
@ -1154,12 +1172,17 @@ INSTANTIATE_TEST_SUITE_P(ResolverTypeValidationTest,
ParamsFor<i32>(4, 3),
ParamsFor<u32>(4, 4),
ParamsFor<vec2<f32>>(2, 2),
ParamsFor<vec2<f16>>(2, 2),
ParamsFor<vec3<i32>>(2, 3),
ParamsFor<vec4<u32>>(2, 4),
ParamsFor<mat2x2<f32>>(3, 2),
ParamsFor<mat3x3<f32>>(3, 3),
ParamsFor<mat4x4<f32>>(3, 4),
ParamsFor<array<2, f32>>(4, 2)));
ParamsFor<mat2x2<f16>>(3, 2),
ParamsFor<mat3x3<f16>>(3, 3),
ParamsFor<mat4x4<f16>>(3, 4),
ParamsFor<array<2, f32>>(4, 2),
ParamsFor<array<2, f16>>(4, 2)));
} // namespace MatrixTests
namespace VectorTests {
@ -1177,6 +1200,9 @@ using ValidVectorTypes = ResolverTestWithParam<Params>;
TEST_P(ValidVectorTypes, Okay) {
// var a : vecN<EL_TY>;
auto& params = GetParam();
Enable(ast::Extension::kF16);
GlobalVar("a", ty.vec(params.elem_ty(*this), params.width), ast::StorageClass::kPrivate);
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
@ -1184,18 +1210,22 @@ INSTANTIATE_TEST_SUITE_P(ResolverTypeValidationTest,
ValidVectorTypes,
testing::Values(ParamsFor<bool>(2),
ParamsFor<f32>(2),
ParamsFor<f16>(2),
ParamsFor<i32>(2),
ParamsFor<u32>(2),
ParamsFor<bool>(3),
ParamsFor<f32>(3),
ParamsFor<f16>(3),
ParamsFor<i32>(3),
ParamsFor<u32>(3),
ParamsFor<bool>(4),
ParamsFor<f32>(4),
ParamsFor<f16>(4),
ParamsFor<i32>(4),
ParamsFor<u32>(4),
ParamsFor<alias<bool>>(4),
ParamsFor<alias<f32>>(4),
ParamsFor<alias<f16>>(4),
ParamsFor<alias<i32>>(4),
ParamsFor<alias<u32>>(4)));
@ -1203,11 +1233,15 @@ using InvalidVectorElementTypes = ResolverTestWithParam<Params>;
TEST_P(InvalidVectorElementTypes, InvalidElementType) {
// var a : vecN<EL_TY>;
auto& params = GetParam();
Enable(ast::Extension::kF16);
GlobalVar("a", ty.vec(Source{{12, 34}}, params.elem_ty(*this), params.width),
ast::StorageClass::kPrivate);
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"12:34 error: vector element type must be 'bool', 'f32', 'i32' or 'u32'");
"12:34 error: vector element type must be 'bool', 'f32', 'f16', 'i32' "
"or 'u32'");
}
INSTANTIATE_TEST_SUITE_P(ResolverTypeValidationTest,
InvalidVectorElementTypes,
@ -1215,7 +1249,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverTypeValidationTest,
ParamsFor<vec3<i32>>(2),
ParamsFor<vec4<u32>>(2),
ParamsFor<mat2x2<f32>>(2),
ParamsFor<mat3x3<f32>>(2),
ParamsFor<mat3x3<f16>>(2),
ParamsFor<mat4x4<f32>>(2),
ParamsFor<array<2, f32>>(2)));
} // namespace VectorTests

4
src/tint/resolver/validator.cc
View File

@ -1865,7 +1865,7 @@ bool Validator::ArrayConstructor(const ast::CallExpression* ctor,
bool Validator::Vector(const sem::Vector* ty, const Source& source) const {
if (!ty->type()->is_scalar()) {
AddError("vector element type must be 'bool', 'f32', 'i32' or 'u32'", source);
AddError("vector element type must be 'bool', 'f32', 'f16', 'i32' or 'u32'", source);
return false;
}
return true;
@ -1873,7 +1873,7 @@ bool Validator::Vector(const sem::Vector* ty, const Source& source) const {
bool Validator::Matrix(const sem::Matrix* ty, const Source& source) const {
if (!ty->is_float_matrix()) {
AddError("matrix element type must be 'f32'", source);
AddError("matrix element type must be 'f32' or 'f16'", source);
return false;
}
return true;

40
src/tint/resolver/validator_is_storeable_test.cc
View File

@ -32,6 +32,7 @@ TEST_F(ValidatorIsStorableTest, Scalar) {
EXPECT_TRUE(v()->IsStorable(create<sem::I32>()));
EXPECT_TRUE(v()->IsStorable(create<sem::U32>()));
EXPECT_TRUE(v()->IsStorable(create<sem::F32>()));
EXPECT_TRUE(v()->IsStorable(create<sem::F16>()));
}
TEST_F(ValidatorIsStorableTest, Vector) {
@ -44,21 +45,36 @@ TEST_F(ValidatorIsStorableTest, Vector) {
EXPECT_TRUE(v()->IsStorable(create<sem::Vector>(create<sem::F32>(), 2u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Vector>(create<sem::F32>(), 3u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Vector>(create<sem::F32>(), 4u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Vector>(create<sem::F16>(), 2u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Vector>(create<sem::F16>(), 3u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Vector>(create<sem::F16>(), 4u)));
}
TEST_F(ValidatorIsStorableTest, Matrix) {
auto* vec2 = create<sem::Vector>(create<sem::F32>(), 2u);
auto* vec3 = create<sem::Vector>(create<sem::F32>(), 3u);
auto* vec4 = create<sem::Vector>(create<sem::F32>(), 4u);
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec2, 2u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec2, 3u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec2, 4u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec3, 2u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec3, 3u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec3, 4u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec4, 2u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec4, 3u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec4, 4u)));
auto* vec2_f32 = create<sem::Vector>(create<sem::F32>(), 2u);
auto* vec3_f32 = create<sem::Vector>(create<sem::F32>(), 3u);
auto* vec4_f32 = create<sem::Vector>(create<sem::F32>(), 4u);
auto* vec2_f16 = create<sem::Vector>(create<sem::F16>(), 2u);
auto* vec3_f16 = create<sem::Vector>(create<sem::F16>(), 3u);
auto* vec4_f16 = create<sem::Vector>(create<sem::F16>(), 4u);
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec2_f32, 2u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec2_f32, 3u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec2_f32, 4u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec3_f32, 2u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec3_f32, 3u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec3_f32, 4u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec4_f32, 2u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec4_f32, 3u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec4_f32, 4u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec2_f16, 2u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec2_f16, 3u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec2_f16, 4u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec3_f16, 2u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec3_f16, 3u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec3_f16, 4u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec4_f16, 2u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec4_f16, 3u)));
EXPECT_TRUE(v()->IsStorable(create<sem::Matrix>(vec4_f16, 4u)));
}
TEST_F(ValidatorIsStorableTest, Pointer) {

27
src/tint/resolver/variable_test.cc
View File

@ -35,17 +35,21 @@ TEST_F(ResolverVariableTest, LocalVar_NoConstructor) {
// var i : i32;
// var u : u32;
// var f : f32;
// var h : f16;
// var b : bool;
// var s : S;
// var a : A;
// }
Enable(ast::Extension::kF16);
auto* S = Structure("S", {Member("i", ty.i32())});
auto* A = Alias("A", ty.Of(S));
auto* i = Var("i", ty.i32(), ast::StorageClass::kNone);
auto* u = Var("u", ty.u32(), ast::StorageClass::kNone);
auto* f = Var("f", ty.f32(), ast::StorageClass::kNone);
auto* h = Var("h", ty.f16(), ast::StorageClass::kNone);
auto* b = Var("b", ty.bool_(), ast::StorageClass::kNone);
auto* s = Var("s", ty.Of(S), ast::StorageClass::kNone);
auto* a = Var("a", ty.Of(A), ast::StorageClass::kNone);
@ -55,6 +59,7 @@ TEST_F(ResolverVariableTest, LocalVar_NoConstructor) {
Decl(i),
Decl(u),
Decl(f),
Decl(h),
Decl(b),
Decl(s),
Decl(a),
@ -66,6 +71,7 @@ TEST_F(ResolverVariableTest, LocalVar_NoConstructor) {
ASSERT_TRUE(TypeOf(i)->Is<sem::Reference>());
ASSERT_TRUE(TypeOf(u)->Is<sem::Reference>());
ASSERT_TRUE(TypeOf(f)->Is<sem::Reference>());
ASSERT_TRUE(TypeOf(h)->Is<sem::Reference>());
ASSERT_TRUE(TypeOf(b)->Is<sem::Reference>());
ASSERT_TRUE(TypeOf(s)->Is<sem::Reference>());
ASSERT_TRUE(TypeOf(a)->Is<sem::Reference>());
@ -73,6 +79,7 @@ TEST_F(ResolverVariableTest, LocalVar_NoConstructor) {
EXPECT_TRUE(TypeOf(i)->As<sem::Reference>()->StoreType()->Is<sem::I32>());
EXPECT_TRUE(TypeOf(u)->As<sem::Reference>()->StoreType()->Is<sem::U32>());
EXPECT_TRUE(TypeOf(f)->As<sem::Reference>()->StoreType()->Is<sem::F32>());
EXPECT_TRUE(TypeOf(h)->As<sem::Reference>()->StoreType()->Is<sem::F16>());
EXPECT_TRUE(TypeOf(b)->As<sem::Reference>()->StoreType()->Is<sem::Bool>());
EXPECT_TRUE(TypeOf(s)->As<sem::Reference>()->StoreType()->Is<sem::Struct>());
EXPECT_TRUE(TypeOf(a)->As<sem::Reference>()->StoreType()->Is<sem::Struct>());
@ -80,6 +87,7 @@ TEST_F(ResolverVariableTest, LocalVar_NoConstructor) {
EXPECT_EQ(Sem().Get(i)->Constructor(), nullptr);
EXPECT_EQ(Sem().Get(u)->Constructor(), nullptr);
EXPECT_EQ(Sem().Get(f)->Constructor(), nullptr);
EXPECT_EQ(Sem().Get(h)->Constructor(), nullptr);
EXPECT_EQ(Sem().Get(b)->Constructor(), nullptr);
EXPECT_EQ(Sem().Get(s)->Constructor(), nullptr);
EXPECT_EQ(Sem().Get(a)->Constructor(), nullptr);
@ -92,17 +100,21 @@ TEST_F(ResolverVariableTest, LocalVar_WithConstructor) {
// var i : i32 = 1i;
// var u : u32 = 1u;
// var f : f32 = 1.f;
// var h : f16 = 1.h;
// var b : bool = true;
// var s : S = S(1);
// var a : A = A(1);
// }
Enable(ast::Extension::kF16);
auto* S = Structure("S", {Member("i", ty.i32())});
auto* A = Alias("A", ty.Of(S));
auto* i_c = Expr(1_i);
auto* u_c = Expr(1_u);
auto* f_c = Expr(1_f);
auto* h_c = Expr(1_h);
auto* b_c = Expr(true);
auto* s_c = Construct(ty.Of(S), Expr(1_i));
auto* a_c = Construct(ty.Of(A), Expr(1_i));
@ -110,6 +122,7 @@ TEST_F(ResolverVariableTest, LocalVar_WithConstructor) {
auto* i = Var("i", ty.i32(), ast::StorageClass::kNone, i_c);
auto* u = Var("u", ty.u32(), ast::StorageClass::kNone, u_c);
auto* f = Var("f", ty.f32(), ast::StorageClass::kNone, f_c);
auto* h = Var("h", ty.f16(), ast::StorageClass::kNone, h_c);
auto* b = Var("b", ty.bool_(), ast::StorageClass::kNone, b_c);
auto* s = Var("s", ty.Of(S), ast::StorageClass::kNone, s_c);
auto* a = Var("a", ty.Of(A), ast::StorageClass::kNone, a_c);
@ -119,6 +132,7 @@ TEST_F(ResolverVariableTest, LocalVar_WithConstructor) {
Decl(i),
Decl(u),
Decl(f),
Decl(h),
Decl(b),
Decl(s),
Decl(a),
@ -130,6 +144,7 @@ TEST_F(ResolverVariableTest, LocalVar_WithConstructor) {
ASSERT_TRUE(TypeOf(i)->Is<sem::Reference>());
ASSERT_TRUE(TypeOf(u)->Is<sem::Reference>());
ASSERT_TRUE(TypeOf(f)->Is<sem::Reference>());
ASSERT_TRUE(TypeOf(h)->Is<sem::Reference>());
ASSERT_TRUE(TypeOf(b)->Is<sem::Reference>());
ASSERT_TRUE(TypeOf(s)->Is<sem::Reference>());
ASSERT_TRUE(TypeOf(a)->Is<sem::Reference>());
@ -144,6 +159,7 @@ TEST_F(ResolverVariableTest, LocalVar_WithConstructor) {
EXPECT_TRUE(TypeOf(i)->As<sem::Reference>()->StoreType()->Is<sem::I32>());
EXPECT_TRUE(TypeOf(u)->As<sem::Reference>()->StoreType()->Is<sem::U32>());
EXPECT_TRUE(TypeOf(f)->As<sem::Reference>()->StoreType()->Is<sem::F32>());
EXPECT_TRUE(TypeOf(h)->As<sem::Reference>()->StoreType()->Is<sem::F16>());
EXPECT_TRUE(TypeOf(b)->As<sem::Reference>()->StoreType()->Is<sem::Bool>());
EXPECT_TRUE(TypeOf(s)->As<sem::Reference>()->StoreType()->Is<sem::Struct>());
EXPECT_TRUE(TypeOf(a)->As<sem::Reference>()->StoreType()->Is<sem::Struct>());
@ -151,6 +167,7 @@ TEST_F(ResolverVariableTest, LocalVar_WithConstructor) {
EXPECT_EQ(Sem().Get(i)->Constructor()->Declaration(), i_c);
EXPECT_EQ(Sem().Get(u)->Constructor()->Declaration(), u_c);
EXPECT_EQ(Sem().Get(f)->Constructor()->Declaration(), f_c);
EXPECT_EQ(Sem().Get(h)->Constructor()->Declaration(), h_c);
EXPECT_EQ(Sem().Get(b)->Constructor()->Declaration(), b_c);
EXPECT_EQ(Sem().Get(s)->Constructor()->Declaration(), s_c);
EXPECT_EQ(Sem().Get(a)->Constructor()->Declaration(), a_c);
@ -396,13 +413,16 @@ TEST_F(ResolverVariableTest, LocalLet) {
// var v : i32;
// let i : i32 = 1i;
// let u : u32 = 1u;
// let f : f32 = 1.;
// let f : f32 = 1.f;
// let h : h32 = 1.h;
// let b : bool = true;
// let s : S = S(1);
// let a : A = A(1);
// let p : pointer<function, i32> = &v;
// }
Enable(ast::Extension::kF16);
auto* S = Structure("S", {Member("i", ty.i32())});
auto* A = Alias("A", ty.Of(S));
auto* v = Var("v", ty.i32(), ast::StorageClass::kNone);
@ -410,6 +430,7 @@ TEST_F(ResolverVariableTest, LocalLet) {
auto* i_c = Expr(1_i);
auto* u_c = Expr(1_u);
auto* f_c = Expr(1_f);
auto* h_c = Expr(1_h);
auto* b_c = Expr(true);
auto* s_c = Construct(ty.Of(S), Expr(1_i));
auto* a_c = Construct(ty.Of(A), Expr(1_i));
@ -418,6 +439,7 @@ TEST_F(ResolverVariableTest, LocalLet) {
auto* i = Let("i", ty.i32(), i_c);
auto* u = Let("u", ty.u32(), u_c);
auto* f = Let("f", ty.f32(), f_c);
auto* h = Let("h", ty.f16(), h_c);
auto* b = Let("b", ty.bool_(), b_c);
auto* s = Let("s", ty.Of(S), s_c);
auto* a = Let("a", ty.Of(A), a_c);
@ -429,6 +451,7 @@ TEST_F(ResolverVariableTest, LocalLet) {
Decl(i),
Decl(u),
Decl(f),
Decl(h),
Decl(b),
Decl(s),
Decl(a),
@ -441,6 +464,7 @@ TEST_F(ResolverVariableTest, LocalLet) {
ASSERT_TRUE(TypeOf(i)->Is<sem::I32>());
ASSERT_TRUE(TypeOf(u)->Is<sem::U32>());
ASSERT_TRUE(TypeOf(f)->Is<sem::F32>());
ASSERT_TRUE(TypeOf(h)->Is<sem::F16>());
ASSERT_TRUE(TypeOf(b)->Is<sem::Bool>());
ASSERT_TRUE(TypeOf(s)->Is<sem::Struct>());
ASSERT_TRUE(TypeOf(a)->Is<sem::Struct>());
@ -450,6 +474,7 @@ TEST_F(ResolverVariableTest, LocalLet) {
EXPECT_EQ(Sem().Get(i)->Constructor()->Declaration(), i_c);
EXPECT_EQ(Sem().Get(u)->Constructor()->Declaration(), u_c);
EXPECT_EQ(Sem().Get(f)->Constructor()->Declaration(), f_c);
EXPECT_EQ(Sem().Get(h)->Constructor()->Declaration(), h_c);
EXPECT_EQ(Sem().Get(b)->Constructor()->Declaration(), b_c);
EXPECT_EQ(Sem().Get(s)->Constructor()->Declaration(), s_c);
EXPECT_EQ(Sem().Get(a)->Constructor()->Declaration(), a_c);

49
test/tint/builtins/gen/literal/arrayLength/8421b9.wgsl Normal file
View File

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

1
test/tint/builtins/gen/literal/arrayLength/8421b9.wgsl.expected.glsl Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/literal/arrayLength/8421b9.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/literal/arrayLength/8421b9.wgsl.expected.msl Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/literal/arrayLength/8421b9.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/literal/arrayLength/8421b9.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1 @@
SKIP

49
test/tint/builtins/gen/literal/arrayLength/cbd6b5.wgsl Normal file
View File

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

1
test/tint/builtins/gen/literal/arrayLength/cbd6b5.wgsl.expected.glsl Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/literal/arrayLength/cbd6b5.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/literal/arrayLength/cbd6b5.wgsl.expected.msl Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/literal/arrayLength/cbd6b5.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/literal/arrayLength/cbd6b5.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1 @@
SKIP

49
test/tint/builtins/gen/var/arrayLength/8421b9.wgsl Normal file
View File

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

1
test/tint/builtins/gen/var/arrayLength/8421b9.wgsl.expected.glsl Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/var/arrayLength/8421b9.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/var/arrayLength/8421b9.wgsl.expected.msl Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/var/arrayLength/8421b9.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/var/arrayLength/8421b9.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1 @@
SKIP

49
test/tint/builtins/gen/var/arrayLength/cbd6b5.wgsl Normal file
View File

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

1
test/tint/builtins/gen/var/arrayLength/cbd6b5.wgsl.expected.glsl Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/var/arrayLength/cbd6b5.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/var/arrayLength/cbd6b5.wgsl.expected.msl Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/var/arrayLength/cbd6b5.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1 @@
SKIP

1
test/tint/builtins/gen/var/arrayLength/cbd6b5.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1 @@
SKIP