blob: 3bae4eb138fd4740b2c056ce4f32852e48704529 [file] [log] [blame]
// Copyright 2021 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
// list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include <unordered_set>
#include "src/tint/lang/wgsl/ast/builtin_texture_helper_test.h"
#include "src/tint/lang/wgsl/resolver/resolver_helper_test.h"
#include "src/tint/lang/wgsl/sem/value_constructor.h"
#include "src/tint/utils/text/string_stream.h"
namespace tint::resolver {
namespace {
using namespace tint::core::fluent_types; // NOLINT
using namespace tint::core::number_suffixes; // NOLINT
using ResolverBuiltinValidationTest = ResolverTest;
TEST_F(ResolverBuiltinValidationTest, FunctionTypeMustMatchReturnStatementType_void_fail) {
// fn func { return workgroupBarrier(); }
Func("func", tint::Empty, ty.void_(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "workgroupBarrier")),
});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"12:34 error: builtin function 'workgroupBarrier' does not return a value");
}
TEST_F(ResolverBuiltinValidationTest, InvalidPipelineStageDirect) {
// @compute @workgroup_size(1) fn func { return dpdx(1.0); }
auto* dpdx = Call(Source{{3, 4}}, "dpdx", 1_f);
Func(Source{{1, 2}}, "func", tint::Empty, ty.void_(),
Vector{
Assign(Phony(), dpdx),
},
Vector{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1_i),
});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), "3:4 error: built-in cannot be used by compute pipeline stage");
}
TEST_F(ResolverBuiltinValidationTest, InvalidPipelineStageIndirect) {
// fn f0 { return dpdx(1.0); }
// fn f1 { f0(); }
// fn f2 { f1(); }
// @compute @workgroup_size(1) fn main { return f2(); }
auto* dpdx = Call(Source{{3, 4}}, "dpdx", 1_f);
Func(Source{{1, 2}}, "f0", tint::Empty, ty.void_(),
Vector{
Assign(Phony(), dpdx),
});
Func(Source{{3, 4}}, "f1", tint::Empty, ty.void_(),
Vector{
CallStmt(Call("f0")),
});
Func(Source{{5, 6}}, "f2", tint::Empty, ty.void_(),
Vector{
CallStmt(Call("f1")),
});
Func(Source{{7, 8}}, "main", tint::Empty, ty.void_(),
Vector{
CallStmt(Call("f2")),
},
Vector{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1_i),
});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
R"(3:4 error: built-in cannot be used by compute pipeline stage
1:2 note: called by function 'f0'
3:4 note: called by function 'f1'
5:6 note: called by function 'f2'
7:8 note: called by entry point 'main')");
}
TEST_F(ResolverBuiltinValidationTest, BuiltinRedeclaredAsFunctionUsedAsFunction) {
auto* mix = Func(Source{{12, 34}}, "mix", tint::Empty, ty.i32(),
Vector{
Return(1_i),
});
auto* use = Call("mix");
WrapInFunction(use);
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get<sem::Call>(use);
ASSERT_NE(sem, nullptr);
EXPECT_EQ(sem->Target(), Sem().Get(mix));
}
TEST_F(ResolverBuiltinValidationTest, BuiltinRedeclaredAsFunctionUsedAsVariable) {
Func(Source{{12, 34}}, "mix", tint::Empty, ty.i32(),
Vector{
Return(1_i),
});
WrapInFunction(Decl(Var("v", Expr(Source{{56, 78}}, "mix"))));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), R"(56:78 error: cannot use function 'mix' as value
12:34 note: function 'mix' declared here
56:78 note: are you missing '()'?)");
}
TEST_F(ResolverBuiltinValidationTest, BuiltinRedeclaredAsGlobalConstUsedAsVariable) {
auto* mix = GlobalConst(Source{{12, 34}}, "mix", ty.i32(), Expr(1_i));
auto* use = Expr("mix");
WrapInFunction(Decl(Var("v", use)));
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get<sem::VariableUser>(use);
ASSERT_NE(sem, nullptr);
EXPECT_EQ(sem->Variable(), Sem().Get(mix));
}
TEST_F(ResolverBuiltinValidationTest, BuiltinRedeclaredAsGlobalVarUsedAsVariable) {
auto* mix =
GlobalVar(Source{{12, 34}}, "mix", ty.i32(), Expr(1_i), core::AddressSpace::kPrivate);
auto* use = Expr("mix");
WrapInFunction(Decl(Var("v", use)));
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().GetVal(use)->UnwrapLoad()->As<sem::VariableUser>();
ASSERT_NE(sem, nullptr);
EXPECT_EQ(sem->Variable(), Sem().Get(mix));
}
TEST_F(ResolverBuiltinValidationTest, BuiltinRedeclaredAsAliasUsedAsFunction) {
Alias(Source{{12, 34}}, "mix", ty.i32());
WrapInFunction(Call(Source{{56, 78}}, "mix", 1_f, 2_f, 3_f));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), R"(56:78 error: no matching constructor for i32(f32, f32, f32)
2 candidate constructors:
i32(i32) -> i32
i32() -> i32
1 candidate conversion:
i32<T>(T) -> i32 where: T is abstract-int, abstract-float, f32, f16, u32 or bool
)");
}
TEST_F(ResolverBuiltinValidationTest, BuiltinRedeclaredAsAliasUsedAsType) {
auto* mix = Alias(Source{{12, 34}}, "mix", ty.i32());
auto* use = Call("mix");
WrapInFunction(use);
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get<sem::Call>(use);
ASSERT_NE(sem, nullptr);
EXPECT_EQ(sem->Type(), Sem().Get(mix));
}
TEST_F(ResolverBuiltinValidationTest, BuiltinRedeclaredAsStructUsedAsFunction) {
Structure("mix", Vector{
Member("m", ty.i32()),
});
WrapInFunction(Call(Source{{12, 34}}, "mix", 1_f, 2_f, 3_f));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
R"(12:34 error: structure constructor has too many inputs: expected 1, found 3)");
}
TEST_F(ResolverBuiltinValidationTest, BuiltinRedeclaredAsStructUsedAsType) {
auto* mix = Structure("mix", Vector{
Member("m", ty.i32()),
});
auto* use = Call("mix");
WrapInFunction(use);
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* sem = Sem().Get<sem::Call>(use);
ASSERT_NE(sem, nullptr);
EXPECT_EQ(sem->Type(), Sem().Get(mix));
}
namespace texture_constexpr_args {
using TextureOverloadCase = ast::test::TextureOverloadCase;
using ValidTextureOverload = ast::test::ValidTextureOverload;
using TextureKind = ast::test::TextureKind;
using TextureDataType = ast::test::TextureDataType;
static std::vector<TextureOverloadCase> TextureCases(
std::unordered_set<ValidTextureOverload> overloads) {
std::vector<TextureOverloadCase> cases;
for (auto c : TextureOverloadCase::ValidCases()) {
if (overloads.count(c.overload)) {
cases.push_back(c);
}
}
return cases;
}
enum class Position {
kFirst,
kLast,
};
struct Parameter {
const char* const name;
const Position position;
int min;
int max;
};
class Constexpr {
public:
enum class Kind {
kScalar,
kVec2,
kVec3,
kVec3_Scalar_Vec2,
kVec3_Vec2_Scalar,
kEmptyVec2,
kEmptyVec3,
};
Constexpr(int32_t invalid_idx, Kind k, int32_t x = 0, int32_t y = 0, int32_t z = 0)
: invalid_index(invalid_idx), kind(k), values{x, y, z} {}
const ast::Expression* operator()(Source src, ProgramBuilder& b) {
switch (kind) {
case Kind::kScalar:
return b.Expr(src, i32(values[0]));
case Kind::kVec2:
return b.Call(src, b.ty.vec2<i32>(), i32(values[0]), i32(values[1]));
case Kind::kVec3:
return b.Call(src, b.ty.vec3<i32>(), i32(values[0]), i32(values[1]),
i32(values[2]));
case Kind::kVec3_Scalar_Vec2:
return b.Call(src, b.ty.vec3<i32>(), i32(values[0]),
b.Call<vec2<i32>>(i32(values[1]), i32(values[2])));
case Kind::kVec3_Vec2_Scalar:
return b.Call(src, b.ty.vec3<i32>(),
b.Call<vec2<i32>>(i32(values[0]), i32(values[1])), i32(values[2]));
case Kind::kEmptyVec2:
return b.Call(src, b.ty.vec2<i32>());
case Kind::kEmptyVec3:
return b.Call(src, b.ty.vec3<i32>());
}
return nullptr;
}
static const constexpr int32_t kValid = -1;
const int32_t invalid_index; // Expected error value, or kValid
const Kind kind;
const std::array<int32_t, 3> values;
};
static std::ostream& operator<<(std::ostream& out, Parameter param) {
return out << param.name;
}
static std::ostream& operator<<(std::ostream& out, Constexpr expr) {
switch (expr.kind) {
case Constexpr::Kind::kScalar:
return out << expr.values[0];
case Constexpr::Kind::kVec2:
return out << "vec2(" << expr.values[0] << ", " << expr.values[1] << ")";
case Constexpr::Kind::kVec3:
return out << "vec3(" << expr.values[0] << ", " << expr.values[1] << ", "
<< expr.values[2] << ")";
case Constexpr::Kind::kVec3_Scalar_Vec2:
return out << "vec3(" << expr.values[0] << ", vec2(" << expr.values[1] << ", "
<< expr.values[2] << "))";
case Constexpr::Kind::kVec3_Vec2_Scalar:
return out << "vec3(vec2(" << expr.values[0] << ", " << expr.values[1] << "), "
<< expr.values[2] << ")";
case Constexpr::Kind::kEmptyVec2:
return out << "vec2()";
case Constexpr::Kind::kEmptyVec3:
return out << "vec3()";
}
return out;
}
using BuiltinTextureConstExprArgValidationTest =
ResolverTestWithParam<std::tuple<TextureOverloadCase, Parameter, Constexpr>>;
TEST_P(BuiltinTextureConstExprArgValidationTest, Immediate) {
auto& p = GetParam();
auto overload = std::get<0>(p);
auto param = std::get<1>(p);
auto expr = std::get<2>(p);
overload.BuildTextureVariable(this);
overload.BuildSamplerVariable(this);
auto args = overload.args(this);
auto*& arg_to_replace = (param.position == Position::kFirst) ? args.Front() : args.Back();
// BuildTextureVariable() uses a Literal for scalars, and a CallExpression for a vector
// constructor.
bool is_vector = arg_to_replace->Is<ast::CallExpression>();
// Make the expression to be replaced, reachable. This keeps the resolver happy.
WrapInFunction(arg_to_replace);
arg_to_replace = expr(Source{{12, 34}}, *this);
auto* call = Call(overload.function, args);
auto* stmt = overload.returns_value ? static_cast<const ast::Statement*>(Assign(Phony(), call))
: static_cast<const ast::Statement*>(CallStmt(call));
// Call the builtin with the constexpr argument replaced
Func("func", tint::Empty, ty.void_(),
Vector{
stmt,
},
Vector{
Stage(ast::PipelineStage::kFragment),
});
if (expr.invalid_index == Constexpr::kValid) {
EXPECT_TRUE(r()->Resolve()) << r()->error();
} else {
EXPECT_FALSE(r()->Resolve());
StringStream err;
if (is_vector) {
err << "12:34 error: each component of the " << param.name
<< " argument must be at least " << param.min << " and at most " << param.max
<< ". " << param.name << " component " << expr.invalid_index << " is "
<< std::to_string(expr.values[static_cast<size_t>(expr.invalid_index)]);
} else {
err << "12:34 error: the " << param.name << " argument must be at least " << param.min
<< " and at most " << param.max << ". " << param.name << " is "
<< std::to_string(expr.values[static_cast<size_t>(expr.invalid_index)]);
}
EXPECT_EQ(r()->error(), err.str());
}
}
TEST_P(BuiltinTextureConstExprArgValidationTest, GlobalConst) {
auto& p = GetParam();
auto overload = std::get<0>(p);
auto param = std::get<1>(p);
auto expr = std::get<2>(p);
// Build the global texture and sampler variables
overload.BuildTextureVariable(this);
overload.BuildSamplerVariable(this);
// Build the module-scope const 'G' with the offset value
GlobalConst("G", expr({}, *this));
auto args = overload.args(this);
auto*& arg_to_replace = (param.position == Position::kFirst) ? args.Front() : args.Back();
// BuildTextureVariable() uses a Literal for scalars, and a CallExpression for a vector
// constructor.
bool is_vector = arg_to_replace->Is<ast::CallExpression>();
// Make the expression to be replaced, reachable. This keeps the resolver happy.
WrapInFunction(arg_to_replace);
arg_to_replace = Expr(Source{{12, 34}}, "G");
auto* call = Call(overload.function, args);
auto* stmt = overload.returns_value ? static_cast<const ast::Statement*>(Assign(Phony(), call))
: static_cast<const ast::Statement*>(CallStmt(call));
// Call the builtin with the constant-expression argument replaced
Func("func", tint::Empty, ty.void_(),
Vector{
stmt,
},
Vector{
Stage(ast::PipelineStage::kFragment),
});
if (expr.invalid_index == Constexpr::kValid) {
EXPECT_TRUE(r()->Resolve()) << r()->error();
} else {
EXPECT_FALSE(r()->Resolve());
StringStream err;
if (is_vector) {
err << "12:34 error: each component of the " << param.name
<< " argument must be at least " << param.min << " and at most " << param.max
<< ". " << param.name << " component " << expr.invalid_index << " is "
<< std::to_string(expr.values[static_cast<size_t>(expr.invalid_index)]);
} else {
err << "12:34 error: the " << param.name << " argument must be at least " << param.min
<< " and at most " << param.max << ". " << param.name << " is "
<< std::to_string(expr.values[static_cast<size_t>(expr.invalid_index)]);
}
EXPECT_EQ(r()->error(), err.str());
}
}
TEST_P(BuiltinTextureConstExprArgValidationTest, GlobalVar) {
auto& p = GetParam();
auto overload = std::get<0>(p);
auto param = std::get<1>(p);
auto expr = std::get<2>(p);
// Build the global texture and sampler variables
overload.BuildTextureVariable(this);
overload.BuildSamplerVariable(this);
// Build the module-scope var 'G' with the offset value
GlobalVar("G", expr({}, *this), core::AddressSpace::kPrivate);
auto args = overload.args(this);
auto*& arg_to_replace = (param.position == Position::kFirst) ? args.Front() : args.Back();
// Make the expression to be replaced, reachable. This keeps the resolver happy.
WrapInFunction(arg_to_replace);
arg_to_replace = Expr(Source{{12, 34}}, "G");
auto* call = Call(overload.function, args);
auto* stmt = overload.returns_value ? static_cast<const ast::Statement*>(Assign(Phony(), call))
: static_cast<const ast::Statement*>(CallStmt(call));
// Call the builtin with the constant-expression argument replaced
Func("func", tint::Empty, ty.void_(),
Vector{
stmt,
},
Vector{
Stage(ast::PipelineStage::kFragment),
});
EXPECT_FALSE(r()->Resolve());
StringStream err;
err << "12:34 error: the " << param.name << " argument must be a const-expression";
EXPECT_EQ(r()->error(), err.str());
}
INSTANTIATE_TEST_SUITE_P(
Offset2D,
BuiltinTextureConstExprArgValidationTest,
testing::Combine(testing::ValuesIn(TextureCases({
ValidTextureOverload::kSample2dOffsetF32,
ValidTextureOverload::kSample2dArrayOffsetF32,
ValidTextureOverload::kSampleDepth2dOffsetF32,
ValidTextureOverload::kSampleDepth2dArrayOffsetF32,
ValidTextureOverload::kSampleBias2dOffsetF32,
ValidTextureOverload::kSampleBias2dArrayOffsetF32,
ValidTextureOverload::kSampleLevel2dOffsetF32,
ValidTextureOverload::kSampleLevel2dArrayOffsetF32,
ValidTextureOverload::kSampleLevelDepth2dOffsetF32,
ValidTextureOverload::kSampleLevelDepth2dArrayOffsetF32,
ValidTextureOverload::kSampleGrad2dOffsetF32,
ValidTextureOverload::kSampleGrad2dArrayOffsetF32,
ValidTextureOverload::kSampleCompareDepth2dOffsetF32,
ValidTextureOverload::kSampleCompareDepth2dArrayOffsetF32,
ValidTextureOverload::kSampleCompareLevelDepth2dOffsetF32,
ValidTextureOverload::kSampleCompareLevelDepth2dArrayOffsetF32,
})),
testing::Values(Parameter{"offset", Position::kLast, -8, 7}),
testing::Values(Constexpr{Constexpr::kValid, Constexpr::Kind::kEmptyVec2},
Constexpr{Constexpr::kValid, Constexpr::Kind::kVec2, -1, 1},
Constexpr{Constexpr::kValid, Constexpr::Kind::kVec2, 7, -8},
Constexpr{0, Constexpr::Kind::kVec2, 8, 0},
Constexpr{1, Constexpr::Kind::kVec2, 0, 8},
Constexpr{0, Constexpr::Kind::kVec2, -9, 0},
Constexpr{1, Constexpr::Kind::kVec2, 0, -9},
Constexpr{0, Constexpr::Kind::kVec2, 8, 8},
Constexpr{0, Constexpr::Kind::kVec2, -9, -9})));
INSTANTIATE_TEST_SUITE_P(
Offset3D,
BuiltinTextureConstExprArgValidationTest,
testing::Combine(testing::ValuesIn(TextureCases({
ValidTextureOverload::kSample3dOffsetF32,
ValidTextureOverload::kSampleBias3dOffsetF32,
ValidTextureOverload::kSampleLevel3dOffsetF32,
ValidTextureOverload::kSampleGrad3dOffsetF32,
})),
testing::Values(Parameter{"offset", Position::kLast, -8, 7}),
testing::Values(Constexpr{Constexpr::kValid, Constexpr::Kind::kEmptyVec3},
Constexpr{Constexpr::kValid, Constexpr::Kind::kVec3, 0, 0, 0},
Constexpr{Constexpr::kValid, Constexpr::Kind::kVec3, 7, -8, 7},
Constexpr{0, Constexpr::Kind::kVec3, 10, 0, 0},
Constexpr{1, Constexpr::Kind::kVec3, 0, 10, 0},
Constexpr{2, Constexpr::Kind::kVec3, 0, 0, 10},
Constexpr{0, Constexpr::Kind::kVec3, 10, 11, 12},
Constexpr{0, Constexpr::Kind::kVec3_Scalar_Vec2, 10, 0, 0},
Constexpr{1, Constexpr::Kind::kVec3_Scalar_Vec2, 0, 10, 0},
Constexpr{2, Constexpr::Kind::kVec3_Scalar_Vec2, 0, 0, 10},
Constexpr{0, Constexpr::Kind::kVec3_Scalar_Vec2, 10, 11, 12},
Constexpr{0, Constexpr::Kind::kVec3_Vec2_Scalar, 10, 0, 0},
Constexpr{1, Constexpr::Kind::kVec3_Vec2_Scalar, 0, 10, 0},
Constexpr{2, Constexpr::Kind::kVec3_Vec2_Scalar, 0, 0, 10},
Constexpr{0, Constexpr::Kind::kVec3_Vec2_Scalar, 10, 11,
12})));
INSTANTIATE_TEST_SUITE_P(
Component,
BuiltinTextureConstExprArgValidationTest,
testing::Combine(
testing::ValuesIn(TextureCases({
ValidTextureOverload::kGather2dF32, ValidTextureOverload::kGather2dOffsetF32,
ValidTextureOverload::kGather2dArrayF32, ValidTextureOverload::kGatherCubeF32,
// The below require mixed integer signedness.
// See https://github.com/gpuweb/gpuweb/issues/3536
// ValidTextureOverload::kGather2dArrayOffsetF32,
// ValidTextureOverload::kGatherCubeArrayF32,
})),
testing::Values(Parameter{"component", Position::kFirst, 0, 3}),
testing::Values(Constexpr{Constexpr::kValid, Constexpr::Kind::kScalar, 0},
Constexpr{Constexpr::kValid, Constexpr::Kind::kScalar, 1},
Constexpr{Constexpr::kValid, Constexpr::Kind::kScalar, 2},
Constexpr{Constexpr::kValid, Constexpr::Kind::kScalar, 3},
Constexpr{0, Constexpr::Kind::kScalar, 4},
Constexpr{0, Constexpr::Kind::kScalar, 123},
Constexpr{0, Constexpr::Kind::kScalar, -1})));
} // namespace texture_constexpr_args
using ResolverPacked4x8IntegerDotProductValidationTest = ResolverTest;
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Dot4I8Packed) {
// fn func { return dot4I8Packed(1u, 2u); }
Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
Func("func", tint::Empty, ty.i32(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "dot4I8Packed",
Vector{Expr(1_u), Expr(2_u)})),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Dot4I8Packed_FeatureDisallowed) {
// fn func { return dot4I8Packed(1u, 2u); }
Func("func", tint::Empty, ty.i32(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "dot4I8Packed",
Vector{Expr(1_u), Expr(2_u)})),
});
auto resolver = Resolver(this, {});
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'dot4I8Packed' requires the "
"packed_4x8_integer_dot_product language feature, which is not allowed in the "
"current environment");
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Dot4U8Packed) {
// fn func { return dot4U8Packed(1u, 2u); }
Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
Func("func", tint::Empty, ty.u32(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "dot4U8Packed",
Vector{Expr(1_u), Expr(2_u)})),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Dot4U8Packed_FeatureDisallowed) {
// fn func { return dot4U8Packed(1u, 2u); }
Func("func", tint::Empty, ty.u32(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "dot4U8Packed",
Vector{Expr(1_u), Expr(2_u)})),
});
auto resolver = Resolver(this, {});
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'dot4U8Packed' requires the "
"packed_4x8_integer_dot_product language feature, which is not allowed in the "
"current environment");
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Pack4xI8) {
// fn func { return pack4xI8(vec4i()); }
Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
Func("func", tint::Empty, ty.u32(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "pack4xI8", Call<vec4<i32>>())),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Pack4xI8_FeatureDisallowed) {
// fn func { return pack4xI8(vec4i()); }
Func("func", tint::Empty, ty.u32(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "pack4xI8", Call<vec4<i32>>())),
});
auto resolver = Resolver(this, {});
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'pack4xI8' requires the "
"packed_4x8_integer_dot_product language feature, which is not allowed in the "
"current environment");
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Pack4xU8) {
// fn func { return pack4xU8(vec4u()); }
Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
Func("func", tint::Empty, ty.u32(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "pack4xU8", Call<vec4<u32>>())),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Pack4xU8_FeatureDisallowed) {
// fn func { return pack4xU8(vec4u()); }
Func("func", tint::Empty, ty.u32(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "pack4xU8", Call<vec4<u32>>())),
});
auto resolver = Resolver(this, {});
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'pack4xU8' requires the "
"packed_4x8_integer_dot_product language feature, which is not allowed in the "
"current environment");
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Pack4xI8Clamp) {
// fn func { return pack4xI8Clamp(vec4i()); }
Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
Func("func", tint::Empty, ty.u32(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "pack4xI8Clamp", Call<vec4<i32>>())),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Pack4xI8Clamp_FeatureDisallowed) {
// fn func { return pack4xI8Clamp(vec4i()); }
Func("func", tint::Empty, ty.u32(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "pack4xI8Clamp", Call<vec4<i32>>())),
});
auto resolver = Resolver(this, {});
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'pack4xI8Clamp' requires the "
"packed_4x8_integer_dot_product language feature, which is not allowed in the "
"current environment");
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Pack4xU8Clamp) {
// fn func { return pack4xU8Clamp(vec4u()); }
Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
Func("func", tint::Empty, ty.u32(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "pack4xU8Clamp", Call<vec4<u32>>())),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Pack4xU8Clamp_FeatureDisallowed) {
// fn func { return pack4xU8Clamp(vec4u()); }
Func("func", tint::Empty, ty.u32(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "pack4xU8Clamp", Call<vec4<u32>>())),
});
auto resolver = Resolver(this, {});
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'pack4xU8Clamp' requires the "
"packed_4x8_integer_dot_product language feature, which is not allowed in the "
"current environment");
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Unpack4xI8) {
// fn func { return unpack4xI8(u32()); }
Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
Func("func", tint::Empty, ty.vec4<i32>(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "unpack4xI8", Call<u32>())),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Unpack4xI8_FeatureDisallowed) {
// fn func { return unpack4xI8(u32()); }
Func("func", tint::Empty, ty.vec4<i32>(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "unpack4xI8", Call<u32>())),
});
auto resolver = Resolver(this, {});
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'unpack4xI8' requires the "
"packed_4x8_integer_dot_product language feature, which is not allowed in the "
"current environment");
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Unpack4xU8) {
// fn func { return unpack4xU8(u32()); }
Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
Func("func", tint::Empty, ty.vec4<u32>(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "unpack4xU8", Call<u32>())),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverPacked4x8IntegerDotProductValidationTest, Unpack4xU8_FeatureDisallowed) {
// fn func { return unpack4xU8(u32()); }
Func("func", tint::Empty, ty.vec4<u32>(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "unpack4xU8", Call<u32>())),
});
auto resolver = Resolver(this, {});
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'unpack4xU8' requires the "
"packed_4x8_integer_dot_product language feature, which is not allowed in the "
"current environment");
}
TEST_F(ResolverBuiltinValidationTest, WorkgroupUniformLoad_WrongAddressSpace) {
// @group(0) @binding(0) var<storage, read_write> v : i32;
// fn foo() {
// workgroupUniformLoad(&v);
// }
GlobalVar("v", ty.i32(), core::AddressSpace::kStorage, core::Access::kReadWrite,
Vector{Group(0_a), Binding(0_a)});
WrapInFunction(CallStmt(Call("workgroupUniformLoad", AddressOf(Source{{12, 34}}, "v"))));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
R"(error: no matching call to workgroupUniformLoad(ptr<storage, i32, read_write>)
1 candidate function:
workgroupUniformLoad(ptr<workgroup, T, read_write>) -> T
)");
}
TEST_F(ResolverBuiltinValidationTest, WorkgroupUniformLoad_Atomic) {
// var<workgroup> v : atomic<i32>;
// fn foo() {
// workgroupUniformLoad(&v);
// }
GlobalVar("v", ty.atomic<i32>(), core::AddressSpace::kWorkgroup);
WrapInFunction(CallStmt(Call("workgroupUniformLoad", AddressOf(Source{{12, 34}}, "v"))));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
R"(12:34 error: workgroupUniformLoad must not be called with an argument that contains an atomic type)");
}
TEST_F(ResolverBuiltinValidationTest, WorkgroupUniformLoad_AtomicInArray) {
// var<workgroup> v : array<atomic<i32>, 4>;
// fn foo() {
// workgroupUniformLoad(&v);
// }
GlobalVar("v", ty.array(ty.atomic<i32>(), 4_a), core::AddressSpace::kWorkgroup);
WrapInFunction(CallStmt(Call("workgroupUniformLoad", AddressOf(Source{{12, 34}}, "v"))));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
R"(12:34 error: workgroupUniformLoad must not be called with an argument that contains an atomic type)");
}
TEST_F(ResolverBuiltinValidationTest, WorkgroupUniformLoad_AtomicInStruct) {
// struct Inner { a : array<atomic<i32, 4> }
// struct S { i : Inner }
// var<workgroup> v : array<S, 4>;
// fn foo() {
// workgroupUniformLoad(&v);
// }
Structure("Inner", Vector{Member("a", ty.array(ty.atomic<i32>(), 4_a))});
Structure("S", Vector{Member("i", ty("Inner"))});
GlobalVar(Source{{12, 34}}, "v", ty.array(ty("S"), 4_a), core::AddressSpace::kWorkgroup);
WrapInFunction(CallStmt(Call("workgroupUniformLoad", AddressOf("v"))));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
R"(error: workgroupUniformLoad must not be called with an argument that contains an atomic type)");
}
TEST_F(ResolverBuiltinValidationTest, SubgroupBallotWithoutExtension) {
// fn func { return subgroupBallot(); }
Func("func", tint::Empty, ty.vec4<u32>(),
Vector{
Return(Call(Source{Source::Location{12, 34}}, "subgroupBallot")),
});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
R"(12:34 error: cannot call built-in function 'subgroupBallot' without extension chromium_experimental_subgroups)");
}
TEST_F(ResolverBuiltinValidationTest, SubgroupBallotWithExtension) {
// enable chromium_experimental_subgroups;
// fn func -> vec4<u32> { return subgroupBallot(); }
Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
Func("func", tint::Empty, ty.vec4<u32>(),
Vector{
Return(Call("subgroupBallot")),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverBuiltinValidationTest, SubgroupBroadcastWithoutExtension) {
// fn func -> i32 { return subgroupBroadcast(1,0); }
Func("func", tint::Empty, ty.i32(),
Vector{
Return(Call(Source{{12, 34}}, "subgroupBroadcast", 1_i, 0_u)),
});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
R"(12:34 error: cannot call built-in function 'subgroupBroadcast' without extension chromium_experimental_subgroups)");
}
TEST_F(ResolverBuiltinValidationTest, SubgroupBroadcastWithExtension) {
// enable chromium_experimental_subgroups;
// fn func -> i32 { return subgroupBroadcast(1,0); }
Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
Func("func", tint::Empty, ty.i32(),
Vector{
Return(Call("subgroupBroadcast", 1_i, 0_u)),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverBuiltinValidationTest, SubroupBroadcastInComputeStage) {
// @vertex fn func { dpdx(1.0); }
Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
auto* call = Call("subgroupBroadcast", 1_f, 0_u);
Func(Source{{1, 2}}, "func", tint::Empty, ty.void_(), Vector{Ignore(call)},
Vector{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1_i),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverBuiltinValidationTest, SubroupBroadcastInVertexStageIsError) {
// @vertex fn func { dpdx(1.0); }
Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
auto* call = Call(Source{{3, 4}}, "subgroupBroadcast", 1_f, 0_u);
Func("func", tint::Empty, ty.vec4<f32>(), Vector{Ignore(call), Return(Call(ty.vec4<f32>()))},
Vector{
Stage(ast::PipelineStage::kVertex),
},
Vector{Builtin(core::BuiltinValue::kPosition)});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), "3:4 error: built-in cannot be used by vertex pipeline stage");
}
TEST_F(ResolverBuiltinValidationTest, SubroupBroadcastInFragmentStageIsError) {
// @vertex fn func { dpdx(1.0); }
Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
auto* call = Call(Source{{3, 4}}, "subgroupBroadcast", 1_f, 0_u);
Func("func",
Vector{Param("pos", ty.vec4<f32>(), Vector{Builtin(core::BuiltinValue::kPosition)})},
ty.void_(), Vector{Ignore(call)},
Vector{
Stage(ast::PipelineStage::kFragment),
});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), "3:4 error: built-in cannot be used by fragment pipeline stage");
}
TEST_F(ResolverBuiltinValidationTest, SubgroupBroadcastValueF32) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
Func("func", tint::Empty, ty.f32(),
Vector{
Return(Call("subgroupBroadcast", 1_f, 0_u)),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverBuiltinValidationTest, SubgroupBroadcastValueI32) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
Func("func", tint::Empty, ty.i32(),
Vector{
Return(Call("subgroupBroadcast", 1_i, 0_u)),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverBuiltinValidationTest, SubgroupBroadcastValueU32) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
Func("func", tint::Empty, ty.u32(),
Vector{
Return(Call("subgroupBroadcast", 1_u, 0_u)),
});
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverBuiltinValidationTest, SubgroupBroadcastLaneArgMustBeConst) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
Func("func", tint::Empty, ty.void_(),
Vector{
Decl(Let("lane", Expr(1_u))),
Ignore(Call("subgroupBroadcast", 1_f, Ident(Source{{12, 34}}, "lane"))),
});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
R"(12:34 error: the sourceLaneIndex argument of subgroupBroadcast must be a const-expression)");
}
TEST_F(ResolverBuiltinValidationTest, TextureBarrier) {
// fn func { textureBarrier(); }
Func("func", tint::Empty, ty.void_(),
Vector{
CallStmt(Call(Source{Source::Location{12, 34}}, "textureBarrier")),
});
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
TEST_F(ResolverBuiltinValidationTest, TextureBarrier_FeatureDisallowed) {
// fn func { textureBarrier(); }
Func("func", tint::Empty, ty.void_(),
Vector{
CallStmt(Call(Source{Source::Location{12, 34}}, "textureBarrier")),
});
auto resolver = Resolver(this, {});
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'textureBarrier' requires the "
"readonly_and_readwrite_storage_textures language feature, which is not allowed in "
"the current environment");
}
} // namespace
} // namespace tint::resolver