spirv-val: Add Mesh/Task to check for LocalSize (#6459)

spencer-lunarg · web-flow · commit 66798c86ac7c · 2025-12-11T21:40:35.000-05:00
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
@@ -27,6 +27,48 @@ namespace spvtools {
 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 @@ spv_result_t ValidateEntryPoint(ValidationState_t& _, const Instruction* inst) {
   }
 
   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) {
-            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.";
-          }
+    // 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 (_.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.";
-          }
+        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.";
         }
-        break;
-      default:
+      } 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 the TileShadingRateQCOM execution mode is used, the "
+                    "TileShadingQCOM capability must be enabled.";
         }
-        if (execution_model != spv::ExecutionModel::Fragment) {
-          if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) {
+      }
+    } 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)
-                   << "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.";
-            }
+                   << "If the NonCoherentTileAttachmentReadQCOM 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:
         break;
     }
   }
diff --git a/test/val/val_builtins_test.cpp b/test/val/val_builtins_test.cpp
@@ -141,7 +141,11 @@ CodeGenerator GetInMainCodeGenerator(const char* const built_in,
     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 @@ CodeGenerator GetInFunctionCodeGenerator(const char* const built_in,
     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 @@ CodeGenerator GetVariableCodeGenerator(const char* const built_in,
     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 @@ CodeGenerator GetArrayedVariableCodeGenerator(const char* const built_in,
     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 @@ OpDecorate %gl_ViewportIndex PerPrimitiveNV
   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 @@ OpDecorate %gl_ViewportIndex PerPrimitiveNV
   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
@@ -234,6 +234,8 @@ TEST_F(ValidateMeshShading, VulkanBasicMeshAndTaskSuccess) {
                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
@@ -73,6 +73,53 @@ OpEntryPoint GLCompute %main "main"
           "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 @@ OpCapability TileShadingQCOM
 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;
