[dawn][vulkan] Handle ResourceTableConfig when compiling to SPIR-V

Bug: 463925499
Change-Id: I27d98d0fb47df5ee641f81e8d3c754063d06def3
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/279895
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Corentin Wallez <cwallez@chromium.org>

src/dawn/native/vulkan/ShaderModuleVk.cpp

diff --git a/src/dawn/native/vulkan/ShaderModuleVk.cpp b/src/dawn/native/vulkan/ShaderModuleVk.cpp
index b248b89..40ce1f6 100644
--- a/src/dawn/native/vulkan/ShaderModuleVk.cpp
+++ b/src/dawn/native/vulkan/ShaderModuleVk.cpp
@@ -132,10 +132,27 @@
 #if TINT_BUILD_SPV_WRITER
     // Creation of module and spirv is deferred to this point when using tint generator
 
+    // The first VkDescriptorSetLayout is the one for the resource table if needed and pushes the
+    // bindings for all other bindgroups by 1.
+    BindGroupIndex startOfBindGroups{0};
+    std::optional<tint::ResourceTableConfig> resourceTableConfig = std::nullopt;
+    if (layout->UsesResourceTable()) {
+        startOfBindGroups = BindGroupIndex(1);
+
+        // TODO(https://issues.chromium.org/435317394): Update to not pass a DynamicBindingKind once
+        // support for dynamic binding arrays is removed.
+        auto bindingTypeOrder = GetDefaultBindingOrder(wgpu::DynamicBindingKind::SampledTexture);
+        resourceTableConfig = tint::ResourceTableConfig{
+            .resource_table_binding = tint::BindingPoint(0, 1),
+            .storage_buffer_binding = tint::BindingPoint(0, 0),
+            .default_binding_type_order = {bindingTypeOrder.begin(), bindingTypeOrder.end()},
+        };
+    }
+
     tint::Bindings bindings =
         GenerateBindingRemapping(layout, stage, [&](BindGroupIndex group, BindingIndex index) {
             return tint::BindingPoint{
-                .group = uint32_t(group),
+                .group = uint32_t(startOfBindGroups + group),
                 .binding = uint32_t(index),
             };
         });
@@ -150,7 +167,7 @@
             const auto& bindingInfo = bgl->GetBindingInfo(index);
 
             if (auto samplerIndex = bgl->GetStaticSamplerIndexForTexture(index)) {
-                tint::BindingPoint wgslBindingPoint = {.group = uint32_t(group),
+                tint::BindingPoint wgslBindingPoint = {.group = uint32_t(startOfBindGroups + group),
                                                        .binding = uint32_t(bindingInfo.binding)};
                 bindings.texture[wgslBindingPoint].binding = uint32_t(samplerIndex.value());
                 staticallyPairedTextureBindingPoints.insert(wgslBindingPoint);
@@ -167,6 +184,7 @@
         if (!bgl->HasDynamicArray()) {
             continue;
         }
+        DAWN_ASSERT(startOfBindGroups == BindGroupIndex(0));
 
         tint::BindingPoint wgslDynamicArrayBindPoint = {
             .group = uint32_t(group), .binding = uint32_t(bgl->GetAPIDynamicArrayStart())};
@@ -217,6 +235,7 @@
     };
     req.tintOptions.bindings = std::move(bindings);
     req.tintOptions.resource_binding = std::move(resourceBindingConfig);
+    req.tintOptions.resource_table = std::move(resourceTableConfig);
 
     req.tintOptions.workarounds.polyfill_unary_f32_negation =
         GetDevice()->IsToggleEnabled(Toggle::VulkanPolyfillF32Negation);

src/dawn/tests/end2end/ResourceTableTests.cpp

diff --git a/src/dawn/tests/end2end/ResourceTableTests.cpp b/src/dawn/tests/end2end/ResourceTableTests.cpp
index b6f4f95..18af22f 100644
--- a/src/dawn/tests/end2end/ResourceTableTests.cpp
+++ b/src/dawn/tests/end2end/ResourceTableTests.cpp
@@ -106,6 +106,46 @@
     MakePipelineLayoutWithTable({testBgl, testBgl, testBgl}, 4);
 }
 
+// Test that creating pipelines that use resource tables doesn't crash in backends.
+TEST_P(ResourceTableTests, ShaderWithResourceTableCreation) {
+    wgpu::ComputePipelineDescriptor csDesc;
+
+    // Test compiling a pipeline using only the resource table.
+    csDesc.compute.module = utils::CreateShaderModule(device, R"(
+        enable chromium_experimental_resource_table;
+        @compute @workgroup_size(1) fn main() {
+            _ = hasResource<texture_2d<f32>>(0);
+        }
+    )");
+    device.CreateComputePipeline(&csDesc);
+
+    // Test compiling a pipeline using the resource table and a bindgroup.
+    csDesc.compute.module = utils::CreateShaderModule(device, R"(
+        enable chromium_experimental_resource_table;
+        @group(0) @binding(0) var t0 : texture_2d<f32>;
+        @compute @workgroup_size(1) fn main() {
+            _ = hasResource<texture_2d<f32>>(0);
+            _ = t0;
+        }
+    )");
+    device.CreateComputePipeline(&csDesc);
+
+    // Test compiling a pipeline using the resource table and many bindgroup.
+    csDesc.compute.module = utils::CreateShaderModule(device, R"(
+        enable chromium_experimental_resource_table;
+        @group(0) @binding(0) var t0 : texture_2d<f32>;
+        @group(1) @binding(0) var t1 : texture_2d<f32>;
+        @group(2) @binding(0) var t2 : texture_2d<f32>;
+        @compute @workgroup_size(1) fn main() {
+            _ = hasResource<texture_2d<f32>>(0);
+            _ = t0;
+            _ = t1;
+            _ = t2;
+        }
+    )");
+    device.CreateComputePipeline(&csDesc);
+}
+
 DAWN_INSTANTIATE_TEST(ResourceTableTests, D3D12Backend(), MetalBackend(), VulkanBackend());
 
 }  // anonymous namespace