spirv-val: Add Mesh/Task to check for LocalSize (#6459)
obvious oversight of Mesh/Task that is being fixed up in the spec here
https://gitlab.khronos.org/vulkan/vulkan/-/merge_requests/7901
(update, we approved the spec language, will be out in the 1.4.336 spec)
diff --git a/source/val/validate_mode_setting.cpp b/source/val/validate_mode_setting.cpp
index f2b43b6..22d464f 100644
--- a/source/val/validate_mode_setting.cpp
+++ b/source/val/validate_mode_setting.cpp
@@ -27,6 +27,48 @@
namespace val {
namespace {
+// TODO - Make a common util if someone else needs it too outside this file
+const char* ExecutionModelToString(spv::ExecutionModel value) {
+ switch (value) {
+ case spv::ExecutionModel::Vertex:
+ return "Vertex";
+ case spv::ExecutionModel::TessellationControl:
+ return "TessellationControl";
+ case spv::ExecutionModel::TessellationEvaluation:
+ return "TessellationEvaluation";
+ case spv::ExecutionModel::Geometry:
+ return "Geometry";
+ case spv::ExecutionModel::Fragment:
+ return "Fragment";
+ case spv::ExecutionModel::GLCompute:
+ return "GLCompute";
+ case spv::ExecutionModel::Kernel:
+ return "Kernel";
+ case spv::ExecutionModel::TaskNV:
+ return "TaskNV";
+ case spv::ExecutionModel::MeshNV:
+ return "MeshNV";
+ case spv::ExecutionModel::RayGenerationKHR:
+ return "RayGenerationKHR";
+ case spv::ExecutionModel::IntersectionKHR:
+ return "IntersectionKHR";
+ case spv::ExecutionModel::AnyHitKHR:
+ return "AnyHitKHR";
+ case spv::ExecutionModel::ClosestHitKHR:
+ return "ClosestHitKHR";
+ case spv::ExecutionModel::MissKHR:
+ return "MissKHR";
+ case spv::ExecutionModel::CallableKHR:
+ return "CallableKHR";
+ case spv::ExecutionModel::TaskEXT:
+ return "TaskEXT";
+ case spv::ExecutionModel::MeshEXT:
+ return "MeshEXT";
+ default:
+ return "Unknown";
+ }
+}
+
spv_result_t ValidateEntryPoint(ValidationState_t& _, const Instruction* inst) {
const auto entry_point_id = inst->GetOperandAs<uint32_t>(1);
auto entry_point = _.FindDef(entry_point_id);
@@ -306,74 +348,79 @@
}
if (spvIsVulkanEnv(_.context()->target_env)) {
- switch (execution_model) {
- case spv::ExecutionModel::GLCompute:
- if (!has_mode(spv::ExecutionMode::LocalSize)) {
- bool ok = has_workgroup_size || has_local_size_id;
- if (!ok && _.HasCapability(spv::Capability::TileShadingQCOM)) {
- ok = has_mode(spv::ExecutionMode::TileShadingRateQCOM);
- }
- if (!ok) {
+ // SPV_QCOM_tile_shading checks
+ if (execution_model == spv::ExecutionModel::GLCompute) {
+ if (_.HasCapability(spv::Capability::TileShadingQCOM)) {
+ if (has_mode(spv::ExecutionMode::TileShadingRateQCOM) &&
+ (has_mode(spv::ExecutionMode::LocalSize) ||
+ has_mode(spv::ExecutionMode::LocalSizeId))) {
+ return _.diag(SPV_ERROR_INVALID_DATA, inst)
+ << "If the TileShadingRateQCOM execution mode is used, "
+ << "LocalSize and LocalSizeId must not be specified.";
+ }
+ if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) {
+ return _.diag(SPV_ERROR_INVALID_DATA, inst)
+ << "The NonCoherentTileAttachmentQCOM execution mode must "
+ "not be used in any stage other than fragment.";
+ }
+ } else {
+ if (has_mode(spv::ExecutionMode::TileShadingRateQCOM)) {
+ return _.diag(SPV_ERROR_INVALID_DATA, inst)
+ << "If the TileShadingRateQCOM execution mode is used, the "
+ "TileShadingQCOM capability must be enabled.";
+ }
+ }
+ } else {
+ if (has_mode(spv::ExecutionMode::TileShadingRateQCOM)) {
+ return _.diag(SPV_ERROR_INVALID_DATA, inst)
+ << "The TileShadingRateQCOM execution mode must not be used "
+ "in any stage other than compute.";
+ }
+ if (execution_model != spv::ExecutionModel::Fragment) {
+ if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) {
+ return _.diag(SPV_ERROR_INVALID_DATA, inst)
+ << "The NonCoherentTileAttachmentQCOM execution mode must "
+ "not be used in any stage other than fragment.";
+ }
+ if (_.HasCapability(spv::Capability::TileShadingQCOM)) {
+ return _.diag(SPV_ERROR_INVALID_CAPABILITY, inst)
+ << "The TileShadingQCOM capability must not be enabled in "
+ "any stage other than compute or fragment.";
+ }
+ } else {
+ if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) {
+ if (!_.HasCapability(spv::Capability::TileShadingQCOM)) {
return _.diag(SPV_ERROR_INVALID_DATA, inst)
- << _.VkErrorID(10685)
- << "In the Vulkan environment, GLCompute execution model "
- "entry points require either the "
- << (_.HasCapability(spv::Capability::TileShadingQCOM)
- ? "TileShadingRateQCOM, "
- : "")
- << "LocalSize or LocalSizeId execution mode or an object "
- "decorated with WorkgroupSize must be specified.";
+ << "If the NonCoherentTileAttachmentReadQCOM execution "
+ "mode is used, the TileShadingQCOM capability must be "
+ "enabled.";
}
}
+ }
+ }
- if (_.HasCapability(spv::Capability::TileShadingQCOM)) {
- if (has_mode(spv::ExecutionMode::TileShadingRateQCOM) &&
- (has_mode(spv::ExecutionMode::LocalSize) ||
- has_mode(spv::ExecutionMode::LocalSizeId))) {
- return _.diag(SPV_ERROR_INVALID_DATA, inst)
- << "If the TileShadingRateQCOM execution mode is used, "
- << "LocalSize and LocalSizeId must not be specified.";
- }
- if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) {
- return _.diag(SPV_ERROR_INVALID_DATA, inst)
- << "The NonCoherentTileAttachmentQCOM execution mode must "
- "not be used in any stage other than fragment.";
- }
- } else {
- if (has_mode(spv::ExecutionMode::TileShadingRateQCOM)) {
- return _.diag(SPV_ERROR_INVALID_DATA, inst)
- << "If the TileShadingRateQCOM execution mode is used, the "
- "TileShadingQCOM capability must be enabled.";
- }
+ switch (execution_model) {
+ case spv::ExecutionModel::GLCompute:
+ case spv::ExecutionModel::MeshEXT:
+ case spv::ExecutionModel::MeshNV:
+ case spv::ExecutionModel::TaskEXT:
+ case spv::ExecutionModel::TaskNV:
+ if (!has_mode(spv::ExecutionMode::LocalSize) && !has_workgroup_size &&
+ !has_local_size_id &&
+ !has_mode(spv::ExecutionMode::TileShadingRateQCOM)) {
+ return _.diag(SPV_ERROR_INVALID_DATA, inst)
+ << _.VkErrorID(10685) << "In the Vulkan environment, "
+ << ExecutionModelToString(execution_model)
+ << " execution model "
+ "entry points require either the "
+ << (_.HasCapability(spv::Capability::TileShadingQCOM)
+ ? "TileShadingRateQCOM, "
+ : "")
+ << "LocalSize or LocalSizeId execution mode or an object "
+ "decorated with WorkgroupSize must be specified.";
}
break;
default:
- if (has_mode(spv::ExecutionMode::TileShadingRateQCOM)) {
- return _.diag(SPV_ERROR_INVALID_DATA, inst)
- << "The TileShadingRateQCOM execution mode must not be used "
- "in any stage other than compute.";
- }
- if (execution_model != spv::ExecutionModel::Fragment) {
- if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) {
- return _.diag(SPV_ERROR_INVALID_DATA, inst)
- << "The NonCoherentTileAttachmentQCOM execution mode must "
- "not be used in any stage other than fragment.";
- }
- if (_.HasCapability(spv::Capability::TileShadingQCOM)) {
- return _.diag(SPV_ERROR_INVALID_CAPABILITY, inst)
- << "The TileShadingQCOM capability must not be enabled in "
- "any stage other than compute or fragment.";
- }
- } else {
- if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) {
- if (!_.HasCapability(spv::Capability::TileShadingQCOM)) {
- return _.diag(SPV_ERROR_INVALID_DATA, inst)
- << "If the NonCoherentTileAttachmentReadQCOM execution "
- "mode is used, the TileShadingQCOM capability must be "
- "enabled.";
- }
- }
- }
break;
}
}
diff --git a/test/val/val_builtins_test.cpp b/test/val/val_builtins_test.cpp
index 65d6950..8d5f362 100644
--- a/test/val/val_builtins_test.cpp
+++ b/test/val/val_builtins_test.cpp
@@ -141,7 +141,11 @@
execution_modes << "OpExecutionMode %" << entry_point.name
<< " OutputPoints\n";
}
- if (0 == std::strcmp(execution_model, "GLCompute")) {
+ if (0 == std::strcmp(execution_model, "GLCompute") ||
+ 0 == std::strcmp(execution_model, "MeshEXT") ||
+ 0 == std::strcmp(execution_model, "MeshNV") ||
+ 0 == std::strcmp(execution_model, "MeshEXT") ||
+ 0 == std::strcmp(execution_model, "TaskNV")) {
execution_modes << "OpExecutionMode %" << entry_point.name
<< " LocalSize 1 1 1\n";
}
@@ -303,7 +307,11 @@
execution_modes << "OpExecutionMode %" << entry_point.name
<< " OutputPoints\n";
}
- if (0 == std::strcmp(execution_model, "GLCompute")) {
+ if (0 == std::strcmp(execution_model, "GLCompute") ||
+ 0 == std::strcmp(execution_model, "MeshEXT") ||
+ 0 == std::strcmp(execution_model, "MeshNV") ||
+ 0 == std::strcmp(execution_model, "MeshEXT") ||
+ 0 == std::strcmp(execution_model, "TaskNV")) {
execution_modes << "OpExecutionMode %" << entry_point.name
<< " LocalSize 1 1 1\n";
}
@@ -452,7 +460,11 @@
execution_modes << "OpExecutionMode %" << entry_point.name
<< " OutputPoints\n";
}
- if (0 == std::strcmp(execution_model, "GLCompute")) {
+ if (0 == std::strcmp(execution_model, "GLCompute") ||
+ 0 == std::strcmp(execution_model, "MeshEXT") ||
+ 0 == std::strcmp(execution_model, "MeshNV") ||
+ 0 == std::strcmp(execution_model, "MeshEXT") ||
+ 0 == std::strcmp(execution_model, "TaskNV")) {
execution_modes << "OpExecutionMode %" << entry_point.name
<< " LocalSize 1 1 1\n";
}
@@ -2703,7 +2715,11 @@
execution_modes << "OpExecutionMode %" << entry_point.name
<< " OutputPoints\n";
}
- if (0 == std::strcmp(execution_model, "GLCompute")) {
+ if (0 == std::strcmp(execution_model, "GLCompute") ||
+ 0 == std::strcmp(execution_model, "MeshEXT") ||
+ 0 == std::strcmp(execution_model, "MeshNV") ||
+ 0 == std::strcmp(execution_model, "MeshEXT") ||
+ 0 == std::strcmp(execution_model, "TaskNV")) {
execution_modes << "OpExecutionMode %" << entry_point.name
<< " LocalSize 1 1 1\n";
}
@@ -3615,6 +3631,7 @@
EntryPoint entry_point;
entry_point.name = "main_d_r";
entry_point.execution_model = "MeshNV";
+ entry_point.execution_modes = "OpExecutionMode %main_d_r LocalSize 1 1 1";
entry_point.interfaces = "%gl_PrimitiveID %gl_Layer %gl_ViewportIndex";
generator.entry_points_.push_back(std::move(entry_point));
@@ -3653,6 +3670,7 @@
EntryPoint entry_point;
entry_point.name = "main_d_r";
entry_point.execution_model = "MeshNV";
+ entry_point.execution_modes = "OpExecutionMode %main_d_r LocalSize 1 1 1";
entry_point.interfaces = "%gl_PrimitiveID %gl_Layer %gl_ViewportIndex";
entry_point.body = "%ref_load = OpLoad %_arr_float_uint_81 %gl_PrimitiveID";
generator.entry_points_.push_back(std::move(entry_point));
diff --git a/test/val/val_mesh_shading_test.cpp b/test/val/val_mesh_shading_test.cpp
index 49d5ebc..1897015 100644
--- a/test/val/val_mesh_shading_test.cpp
+++ b/test/val/val_mesh_shading_test.cpp
@@ -234,6 +234,8 @@
OpMemoryModel Logical GLSL450
OpEntryPoint MeshEXT %mainMesh "mainMesh"
OpEntryPoint TaskEXT %mainTask "mainTask"
+ OpExecutionMode %mainMesh LocalSize 1 1 1
+ OpExecutionMode %mainTask LocalSize 1 1 1
OpExecutionMode %mainMesh OutputVertices 1
OpExecutionMode %mainMesh OutputPrimitivesEXT 1
OpExecutionMode %mainMesh OutputTrianglesEXT
diff --git a/test/val/val_modes_test.cpp b/test/val/val_modes_test.cpp
index 58e41ed..01e215c 100644
--- a/test/val/val_modes_test.cpp
+++ b/test/val/val_modes_test.cpp
@@ -73,6 +73,53 @@
"or an object decorated with WorkgroupSize must be specified."));
}
+TEST_F(ValidateMode, MeshNoModeVulkan) {
+ const std::string spirv = R"(
+OpCapability Shader
+OpCapability MeshShadingEXT
+OpExtension "SPV_EXT_mesh_shader"
+OpMemoryModel Logical GLSL450
+OpEntryPoint MeshEXT %main "main"
+OpExecutionMode %main OutputVertices 81
+OpExecutionMode %main OutputPrimitivesEXT 16
+OpExecutionMode %main OutputPoints
+)" + kVoidFunction;
+
+ spv_target_env env = SPV_ENV_VULKAN_1_3;
+ CompileSuccessfully(spirv, env);
+ EXPECT_THAT(SPV_ERROR_INVALID_DATA, ValidateInstructions(env));
+ EXPECT_THAT(getDiagnosticString(),
+ AnyVUID("VUID-StandaloneSpirv-None-10685"));
+ EXPECT_THAT(
+ getDiagnosticString(),
+ HasSubstr(
+ "In the Vulkan environment, MeshEXT execution model entry "
+ "points require either the LocalSize or LocalSizeId execution mode "
+ "or an object decorated with WorkgroupSize must be specified."));
+}
+
+TEST_F(ValidateMode, TaskNoModeVulkan) {
+ const std::string spirv = R"(
+OpCapability Shader
+OpCapability MeshShadingEXT
+OpExtension "SPV_EXT_mesh_shader"
+OpMemoryModel Logical GLSL450
+OpEntryPoint TaskEXT %main "main"
+)" + kVoidFunction;
+
+ spv_target_env env = SPV_ENV_VULKAN_1_3;
+ CompileSuccessfully(spirv, env);
+ EXPECT_THAT(SPV_ERROR_INVALID_DATA, ValidateInstructions(env));
+ EXPECT_THAT(getDiagnosticString(),
+ AnyVUID("VUID-StandaloneSpirv-None-10685"));
+ EXPECT_THAT(
+ getDiagnosticString(),
+ HasSubstr(
+ "In the Vulkan environment, TaskEXT execution model entry "
+ "points require either the LocalSize or LocalSizeId execution mode "
+ "or an object decorated with WorkgroupSize must be specified."));
+}
+
TEST_F(ValidateMode, GLComputeNoModeVulkanWorkgroupSize) {
const std::string spirv = R"(
OpCapability Shader
@@ -2926,7 +2973,6 @@
OpExtension "SPV_QCOM_tile_shading"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
-OpExecutionMode %main NonCoherentTileAttachmentReadQCOM
)" + kVoidFunction;
spv_target_env env = SPV_ENV_VULKAN_1_4;
