[tint][ir][val] Implement `CheckBuiltIn`
This adds validation code for builtins, specifically it validates the
shader stage, IO direction, and data type it is annotating.
All of the existing tests have been updated to pass using the new
rules.
Fixes: 372285196
Change-Id: I14a6750fa37470ea12d91518d100cf8bc025584b
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/211195
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/core/ir/transform/zero_init_workgroup_memory_test.cc b/src/tint/lang/core/ir/transform/zero_init_workgroup_memory_test.cc
index 45f122f..b1e175b 100644
--- a/src/tint/lang/core/ir/transform/zero_init_workgroup_memory_test.cc
+++ b/src/tint/lang/core/ir/transform/zero_init_workgroup_memory_test.cc
@@ -1475,7 +1475,7 @@
auto* var = MakeVar("wgvar", ty.bool_());
auto* func = MakeEntryPoint("main", 1, 1, 1);
- auto* global_id = b.FunctionParam("global_id", ty.vec4<u32>());
+ auto* global_id = b.FunctionParam("global_id", ty.vec3<u32>());
global_id->SetBuiltin(BuiltinValue::kGlobalInvocationId);
auto* index = b.FunctionParam("index", ty.u32());
index->SetBuiltin(BuiltinValue::kLocalInvocationIndex);
@@ -1490,7 +1490,7 @@
%wgvar:ptr<workgroup, bool, read_write> = var
}
-%main = @compute @workgroup_size(1, 1, 1) func(%global_id:vec4<u32> [@global_invocation_id], %index:u32 [@local_invocation_index]):void {
+%main = @compute @workgroup_size(1, 1, 1) func(%global_id:vec3<u32> [@global_invocation_id], %index:u32 [@local_invocation_index]):void {
$B2: {
%5:bool = load %wgvar
ret
@@ -1504,7 +1504,7 @@
%wgvar:ptr<workgroup, bool, read_write> = var
}
-%main = @compute @workgroup_size(1, 1, 1) func(%global_id:vec4<u32> [@global_invocation_id], %index:u32 [@local_invocation_index]):void {
+%main = @compute @workgroup_size(1, 1, 1) func(%global_id:vec3<u32> [@global_invocation_id], %index:u32 [@local_invocation_index]):void {
$B2: {
%5:bool = eq %index, 0u
if %5 [t: $B3] { # if_1
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 9766d25..d1bfa1e 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -83,6 +83,7 @@
#include "src/tint/lang/core/type/pointer.h"
#include "src/tint/lang/core/type/reference.h"
#include "src/tint/lang/core/type/type.h"
+#include "src/tint/lang/core/type/u32.h"
#include "src/tint/lang/core/type/u8.h"
#include "src/tint/lang/core/type/vector.h"
#include "src/tint/lang/core/type/void.h"
@@ -161,6 +162,23 @@
ty->Is<type::Sampler>();
}
+/// @returns true if @p ty is a non-struct and decorated with @builtin(position), or if it is a
+/// struct and one of its members is decorated, otherwise false.
+/// @param attr attributes attached to data
+/// @param ty type of the data being tested
+bool IsPositionPresent(const IOAttributes& attr, const core::type::Type* ty) {
+ if (auto* ty_struct = ty->As<core::type::Struct>()) {
+ for (const auto* mem : ty_struct->Members()) {
+ if (mem->Attributes().builtin == BuiltinValue::kPosition) {
+ return true;
+ }
+ }
+ return false;
+ }
+
+ return attr.builtin == BuiltinValue::kPosition;
+}
+
/// Utility for running checks on attributes.
/// If the type that the attributes are attached to is a struct, the check is run over the members,
/// otherwise it run on the attributes directly.
@@ -239,9 +257,21 @@
}
}
-/// Helper for calling CheckIOAttributesAndType on a function return
-/// @param func function whose return is to be tested
-/// See @ref CheckIOAttributesAndType for more details
+/// Helper for calling IOAttributesAndType on a function param
+/// @param param function param to be tested
+/// See @ref IOAttributesAndType for more details
+template <typename IS_NOT_STRUCT, typename IS_STRUCT>
+void CheckFunctionParamAttributesAndType(const FunctionParam* param,
+ IS_NOT_STRUCT&& is_not_struct_impl,
+ IS_STRUCT&& is_struct_impl) {
+ CheckIOAttributesAndType(param, param->Attributes(), param->Type(),
+ std::forward<IS_NOT_STRUCT>(is_not_struct_impl),
+ std::forward<IS_STRUCT>(is_struct_impl));
+}
+
+/// Helper for calling IOAttributesAndType on a function return
+/// @param func function's return to be tested
+/// See @ref IOAttributesAndType for more details
template <typename IS_NOT_STRUCT, typename IS_STRUCT>
void CheckFunctionReturnAttributesAndType(const Function* func,
IS_NOT_STRUCT&& is_not_struct_impl,
@@ -251,6 +281,321 @@
std::forward<IS_STRUCT>(is_struct_impl));
}
+/// A BuiltinChecker is the interface used to check that a usage of a builtin attribute meets the
+/// basic spec rules, i.e. correct shader stage, data type, and IO direction.
+/// It does not test more sophisticated rules like location and builtins being mutually exclusive or
+/// the correct capabilities are enabled.
+struct BuiltinChecker {
+ /// User friendly name to print in logging messages
+ const char* name;
+
+ /// What type of entry point is this builtin legal for
+ EnumSet<Function::PipelineStage> stages;
+
+ enum IODirection : uint8_t { kInput, kOutput };
+ /// Is this expected to be a param going into the entry point or a result coming out
+ IODirection direction;
+
+ /// Implements logic for checking if the given type is valid or not
+ using TypeCheckFn = bool(const core::type::Type* type);
+
+ /// @see #TypeCheckFn
+ TypeCheckFn* const type_check;
+
+ /// Message that should logged if the type check fails
+ const char* type_error;
+};
+
+std::string_view ToString(BuiltinChecker::IODirection value) {
+ switch (value) {
+ case BuiltinChecker::IODirection::kInput:
+ return "input";
+ case BuiltinChecker::IODirection::kOutput:
+ return "output";
+ }
+ TINT_ICE() << "Unknown enum passed to ToString(BuiltinChecker::IODirection)";
+}
+
+constexpr BuiltinChecker kPointSizeChecker{
+ /* name */ "__point_size",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kVertex),
+ /* direction */ BuiltinChecker::IODirection::kOutput,
+ /* type_check */ [](const core::type::Type* ty) -> bool { return ty->Is<core::type::F32>(); },
+ /* type_error */ "__point_size must be a f32",
+};
+
+constexpr BuiltinChecker kClipDistancesChecker{
+ /* name */ "clip_distances",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kVertex),
+ /* direction */ BuiltinChecker::IODirection::kOutput,
+ /* type_check */
+ [](const core::type::Type* ty) -> bool {
+ auto elems = ty->Elements();
+ return elems.type && elems.type->Is<core::type::F32>() && elems.count <= 8;
+ },
+ /* type_error */ "clip_distances must be an array<f32, N>, where N <= 8",
+};
+
+constexpr BuiltinChecker kFragDepthChecker{
+ /* name */ "frag_depth",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kFragment),
+ /* direction */ BuiltinChecker::IODirection::kOutput,
+ /* type_check */ [](const core::type::Type* ty) -> bool { return ty->Is<core::type::F32>(); },
+ /* type_error */ "frag_depth must be a f32",
+};
+
+constexpr BuiltinChecker kFrontFacingChecker{
+ /* name */ "front_facing",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kFragment),
+ /* direction */ BuiltinChecker::IODirection::kInput,
+ /* type_check */ [](const core::type::Type* ty) -> bool { return ty->Is<core::type::Bool>(); },
+ /* type_error */ "front_facing must be a bool",
+};
+
+constexpr BuiltinChecker kGlobalInvocationIdChecker{
+ /* name */ "global_invocation_id",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kCompute),
+ /* direction */ BuiltinChecker::IODirection::kInput,
+ /* type_check */
+ [](const core::type::Type* ty) -> bool {
+ return ty->IsUnsignedIntegerVector() && ty->Elements().count == 3;
+ },
+ /* type_error */ "global_invocation_id must be an vec3<u32>",
+};
+
+constexpr BuiltinChecker kInstanceIndexChecker{
+ /* name */ "instance_index",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kVertex),
+ /* direction */ BuiltinChecker::IODirection::kInput,
+ /* type_check */ [](const core::type::Type* ty) -> bool { return ty->Is<core::type::U32>(); },
+ /* type_error */ "instance_index must be an u32",
+};
+
+constexpr BuiltinChecker kLocalInvocationIdChecker{
+ /* name */ "local_invocation_id",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kCompute),
+ /* direction */ BuiltinChecker::IODirection::kInput,
+ /* type_check */
+ [](const core::type::Type* ty) -> bool {
+ return ty->IsUnsignedIntegerVector() && ty->Elements().count == 3;
+ },
+ /* type_error */ "local_invocation_id must be an vec3<u32>",
+};
+
+constexpr BuiltinChecker kLocalInvocationIndexChecker{
+ /* name */ "local_invocation_index",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kCompute),
+ /* direction */ BuiltinChecker::IODirection::kInput,
+ /* type_check */ [](const core::type::Type* ty) -> bool { return ty->Is<core::type::U32>(); },
+ /* type_error */ "local_invocation_index must be an u32",
+};
+
+constexpr BuiltinChecker kNumWorkgroupsChecker{
+ /* name */ "num_workgroups",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kCompute),
+ /* direction */ BuiltinChecker::IODirection::kInput,
+ /* type_check */
+ [](const core::type::Type* ty) -> bool {
+ return ty->IsUnsignedIntegerVector() && ty->Elements().count == 3;
+ },
+ /* type_error */ "num_workgroups must be an vec3<u32>",
+};
+
+constexpr BuiltinChecker kSampleIndexChecker{
+ /* name */ "sample_index",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kFragment),
+ /* direction */ BuiltinChecker::IODirection::kInput,
+ /* type_check */ [](const core::type::Type* ty) -> bool { return ty->Is<core::type::U32>(); },
+ /* type_error */ "sample_index must be an u32",
+};
+
+constexpr BuiltinChecker kSubgroupInvocationIdChecker{
+ /* name */ "subgroup_invocation_id",
+ /* stages */
+ EnumSet<Function::PipelineStage>(Function::PipelineStage::kFragment,
+ Function::PipelineStage::kCompute),
+ /* direction */ BuiltinChecker::IODirection::kInput,
+ /* type_check */ [](const core::type::Type* ty) -> bool { return ty->Is<core::type::U32>(); },
+ /* type_error */ "subgroup_invocation_id must be an u32",
+};
+
+constexpr BuiltinChecker kSubgroupSizeChecker{
+ /* name */ "subgroup_size",
+ /* stages */
+ EnumSet<Function::PipelineStage>(Function::PipelineStage::kFragment,
+ Function::PipelineStage::kCompute),
+ /* direction */ BuiltinChecker::IODirection::kInput,
+ /* type_check */ [](const core::type::Type* ty) -> bool { return ty->Is<core::type::U32>(); },
+ /* type_error */ "subgroup_size must be an u32",
+};
+
+constexpr BuiltinChecker kVertexIndexChecker{
+ /* name */ "vertex_index",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kVertex),
+ /* direction */ BuiltinChecker::IODirection::kInput,
+ /* type_check */ [](const core::type::Type* ty) -> bool { return ty->Is<core::type::U32>(); },
+ /* type_error */ "vertex_index must be an u32",
+};
+
+constexpr BuiltinChecker kWorkgroupIdChecker{
+ /* name */ "workgroup_id",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kCompute),
+ /* direction */ BuiltinChecker::IODirection::kInput,
+ /* type_check */
+ [](const core::type::Type* ty) -> bool {
+ return ty->IsUnsignedIntegerVector() && ty->Elements().count == 3;
+ },
+ /* type_error */ "workgroup_id must be an vec3<u32>",
+};
+
+/// @returns an appropriate BuiltInCheck for @p builtin, ICEs when one isn't defined
+const BuiltinChecker& BuiltinCheckerFor(BuiltinValue builtin) {
+ switch (builtin) {
+ case BuiltinValue::kPointSize:
+ return kPointSizeChecker;
+ case BuiltinValue::kClipDistances:
+ return kClipDistancesChecker;
+ case BuiltinValue::kFragDepth:
+ return kFragDepthChecker;
+ case BuiltinValue::kFrontFacing:
+ return kFrontFacingChecker;
+ case BuiltinValue::kGlobalInvocationId:
+ return kGlobalInvocationIdChecker;
+ case BuiltinValue::kInstanceIndex:
+ return kInstanceIndexChecker;
+ case BuiltinValue::kLocalInvocationId:
+ return kLocalInvocationIdChecker;
+ case BuiltinValue::kLocalInvocationIndex:
+ return kLocalInvocationIndexChecker;
+ case BuiltinValue::kNumWorkgroups:
+ return kNumWorkgroupsChecker;
+ case BuiltinValue::kSampleIndex:
+ return kSampleIndexChecker;
+ case BuiltinValue::kSubgroupInvocationId:
+ return kSubgroupInvocationIdChecker;
+ case BuiltinValue::kSubgroupSize:
+ return kSubgroupSizeChecker;
+ case BuiltinValue::kVertexIndex:
+ return kVertexIndexChecker;
+ case BuiltinValue::kWorkgroupId:
+ return kWorkgroupIdChecker;
+ case BuiltinValue::kPosition:
+ TINT_ICE() << "BuiltinValue::kPosition requires special handling, so does not have a "
+ "checker defined";
+ case BuiltinValue::kSampleMask:
+ TINT_ICE() << "BuiltinValue::kSampleMask requires special handling, so does not have a "
+ "checker defined";
+ default:
+ TINT_ICE() << builtin << " is does not have a checker defined for it";
+ }
+}
+
+/// Validates the basic spec rules for @builtin(position) usage
+/// @param stage the shader stage the builtin is being used
+/// @param is_input the IO direction of usage, true if input, false if output
+/// @param ty the data type being decorated by the builtin
+/// @returns Success if a valid usage, or reason for invalidity in Failure
+Result<SuccessType, std::string> ValidatePositionBuiltIn(Function::PipelineStage stage,
+ bool is_input,
+ const core::type::Type* ty) {
+ if (stage != Function::PipelineStage::kVertex && stage != Function::PipelineStage::kFragment) {
+ return std::string("position must be used in a fragment or vertex shader entry point");
+ }
+
+ if (stage == Function::PipelineStage::kVertex && is_input) {
+ return std::string("position must be an output for a vertex entry point");
+ }
+
+ if (stage == Function::PipelineStage::kFragment && !is_input) {
+ return std::string("position must be an input for a fragment entry point");
+ }
+
+ if (!ty->IsFloatVector() || ty->Elements().count != 4 ||
+ !ty->Element(0)->Is<core::type::F32>()) {
+ return std::string("position must be an vec4<f32>");
+ }
+
+ return Success;
+}
+
+/// Validates the basic spec rules for @builtin(sample_mask) usage
+/// @param stage the shader stage the builtin is being used
+/// @param ty the data type being decorated by the builtin
+/// @returns Success if a valid usage, or reason for invalidity in Failure
+Result<SuccessType, std::string> ValidateSampleMaskBuiltIn(Function::PipelineStage stage,
+ const core::type::Type* ty) {
+ if (stage != Function::PipelineStage::kFragment) {
+ return std::string("sample_mask must be used in a fragment entry point");
+ }
+
+ if (!ty->Is<core::type::U32>()) {
+ return std::string("sample_mask must be an u32");
+ }
+
+ return Success;
+}
+
+/// Validates the basic spec rules for builtin usage
+/// @param builtin the builtin to test
+/// @param stage the shader stage the builtin is being used
+/// @param is_input the IO direction of usage, true if input, false if output
+/// @param ty the data type being decorated by the builtin
+/// @returns Success if a valid usage, or reason for invalidity in Failure
+Result<SuccessType, std::string> ValidateBuiltIn(BuiltinValue builtin,
+ Function::PipelineStage stage,
+ bool is_input,
+ const core::type::Type* ty) {
+ // This is not an entry point function, either it is dead code and thus never called, or any
+ // issues will be detected when validating the calling entry point.
+ if (stage == Function::PipelineStage::kUndefined) {
+ return Success;
+ }
+
+ // Some builtins have multiple contexts that they are valid in, so have special handling
+ // instead of making the checker/lookup table more complex.
+ switch (builtin) {
+ case BuiltinValue::kPosition:
+ return ValidatePositionBuiltIn(stage, is_input, ty);
+ case BuiltinValue::kSampleMask:
+ return ValidateSampleMaskBuiltIn(stage, ty);
+ default: {
+ }
+ }
+
+ const auto& checker = BuiltinCheckerFor(builtin);
+ std::stringstream msg;
+ if (!checker.stages.Contains(stage)) {
+ auto stages_size = checker.stages.Size();
+ switch (stages_size) {
+ case 1:
+ msg << checker.name << " must be used in a " << ToString(*checker.stages.begin())
+ << " shader entry point";
+ break;
+ case 2:
+ msg << checker.name << " must be used in a " << ToString(*checker.stages.begin())
+ << " or " << ToString(*(++checker.stages.begin())) << " shader entry point";
+ break;
+ default:
+ TINT_ICE() << "Unexpected number of stages set, " << stages_size;
+ }
+ return msg.str();
+ }
+
+ auto io_direction =
+ is_input ? BuiltinChecker::IODirection::kInput : BuiltinChecker::IODirection::kOutput;
+ if (io_direction != checker.direction) {
+ msg << checker.name << " must be an " << ToString(checker.direction)
+ << " of a shader entry point";
+ return msg.str();
+ }
+
+ if (!checker.type_check(ty)) {
+ return std::string(checker.type_error);
+ }
+
+ return Success;
+}
+
/// The core IR validator.
class Validator {
public:
@@ -498,6 +843,34 @@
};
}
+ /// @returns a function that validates builtins on function params
+ auto CheckBuiltinFunctionParam(const std::string& err) {
+ return [this, err](const FunctionParam* param, const IOAttributes& attr,
+ const type::Type* ty) {
+ if (!attr.builtin.has_value()) {
+ return;
+ }
+ auto result =
+ ValidateBuiltIn(attr.builtin.value(), param->Function()->Stage(), true, ty);
+ if (result != Success) {
+ AddError(param) << err << result.Failure();
+ }
+ };
+ }
+
+ /// @returns a function that validates builtins on function returns
+ auto CheckBuiltinFunctionReturn(const std::string& err) {
+ return [this, err](const Function* func, const IOAttributes& attr, const type::Type* ty) {
+ if (!attr.builtin.has_value()) {
+ return;
+ }
+ auto result = ValidateBuiltIn(attr.builtin.value(), func->Stage(), false, ty);
+ if (result != Success) {
+ AddError(func) << err << result.Failure();
+ }
+ };
+ }
+
/// @returns a function that validates that location and builtin attributes are not present at
/// the same time
/// @param err error message to log when check fails
@@ -521,16 +894,6 @@
};
}
- /// Validates that the type annotated with @builtin(position) is correct
- /// @param ep the entry point to associate errors with
- /// @param type the type to validate
- void CheckBuiltinPosition(const Function* ep, const core::type::Type* type);
-
- /// Validates that the type annotated with @builtin(clip_distances) is correct
- /// @param ep the entry point to associate errors with
- /// @param type the type to validate
- void CheckBuiltinClipDistances(const Function* ep, const core::type::Type* type);
-
/// Validates the given instruction
/// @param inst the instruction to validate
void CheckInstruction(const Instruction* inst);
@@ -1353,6 +1716,9 @@
}
}
+ CheckFunctionParamAttributesAndType(param, CheckBuiltinFunctionParam(""),
+ CheckBuiltinFunctionParam(""));
+
CheckFunctionParamAttributes(
param,
CheckInvariantFunc<FunctionParam>(
@@ -1375,6 +1741,9 @@
func->ReturnType(), [&]() -> diag::Diagnostic& { return AddError(func); },
Capabilities{Capability::kAllowRefTypes});
+ CheckFunctionReturnAttributesAndType(func, CheckBuiltinFunctionReturn(""),
+ CheckBuiltinFunctionReturn(""));
+
CheckFunctionReturnAttributes(
func,
CheckInvariantFunc<Function>(
@@ -1382,6 +1751,13 @@
CheckInvariantFunc<Function>(
"invariant can only decorate output members iff they are also position builtins"));
+ CheckFunctionReturnAttributes(
+ func,
+ CheckDoesNotHaveBothLocationAndBuiltinFunc<Function>(
+ "a builtin and location cannot be both declared for a function return"),
+ CheckDoesNotHaveBothLocationAndBuiltinFunc<Function>(
+ "a builtin and location cannot be both declared for a struct member"));
+
// void needs to be filtered out, since it isn't constructible, but used in the IR when no
// return is specified.
if (DAWN_UNLIKELY(!func->ReturnType()->Is<core::type::Void>() &&
@@ -1403,20 +1779,6 @@
if (DAWN_UNLIKELY(func->ReturnType() && !func->ReturnType()->Is<core::type::Void>())) {
AddError(func) << "compute entry point must not have a return type";
}
- } else {
- CheckFunctionReturnAttributes(
- func,
- CheckDoesNotHaveBothLocationAndBuiltinFunc<Function>(
- "a builtin and location cannot be both declared for a function return"),
- CheckDoesNotHaveBothLocationAndBuiltinFunc<Function>(
- "a builtin and location cannot be both declared for a struct member"));
- }
-
- if (func->Stage() != Function::PipelineStage::kFragment) {
- if (DAWN_UNLIKELY(func->ReturnBuiltin().has_value() &&
- func->ReturnBuiltin().value() == BuiltinValue::kFragDepth)) {
- AddError(func) << "frag_depth can only be declared for fragment entry points";
- }
}
if (func->Stage() == Function::PipelineStage::kFragment) {
@@ -1441,38 +1803,28 @@
}
void Validator::CheckVertexEntryPoint(const Function* ep) {
- bool contains_position = false;
- auto check_position = [&](const Function* func, const IOAttributes& attr,
- const core::type::Type* ty) {
- if (attr.builtin == BuiltinValue::kPosition) {
- contains_position = true;
- CheckBuiltinPosition(func, ty);
- }
- };
-
- auto check_clip_distances = [&](const Function* func, const IOAttributes& attr,
- const core::type::Type* ty) {
- if (attr.builtin == BuiltinValue::kClipDistances) {
- CheckBuiltinClipDistances(func, ty);
- }
- };
- auto check_clip_distances_noop = [](const Function*, const IOAttributes&,
- const core::type::Type*) {};
-
- CheckFunctionReturnAttributesAndType(ep, check_position, check_position);
- CheckFunctionReturnAttributesAndType(ep, check_clip_distances_noop, check_clip_distances);
+ bool contains_position = IsPositionPresent(ep->ReturnAttributes(), ep->ReturnType());
for (auto var : referenced_module_vars_.TransitiveReferences(ep)) {
const auto* ty = var->Result(0)->Type()->UnwrapPtrOrRef();
const auto attr = var->Attributes();
+ if (!ty) {
+ continue;
+ }
+
+ if (!contains_position) {
+ contains_position = IsPositionPresent(attr, ty);
+ }
+
CheckIOAttributes(
ep, attr, ty,
CheckInvariantFunc<Function>(
"invariant can only decorate vars iff they are also position builtins"),
CheckInvariantFunc<Function>(
"invariant can only decorate members iff they are also position builtins"));
- CheckIOAttributesAndType(ep, attr, ty, check_position, check_position);
- CheckIOAttributesAndType(ep, attr, ty, check_clip_distances_noop, check_clip_distances);
+
+ // Builtin rules are not checked on module-scope variables, because they are often generated
+ // as part of the backend transforms, and have different rules for correctness.
}
if (DAWN_UNLIKELY(!contains_position)) {
@@ -1480,20 +1832,6 @@
}
}
-void Validator::CheckBuiltinPosition(const Function* ep, const core::type::Type* type) {
- auto elems = type->Elements();
- if (!type->IsFloatVector() || !elems.type->Is<core::type::F32>() || elems.count != 4) {
- AddError(ep) << "position must be a vec4<f32>";
- }
-}
-
-void Validator::CheckBuiltinClipDistances(const Function* ep, const core::type::Type* type) {
- const auto elems = type->Elements();
- if (!elems.type || !elems.type->Is<core::type::F32>() || elems.count > 8) {
- AddError(ep) << "clip_distances must be an array<f32, N>, where N <= 8";
- }
-}
-
void Validator::ProcessTasks() {
while (!tasks_.IsEmpty()) {
tasks_.Pop()();
diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc
index d9724b3..c92477e 100644
--- a/src/tint/lang/core/ir/validator_test.cc
+++ b/src/tint/lang/core/ir/validator_test.cc
@@ -52,7 +52,75 @@
using namespace tint::core::fluent_types; // NOLINT
using namespace tint::core::number_suffixes; // NOLINT
-using IR_ValidatorTest = IRTestHelper;
+class IR_ValidatorTest : public IRTestHelper {
+ public:
+ /// Builds and returns a basic 'compute' entry point function, named @p name
+ Function* ComputeEntryPoint(const std::string& name = "f") {
+ return b.Function(name, ty.void_(), Function::PipelineStage::kCompute,
+ std::array<uint32_t, 3>({0, 0, 0}));
+ }
+
+ /// Builds and returns a basic 'fragment' entry point function, named @p name
+ Function* FragmentEntryPoint(const std::string& name = "f") {
+ return b.Function(name, ty.void_(), Function::PipelineStage::kFragment);
+ }
+
+ /// Builds and returns a basic 'vertex' entry point function, named @p name
+ Function* VertexEntryPoint(const std::string& name = "f") {
+ auto* f = b.Function(name, ty.vec4<f32>(), Function::PipelineStage::kVertex);
+ IOAttributes attr;
+ attr.builtin = BuiltinValue::kPosition;
+ f->SetReturnAttributes(attr);
+ return f;
+ }
+
+ /// Adds to a function an input param named @p name of type @p type, and decorated with @p
+ /// builtin
+ void AddBuiltinParam(Function* func,
+ const std::string& name,
+ BuiltinValue builtin,
+ const core::type::Type* type) {
+ IOAttributes attr;
+ attr.builtin = builtin;
+ auto* p = b.FunctionParam(name, type);
+ p->SetAttributes(attr);
+ func->AppendParam(p);
+ }
+
+ /// Adds to a function an return value of type @p type, and decorated with @p builtin.
+ /// If there is an already existing non-structured return, both values are moved into a
+ /// structured return using @p name as the name.
+ /// If there is an already existing structured return, then this ICEs, since that is beyond the
+ /// scope of this implementation.
+ void AddBuiltinReturn(Function* func,
+ const std::string& name,
+ BuiltinValue builtin,
+ const core::type::Type* type) {
+ if (func->ReturnType()->Is<core::type::Struct>()) {
+ TINT_ICE() << "AddBuiltinReturn does not support adding to structured returns";
+ }
+
+ IOAttributes attr;
+ attr.builtin = builtin;
+ if (func->ReturnType() == ty.void_()) {
+ func->SetReturnAttributes(attr);
+ func->SetReturnType(type);
+ return;
+ }
+
+ std::string old_name =
+ func->ReturnAttributes().builtin == BuiltinValue::kPosition ? "pos" : "old_ret";
+ auto* str_ty =
+ ty.Struct(mod.symbols.New("OutputStruct"),
+ {
+ {mod.symbols.New(old_name), func->ReturnType(), func->ReturnAttributes()},
+ {mod.symbols.New(name), type, attr},
+ });
+
+ func->SetReturnAttributes({});
+ func->SetReturnType(str_ty);
+ }
+};
TEST_F(IR_ValidatorTest, RootBlock_Var) {
mod.root_block->Append(b.Var(ty.ptr<private_, i32>()));
@@ -321,8 +389,9 @@
}
TEST_F(IR_ValidatorTest, Function_Param_BothLocationAndBuiltin) {
- auto* f = b.Function("my_func", ty.void_());
- auto* p = b.FunctionParam("my_param", ty.f32());
+ auto* f = b.Function("my_func", ty.void_(), Function::PipelineStage::kFragment);
+
+ auto* p = b.FunctionParam("my_param", ty.vec4<f32>());
IOAttributes attr;
attr.builtin = BuiltinValue::kPosition;
attr.location = 0;
@@ -334,12 +403,12 @@
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
- R"(:1:17 error: a builtin and location cannot be both declared for a param
-%my_func = func(%my_param:f32 [@location(0), @position]):void {
- ^^^^^^^^^^^^^
+ R"(:1:27 error: a builtin and location cannot be both declared for a param
+%my_func = @fragment func(%my_param:vec4<f32> [@location(0), @position]):void {
+ ^^^^^^^^^^^^^^^^^^^
note: # Disassembly
-%my_func = func(%my_param:f32 [@location(0), @position]):void {
+%my_func = @fragment func(%my_param:vec4<f32> [@location(0), @position]):void {
$B1: {
ret
}
@@ -348,16 +417,15 @@
}
TEST_F(IR_ValidatorTest, Function_Param_Struct_BothLocationAndBuiltin) {
+ auto* f = b.Function("my_func", ty.void_(), Function::PipelineStage::kFragment);
+
IOAttributes attr;
attr.builtin = BuiltinValue::kPosition;
attr.location = 0;
-
auto* str_ty =
ty.Struct(mod.symbols.New("MyStruct"), {
- {mod.symbols.New("a"), ty.f32(), attr},
+ {mod.symbols.New("a"), ty.vec4<f32>(), attr},
});
-
- auto* f = b.Function("my_func", ty.void_());
auto* p = b.FunctionParam("my_param", str_ty);
f->SetParams({p});
@@ -366,16 +434,16 @@
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
- R"(:5:17 error: a builtin and location cannot be both declared for a struct member
-%my_func = func(%my_param:MyStruct):void {
- ^^^^^^^^^^^^^^^^^^
+ R"(:5:27 error: a builtin and location cannot be both declared for a struct member
+%my_func = @fragment func(%my_param:MyStruct):void {
+ ^^^^^^^^^^^^^^^^^^
note: # Disassembly
-MyStruct = struct @align(4) {
- a:f32 @offset(0), @location(0), @builtin(position)
+MyStruct = struct @align(16) {
+ a:vec4<f32> @offset(0), @location(0), @builtin(position)
}
-%my_func = func(%my_param:MyStruct):void {
+%my_func = @fragment func(%my_param:MyStruct):void {
$B1: {
ret
}
@@ -447,7 +515,8 @@
}
TEST_F(IR_ValidatorTest, Function_Param_InvariantWithPosition) {
- auto* f = b.Function("my_func", ty.void_());
+ auto* f = b.Function("my_func", ty.void_(), Function::PipelineStage::kFragment);
+
auto* p = b.FunctionParam("my_param", ty.vec4<f32>());
IOAttributes attr;
attr.builtin = BuiltinValue::kPosition;
@@ -489,16 +558,15 @@
}
TEST_F(IR_ValidatorTest, Function_Param_Struct_InvariantWithPosition) {
+ auto* f = b.Function("my_func", ty.void_(), Function::PipelineStage::kFragment);
+
IOAttributes attr;
attr.invariant = true;
attr.builtin = BuiltinValue::kPosition;
-
auto* str_ty =
ty.Struct(mod.symbols.New("MyStruct"), {
{mod.symbols.New("pos"), ty.vec4<f32>(), attr},
});
-
- auto* f = b.Function("my_func", ty.void_());
auto* p = b.FunctionParam("my_param", str_ty);
f->SetParams({p});
@@ -545,8 +613,7 @@
}
TEST_F(IR_ValidatorTest, Function_Return_BothLocationAndBuiltin) {
- auto* f = b.Function("my_func", ty.f32());
-
+ auto* f = b.Function("my_func", ty.vec4<f32>(), Function::PipelineStage::kVertex);
IOAttributes attr;
attr.builtin = BuiltinValue::kPosition;
attr.location = 0;
@@ -558,11 +625,11 @@
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
R"(:1:1 error: a builtin and location cannot be both declared for a function return
-%my_func = func():f32 [@location(0), @position] {
+%my_func = @vertex func():vec4<f32> [@location(0), @position] {
^^^^^^^^
note: # Disassembly
-%my_func = func():f32 [@location(0), @position] {
+%my_func = @vertex func():vec4<f32> [@location(0), @position] {
$B1: {
unreachable
}
@@ -574,29 +641,26 @@
IOAttributes attr;
attr.builtin = BuiltinValue::kPosition;
attr.location = 0;
-
auto* str_ty =
ty.Struct(mod.symbols.New("MyStruct"), {
- {mod.symbols.New("a"), ty.f32(), attr},
+ {mod.symbols.New("a"), ty.vec4<f32>(), attr},
});
-
- auto* f = b.Function("my_func", str_ty);
-
+ auto* f = b.Function("my_func", str_ty, Function::PipelineStage::kVertex);
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
R"(:5:1 error: a builtin and location cannot be both declared for a struct member
-%my_func = func():MyStruct {
+%my_func = @vertex func():MyStruct {
^^^^^^^^
note: # Disassembly
-MyStruct = struct @align(4) {
- a:f32 @offset(0), @location(0), @builtin(position)
+MyStruct = struct @align(16) {
+ a:vec4<f32> @offset(0), @location(0), @builtin(position)
}
-%my_func = func():MyStruct {
+%my_func = @vertex func():MyStruct {
$B1: {
unreachable
}
@@ -605,8 +669,7 @@
}
TEST_F(IR_ValidatorTest, Function_Return_NonVoid_MissingLocationAndBuiltin) {
- auto* f = b.Function("my_func", ty.f32());
- f->SetStage(Function::PipelineStage::kFragment);
+ auto* f = b.Function("my_func", ty.f32(), Function::PipelineStage::kFragment);
b.Append(f->Block(), [&] { b.Unreachable(); });
@@ -632,9 +695,7 @@
{mod.symbols.New("a"), ty.f32(), {}},
});
- auto* f = b.Function("my_func", str_ty);
- f->SetStage(Function::PipelineStage::kFragment);
-
+ auto* f = b.Function("my_func", str_ty, Function::PipelineStage::kFragment);
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
@@ -663,7 +724,7 @@
attr.builtin = BuiltinValue::kPosition;
attr.invariant = true;
- auto* f = b.Function("my_func", ty.vec4<f32>());
+ auto* f = b.Function("my_func", ty.vec4<f32>(), Function::PipelineStage::kVertex);
f->SetReturnAttributes(attr);
b.Append(f->Block(), [&] { b.Unreachable(); });
@@ -701,14 +762,12 @@
IOAttributes attr;
attr.invariant = true;
attr.builtin = BuiltinValue::kPosition;
-
auto* str_ty =
ty.Struct(mod.symbols.New("MyStruct"), {
{mod.symbols.New("pos"), ty.vec4<f32>(), attr},
});
- auto* f = b.Function("my_func", str_ty);
-
+ auto* f = b.Function("my_func", str_ty, Function::PipelineStage::kVertex);
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
@@ -725,7 +784,6 @@
});
auto* f = b.Function("my_func", str_ty);
-
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
@@ -770,9 +828,8 @@
}
TEST_F(IR_ValidatorTest, Function_UnnamedEntryPoint) {
- auto* f = b.Function(ty.void_());
- f->SetWorkgroupSize(0, 0, 0);
- f->SetStage(Function::PipelineStage::kCompute);
+ auto* f = b.Function(ty.void_(), Function::PipelineStage::kCompute,
+ std::array<uint32_t, 3>({0, 0, 0}));
b.Append(f->Block(), [&] { b.Return(f); });
@@ -859,42 +916,21 @@
)");
}
-TEST_F(IR_ValidatorTest, Function_UnexpectedFragDepth) {
- auto* f = b.Function("my_func", ty.void_());
- f->SetReturnBuiltin(BuiltinValue::kFragDepth);
- b.Append(f->Block(), [&] { b.Return(f); });
+TEST_F(IR_ValidatorTest, Function_Compute_NonVoidReturn) {
+ auto* f = b.Function("my_func", ty.f32(), Function::PipelineStage::kCompute,
+ std::array<uint32_t, 3>({0, 0, 0}));
- auto res = ir::Validate(mod);
- ASSERT_NE(res, Success);
- EXPECT_EQ(res.Failure().reason.Str(),
- R"(:1:1 error: frag_depth can only be declared for fragment entry points
-%my_func = func():void [@frag_depth] {
-^^^^^^^^
-
-note: # Disassembly
-%my_func = func():void [@frag_depth] {
- $B1: {
- ret
- }
-}
-)");
-}
-
-TEST_F(IR_ValidatorTest, Function_ComputeNonVoidReturn) {
- auto* f = b.Function("my_func", ty.f32());
- f->SetStage(Function::PipelineStage::kCompute);
- f->SetWorkgroupSize(1, 1, 1);
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
R"(:1:1 error: compute entry point must not have a return type
-%my_func = @compute @workgroup_size(1, 1, 1) func():f32 {
+%my_func = @compute @workgroup_size(0, 0, 0) func():f32 {
^^^^^^^^
note: # Disassembly
-%my_func = @compute @workgroup_size(1, 1, 1) func():f32 {
+%my_func = @compute @workgroup_size(0, 0, 0) func():f32 {
$B1: {
unreachable
}
@@ -902,9 +938,8 @@
)");
}
-TEST_F(IR_ValidatorTest, Function_VertexBasicPosition) {
- auto* f = b.Function("my_func", ty.vec4<f32>());
- f->SetStage(Function::PipelineStage::kVertex);
+TEST_F(IR_ValidatorTest, Function_Vertex_BasicPosition) {
+ auto* f = b.Function("my_func", ty.vec4<f32>(), Function::PipelineStage::kVertex);
f->SetReturnBuiltin(BuiltinValue::kPosition);
b.Append(f->Block(), [&] { b.Unreachable(); });
@@ -912,7 +947,7 @@
ASSERT_EQ(res, Success);
}
-TEST_F(IR_ValidatorTest, Function_VertexStructPosition) {
+TEST_F(IR_ValidatorTest, Function_Vertex_StructPosition) {
auto pos_ty = ty.vec4<f32>();
auto pos_attr = IOAttributes();
pos_attr.builtin = BuiltinValue::kPosition;
@@ -922,15 +957,14 @@
{mod.symbols.New("pos"), pos_ty, pos_attr},
});
- auto* f = b.Function("my_func", str_ty);
- f->SetStage(Function::PipelineStage::kVertex);
+ auto* f = b.Function("my_func", str_ty, Function::PipelineStage::kVertex);
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
ASSERT_EQ(res, Success);
}
-TEST_F(IR_ValidatorTest, Function_VertexStructPositionAndClipDistances) {
+TEST_F(IR_ValidatorTest, Function_Vertex_StructPositionAndClipDistances) {
auto pos_ty = ty.vec4<f32>();
auto pos_attr = IOAttributes();
pos_attr.builtin = BuiltinValue::kPosition;
@@ -945,15 +979,14 @@
{mod.symbols.New("clip"), clip_ty, clip_attr},
});
- auto* f = b.Function("my_func", str_ty);
- f->SetStage(Function::PipelineStage::kVertex);
+ auto* f = b.Function("my_func", str_ty, Function::PipelineStage::kVertex);
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
ASSERT_EQ(res, Success);
}
-TEST_F(IR_ValidatorTest, Function_VertexStructOnlyClipDistances) {
+TEST_F(IR_ValidatorTest, Function_Vertex_StructOnlyClipDistances) {
auto clip_ty = ty.array<f32, 4>();
auto clip_attr = IOAttributes();
clip_attr.builtin = BuiltinValue::kClipDistances;
@@ -963,8 +996,7 @@
{mod.symbols.New("clip"), clip_ty, clip_attr},
});
- auto* f = b.Function("my_func", str_ty);
- f->SetStage(Function::PipelineStage::kVertex);
+ auto* f = b.Function("my_func", str_ty, Function::PipelineStage::kVertex);
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
@@ -987,9 +1019,8 @@
)");
}
-TEST_F(IR_ValidatorTest, Function_VertexMissingPosition) {
- auto* f = b.Function("my_func", ty.vec4<f32>());
- f->SetStage(Function::PipelineStage::kVertex);
+TEST_F(IR_ValidatorTest, Function_Vertex_MissingPosition) {
+ auto* f = b.Function("my_func", ty.vec4<f32>(), Function::PipelineStage::kVertex);
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
@@ -1008,20 +1039,1126 @@
)");
}
-TEST_F(IR_ValidatorTest, Function_VertexPositionWrongType) {
- auto* f = b.Function("my_func", ty.void_());
- f->SetStage(Function::PipelineStage::kVertex);
+TEST_F(IR_ValidatorTest, Builtin_PointSize_WrongStage) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinReturn(f, "size", BuiltinValue::kPointSize, ty.f32());
+
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
- R"(:1:1 error: position must be declared for vertex entry point output
-%my_func = @vertex func():void {
-^^^^^^^^
+ R"(:1:1 error: __point_size must be used in a vertex shader entry point
+%f = @fragment func():f32 [@__point_size] {
+^^
note: # Disassembly
-%my_func = @vertex func():void {
+%f = @fragment func():f32 [@__point_size] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_PointSize_WrongIODirection) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinParam(f, "size", BuiltinValue::kPointSize, ty.f32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:19 error: __point_size must be an output of a shader entry point
+%f = @vertex func(%size:f32 [@__point_size]):vec4<f32> [@position] {
+ ^^^^^^^^^
+
+note: # Disassembly
+%f = @vertex func(%size:f32 [@__point_size]):vec4<f32> [@position] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_PointSize_WrongType) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinReturn(f, "size", BuiltinValue::kPointSize, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:6:1 error: __point_size must be a f32
+%f = @vertex func():OutputStruct {
+^^
+
+note: # Disassembly
+OutputStruct = struct @align(16) {
+ pos:vec4<f32> @offset(0), @builtin(position)
+ size:u32 @offset(16), @builtin(__point_size)
+}
+
+%f = @vertex func():OutputStruct {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_ClipDistances_WrongStage) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinReturn(f, "distances", BuiltinValue::kClipDistances, ty.array<f32, 2>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: clip_distances must be used in a vertex shader entry point
+%f = @fragment func():array<f32, 2> [@clip_distances] {
+^^
+
+note: # Disassembly
+%f = @fragment func():array<f32, 2> [@clip_distances] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_ClipDistances_WrongIODirection) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinParam(f, "distances", BuiltinValue::kClipDistances, ty.array<f32, 2>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:19 error: clip_distances must be an output of a shader entry point
+%f = @vertex func(%distances:array<f32, 2> [@clip_distances]):vec4<f32> [@position] {
+ ^^^^^^^^^^^^^^^^^^^^^^^^
+
+note: # Disassembly
+%f = @vertex func(%distances:array<f32, 2> [@clip_distances]):vec4<f32> [@position] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_ClipDistances_WrongType) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinReturn(f, "distances", BuiltinValue::kClipDistances, ty.f32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:6:1 error: clip_distances must be an array<f32, N>, where N <= 8
+%f = @vertex func():OutputStruct {
+^^
+
+note: # Disassembly
+OutputStruct = struct @align(16) {
+ pos:vec4<f32> @offset(0), @builtin(position)
+ distances:f32 @offset(16), @builtin(clip_distances)
+}
+
+%f = @vertex func():OutputStruct {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_FragDepth_WrongStage) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinReturn(f, "depth", BuiltinValue::kFragDepth, ty.f32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:6:1 error: frag_depth must be used in a fragment shader entry point
+%f = @vertex func():OutputStruct {
+^^
+
+note: # Disassembly
+OutputStruct = struct @align(16) {
+ pos:vec4<f32> @offset(0), @builtin(position)
+ depth:f32 @offset(16), @builtin(frag_depth)
+}
+
+%f = @vertex func():OutputStruct {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_FragDepth_WrongIODirection) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "depth", BuiltinValue::kFragDepth, ty.f32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:21 error: frag_depth must be an output of a shader entry point
+%f = @fragment func(%depth:f32 [@frag_depth]):void {
+ ^^^^^^^^^^
+
+note: # Disassembly
+%f = @fragment func(%depth:f32 [@frag_depth]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_FragDepth_WrongType) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinReturn(f, "depth", BuiltinValue::kFragDepth, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: frag_depth must be a f32
+%f = @fragment func():u32 [@frag_depth] {
+^^
+
+note: # Disassembly
+%f = @fragment func():u32 [@frag_depth] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_FrontFacing_WrongStage) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinParam(f, "facing", BuiltinValue::kFrontFacing, ty.bool_());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:19 error: front_facing must be used in a fragment shader entry point
+%f = @vertex func(%facing:bool [@front_facing]):vec4<f32> [@position] {
+ ^^^^^^^^^^^^
+
+note: # Disassembly
+%f = @vertex func(%facing:bool [@front_facing]):vec4<f32> [@position] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_FrontFacing_WrongIODirection) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinReturn(f, "facing", BuiltinValue::kFrontFacing, ty.bool_());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: front_facing must be an input of a shader entry point
+%f = @fragment func():bool [@front_facing] {
+^^
+
+note: # Disassembly
+%f = @fragment func():bool [@front_facing] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_FrontFacing_WrongType) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "facing", BuiltinValue::kFrontFacing, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:21 error: front_facing must be a bool
+%f = @fragment func(%facing:u32 [@front_facing]):void {
+ ^^^^^^^^^^^
+
+note: # Disassembly
+%f = @fragment func(%facing:u32 [@front_facing]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_GlobalInvocationId_WrongStage) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "invocation", BuiltinValue::kGlobalInvocationId, ty.vec3<u32>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:21 error: global_invocation_id must be used in a compute shader entry point
+%f = @fragment func(%invocation:vec3<u32> [@global_invocation_id]):void {
+ ^^^^^^^^^^^^^^^^^^^^^
+
+note: # Disassembly
+%f = @fragment func(%invocation:vec3<u32> [@global_invocation_id]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_GlobalInvocationId_WrongIODirection) {
+ // This will also trigger the compute entry points should have void returns check
+ auto* f = ComputeEntryPoint();
+ AddBuiltinReturn(f, "invocation", BuiltinValue::kGlobalInvocationId, ty.vec3<u32>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: global_invocation_id must be an input of a shader entry point
+%f = @compute @workgroup_size(0, 0, 0) func():vec3<u32> [@global_invocation_id] {
+^^
+
+:1:1 error: compute entry point must not have a return type
+%f = @compute @workgroup_size(0, 0, 0) func():vec3<u32> [@global_invocation_id] {
+^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func():vec3<u32> [@global_invocation_id] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_GlobalInvocationId_WrongType) {
+ auto* f = ComputeEntryPoint();
+ AddBuiltinParam(f, "invocation", BuiltinValue::kGlobalInvocationId, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:45 error: global_invocation_id must be an vec3<u32>
+%f = @compute @workgroup_size(0, 0, 0) func(%invocation:u32 [@global_invocation_id]):void {
+ ^^^^^^^^^^^^^^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func(%invocation:u32 [@global_invocation_id]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_InstanceIndex_WrongStage) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "instance", BuiltinValue::kInstanceIndex, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:21 error: instance_index must be used in a vertex shader entry point
+%f = @fragment func(%instance:u32 [@instance_index]):void {
+ ^^^^^^^^^^^^^
+
+note: # Disassembly
+%f = @fragment func(%instance:u32 [@instance_index]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_InstanceIndex_WrongIODirection) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinReturn(f, "instance", BuiltinValue::kInstanceIndex, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:6:1 error: instance_index must be an input of a shader entry point
+%f = @vertex func():OutputStruct {
+^^
+
+note: # Disassembly
+OutputStruct = struct @align(16) {
+ pos:vec4<f32> @offset(0), @builtin(position)
+ instance:u32 @offset(16), @builtin(instance_index)
+}
+
+%f = @vertex func():OutputStruct {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_InstanceIndex_WrongType) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinParam(f, "instance", BuiltinValue::kInstanceIndex, ty.i32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:19 error: instance_index must be an u32
+%f = @vertex func(%instance:i32 [@instance_index]):vec4<f32> [@position] {
+ ^^^^^^^^^^^^^
+
+note: # Disassembly
+%f = @vertex func(%instance:i32 [@instance_index]):vec4<f32> [@position] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_LocalInvocationId_WrongStage) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "id", BuiltinValue::kLocalInvocationId, ty.vec3<u32>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:21 error: local_invocation_id must be used in a compute shader entry point
+%f = @fragment func(%id:vec3<u32> [@local_invocation_id]):void {
+ ^^^^^^^^^^^^^
+
+note: # Disassembly
+%f = @fragment func(%id:vec3<u32> [@local_invocation_id]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_LocalInvocationId_WrongIODirection) {
+ // This will also trigger the compute entry points should have void returns check
+ auto* f = ComputeEntryPoint();
+ AddBuiltinReturn(f, "id", BuiltinValue::kLocalInvocationId, ty.vec3<u32>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: local_invocation_id must be an input of a shader entry point
+%f = @compute @workgroup_size(0, 0, 0) func():vec3<u32> [@local_invocation_id] {
+^^
+
+:1:1 error: compute entry point must not have a return type
+%f = @compute @workgroup_size(0, 0, 0) func():vec3<u32> [@local_invocation_id] {
+^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func():vec3<u32> [@local_invocation_id] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_LocalInvocationId_WrongType) {
+ auto* f = ComputeEntryPoint();
+ AddBuiltinParam(f, "id", BuiltinValue::kLocalInvocationId, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:45 error: local_invocation_id must be an vec3<u32>
+%f = @compute @workgroup_size(0, 0, 0) func(%id:u32 [@local_invocation_id]):void {
+ ^^^^^^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func(%id:u32 [@local_invocation_id]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_LocalInvocationIndex_WrongStage) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "index", BuiltinValue::kLocalInvocationIndex, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:21 error: local_invocation_index must be used in a compute shader entry point
+%f = @fragment func(%index:u32 [@local_invocation_index]):void {
+ ^^^^^^^^^^
+
+note: # Disassembly
+%f = @fragment func(%index:u32 [@local_invocation_index]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_LocalInvocationIndex_WrongIODirection) {
+ // This will also trigger the compute entry points should have void returns check
+ auto* f = ComputeEntryPoint();
+ AddBuiltinReturn(f, "index", BuiltinValue::kLocalInvocationIndex, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: local_invocation_index must be an input of a shader entry point
+%f = @compute @workgroup_size(0, 0, 0) func():u32 [@local_invocation_index] {
+^^
+
+:1:1 error: compute entry point must not have a return type
+%f = @compute @workgroup_size(0, 0, 0) func():u32 [@local_invocation_index] {
+^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func():u32 [@local_invocation_index] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_LocalInvocationIndex_WrongType) {
+ auto* f = ComputeEntryPoint();
+ AddBuiltinParam(f, "index", BuiltinValue::kLocalInvocationIndex, ty.i32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:45 error: local_invocation_index must be an u32
+%f = @compute @workgroup_size(0, 0, 0) func(%index:i32 [@local_invocation_index]):void {
+ ^^^^^^^^^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func(%index:i32 [@local_invocation_index]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_NumWorkgroups_WrongStage) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "num", BuiltinValue::kNumWorkgroups, ty.vec3<u32>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:21 error: num_workgroups must be used in a compute shader entry point
+%f = @fragment func(%num:vec3<u32> [@num_workgroups]):void {
+ ^^^^^^^^^^^^^^
+
+note: # Disassembly
+%f = @fragment func(%num:vec3<u32> [@num_workgroups]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_NumWorkgroups_WrongIODirection) {
+ // This will also trigger the compute entry points should have void returns check
+ auto* f = ComputeEntryPoint();
+ AddBuiltinReturn(f, "num", BuiltinValue::kNumWorkgroups, ty.vec3<u32>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: num_workgroups must be an input of a shader entry point
+%f = @compute @workgroup_size(0, 0, 0) func():vec3<u32> [@num_workgroups] {
+^^
+
+:1:1 error: compute entry point must not have a return type
+%f = @compute @workgroup_size(0, 0, 0) func():vec3<u32> [@num_workgroups] {
+^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func():vec3<u32> [@num_workgroups] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_NumWorkgroups_WrongType) {
+ auto* f = ComputeEntryPoint();
+ AddBuiltinParam(f, "num", BuiltinValue::kNumWorkgroups, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:45 error: num_workgroups must be an vec3<u32>
+%f = @compute @workgroup_size(0, 0, 0) func(%num:u32 [@num_workgroups]):void {
+ ^^^^^^^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func(%num:u32 [@num_workgroups]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SampleIndex_WrongStage) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinParam(f, "index", BuiltinValue::kSampleIndex, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:19 error: sample_index must be used in a fragment shader entry point
+%f = @vertex func(%index:u32 [@sample_index]):vec4<f32> [@position] {
+ ^^^^^^^^^^
+
+note: # Disassembly
+%f = @vertex func(%index:u32 [@sample_index]):vec4<f32> [@position] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SampleIndex_WrongIODirection) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinReturn(f, "index", BuiltinValue::kSampleIndex, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: sample_index must be an input of a shader entry point
+%f = @fragment func():u32 [@sample_index] {
+^^
+
+note: # Disassembly
+%f = @fragment func():u32 [@sample_index] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SampleIndex_WrongType) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "index", BuiltinValue::kSampleIndex, ty.f32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:21 error: sample_index must be an u32
+%f = @fragment func(%index:f32 [@sample_index]):void {
+ ^^^^^^^^^^
+
+note: # Disassembly
+%f = @fragment func(%index:f32 [@sample_index]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_VertexIndex_WrongStage) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "index", BuiltinValue::kVertexIndex, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:21 error: vertex_index must be used in a vertex shader entry point
+%f = @fragment func(%index:u32 [@vertex_index]):void {
+ ^^^^^^^^^^
+
+note: # Disassembly
+%f = @fragment func(%index:u32 [@vertex_index]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_VertexIndex_WrongIODirection) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinReturn(f, "index", BuiltinValue::kVertexIndex, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:6:1 error: vertex_index must be an input of a shader entry point
+%f = @vertex func():OutputStruct {
+^^
+
+note: # Disassembly
+OutputStruct = struct @align(16) {
+ pos:vec4<f32> @offset(0), @builtin(position)
+ index:u32 @offset(16), @builtin(vertex_index)
+}
+
+%f = @vertex func():OutputStruct {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_VertexIndex_WrongType) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinParam(f, "index", BuiltinValue::kVertexIndex, ty.f32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:19 error: vertex_index must be an u32
+%f = @vertex func(%index:f32 [@vertex_index]):vec4<f32> [@position] {
+ ^^^^^^^^^^
+
+note: # Disassembly
+%f = @vertex func(%index:f32 [@vertex_index]):vec4<f32> [@position] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_WorkgroupId_WrongStage) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "id", BuiltinValue::kWorkgroupId, ty.vec3<u32>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:21 error: workgroup_id must be used in a compute shader entry point
+%f = @fragment func(%id:vec3<u32> [@workgroup_id]):void {
+ ^^^^^^^^^^^^^
+
+note: # Disassembly
+%f = @fragment func(%id:vec3<u32> [@workgroup_id]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_WorkgroupId_WrongIODirection) {
+ // This will also trigger the compute entry points should have void returns check
+ auto* f = ComputeEntryPoint();
+ AddBuiltinReturn(f, "id", BuiltinValue::kWorkgroupId, ty.vec3<u32>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: workgroup_id must be an input of a shader entry point
+%f = @compute @workgroup_size(0, 0, 0) func():vec3<u32> [@workgroup_id] {
+^^
+
+:1:1 error: compute entry point must not have a return type
+%f = @compute @workgroup_size(0, 0, 0) func():vec3<u32> [@workgroup_id] {
+^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func():vec3<u32> [@workgroup_id] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_WorkgroupId_WrongType) {
+ auto* f = ComputeEntryPoint();
+ AddBuiltinParam(f, "id", BuiltinValue::kWorkgroupId, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:45 error: workgroup_id must be an vec3<u32>
+%f = @compute @workgroup_size(0, 0, 0) func(%id:u32 [@workgroup_id]):void {
+ ^^^^^^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func(%id:u32 [@workgroup_id]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_Position_WrongStage) {
+ auto* f = ComputeEntryPoint();
+ AddBuiltinParam(f, "pos", BuiltinValue::kPosition, ty.vec4<f32>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:45 error: position must be used in a fragment or vertex shader entry point
+%f = @compute @workgroup_size(0, 0, 0) func(%pos:vec4<f32> [@position]):void {
+ ^^^^^^^^^^^^^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func(%pos:vec4<f32> [@position]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_Position_WrongIODirectionForVertex) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinParam(f, "pos", BuiltinValue::kPosition, ty.vec4<f32>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:19 error: position must be an output for a vertex entry point
+%f = @vertex func(%pos:vec4<f32> [@position]):vec4<f32> [@position] {
+ ^^^^^^^^^^^^^^
+
+note: # Disassembly
+%f = @vertex func(%pos:vec4<f32> [@position]):vec4<f32> [@position] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_Position_WrongIODirectionForFragment) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinReturn(f, "pos", BuiltinValue::kPosition, ty.vec4<f32>());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: position must be an input for a fragment entry point
+%f = @fragment func():vec4<f32> [@position] {
+^^
+
+note: # Disassembly
+%f = @fragment func():vec4<f32> [@position] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_Position_WrongType) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "pos", BuiltinValue::kPosition, ty.f32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:21 error: position must be an vec4<f32>
+%f = @fragment func(%pos:f32 [@position]):void {
+ ^^^^^^^^
+
+note: # Disassembly
+%f = @fragment func(%pos:f32 [@position]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SampleMask_WrongStage) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinParam(f, "mask", BuiltinValue::kSampleMask, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:19 error: sample_mask must be used in a fragment entry point
+%f = @vertex func(%mask:u32 [@sample_mask]):vec4<f32> [@position] {
+ ^^^^^^^^^
+
+note: # Disassembly
+%f = @vertex func(%mask:u32 [@sample_mask]):vec4<f32> [@position] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SampleMask_InputValid) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "mask", BuiltinValue::kSampleMask, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_EQ(res, Success);
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SampleMask_OutputValid) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinReturn(f, "mask", BuiltinValue::kSampleMask, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_EQ(res, Success);
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SampleMask_WrongType) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinParam(f, "mask", BuiltinValue::kSampleMask, ty.f32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:21 error: sample_mask must be an u32
+%f = @fragment func(%mask:f32 [@sample_mask]):void {
+ ^^^^^^^^^
+
+note: # Disassembly
+%f = @fragment func(%mask:f32 [@sample_mask]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SubgroupSize_WrongStage) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinParam(f, "size", BuiltinValue::kSubgroupSize, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:19 error: subgroup_size must be used in a compute or fragment shader entry point
+%f = @vertex func(%size:u32 [@subgroup_size]):vec4<f32> [@position] {
+ ^^^^^^^^^
+
+note: # Disassembly
+%f = @vertex func(%size:u32 [@subgroup_size]):vec4<f32> [@position] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SubgroupSize_WrongIODirection) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinReturn(f, "size", BuiltinValue::kSubgroupSize, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: subgroup_size must be an input of a shader entry point
+%f = @fragment func():u32 [@subgroup_size] {
+^^
+
+note: # Disassembly
+%f = @fragment func():u32 [@subgroup_size] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SubgroupSize_WrongType) {
+ auto* f = ComputeEntryPoint();
+ AddBuiltinParam(f, "size", BuiltinValue::kSubgroupSize, ty.i32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:45 error: subgroup_size must be an u32
+%f = @compute @workgroup_size(0, 0, 0) func(%size:i32 [@subgroup_size]):void {
+ ^^^^^^^^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func(%size:i32 [@subgroup_size]):void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SubgroupInvocationId_WrongStage) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinParam(f, "id", BuiltinValue::kSubgroupInvocationId, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(
+ res.Failure().reason.Str(),
+ R"(:1:19 error: subgroup_invocation_id must be used in a compute or fragment shader entry point
+%f = @vertex func(%id:u32 [@subgroup_invocation_id]):vec4<f32> [@position] {
+ ^^^^^^^
+
+note: # Disassembly
+%f = @vertex func(%id:u32 [@subgroup_invocation_id]):vec4<f32> [@position] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SubgroupInvocationId_WrongIODirection) {
+ auto* f = FragmentEntryPoint();
+ AddBuiltinReturn(f, "id", BuiltinValue::kSubgroupInvocationId, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: subgroup_invocation_id must be an input of a shader entry point
+%f = @fragment func():u32 [@subgroup_invocation_id] {
+^^
+
+note: # Disassembly
+%f = @fragment func():u32 [@subgroup_invocation_id] {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Builtin_SubgroupInvocationId_WrongType) {
+ auto* f = ComputeEntryPoint();
+ AddBuiltinParam(f, "id", BuiltinValue::kSubgroupInvocationId, ty.i32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:45 error: subgroup_invocation_id must be an u32
+%f = @compute @workgroup_size(0, 0, 0) func(%id:i32 [@subgroup_invocation_id]):void {
+ ^^^^^^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(0, 0, 0) func(%id:i32 [@subgroup_invocation_id]):void {
$B1: {
unreachable
}
@@ -1063,8 +2200,8 @@
TEST_F(IR_ValidatorTest, CallToEntryPointFunction) {
auto* f = b.Function("f", ty.void_());
- auto* g = b.Function("g", ty.void_(), Function::PipelineStage::kCompute);
- g->SetWorkgroupSize(1, 1, 1);
+ auto* g = b.Function("g", ty.void_(), Function::PipelineStage::kCompute,
+ std::array<uint32_t, 3>({0, 0, 0}));
b.Append(f->Block(), [&] {
b.Call(g);
@@ -1090,7 +2227,7 @@
ret
}
}
-%g = @compute @workgroup_size(1, 1, 1) func():void {
+%g = @compute @workgroup_size(0, 0, 0) func():void {
$B2: {
ret
}
@@ -2009,8 +3146,7 @@
b.Return(func);
});
- auto* ep = b.Function("ep", ty.void_());
- ep->SetStage(Function::PipelineStage::kFragment);
+ auto* ep = b.Function("ep", ty.void_(), Function::PipelineStage::kFragment);
b.Append(ep->Block(), [&] {
b.Call(func);
b.Return(ep);
@@ -2051,8 +3187,7 @@
b.Return(func);
});
- auto* ep = b.Function("ep", ty.void_());
- ep->SetStage(Function::PipelineStage::kFragment);
+ auto* ep = b.Function("ep", ty.void_(), Function::PipelineStage::kFragment);
b.Append(ep->Block(), [&] {
b.Call(func);
b.Return(ep);
@@ -2092,9 +3227,9 @@
b.Return(func);
});
- auto* ep = b.Function("ep", ty.void_());
- ep->SetStage(Function::PipelineStage::kCompute);
- ep->SetWorkgroupSize(0, 0, 0);
+ auto* ep = b.Function("ep", ty.void_(), Function::PipelineStage::kCompute,
+ std::array<uint32_t, 3>({0, 0, 0}));
+
b.Append(ep->Block(), [&] {
b.Call(func);
b.Return(ep);
diff --git a/src/tint/lang/hlsl/writer/raise/pixel_local_test.cc b/src/tint/lang/hlsl/writer/raise/pixel_local_test.cc
index 234da63..59c5d2e 100644
--- a/src/tint/lang/hlsl/writer/raise/pixel_local_test.cc
+++ b/src/tint/lang/hlsl/writer/raise/pixel_local_test.cc
@@ -64,7 +64,7 @@
if (multiple_builtins) {
attrs.builtin = core::BuiltinValue::kFrontFacing;
members.Emplace(mod.symbols.New("front_facing"), ty.bool_(), attrs);
- attrs.builtin = core::BuiltinValue::kFragDepth;
+ attrs.builtin = core::BuiltinValue::kSampleIndex;
members.Emplace(mod.symbols.New("sample_index"), ty.u32(), attrs);
}
auto* param_struct_ty = ty.Struct(mod.symbols.New("params"), members);
@@ -445,7 +445,7 @@
params = struct @align(16) {
pos:vec4<f32> @offset(0), @builtin(position)
front_facing:bool @offset(16), @builtin(front_facing)
- sample_index:u32 @offset(20), @builtin(frag_depth)
+ sample_index:u32 @offset(20), @builtin(sample_index)
}
$B1: { # root
@@ -474,7 +474,7 @@
params = struct @align(16) {
pos:vec4<f32> @offset(0), @builtin(position)
front_facing:bool @offset(16), @builtin(front_facing)
- sample_index:u32 @offset(20), @builtin(frag_depth)
+ sample_index:u32 @offset(20), @builtin(sample_index)
}
$B1: { # root
diff --git a/src/tint/utils/containers/enum_set.h b/src/tint/utils/containers/enum_set.h
index 5736c13..921cfda 100644
--- a/src/tint/utils/containers/enum_set.h
+++ b/src/tint/utils/containers/enum_set.h
@@ -148,6 +148,18 @@
/// @return true if the set is empty
inline bool Empty() const { return set == 0; }
+ /// @returns number of enums currently in the set
+ /// This is an O(N) operation, where N can be upto 64
+ inline size_t Size() const {
+ size_t result = 0;
+ uint64_t bits = set;
+ while (bits) {
+ result += bits & 1;
+ bits >>= 1;
+ }
+ return result;
+ }
+
/// @return the hash value of this object
tint::HashCode HashCode() const { return Hash(Value()); }
