Overridable constants vulkan implementation and tests

Add vulkan backend implementations and tests.
Disabled some tests that fail due to tint missing
features and issues.

Bug: dawn:1041, tint:1155
Change-Id: Iac161317450cff59627e08b1228bffde4cef71da
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/64981
Commit-Queue: Shrek Shao <shrekshao@google.com>
Reviewed-by: Corentin Wallez <cwallez@chromium.org>

dawn.json

@@ -438,7 +438,7 @@
         "category": "structure",
         "extensible": "in",
         "members": [
-            {"name": "key", "type": "char", "annotation": "const*"},
+            {"name": "key", "type": "char", "annotation": "const*", "length": "strlen"},
             {"name": "value", "type": "double"}
         ]
     },
@@ -1348,8 +1348,8 @@
         "members": [
             {"name": "module", "type": "shader module"},
             {"name": "entry point", "type": "char", "annotation": "const*", "length": "strlen"},
-            {"name": "constant count", "type": "uint32_t", "tags": ["upstream"]},
+            {"name": "constant count", "type": "uint32_t", "default": 0},
-            {"name": "constants", "type": "constant entry", "annotation": "const*", "length": "constant count", "tags": ["upstream"]}
+            {"name": "constants", "type": "constant entry", "annotation": "const*", "length": "constant count"}
         ]
     },
     "primitive topology": {
@@ -1817,8 +1817,8 @@
         "members": [
             {"name": "module", "type": "shader module"},
             {"name": "entry point", "type": "char", "annotation": "const*", "length": "strlen"},
-            {"name": "constant count", "type": "uint32_t", "tags": ["upstream"]},
+            {"name": "constant count", "type": "uint32_t", "default": 0},
-            {"name": "constants", "type": "constant entry", "annotation": "const*", "length": "constant count", "tags": ["upstream"]},
+            {"name": "constants", "type": "constant entry", "annotation": "const*", "length": "constant count"},
             {"name": "buffer count", "type": "uint32_t", "default": 0},
             {"name": "buffers", "type": "vertex buffer layout", "annotation": "const*", "length": "buffer count"}
         ]
@@ -1876,8 +1876,8 @@
         "members": [
             {"name": "module", "type": "shader module"},
             {"name": "entry point", "type": "char", "annotation": "const*", "length": "strlen"},
-            {"name": "constant count", "type": "uint32_t", "tags": ["upstream"]},
+            {"name": "constant count", "type": "uint32_t", "default": 0},
-            {"name": "constants", "type": "constant entry", "annotation": "const*", "length": "constant count", "tags": ["upstream"]},
+            {"name": "constants", "type": "constant entry", "annotation": "const*", "length": "constant count"},
             {"name": "target count", "type": "uint32_t"},
             {"name": "targets", "type": "color target state", "annotation": "const*", "length": "target count"}
         ]

src/dawn_native/ComputePipeline.cpp

@@ -30,9 +30,10 @@ namespace dawn_native {
             DAWN_TRY(device->ValidateObject(descriptor->layout));
         }
 
-        return ValidateProgrammableStage(device, descriptor->compute.module,
+        return ValidateProgrammableStage(
-                                         descriptor->compute.entryPoint, descriptor->layout,
+            device, descriptor->compute.module, descriptor->compute.entryPoint,
-                                         SingleShaderStage::Compute);
+            descriptor->compute.constantCount, descriptor->compute.constants, descriptor->layout,
+            SingleShaderStage::Compute);
     }
 
     // ComputePipelineBase
@@ -43,7 +44,8 @@ namespace dawn_native {
                        descriptor->layout,
                        descriptor->label,
                        {{SingleShaderStage::Compute, descriptor->compute.module,
-                         descriptor->compute.entryPoint}}) {
+                         descriptor->compute.entryPoint, descriptor->compute.constantCount,
+                         descriptor->compute.constants}}) {
     }
 
     ComputePipelineBase::ComputePipelineBase(DeviceBase* device, ObjectBase::ErrorTag tag)

src/dawn_native/Device.cpp

@@ -135,9 +135,13 @@ namespace dawn_native {
 
             if (outDescriptor->layout == nullptr) {
                 DAWN_TRY_ASSIGN(layoutRef, PipelineLayoutBase::CreateDefault(
-                                               device, {{SingleShaderStage::Compute,
+                                               device, {{
-                                                         outDescriptor->compute.module,
+                                                           SingleShaderStage::Compute,
-                                                         outDescriptor->compute.entryPoint}}));
+                                                           outDescriptor->compute.module,
+                                                           outDescriptor->compute.entryPoint,
+                                                           outDescriptor->compute.constantCount,
+                                                           outDescriptor->compute.constants,
+                                                       }}));
                 outDescriptor->layout = layoutRef.Get();
             }
 

src/dawn_native/Pipeline.cpp

@@ -26,6 +26,8 @@ namespace dawn_native {
     MaybeError ValidateProgrammableStage(DeviceBase* device,
                                          const ShaderModuleBase* module,
                                          const std::string& entryPoint,
+                                         uint32_t constantCount,
+                                         const ConstantEntry* constants,
                                          const PipelineLayoutBase* layout,
                                          SingleShaderStage stage) {
         DAWN_TRY(device->ValidateObject(module));
@@ -44,6 +46,14 @@ namespace dawn_native {
             DAWN_TRY(ValidateCompatibilityWithPipelineLayout(device, metadata, layout));
         }
 
+        // Validate if overridable constants exist in shader module
+        // pipelineBase is not yet constructed at this moment so iterate constants from descriptor
+        for (uint32_t i = 0; i < constantCount; i++) {
+            DAWN_INVALID_IF(metadata.overridableConstants.count(constants[i].key) == 0,
+                            "Pipeline overridable constant \"%s\" not found in shader module %s.",
+                            constants[i].key, module);
+        }
+
         return {};
     }
 
@@ -68,7 +78,12 @@ namespace dawn_native {
             // Record them internally.
             bool isFirstStage = mStageMask == wgpu::ShaderStage::None;
             mStageMask |= StageBit(shaderStage);
-            mStages[shaderStage] = {module, entryPointName, &metadata};
+            mStages[shaderStage] = {module, entryPointName, &metadata,
+                                    std::vector<PipelineConstantEntry>()};
+            auto& constants = mStages[shaderStage].constants;
+            for (uint32_t i = 0; i < stage.constantCount; i++) {
+                constants.emplace_back(stage.constants[i].key, stage.constants[i].value);
+            }
 
             // Compute the max() of all minBufferSizes across all stages.
             RequiredBufferSizes stageMinBufferSizes =

src/dawn_native/Pipeline.h

@@ -32,15 +32,20 @@ namespace dawn_native {
     MaybeError ValidateProgrammableStage(DeviceBase* device,
                                          const ShaderModuleBase* module,
                                          const std::string& entryPoint,
+                                         uint32_t constantCount,
+                                         const ConstantEntry* constants,
                                          const PipelineLayoutBase* layout,
                                          SingleShaderStage stage);
 
+    using PipelineConstantEntry = std::pair<std::string, double>;
     struct ProgrammableStage {
         Ref<ShaderModuleBase> module;
         std::string entryPoint;
 
         // The metadata lives as long as module, that's ref-ed in the same structure.
         const EntryPointMetadata* metadata = nullptr;
+
+        std::vector<PipelineConstantEntry> constants;
     };
 
     class PipelineBase : public ApiObjectBase, public CachedObject {

src/dawn_native/PipelineLayout.h

@@ -44,6 +44,8 @@ namespace dawn_native {
         SingleShaderStage shaderStage;
         ShaderModuleBase* module;
         std::string entryPoint;
+        uint32_t constantCount = 0u;
+        ConstantEntry const* constants = nullptr;
     };
 
     class PipelineLayoutBase : public ApiObjectBase, public CachedObject {

src/dawn_native/RenderPipeline.cpp

@@ -114,6 +114,7 @@ namespace dawn_native {
             }
 
             DAWN_TRY(ValidateProgrammableStage(device, descriptor->module, descriptor->entryPoint,
+                                               descriptor->constantCount, descriptor->constants,
                                                layout, SingleShaderStage::Vertex));
             const EntryPointMetadata& vertexMetadata =
                 descriptor->module->GetEntryPoint(descriptor->entryPoint);
@@ -302,6 +303,7 @@ namespace dawn_native {
             }
 
             DAWN_TRY(ValidateProgrammableStage(device, descriptor->module, descriptor->entryPoint,
+                                               descriptor->constantCount, descriptor->constants,
                                                layout, SingleShaderStage::Fragment));
 
             if (descriptor->targetCount > kMaxColorAttachments) {
@@ -422,17 +424,20 @@ namespace dawn_native {
         DeviceBase* device,
         const RenderPipelineDescriptor* descriptor) {
         std::vector<StageAndDescriptor> stages;
-        stages.push_back(
+        stages.push_back({SingleShaderStage::Vertex, descriptor->vertex.module,
-            {SingleShaderStage::Vertex, descriptor->vertex.module, descriptor->vertex.entryPoint});
+                          descriptor->vertex.entryPoint, descriptor->vertex.constantCount,
+                          descriptor->vertex.constants});
         if (descriptor->fragment != nullptr) {
             stages.push_back({SingleShaderStage::Fragment, descriptor->fragment->module,
-                              descriptor->fragment->entryPoint});
+                              descriptor->fragment->entryPoint, descriptor->fragment->constantCount,
+                              descriptor->fragment->constants});
         } else if (device->IsToggleEnabled(Toggle::UseDummyFragmentInVertexOnlyPipeline)) {
             InternalPipelineStore* store = device->GetInternalPipelineStore();
             // The dummy fragment shader module should already be initialized
             DAWN_ASSERT(store->dummyFragmentShader != nullptr);
             ShaderModuleBase* dummyFragmentShader = store->dummyFragmentShader.Get();
-            stages.push_back({SingleShaderStage::Fragment, dummyFragmentShader, "fs_empty_main"});
+            stages.push_back(
+                {SingleShaderStage::Fragment, dummyFragmentShader, "fs_empty_main", 0, nullptr});
         }
         return stages;
     }

src/dawn_native/ShaderModule.cpp

@@ -33,6 +33,22 @@
 
 namespace dawn_native {
 
+    EntryPointMetadata::OverridableConstant::Type GetDawnOverridableConstantType(
+        tint::inspector::OverridableConstant::Type type) {
+        switch (type) {
+            case tint::inspector::OverridableConstant::Type::kBool:
+                return EntryPointMetadata::OverridableConstant::Type::Boolean;
+            case tint::inspector::OverridableConstant::Type::kFloat32:
+                return EntryPointMetadata::OverridableConstant::Type::Float32;
+            case tint::inspector::OverridableConstant::Type::kInt32:
+                return EntryPointMetadata::OverridableConstant::Type::Int32;
+            case tint::inspector::OverridableConstant::Type::kUint32:
+                return EntryPointMetadata::OverridableConstant::Type::Uint32;
+            default:
+                UNREACHABLE();
+        }
+    }
+
     namespace {
 
         std::string GetShaderDeclarationString(BindGroupIndex group, BindingNumber binding) {
@@ -629,13 +645,22 @@ namespace dawn_native {
             for (auto& entryPoint : entryPoints) {
                 ASSERT(result.count(entryPoint.name) == 0);
 
-                if (!entryPoint.overridable_constants.empty()) {
-                    return DAWN_VALIDATION_ERROR(
-                        "Pipeline overridable constants are not implemented yet");
-                }
-
                 auto metadata = std::make_unique<EntryPointMetadata>();
 
+                if (!entryPoint.overridable_constants.empty()) {
+                    const auto& name2Id = inspector.GetConstantNameToIdMap();
+
+                    for (auto& c : entryPoint.overridable_constants) {
+                        EntryPointMetadata::OverridableConstant constant = {
+                            name2Id.at(c.name), GetDawnOverridableConstantType(c.type)};
+                        metadata->overridableConstants[c.name] = constant;
+                        // TODO(tint:1155) tint needs ways to differentiate whether a pipeline
+                        // constant id is specified explicitly. Now we just store numeric id and
+                        // variable name in the index at the same time
+                        metadata->overridableConstants[std::to_string(constant.id)] = constant;
+                    }
+                }
+
                 DAWN_TRY_ASSIGN(metadata->stage, TintPipelineStageToShaderStage(entryPoint.stage));
 
                 if (metadata->stage == SingleShaderStage::Compute) {

src/dawn_native/ShaderModule.h

@@ -194,6 +194,16 @@ namespace dawn_native {
 
         // The shader stage for this binding.
         SingleShaderStage stage;
+
+        struct OverridableConstant {
+            uint32_t id;
+            // Match tint::inspector::OverridableConstant::Type
+            // Bool is defined as a macro on linux X11 and cannot compile
+            enum class Type { Boolean, Float32, Uint32, Int32 } type;
+        };
+
+        // Store overridableConstants from tint program
+        std::unordered_map<std::string, OverridableConstant> overridableConstants;
     };
 
     class ShaderModuleBase : public ApiObjectBase, public CachedObject {

src/dawn_native/vulkan/ComputePipelineVk.cpp

@@ -52,8 +52,15 @@ namespace dawn_native { namespace vulkan {
                         ToBackend(computeStage.module.Get())
                             ->GetTransformedModuleHandle(computeStage.entryPoint.c_str(),
                                                          ToBackend(GetLayout())));
+
         createInfo.stage.pName = computeStage.entryPoint.c_str();
-        createInfo.stage.pSpecializationInfo = nullptr;
+
+        std::vector<SpecializationDataEntry> specializationDataEntries;
+        std::vector<VkSpecializationMapEntry> specializationMapEntries;
+        VkSpecializationInfo specializationInfo{};
+        createInfo.stage.pSpecializationInfo =
+            GetVkSpecializationInfo(computeStage, &specializationInfo, &specializationDataEntries,
+                                    &specializationMapEntries);
 
         Device* device = ToBackend(GetDevice());
 

src/dawn_native/vulkan/RenderPipelineVk.cpp

@@ -341,21 +341,25 @@ namespace dawn_native { namespace vulkan {
 
         // There are at most 2 shader stages in render pipeline, i.e. vertex and fragment
         std::array<VkPipelineShaderStageCreateInfo, 2> shaderStages;
+        std::array<std::vector<SpecializationDataEntry>, 2> specializationDataEntriesPerStages;
+        std::array<std::vector<VkSpecializationMapEntry>, 2> specializationMapEntriesPerStages;
+        std::array<VkSpecializationInfo, 2> specializationInfoPerStages;
         uint32_t stageCount = 0;
 
         for (auto stage : IterateStages(this->GetStageMask())) {
             VkPipelineShaderStageCreateInfo shaderStage;
 
+            const ProgrammableStage& programmableStage = GetStage(stage);
             DAWN_TRY_ASSIGN(shaderStage.module,
-                            ToBackend(GetStage(stage).module)
+                            ToBackend(programmableStage.module)
-                                ->GetTransformedModuleHandle(GetStage(stage).entryPoint.c_str(),
+                                ->GetTransformedModuleHandle(programmableStage.entryPoint.c_str(),
                                                              ToBackend(GetLayout())));
 
             shaderStage.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO;
             shaderStage.pNext = nullptr;
             shaderStage.flags = 0;
             shaderStage.pSpecializationInfo = nullptr;
-            shaderStage.pName = GetStage(stage).entryPoint.c_str();
+            shaderStage.pName = programmableStage.entryPoint.c_str();
 
             switch (stage) {
                 case dawn_native::SingleShaderStage::Vertex: {
@@ -373,6 +377,11 @@ namespace dawn_native { namespace vulkan {
                 }
             }
 
+            shaderStage.pSpecializationInfo =
+                GetVkSpecializationInfo(programmableStage, &specializationInfoPerStages[stageCount],
+                                        &specializationDataEntriesPerStages[stageCount],
+                                        &specializationMapEntriesPerStages[stageCount]);
+
             DAWN_ASSERT(stageCount < 2);
             shaderStages[stageCount] = shaderStage;
             stageCount++;

src/dawn_native/vulkan/UtilsVulkan.cpp

@@ -17,6 +17,8 @@
 #include "common/Assert.h"
 #include "dawn_native/EnumMaskIterator.h"
 #include "dawn_native/Format.h"
+#include "dawn_native/Pipeline.h"
+#include "dawn_native/ShaderModule.h"
 #include "dawn_native/vulkan/DeviceVk.h"
 #include "dawn_native/vulkan/Forward.h"
 #include "dawn_native/vulkan/TextureVk.h"
@@ -195,4 +197,63 @@ namespace dawn_native { namespace vulkan {
             device->fn.SetDebugUtilsObjectNameEXT(device->GetVkDevice(), &objectNameInfo);
         }
     }
+
+    VkSpecializationInfo* GetVkSpecializationInfo(
+        const ProgrammableStage& programmableStage,
+        VkSpecializationInfo* specializationInfo,
+        std::vector<SpecializationDataEntry>* specializationDataEntries,
+        std::vector<VkSpecializationMapEntry>* specializationMapEntries) {
+        ASSERT(specializationInfo);
+        ASSERT(specializationDataEntries);
+        ASSERT(specializationMapEntries);
+
+        if (programmableStage.constants.size() == 0) {
+            return nullptr;
+        }
+
+        const EntryPointMetadata& entryPointMetaData =
+            programmableStage.module->GetEntryPoint(programmableStage.entryPoint);
+
+        for (const auto& pipelineConstant : programmableStage.constants) {
+            const std::string& name = pipelineConstant.first;
+            double value = pipelineConstant.second;
+
+            // This is already validated so `name` must exist
+            const auto& moduleConstant = entryPointMetaData.overridableConstants.at(name);
+
+            specializationMapEntries->push_back(
+                VkSpecializationMapEntry{moduleConstant.id,
+                                         static_cast<uint32_t>(specializationDataEntries->size() *
+                                                               sizeof(SpecializationDataEntry)),
+                                         sizeof(SpecializationDataEntry)});
+
+            SpecializationDataEntry entry;
+            switch (moduleConstant.type) {
+                case EntryPointMetadata::OverridableConstant::Type::Boolean:
+                    entry.b = static_cast<bool>(value);
+                    break;
+                case EntryPointMetadata::OverridableConstant::Type::Float32:
+                    entry.f32 = static_cast<float>(value);
+                    break;
+                case EntryPointMetadata::OverridableConstant::Type::Int32:
+                    entry.i32 = static_cast<int32_t>(value);
+                    break;
+                case EntryPointMetadata::OverridableConstant::Type::Uint32:
+                    entry.u32 = static_cast<uint32_t>(value);
+                    break;
+                default:
+                    UNREACHABLE();
+            }
+            specializationDataEntries->push_back(entry);
+        }
+
+        specializationInfo->mapEntryCount = static_cast<uint32_t>(specializationMapEntries->size());
+        specializationInfo->pMapEntries = specializationMapEntries->data();
+        specializationInfo->dataSize =
+            specializationDataEntries->size() * sizeof(SpecializationDataEntry);
+        specializationInfo->pData = specializationDataEntries->data();
+
+        return specializationInfo;
+    }
+
 }}  // namespace dawn_native::vulkan

src/dawn_native/vulkan/UtilsVulkan.h

@@ -19,6 +19,10 @@
 #include "dawn_native/Commands.h"
 #include "dawn_native/dawn_platform.h"
 
+namespace dawn_native {
+    struct ProgrammableStage;
+}  // namespace dawn_native
+
 namespace dawn_native { namespace vulkan {
 
     class Device;
@@ -107,6 +111,24 @@ namespace dawn_native { namespace vulkan {
                       const char* prefix,
                       std::string label = "");
 
+    // Helpers for creating VkSpecializationInfo
+    // The WebGPU overridable constants only support these scalar types
+    union SpecializationDataEntry {
+        bool b;
+        float f32;
+        int32_t i32;
+        uint32_t u32;
+    };
+
+    // Returns nullptr or &specializationInfo
+    // specializationInfo, specializationDataEntries, specializationMapEntries needs to
+    // be alive at least until VkSpecializationInfo is passed into Vulkan Create*Pipelines
+    VkSpecializationInfo* GetVkSpecializationInfo(
+        const ProgrammableStage& programmableStage,
+        VkSpecializationInfo* specializationInfo,
+        std::vector<SpecializationDataEntry>* specializationDataEntries,
+        std::vector<VkSpecializationMapEntry>* specializationMapEntries);
+
 }}  // namespace dawn_native::vulkan
 
 #endif  // DAWNNATIVE_VULKAN_UTILSVULKAN_H_

src/tests/BUILD.gn

@@ -207,6 +207,7 @@ test("dawn_unittests") {
     "unittests/validation/LabelTests.cpp",
     "unittests/validation/MinimumBufferSizeValidationTests.cpp",
     "unittests/validation/MultipleDeviceTests.cpp",
+    "unittests/validation/OverridableConstantsValidationTests.cpp",
     "unittests/validation/QueryValidationTests.cpp",
     "unittests/validation/QueueOnSubmittedWorkDoneValidationTests.cpp",
     "unittests/validation/QueueSubmitValidationTests.cpp",

src/tests/end2end/ShaderTests.cpp

@@ -17,18 +17,37 @@
 #include "utils/ComboRenderPipelineDescriptor.h"
 #include "utils/WGPUHelpers.h"
 
+#include <numeric>
 #include <vector>
 
-class ShaderTests : public DawnTest {};
+class ShaderTests : public DawnTest {
+  public:
+    wgpu::Buffer CreateBuffer(const uint32_t count) {
+        std::vector<uint32_t> data(count, 0);
+        uint64_t bufferSize = static_cast<uint64_t>(data.size() * sizeof(uint32_t));
+        return utils::CreateBufferFromData(device, data.data(), bufferSize,
+                                           wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc);
+    }
+    wgpu::ComputePipeline CreateComputePipeline(
+        const std::string& shader,
+        const char* entryPoint,
+        const std::vector<wgpu::ConstantEntry>* constants = nullptr) {
+        wgpu::ComputePipelineDescriptor csDesc;
+        csDesc.compute.module = utils::CreateShaderModule(device, shader.c_str());
+        csDesc.compute.entryPoint = entryPoint;
+        if (constants) {
+            csDesc.compute.constants = constants->data();
+            csDesc.compute.constantCount = constants->size();
+        }
+        return device.CreateComputePipeline(&csDesc);
+    }
+};
 
 // Test that log2 is being properly calculated, base on crbug.com/1046622
 TEST_P(ShaderTests, ComputeLog2) {
     uint32_t const kSteps = 19;
-    std::vector<uint32_t> data(kSteps, 0);
     std::vector<uint32_t> expected{0, 1, 1, 2, 2, 3, 3, 4, 4, 5, 5, 6, 6, 7, 7, 8, 8, 9, 32};
-    uint64_t bufferSize = static_cast<uint64_t>(data.size() * sizeof(uint32_t));
+    wgpu::Buffer buffer = CreateBuffer(kSteps);
-    wgpu::Buffer buffer = utils::CreateBufferFromData(
-        device, data.data(), bufferSize, wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc);
 
     std::string shader = R"(
 [[block]] struct Buf {
@@ -61,10 +80,7 @@ TEST_P(ShaderTests, ComputeLog2) {
     buf.data[18] = u32(log2(4294967295.0 * factor));
 })";
 
-    wgpu::ComputePipelineDescriptor csDesc;
+    wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main");
-    csDesc.compute.module = utils::CreateShaderModule(device, shader.c_str());
-    csDesc.compute.entryPoint = "main";
-    wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc);
 
     wgpu::BindGroup bindGroup =
         utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}});
@@ -305,23 +321,6 @@ fn fragmentMain(input : VertexOut) -> [[location(0)]] vec4<f32> {
     wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&rpDesc);
 }
 
-// Feature currently not implemented in Tint, so should fail validation.
-TEST_P(ShaderTests, PipelineOverridableUsed) {
-    DAWN_TEST_UNSUPPORTED_IF(HasToggleEnabled("skip_validation"));
-    DAWN_TEST_UNSUPPORTED_IF(!HasToggleEnabled("use_tint_generator"));
-
-    std::string shader = R"(
-[[override]] let foo : f32;
-
-[[stage(compute), workgroup_size(1)]]
-fn ep_func() {
-  var local_foo : f32;
-  local_foo = foo;
-  return;
-})";
-    ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, shader.c_str()));
-}
-
 // This is a regression test for an issue caused by the FirstIndexOffset transfrom being done before
 // the BindingRemapper, causing an intermediate AST to be invalid (and fail the overall
 // compilation).
@@ -393,6 +392,299 @@ fn main([[location(0)]] pos : vec4<f32>) -> [[builtin(position)]] vec4<f32> {
     device.CreateRenderPipeline(&descriptor);
 }
 
+// Test overridable constants without numeric identifiers
+// TODO(tint:1155): Implicit numeric ID is undetermined in tint
+TEST_P(ShaderTests, DISABLED_OverridableConstants) {
+    // TODO(dawn:1041): Only Vulkan backend is implemented
+    DAWN_TEST_UNSUPPORTED_IF(!IsVulkan());
+
+    uint32_t const kCount = 15;
+    std::vector<uint32_t> expected(kCount);
+    std::iota(expected.begin(), expected.end(), 0);
+    // Test last entry with unspecified default value
+    expected[kCount - 1] = 0u;
+    expected[kCount - 2] = 0u;
+    expected[kCount - 3] = 0u;
+    expected[kCount - 4] = 0u;
+    wgpu::Buffer buffer = CreateBuffer(kCount);
+
+    std::string shader = R"(
+[[override]] let c0: bool;              // type: bool
+[[override]] let c1: bool = false;      // default override
+[[override]] let c2: f32;               // type: float32
+[[override]] let c3: f32 = 0.0;         // default override
+[[override]] let c4: f32 = 4.0;         // default
+[[override]] let c5: i32;               // type: int32
+[[override]] let c6: i32 = 0;           // default override
+[[override]] let c7: i32 = 7;           // default
+[[override]] let c8: u32;               // type: uint32
+[[override]] let c9: u32 = 0u;          // default override
+[[override]] let c10: u32 = 10u;        // default
+[[override]] let c11: bool;             // default unspecified
+[[override]] let c12: f32;              // default unspecified
+[[override]] let c13: i32;              // default unspecified
+[[override]] let c14: u32;              // default unspecified
+
+[[block]] struct Buf {
+    data : array<u32, 15>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> buf : Buf;
+
+[[stage(compute), workgroup_size(1)]] fn main() {
+    buf.data[0] = u32(c0);
+    buf.data[1] = u32(c1);
+    buf.data[2] = u32(c2);
+    buf.data[3] = u32(c3);
+    buf.data[4] = u32(c4);
+    buf.data[5] = u32(c5);
+    buf.data[6] = u32(c6);
+    buf.data[7] = u32(c7);
+    buf.data[8] = u32(c8);
+    buf.data[9] = u32(c9);
+    buf.data[10] = u32(c10);
+    buf.data[11] = u32(c11);
+    buf.data[12] = u32(c12);
+    buf.data[13] = u32(c13);
+    buf.data[14] = u32(c14);
+})";
+
+    std::vector<wgpu::ConstantEntry> constants;
+    constants.push_back({nullptr, "c0", 0});
+    constants.push_back({nullptr, "c1", 1});
+    constants.push_back({nullptr, "c2", 2});
+    constants.push_back({nullptr, "c3", 3});
+    // c4 is not assigned, testing default value
+    constants.push_back({nullptr, "c5", 5});
+    constants.push_back({nullptr, "c6", 6});
+    // c7 is not assigned, testing default value
+    constants.push_back({nullptr, "c8", 8});
+    constants.push_back({nullptr, "c9", 9});
+    // c10 is not assigned, testing default value
+    // c11 is not assigned, testing unspecified default value
+    // c12 is not assigned, testing unspecified default value
+    // c13 is not assigned, testing unspecified default value
+    // c14 is not assigned, testing unspecified default value
+
+    wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main", &constants);
+
+    wgpu::BindGroup bindGroup =
+        utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}});
+
+    wgpu::CommandBuffer commands;
+    {
+        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+        pass.SetPipeline(pipeline);
+        pass.SetBindGroup(0, bindGroup);
+        pass.Dispatch(1);
+        pass.EndPass();
+
+        commands = encoder.Finish();
+    }
+
+    queue.Submit(1, &commands);
+
+    EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), buffer, 0, kCount);
+}
+
+// Test overridable constants with numeric identifiers
+TEST_P(ShaderTests, OverridableConstantsNumericIdentifiers) {
+    // TODO(dawn:1041): Only Vulkan backend is implemented
+    DAWN_TEST_UNSUPPORTED_IF(!IsVulkan());
+
+    uint32_t const kCount = 4;
+    std::vector<uint32_t> expected{1u, 2u, 3u, 0u};
+    wgpu::Buffer buffer = CreateBuffer(kCount);
+
+    std::string shader = R"(
+[[override(1001)]] let c1: u32;            // some big numeric id
+[[override(1)]] let c2: u32 = 0u;          // id == 1 might collide with some generated constant id
+[[override(1003)]] let c3: u32 = 3u;       // default
+[[override(1004)]] let c4: u32;            // default unspecified
+
+[[block]] struct Buf {
+    data : array<u32, 4>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> buf : Buf;
+
+[[stage(compute), workgroup_size(1)]] fn main() {
+    buf.data[0] = c1;
+    buf.data[1] = c2;
+    buf.data[2] = c3;
+    buf.data[3] = c4;
+})";
+
+    std::vector<wgpu::ConstantEntry> constants;
+    constants.push_back({nullptr, "1001", 1});
+    constants.push_back({nullptr, "1", 2});
+    // c3 is not assigned, testing default value
+    // c4 is not assigned, testing unspecified default value
+
+    wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main", &constants);
+
+    wgpu::BindGroup bindGroup =
+        utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}});
+
+    wgpu::CommandBuffer commands;
+    {
+        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+        pass.SetPipeline(pipeline);
+        pass.SetBindGroup(0, bindGroup);
+        pass.Dispatch(1);
+        pass.EndPass();
+
+        commands = encoder.Finish();
+    }
+
+    queue.Submit(1, &commands);
+
+    EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), buffer, 0, kCount);
+}
+
+// Test overridable constants for different entry points
+TEST_P(ShaderTests, OverridableConstantsMultipleEntryPoints) {
+    // TODO(dawn:1041): Only Vulkan backend is implemented
+    DAWN_TEST_UNSUPPORTED_IF(!IsVulkan());
+
+    uint32_t const kCount = 1;
+    std::vector<uint32_t> expected1{1u};
+    std::vector<uint32_t> expected2{2u};
+
+    wgpu::Buffer buffer1 = CreateBuffer(kCount);
+    wgpu::Buffer buffer2 = CreateBuffer(kCount);
+
+    std::string shader = R"(
+[[override(1001)]] let c1: u32;
+[[override(1002)]] let c2: u32;
+
+[[block]] struct Buf {
+    data : array<u32, 1>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> buf : Buf;
+
+[[stage(compute), workgroup_size(1)]] fn main1() {
+    buf.data[0] = c1;
+}
+
+[[stage(compute), workgroup_size(1)]] fn main2() {
+    buf.data[0] = c2;
+}
+)";
+
+    std::vector<wgpu::ConstantEntry> constants1;
+    constants1.push_back({nullptr, "1001", 1});
+    std::vector<wgpu::ConstantEntry> constants2;
+    constants2.push_back({nullptr, "1002", 2});
+
+    wgpu::ShaderModule shaderModule = utils::CreateShaderModule(device, shader.c_str());
+
+    wgpu::ComputePipelineDescriptor csDesc1;
+    csDesc1.compute.module = shaderModule;
+    csDesc1.compute.entryPoint = "main1";
+    csDesc1.compute.constants = constants1.data();
+    csDesc1.compute.constantCount = constants1.size();
+    wgpu::ComputePipeline pipeline1 = device.CreateComputePipeline(&csDesc1);
+
+    wgpu::ComputePipelineDescriptor csDesc2;
+    csDesc2.compute.module = shaderModule;
+    csDesc2.compute.entryPoint = "main2";
+    csDesc2.compute.constants = constants2.data();
+    csDesc2.compute.constantCount = constants2.size();
+    wgpu::ComputePipeline pipeline2 = device.CreateComputePipeline(&csDesc2);
+
+    wgpu::BindGroup bindGroup1 =
+        utils::MakeBindGroup(device, pipeline1.GetBindGroupLayout(0), {{0, buffer1}});
+    wgpu::BindGroup bindGroup2 =
+        utils::MakeBindGroup(device, pipeline2.GetBindGroupLayout(0), {{0, buffer2}});
+
+    wgpu::CommandBuffer commands;
+    {
+        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+        pass.SetPipeline(pipeline1);
+        pass.SetBindGroup(0, bindGroup1);
+        pass.Dispatch(1);
+
+        pass.SetPipeline(pipeline2);
+        pass.SetBindGroup(0, bindGroup2);
+        pass.Dispatch(1);
+
+        pass.EndPass();
+
+        commands = encoder.Finish();
+    }
+
+    queue.Submit(1, &commands);
+
+    EXPECT_BUFFER_U32_RANGE_EQ(expected1.data(), buffer1, 0, kCount);
+    EXPECT_BUFFER_U32_RANGE_EQ(expected2.data(), buffer2, 0, kCount);
+}
+
+// Test overridable constants with render pipeline
+// Draw a triangle covering the render target, with vertex position and color values from
+// overridable constants
+TEST_P(ShaderTests, OverridableConstantsRenderPipeline) {
+    // TODO(dawn:1041): Only Vulkan backend is implemented
+    DAWN_TEST_UNSUPPORTED_IF(!IsVulkan());
+
+    wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"(
+[[override(1111)]] let xright: f32;
+[[override(2222)]] let ytop: f32;
+[[stage(vertex)]]
+fn main([[builtin(vertex_index)]] VertexIndex : u32)
+     -> [[builtin(position)]] vec4<f32> {
+  var pos = array<vec2<f32>, 3>(
+      vec2<f32>(-1.0, ytop),
+      vec2<f32>(-1.0, -ytop),
+      vec2<f32>(xright, 0.0));
+
+  return vec4<f32>(pos[VertexIndex], 0.0, 1.0);
+})");
+
+    wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"(
+[[override(1000)]] let intensity: f32 = 0.0;
+[[stage(fragment)]] fn main()
+    -> [[location(0)]] vec4<f32> {
+    return vec4<f32>(intensity, intensity, intensity, 1.0);
+})");
+
+    utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1);
+
+    utils::ComboRenderPipelineDescriptor descriptor;
+    descriptor.vertex.module = vsModule;
+    descriptor.cFragment.module = fsModule;
+    descriptor.primitive.topology = wgpu::PrimitiveTopology::TriangleList;
+    descriptor.cTargets[0].format = renderPass.colorFormat;
+
+    std::vector<wgpu::ConstantEntry> vertexConstants;
+    vertexConstants.push_back({nullptr, "1111", 3.0});  // x right
+    vertexConstants.push_back({nullptr, "2222", 3.0});  // y top
+    descriptor.vertex.constants = vertexConstants.data();
+    descriptor.vertex.constantCount = vertexConstants.size();
+    std::vector<wgpu::ConstantEntry> fragmentConstants;
+    fragmentConstants.push_back({nullptr, "1000", 1.0});  // color intensity
+    descriptor.cFragment.constants = fragmentConstants.data();
+    descriptor.cFragment.constantCount = fragmentConstants.size();
+
+    wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&descriptor);
+
+    wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+    wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo);
+    pass.SetPipeline(pipeline);
+    pass.Draw(3);
+    pass.EndPass();
+    wgpu::CommandBuffer commands = encoder.Finish();
+    queue.Submit(1, &commands);
+
+    EXPECT_PIXEL_RGBA8_EQ(RGBA8(255, 255, 255, 255), renderPass.color, 0, 0);
+}
+
+// TODO(tint:1155): Test overridable constants used for workgroup size
+
 DAWN_INSTANTIATE_TEST(ShaderTests,
                       D3D12Backend(),
                       MetalBackend(),

src/tests/unittests/validation/OverridableConstantsValidationTests.cpp

@@ -0,0 +1,147 @@
+// Copyright 2021 The Dawn Authors
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "common/Constants.h"
+#include "tests/unittests/validation/ValidationTest.h"
+#include "utils/WGPUHelpers.h"
+
+class ComputePipelineOverridableConstantsValidationTest : public ValidationTest {
+  protected:
+    void SetUp() override {
+        ValidationTest::SetUp();
+
+        computeModule = utils::CreateShaderModule(device, R"(
+[[override]] let c0: bool;              // type: bool
+[[override]] let c1: bool = false;      // default override
+[[override]] let c2: f32;               // type: float32
+[[override]] let c3: f32 = 0.0;         // default override
+[[override]] let c4: f32 = 4.0;         // default
+[[override]] let c5: i32;               // type: int32
+[[override]] let c6: i32 = 0;           // default override
+[[override]] let c7: i32 = 7;           // default
+[[override]] let c8: u32;               // type: uint32
+[[override]] let c9: u32 = 0u;          // default override
+[[override(1000)]] let c10: u32 = 10u;  // default
+
+[[block]] struct Buf {
+    data : array<u32, 11>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> buf : Buf;
+
+[[stage(compute), workgroup_size(1)]] fn main() {
+    // make sure the overridable constants are not optimized out
+    buf.data[0] = u32(c0);
+    buf.data[1] = u32(c1);
+    buf.data[2] = u32(c2);
+    buf.data[3] = u32(c3);
+    buf.data[4] = u32(c4);
+    buf.data[5] = u32(c5);
+    buf.data[6] = u32(c6);
+    buf.data[7] = u32(c7);
+    buf.data[8] = u32(c8);
+    buf.data[9] = u32(c9);
+    buf.data[10] = u32(c10);
+})");
+    }
+
+    void TestCreatePipeline(const std::vector<wgpu::ConstantEntry>& constants) {
+        wgpu::ComputePipelineDescriptor csDesc;
+        csDesc.compute.module = computeModule;
+        csDesc.compute.entryPoint = "main";
+        csDesc.compute.constants = constants.data();
+        csDesc.compute.constantCount = constants.size();
+        wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc);
+    }
+
+    wgpu::ShaderModule computeModule;
+    wgpu::Buffer buffer;
+};
+
+// Basic constants lookup tests
+TEST_F(ComputePipelineOverridableConstantsValidationTest, ConstantsIdentifierLookUp) {
+    {
+        // Valid: no constants specified
+        std::vector<wgpu::ConstantEntry> constants;
+        TestCreatePipeline(constants);
+    }
+    {
+        // Valid: find by constant name
+        std::vector<wgpu::ConstantEntry> constants{{nullptr, "c0", 0}};
+        TestCreatePipeline(constants);
+    }
+    {
+        // Valid: set the same constant twice
+        std::vector<wgpu::ConstantEntry> constants{
+            {nullptr, "c0", 0},
+            {nullptr, "c0", 1},
+        };
+        TestCreatePipeline(constants);
+    }
+    {
+        // Valid: find by constant numeric id
+        std::vector<wgpu::ConstantEntry> constants{{nullptr, "1000", 0}};
+        TestCreatePipeline(constants);
+    }
+    {
+        // Error: constant numeric id not specified
+        std::vector<wgpu::ConstantEntry> constants{{nullptr, "9999", 0}};
+        ASSERT_DEVICE_ERROR(TestCreatePipeline(constants));
+    }
+    {
+        // Error: constant name doesn't exit
+        std::vector<wgpu::ConstantEntry> constants{{nullptr, "c99", 0}};
+        ASSERT_DEVICE_ERROR(TestCreatePipeline(constants));
+    }
+}
+
+// Test that only explicitly specified numeric ID can be referenced
+// TODO(tint:1155): missing feature in tint to differentiate explicitly specified numeric ID
+TEST_F(ComputePipelineOverridableConstantsValidationTest,
+       DISABLED_ConstantsIdentifierExplicitNumericID) {
+    {
+        // Error: constant numeric id not explicitly specified
+        // But could be impliciltly assigned to one of the constants
+        std::vector<wgpu::ConstantEntry> constants{{nullptr, "0", 0}};
+        ASSERT_DEVICE_ERROR(TestCreatePipeline(constants));
+    }
+    {
+        // Error: constant numeric id not explicitly specified
+        // But could be impliciltly assigned to one of the constants
+        std::vector<wgpu::ConstantEntry> constants{{nullptr, "1", 0}};
+        ASSERT_DEVICE_ERROR(TestCreatePipeline(constants));
+    }
+    {
+        // Error: constant numeric id not explicitly specified
+        // But could be impliciltly assigned to one of the constants
+        std::vector<wgpu::ConstantEntry> constants{{nullptr, "2", 0}};
+        ASSERT_DEVICE_ERROR(TestCreatePipeline(constants));
+    }
+    {
+        // Error: constant numeric id not explicitly specified
+        // But could be impliciltly assigned to one of the constants
+        std::vector<wgpu::ConstantEntry> constants{{nullptr, "3", 0}};
+        ASSERT_DEVICE_ERROR(TestCreatePipeline(constants));
+    }
+}
+
+// Test that identifiers are unique
+// TODO(tint:1155): missing feature in tint to differentiate explicitly specified numeric ID
+TEST_F(ComputePipelineOverridableConstantsValidationTest, DISABLED_ConstantsIdentifierUnique) {
+    {
+        // Error: constant with numeric id cannot be referenced with variable name
+        std::vector<wgpu::ConstantEntry> constants{{nullptr, "c10", 0}};
+        ASSERT_DEVICE_ERROR(TestCreatePipeline(constants));
+    }
+}

src/tests/unittests/validation/ShaderModuleValidationTests.cpp

@@ -493,3 +493,22 @@ TEST_F(ShaderModuleValidationTest, ComputeWorkgroupStorageSizeLimits) {
     ASSERT_DEVICE_ERROR(MakeShaderWithWorkgroupStorage(0, kMaxMat4Count + 1));
     ASSERT_DEVICE_ERROR(MakeShaderWithWorkgroupStorage(4, kMaxMat4Count));
 }
+
+// Test that numeric ID must be unique
+TEST_F(ShaderModuleValidationTest, OverridableConstantsNumericIDConflicts) {
+    ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, R"(
+[[override(1234)]] let c0: u32;
+[[override(1234)]] let c1: u32;
+
+[[block]] struct Buf {
+    data : array<u32, 2>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> buf : Buf;
+
+[[stage(compute), workgroup_size(1)]] fn main() {
+    // make sure the overridable constants are not optimized out
+    buf.data[0] = c0;
+    buf.data[1] = c1;
+})"));
+}