Index: mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp
===================================================================
--- mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp
+++ mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp
@@ -131,6 +131,22 @@
               deducedCapabilities)))
         return WalkResult::interrupt();
 
+    // Builtin attach to the op may have capability requirements.
+    auto builtInName = llvm::convertToSnakeFromCamelCase(
+        stringifyDecoration(spirv::Decoration::BuiltIn));
+    if (auto builtin = op->getAttrOfType<StringAttr>(builtInName)) {
+      if (auto varBuiltIn = spirv::symbolizeBuiltIn(builtin.getValue())) {
+        if (auto cap = spirv::getCapabilities(varBuiltIn.getValue())) {
+          if (failed(checkAndUpdateCapabilityRequirements(
+                  op, targetEnv,
+                  SmallVector<ArrayRef<spirv::Capability>, 1>(1,
+                                                              cap.getValue()),
+                  deducedCapabilities)))
+            return WalkResult::interrupt();
+        }
+      }
+    }
+
     SmallVector<Type, 4> valueTypes;
     valueTypes.append(op->operand_type_begin(), op->operand_type_end());
     valueTypes.append(op->result_type_begin(), op->result_type_end());
@@ -157,6 +173,21 @@
         return WalkResult::interrupt();
     }
 
+    // Capabilities may require a minimum version as well.
+    for (auto capability : deducedCapabilities) {
+      if (auto capMinVersion = spirv::getMinVersion(capability)) {
+        deducedVersion = std::max(deducedVersion, capMinVersion.getValue());
+        if (deducedVersion > allowedVersion) {
+          return op->emitError("'") << op->getName() << "' requires "
+                                    << spirv::stringifyCapability(capability)
+                                    << " which requires min version "
+                                    << spirv::stringifyVersion(deducedVersion)
+                                    << " but target environment allows up to "
+                                    << spirv::stringifyVersion(allowedVersion);
+        }
+      }
+    }
+
     return WalkResult::advance();
   });
 
Index: mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir
===================================================================
--- mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir
+++ mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir
@@ -137,6 +137,18 @@
   }
 }
 
+// Test Builtin capabilities as well as capabilities min version.
+// NumSubgroups requires GroupNonUniform which requires v1.3.
+// CHECK: requires #spv.vce<v1.3, [GroupNonUniform, Shader], []>
+spv.module Logical GLSL450 attributes {
+  spv.target_env = #spv.target_env<
+    #spv.vce<v1.5, [Shader, UniformAndStorageBuffer8BitAccess, StorageBuffer16BitAccess, StorageUniform16, Int16, ImageBuffer, StorageImageExtendedFormats, GroupNonUniform], []>,
+    {max_compute_workgroup_invocations = 128 : i32,
+     max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+} {
+   spv.globalVariable @__builtin_var_NumSubgroups__ built_in("NumSubgroups") : !spv.ptr<i32, Input>
+}
+
 //===----------------------------------------------------------------------===//
 // Extension
 //===----------------------------------------------------------------------===//
@@ -162,7 +174,7 @@
 // Vulkan memory model requires SPV_KHR_vulkan_memory_model, which is enabled
 // implicitly by v1.5.
 
-// CHECK: requires #spv.vce<v1.0, [VulkanMemoryModel], [SPV_KHR_vulkan_memory_model]>
+// CHECK: requires #spv.vce<v1.5, [VulkanMemoryModel], [SPV_KHR_vulkan_memory_model]>
 spv.module Logical Vulkan attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.5, [Shader, VulkanMemoryModel], []>,