Adds pipeline cache and implementation for Vulkan backend.

- Adds testing for Vulkan pipeline caching.
- Removed redundant VK_NULL_HANDLE and use explicit {} initialization for 0 handles when necessary.
- Adds some const qualifiers where applicable and useful.
- Removes overloaded GetCacheKey (const/non-const) versions and exposed the cache key member directly for modifiers in derived classes.

Bug: dawn:549
Change-Id: I5e8ab9716eebc916b813c9d032f8dc1f3f5261bc
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/86581
Commit-Queue: Loko Kung <lokokung@google.com>
Reviewed-by: Austin Eng <enga@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>

src/dawn/common/vulkan_platform.h

@@ -117,8 +117,6 @@ class alignas(detail::kNativeVkHandleAlignment) VkHandle {
 };
 }  // namespace detail
 
-static constexpr std::nullptr_t VK_NULL_HANDLE = nullptr;
-
 template <typename Tag, typename HandleType>
 HandleType* AsVkArray(detail::VkHandle<Tag, HandleType>* handle) {
     return reinterpret_cast<HandleType*>(handle);
@@ -179,12 +177,6 @@ HandleType* AsVkArray(detail::VkHandle<Tag, HandleType>* handle) {
 
 // Redefine VK_NULL_HANDLE for better type safety where possible.
 #undef VK_NULL_HANDLE
-#if defined(DAWN_PLATFORM_64_BIT)
 static constexpr std::nullptr_t VK_NULL_HANDLE = nullptr;
-#elif defined(DAWN_PLATFORM_32_BIT)
-static constexpr uint64_t VK_NULL_HANDLE = 0;
-#else
-#error "Unsupported platform"
-#endif
 
 #endif  // SRC_DAWN_COMMON_VULKAN_PLATFORM_H_

src/dawn/native/BUILD.gn

@@ -260,6 +260,8 @@ source_set("sources") {
     "PerStage.h",
     "Pipeline.cpp",
     "Pipeline.h",
+    "PipelineCache.cpp",
+    "PipelineCache.h",
     "PipelineLayout.cpp",
     "PipelineLayout.h",
     "PooledResourceMemoryAllocator.cpp",
@@ -586,6 +588,8 @@ source_set("sources") {
       "vulkan/Forward.h",
       "vulkan/NativeSwapChainImplVk.cpp",
       "vulkan/NativeSwapChainImplVk.h",
+      "vulkan/PipelineCacheVk.cpp",
+      "vulkan/PipelineCacheVk.h",
       "vulkan/PipelineLayoutVk.cpp",
       "vulkan/PipelineLayoutVk.h",
       "vulkan/QuerySetVk.cpp",

src/dawn/native/CMakeLists.txt

@@ -126,6 +126,8 @@ target_sources(dawn_native PRIVATE
     "PerStage.h"
     "Pipeline.cpp"
     "Pipeline.h"
+    "PipelineCache.cpp"
+    "PipelineCache.h"
     "PipelineLayout.cpp"
     "PipelineLayout.h"
     "PooledResourceMemoryAllocator.cpp"
@@ -477,6 +479,8 @@ if (DAWN_ENABLE_VULKAN)
         "vulkan/Forward.h"
         "vulkan/NativeSwapChainImplVk.cpp"
         "vulkan/NativeSwapChainImplVk.h"
+        "vulkan/PipelineCacheVk.cpp"
+        "vulkan/PipelineCacheVk.h"
         "vulkan/PipelineLayoutVk.cpp"
         "vulkan/PipelineLayoutVk.h"
         "vulkan/QuerySetVk.cpp"

src/dawn/native/CachedObject.cpp

@@ -46,8 +46,4 @@ const CacheKey& CachedObject::GetCacheKey() const {
     return mCacheKey;
 }
 
-CacheKey* CachedObject::GetCacheKey() {
-    return &mCacheKey;
-}
-
 }  // namespace dawn::native

src/dawn/native/CachedObject.h

@@ -43,8 +43,8 @@ class CachedObject {
     const CacheKey& GetCacheKey() const;
 
   protected:
-    // Protected accessor for derived classes to access and modify the key.
-    CacheKey* GetCacheKey();
+    // Cache key member is protected so that derived classes can modify it.
+    CacheKey mCacheKey;
 
   private:
     friend class DeviceBase;
@@ -57,7 +57,6 @@ class CachedObject {
 
     size_t mContentHash = 0;
     bool mIsContentHashInitialized = false;
-    CacheKey mCacheKey;
 };
 
 }  // namespace dawn::native

src/dawn/native/ComputePipeline.cpp

@@ -50,7 +50,7 @@ ComputePipelineBase::ComputePipelineBase(DeviceBase* device,
     TrackInDevice();
 
     // Initialize the cache key to include the cache type and device information.
-    GetCacheKey()->Record(CacheKey::Type::ComputePipeline, device->GetCacheKey());
+    mCacheKey.Record(CacheKey::Type::ComputePipeline, device->GetCacheKey());
 }
 
 ComputePipelineBase::ComputePipelineBase(DeviceBase* device) : PipelineBase(device) {

src/dawn/native/Device.cpp

@@ -40,6 +40,7 @@
 #include "dawn/native/Instance.h"
 #include "dawn/native/InternalPipelineStore.h"
 #include "dawn/native/ObjectType_autogen.h"
+#include "dawn/native/PipelineCache.h"
 #include "dawn/native/QuerySet.h"
 #include "dawn/native/Queue.h"
 #include "dawn/native/RenderBundleEncoder.h"
@@ -968,6 +969,10 @@ void DeviceBase::UncacheAttachmentState(AttachmentState* obj) {
     ASSERT(removedCount == 1);
 }
 
+Ref<PipelineCacheBase> DeviceBase::GetOrCreatePipelineCache(const CacheKey& key) {
+    return GetOrCreatePipelineCacheImpl(key);
+}
+
 // Object creation API methods
 
 BindGroupBase* DeviceBase::APICreateBindGroup(const BindGroupDescriptor* descriptor) {
@@ -1377,6 +1382,11 @@ ResultOrError<Ref<CommandEncoder>> DeviceBase::CreateCommandEncoder(
     return CommandEncoder::Create(this, descriptor);
 }
 
+// Overwritten on the backends to return pipeline caches if supported.
+Ref<PipelineCacheBase> DeviceBase::GetOrCreatePipelineCacheImpl(const CacheKey& key) {
+    UNREACHABLE();
+}
+
 MaybeError DeviceBase::CreateComputePipelineAsync(const ComputePipelineDescriptor* descriptor,
                                                   WGPUCreateComputePipelineAsyncCallback callback,
                                                   void* userdata) {

src/dawn/native/Device.h

@@ -195,6 +195,8 @@ class DeviceBase : public RefCounted {
     Ref<AttachmentState> GetOrCreateAttachmentState(const RenderPassDescriptor* descriptor);
     void UncacheAttachmentState(AttachmentState* obj);
 
+    Ref<PipelineCacheBase> GetOrCreatePipelineCache(const CacheKey& key);
+
     // Object creation methods that be used in a reentrant manner.
     ResultOrError<Ref<BindGroupBase>> CreateBindGroup(const BindGroupDescriptor* descriptor);
     ResultOrError<Ref<BindGroupLayoutBase>> CreateBindGroupLayout(
@@ -441,6 +443,7 @@ class DeviceBase : public RefCounted {
     Ref<ComputePipelineBase> AddOrGetCachedComputePipeline(
         Ref<ComputePipelineBase> computePipeline);
     Ref<RenderPipelineBase> AddOrGetCachedRenderPipeline(Ref<RenderPipelineBase> renderPipeline);
+    virtual Ref<PipelineCacheBase> GetOrCreatePipelineCacheImpl(const CacheKey& key);
     virtual void InitializeComputePipelineAsyncImpl(Ref<ComputePipelineBase> computePipeline,
                                                     WGPUCreateComputePipelineAsyncCallback callback,
                                                     void* userdata);

src/dawn/native/Forward.h

@@ -35,6 +35,7 @@ class ComputePassEncoder;
 class ExternalTextureBase;
 class InstanceBase;
 class PipelineBase;
+class PipelineCacheBase;
 class PipelineLayoutBase;
 class QuerySetBase;
 class QueueBase;

src/dawn/native/PipelineCache.cpp

@@ -0,0 +1,55 @@
+// Copyright 2022 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 "dawn/native/PipelineCache.h"
+
+namespace dawn::native {
+
+PipelineCacheBase::PipelineCacheBase(BlobCache* cache, const CacheKey& key)
+    : mCache(cache), mKey(key) {}
+
+CachedBlob PipelineCacheBase::Initialize() {
+    ASSERT(!mInitialized);
+    CachedBlob blob = mCache->Load(mKey);
+    mCacheHit = !blob.Empty();
+    mInitialized = true;
+    return blob;
+}
+
+bool PipelineCacheBase::CacheHit() const {
+    ASSERT(mInitialized);
+    return mCacheHit;
+}
+
+MaybeError PipelineCacheBase::Flush() {
+    // Try to write the data out to the persistent cache.
+    CachedBlob blob;
+    DAWN_TRY_ASSIGN(blob, SerializeToBlobImpl());
+    if (blob.Size() > 0) {
+        // Using a simple heuristic to decide whether to write out the blob right now. May need
+        // smarter tracking when we are dealing with monolithic caches.
+        mCache->Store(mKey, blob);
+    }
+    return {};
+}
+
+MaybeError PipelineCacheBase::FlushIfNeeded() {
+    ASSERT(mInitialized);
+    if (!CacheHit()) {
+        return Flush();
+    }
+    return {};
+}
+
+}  // namespace dawn::native

src/dawn/native/PipelineCache.h

@@ -0,0 +1,63 @@
+// Copyright 2022 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.
+
+#ifndef SRC_DAWN_NATIVE_PIPELINECACHE_H_
+#define SRC_DAWN_NATIVE_PIPELINECACHE_H_
+
+#include "dawn/common/RefCounted.h"
+#include "dawn/native/BlobCache.h"
+#include "dawn/native/CacheKey.h"
+#include "dawn/native/Error.h"
+
+namespace dawn::native {
+
+// Abstraction layer for backend dependent pipeline caching.
+class PipelineCacheBase : public RefCounted {
+  public:
+    // Returns whether or not we got a cache hit when initializing.
+    bool CacheHit() const;
+
+    // Serializes and writes the current contents of the backend cache object into the backing
+    // blob cache, potentially overwriting what is already there. Useful when we are working
+    // with more monolithic-like caches where we expect overwriting sometimes.
+    MaybeError Flush();
+
+    // Serializes and writes the current contents of the backend cache object into the backing
+    // blob cache iff the initial read from the backend cache did not result in a hit.
+    MaybeError FlushIfNeeded();
+
+  protected:
+    PipelineCacheBase(BlobCache* cache, const CacheKey& key);
+
+    // Initializes and returns the cached blob given the cache and keys. Used by backend
+    // implementations to get the cache and set the cache hit state. Should only be called once.
+    CachedBlob Initialize();
+
+  private:
+    // Backend implementation of serialization of the cache into a blob. Note that an empty
+    // blob may be returned.
+    virtual ResultOrError<CachedBlob> SerializeToBlobImpl() = 0;
+
+    // The blob cache is owned by the Adapter and pipeline caches are owned/created by devices
+    // or adapters. Since the device owns a reference to the Instance which owns the Adapter,
+    // the blob cache is guaranteed to be valid throughout the lifetime of the object.
+    BlobCache* mCache;
+    CacheKey mKey;
+    bool mInitialized = false;
+    bool mCacheHit = false;
+};
+
+}  // namespace dawn::native
+
+#endif  // SRC_DAWN_NATIVE_PIPELINECACHE_H_

src/dawn/native/RenderPipeline.cpp

@@ -613,7 +613,7 @@ RenderPipelineBase::RenderPipelineBase(DeviceBase* device,
     TrackInDevice();
 
     // Initialize the cache key to include the cache type and device information.
-    GetCacheKey()->Record(CacheKey::Type::RenderPipeline, device->GetCacheKey());
+    mCacheKey.Record(CacheKey::Type::RenderPipeline, device->GetCacheKey());
 }
 
 RenderPipelineBase::RenderPipelineBase(DeviceBase* device) : PipelineBase(device) {

src/dawn/native/ShaderModule.h

@@ -75,7 +75,7 @@ enum class InterpolationSampling {
     Sample,
 };
 
-using PipelineLayoutEntryPointPair = std::pair<PipelineLayoutBase*, std::string>;
+using PipelineLayoutEntryPointPair = std::pair<const PipelineLayoutBase*, std::string>;
 struct PipelineLayoutEntryPointPairHashFunc {
     size_t operator()(const PipelineLayoutEntryPointPair& pair) const;
 };

src/dawn/native/ToBackend.h

@@ -58,6 +58,11 @@ struct ToBackendTraits<DeviceBase, BackendTraits> {
     using BackendType = typename BackendTraits::DeviceType;
 };
 
+template <typename BackendTraits>
+struct ToBackendTraits<PipelineCacheBase, BackendTraits> {
+    using BackendType = typename BackendTraits::PipelineCacheType;
+};
+
 template <typename BackendTraits>
 struct ToBackendTraits<PipelineLayoutBase, BackendTraits> {
     using BackendType = typename BackendTraits::PipelineLayoutType;

src/dawn/native/vulkan/BindGroupLayoutVk.cpp

@@ -117,7 +117,7 @@ MaybeError BindGroupLayout::Initialize() {
     createInfo.pBindings = bindings.data();
 
     // Record cache key information now since the createInfo is not stored.
-    GetCacheKey()->Record(createInfo);
+    mCacheKey.Record(createInfo);
 
     Device* device = ToBackend(GetDevice());
     DAWN_TRY(CheckVkSuccess(device->fn.CreateDescriptorSetLayout(device->GetVkDevice(), &createInfo,

src/dawn/native/vulkan/ComputePipelineVk.cpp

@@ -21,6 +21,7 @@
 #include "dawn/native/CreatePipelineAsyncTask.h"
 #include "dawn/native/vulkan/DeviceVk.h"
 #include "dawn/native/vulkan/FencedDeleter.h"
+#include "dawn/native/vulkan/PipelineCacheVk.h"
 #include "dawn/native/vulkan/PipelineLayoutVk.h"
 #include "dawn/native/vulkan/ShaderModuleVk.h"
 #include "dawn/native/vulkan/UtilsVulkan.h"
@@ -36,12 +37,18 @@ Ref<ComputePipeline> ComputePipeline::CreateUninitialized(
 }
 
 MaybeError ComputePipeline::Initialize() {
+    Device* device = ToBackend(GetDevice());
+    const PipelineLayout* layout = ToBackend(GetLayout());
+
+    // Vulkan devices need cache UUID field to be serialized into pipeline cache keys.
+    mCacheKey.Record(device->GetDeviceInfo().properties.pipelineCacheUUID);
+
     VkComputePipelineCreateInfo createInfo;
     createInfo.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;
     createInfo.pNext = nullptr;
     createInfo.flags = 0;
-    createInfo.layout = ToBackend(GetLayout())->GetHandle();
-    createInfo.basePipelineHandle = ::VK_NULL_HANDLE;
+    createInfo.layout = layout->GetHandle();
+    createInfo.basePipelineHandle = VkPipeline{};
     createInfo.basePipelineIndex = -1;
 
     createInfo.stage.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO;
@@ -51,7 +58,6 @@ MaybeError ComputePipeline::Initialize() {
     // Generate a new VkShaderModule with BindingRemapper tint transform for each pipeline
     const ProgrammableStage& computeStage = GetStage(SingleShaderStage::Compute);
     ShaderModule* module = ToBackend(computeStage.module.Get());
-    PipelineLayout* layout = ToBackend(GetLayout());
     const ShaderModule::Spirv* spirv;
     DAWN_TRY_ASSIGN((std::tie(createInfo.stage.module, spirv)),
                     module->GetHandleAndSpirv(computeStage.entryPoint.c_str(), layout));
@@ -64,8 +70,6 @@ MaybeError ComputePipeline::Initialize() {
     createInfo.stage.pSpecializationInfo = GetVkSpecializationInfo(
         computeStage, &specializationInfo, &specializationDataEntries, &specializationMapEntries);
 
-    Device* device = ToBackend(GetDevice());
-
     PNextChainBuilder stageExtChain(&createInfo.stage);
 
     VkPipelineShaderStageRequiredSubgroupSizeCreateInfoEXT subgroupSizeInfo = {};
@@ -79,14 +83,16 @@ MaybeError ComputePipeline::Initialize() {
     }
 
     // Record cache key information now since the createInfo is not stored.
-    GetCacheKey()
-        ->Record(createInfo, static_cast<const ComputePipeline*>(this)->GetLayout())
-        .RecordIterable(*spirv);
+    mCacheKey.Record(createInfo, layout).RecordIterable(*spirv);
 
+    // Try to see if we have anything in the blob cache.
+    Ref<PipelineCache> cache = ToBackend(GetDevice()->GetOrCreatePipelineCache(GetCacheKey()));
     DAWN_TRY(
-        CheckVkSuccess(device->fn.CreateComputePipelines(device->GetVkDevice(), ::VK_NULL_HANDLE, 1,
-                                                         &createInfo, nullptr, &*mHandle),
+        CheckVkSuccess(device->fn.CreateComputePipelines(device->GetVkDevice(), cache->GetHandle(),
+                                                         1, &createInfo, nullptr, &*mHandle),
                        "CreateComputePipeline"));
+    // TODO(dawn:549): Flush is currently in the same thread, but perhaps deferrable.
+    DAWN_TRY(cache->FlushIfNeeded());
 
     SetLabelImpl();
 

src/dawn/native/vulkan/DeviceVk.cpp

@@ -28,6 +28,7 @@
 #include "dawn/native/vulkan/CommandBufferVk.h"
 #include "dawn/native/vulkan/ComputePipelineVk.h"
 #include "dawn/native/vulkan/FencedDeleter.h"
+#include "dawn/native/vulkan/PipelineCacheVk.h"
 #include "dawn/native/vulkan/PipelineLayoutVk.h"
 #include "dawn/native/vulkan/QuerySetVk.h"
 #include "dawn/native/vulkan/QueueVk.h"
@@ -167,6 +168,9 @@ ResultOrError<Ref<TextureViewBase>> Device::CreateTextureViewImpl(
     const TextureViewDescriptor* descriptor) {
     return TextureView::Create(texture, descriptor);
 }
+Ref<PipelineCacheBase> Device::GetOrCreatePipelineCacheImpl(const CacheKey& key) {
+    return PipelineCache::Create(this, key);
+}
 void Device::InitializeComputePipelineAsyncImpl(Ref<ComputePipelineBase> computePipeline,
                                                 WGPUCreateComputePipelineAsyncCallback callback,
                                                 void* userdata) {

src/dawn/native/vulkan/DeviceVk.h

@@ -35,8 +35,6 @@
 
 namespace dawn::native::vulkan {
 
-class Adapter;
-class BindGroupLayout;
 class BufferUploader;
 class FencedDeleter;
 class RenderPassCache;
@@ -138,6 +136,7 @@ class Device final : public DeviceBase {
         const ComputePipelineDescriptor* descriptor) override;
     Ref<RenderPipelineBase> CreateUninitializedRenderPipelineImpl(
         const RenderPipelineDescriptor* descriptor) override;
+    Ref<PipelineCacheBase> GetOrCreatePipelineCacheImpl(const CacheKey& key) override;
     void InitializeComputePipelineAsyncImpl(Ref<ComputePipelineBase> computePipeline,
                                             WGPUCreateComputePipelineAsyncCallback callback,
                                             void* userdata) override;

src/dawn/native/vulkan/Forward.h

@@ -26,6 +26,7 @@ class Buffer;
 class CommandBuffer;
 class ComputePipeline;
 class Device;
+class PipelineCache;
 class PipelineLayout;
 class QuerySet;
 class Queue;
@@ -46,6 +47,7 @@ struct VulkanBackendTraits {
     using CommandBufferType = CommandBuffer;
     using ComputePipelineType = ComputePipeline;
     using DeviceType = Device;
+    using PipelineCacheType = PipelineCache;
     using PipelineLayoutType = PipelineLayout;
     using QuerySetType = QuerySet;
     using QueueType = Queue;

src/dawn/native/vulkan/PipelineCacheVk.cpp

@@ -0,0 +1,89 @@
+// Copyright 2022 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 "dawn/native/vulkan/PipelineCacheVk.h"
+
+#include "dawn/common/Assert.h"
+#include "dawn/native/Device.h"
+#include "dawn/native/Error.h"
+#include "dawn/native/vulkan/DeviceVk.h"
+#include "dawn/native/vulkan/FencedDeleter.h"
+#include "dawn/native/vulkan/VulkanError.h"
+
+namespace dawn::native::vulkan {
+
+// static
+Ref<PipelineCache> PipelineCache::Create(DeviceBase* device, const CacheKey& key) {
+    Ref<PipelineCache> cache = AcquireRef(new PipelineCache(device, key));
+    cache->Initialize();
+    return cache;
+}
+
+PipelineCache::PipelineCache(DeviceBase* device, const CacheKey& key)
+    : PipelineCacheBase(device->GetBlobCache(), key), mDevice(device) {}
+
+PipelineCache::~PipelineCache() {
+    if (mHandle == VK_NULL_HANDLE) {
+        return;
+    }
+    Device* device = ToBackend(GetDevice());
+    device->fn.DestroyPipelineCache(device->GetVkDevice(), mHandle, nullptr);
+    mHandle = VK_NULL_HANDLE;
+}
+
+DeviceBase* PipelineCache::GetDevice() const {
+    return mDevice;
+}
+
+VkPipelineCache PipelineCache::GetHandle() const {
+    return mHandle;
+}
+
+ResultOrError<CachedBlob> PipelineCache::SerializeToBlobImpl() {
+    CachedBlob emptyBlob;
+    if (mHandle == VK_NULL_HANDLE) {
+        return emptyBlob;
+    }
+
+    size_t bufferSize;
+    Device* device = ToBackend(GetDevice());
+    DAWN_TRY(CheckVkSuccess(
+        device->fn.GetPipelineCacheData(device->GetVkDevice(), mHandle, &bufferSize, nullptr),
+        "GetPipelineCacheData"));
+
+    CachedBlob blob(bufferSize);
+    DAWN_TRY(CheckVkSuccess(
+        device->fn.GetPipelineCacheData(device->GetVkDevice(), mHandle, &bufferSize, blob.Data()),
+        "GetPipelineCacheData"));
+    return blob;
+}
+
+void PipelineCache::Initialize() {
+    CachedBlob blob = PipelineCacheBase::Initialize();
+
+    VkPipelineCacheCreateInfo createInfo;
+    createInfo.flags = 0;
+    createInfo.sType = VK_STRUCTURE_TYPE_PIPELINE_CACHE_CREATE_INFO;
+    createInfo.pNext = nullptr;
+    createInfo.initialDataSize = blob.Size();
+    createInfo.pInitialData = blob.Data();
+
+    Device* device = ToBackend(GetDevice());
+    mHandle = VK_NULL_HANDLE;
+    GetDevice()->ConsumedError(CheckVkSuccess(
+        device->fn.CreatePipelineCache(device->GetVkDevice(), &createInfo, nullptr, &*mHandle),
+        "CreatePipelineCache"));
+}
+
+}  // namespace dawn::native::vulkan

src/dawn/native/vulkan/PipelineCacheVk.h

@@ -0,0 +1,49 @@
+// Copyright 2022 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.
+
+#ifndef SRC_DAWN_NATIVE_VULKAN_PIPELINECACHEVK_H_
+#define SRC_DAWN_NATIVE_VULKAN_PIPELINECACHEVK_H_
+
+#include "dawn/native/ObjectBase.h"
+#include "dawn/native/PipelineCache.h"
+
+#include "dawn/common/vulkan_platform.h"
+
+namespace dawn::native {
+class DeviceBase;
+}
+
+namespace dawn::native::vulkan {
+
+class PipelineCache final : public PipelineCacheBase {
+  public:
+    static Ref<PipelineCache> Create(DeviceBase* device, const CacheKey& key);
+
+    DeviceBase* GetDevice() const;
+    VkPipelineCache GetHandle() const;
+
+  private:
+    explicit PipelineCache(DeviceBase* device, const CacheKey& key);
+    ~PipelineCache() override;
+
+    void Initialize();
+    ResultOrError<CachedBlob> SerializeToBlobImpl() override;
+
+    DeviceBase* mDevice;
+    VkPipelineCache mHandle = VK_NULL_HANDLE;
+};
+
+}  // namespace dawn::native::vulkan
+
+#endif  // SRC_DAWN_NATIVE_VULKAN_PIPELINECACHEVK_H_

src/dawn/native/vulkan/PipelineLayoutVk.cpp

@@ -56,7 +56,7 @@ MaybeError PipelineLayout::Initialize() {
     createInfo.pPushConstantRanges = nullptr;
 
     // Record cache key information now since the createInfo is not stored.
-    GetCacheKey()->RecordIterable(cachedObjects.data(), numSetLayouts).Record(createInfo);
+    mCacheKey.RecordIterable(cachedObjects.data(), numSetLayouts).Record(createInfo);
 
     Device* device = ToBackend(GetDevice());
     DAWN_TRY(CheckVkSuccess(

src/dawn/native/vulkan/RenderPipelineVk.cpp

@@ -21,6 +21,7 @@
 #include "dawn/native/CreatePipelineAsyncTask.h"
 #include "dawn/native/vulkan/DeviceVk.h"
 #include "dawn/native/vulkan/FencedDeleter.h"
+#include "dawn/native/vulkan/PipelineCacheVk.h"
 #include "dawn/native/vulkan/PipelineLayoutVk.h"
 #include "dawn/native/vulkan/RenderPassCache.h"
 #include "dawn/native/vulkan/ShaderModuleVk.h"
@@ -331,7 +332,10 @@ Ref<RenderPipeline> RenderPipeline::CreateUninitialized(
 
 MaybeError RenderPipeline::Initialize() {
     Device* device = ToBackend(GetDevice());
-    PipelineLayout* layout = ToBackend(GetLayout());
+    const PipelineLayout* layout = ToBackend(GetLayout());
+
+    // Vulkan devices need cache UUID field to be serialized into pipeline cache keys.
+    mCacheKey.Record(device->GetDeviceInfo().properties.pipelineCacheUUID);
 
     // There are at most 2 shader stages in render pipeline, i.e. vertex and fragment
     std::array<VkPipelineShaderStageCreateInfo, 2> shaderStages;
@@ -381,7 +385,7 @@ MaybeError RenderPipeline::Initialize() {
         stageCount++;
 
         // Record cache key for each shader since it will become inaccessible later on.
-        GetCacheKey()->Record(stage).RecordIterable(*spirv);
+        mCacheKey.Record(stage).RecordIterable(*spirv);
     }
 
     PipelineVertexInputStateCreateInfoTemporaryAllocations tempAllocations;
@@ -528,7 +532,7 @@ MaybeError RenderPipeline::Initialize() {
 
         query.SetSampleCount(GetSampleCount());
 
-        GetCacheKey()->Record(query);
+        mCacheKey.Record(query);
         DAWN_TRY_ASSIGN(renderPass, device->GetRenderPassCache()->GetRenderPass(query));
     }
 
@@ -557,13 +561,16 @@ MaybeError RenderPipeline::Initialize() {
     createInfo.basePipelineIndex = -1;
 
     // Record cache key information now since createInfo is not stored.
-    GetCacheKey()->Record(createInfo,
-                          static_cast<const RenderPipeline*>(this)->GetLayout()->GetCacheKey());
+    mCacheKey.Record(createInfo, layout->GetCacheKey());
 
+    // Try to see if we have anything in the blob cache.
+    Ref<PipelineCache> cache = ToBackend(GetDevice()->GetOrCreatePipelineCache(GetCacheKey()));
     DAWN_TRY(
-        CheckVkSuccess(device->fn.CreateGraphicsPipelines(device->GetVkDevice(), VkPipelineCache{},
+        CheckVkSuccess(device->fn.CreateGraphicsPipelines(device->GetVkDevice(), cache->GetHandle(),
                                                           1, &createInfo, nullptr, &*mHandle),
-                       "CreateGraphicsPipeline"));
+                       "CreateGraphicsPipelines"));
+    // TODO(dawn:549): Flush is currently in the same thread, but perhaps deferrable.
+    DAWN_TRY(cache->FlushIfNeeded());
 
     SetLabelImpl();
 

src/dawn/native/vulkan/ShaderModuleVk.cpp

@@ -113,7 +113,7 @@ ShaderModule::~ShaderModule() = default;
 
 ResultOrError<ShaderModule::ModuleAndSpirv> ShaderModule::GetHandleAndSpirv(
     const char* entryPointName,
-    PipelineLayout* layout) {
+    const PipelineLayout* layout) {
     TRACE_EVENT0(GetDevice()->GetPlatform(), General, "ShaderModuleVk::GetHandleAndSpirv");
 
     // If the shader was destroyed, we should never call this function.
@@ -170,7 +170,7 @@ ResultOrError<ShaderModule::ModuleAndSpirv> ShaderModule::GetHandleAndSpirv(
     // TODO(dawn:1082): Replace this block with ShaderModuleBase::AddExternalTextureTransform.
     tint::transform::MultiplanarExternalTexture::BindingsMap newBindingsMap;
     for (BindGroupIndex i : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
-        BindGroupLayoutBase* bgl = layout->GetBindGroupLayout(i);
+        const BindGroupLayoutBase* bgl = layout->GetBindGroupLayout(i);
 
         ExternalTextureBindingExpansionMap expansions =
             bgl->GetExternalTextureBindingExpansionMap();

src/dawn/native/vulkan/ShaderModuleVk.h

@@ -43,7 +43,7 @@ class ShaderModule final : public ShaderModuleBase {
                                                    ShaderModuleParseResult* parseResult);
 
     ResultOrError<ModuleAndSpirv> GetHandleAndSpirv(const char* entryPointName,
-                                                    PipelineLayout* layout);
+                                                    const PipelineLayout* layout);
 
   private:
     ShaderModule(Device* device, const ShaderModuleDescriptor* descriptor);

src/dawn/tests/BUILD.gn

@@ -446,6 +446,7 @@ source_set("end2end_tests_sources") {
     "end2end/NonzeroTextureCreationTests.cpp",
     "end2end/ObjectCachingTests.cpp",
     "end2end/OpArrayLengthTests.cpp",
+    "end2end/PipelineCachingTests.cpp",
     "end2end/PipelineLayoutTests.cpp",
     "end2end/PrimitiveStateTests.cpp",
     "end2end/PrimitiveTopologyTests.cpp",

src/dawn/tests/end2end/PipelineCachingTests.cpp

@@ -0,0 +1,269 @@
+// Copyright 2022 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 <memory>
+#include <string_view>
+
+#include "dawn/tests/DawnTest.h"
+#include "dawn/tests/end2end/mocks/CachingInterfaceMock.h"
+#include "dawn/utils/ComboRenderPipelineDescriptor.h"
+#include "dawn/utils/WGPUHelpers.h"
+
+namespace {
+
+using ::testing::NiceMock;
+
+// TODO(dawn:549) Add some sort of pipeline descriptor repository to test more caching.
+
+static constexpr std::string_view kComputeShader = R"(
+        @stage(compute) @workgroup_size(1) fn main() {}
+    )";
+
+static constexpr std::string_view kVertexShader = R"(
+        @stage(vertex) fn main() -> @builtin(position) vec4<f32> {
+            return vec4<f32>(0.0, 0.0, 0.0, 0.0);
+        }
+    )";
+
+static constexpr std::string_view kFragmentShader = R"(
+        @stage(fragment) fn main() {}
+    )";
+
+class PipelineCachingTests : public DawnTest {
+  protected:
+    std::unique_ptr<dawn::platform::Platform> CreateTestPlatform() override {
+        return std::make_unique<DawnCachingMockPlatform>(&mMockCache);
+    }
+
+    NiceMock<CachingInterfaceMock> mMockCache;
+};
+
+class SinglePipelineCachingTests : public PipelineCachingTests {};
+
+// Tests that pipeline creation works fine even if the cache is disabled.
+// Note: This tests needs to use more than 1 device since the frontend cache on each device
+//   will prevent going out to the blob cache.
+TEST_P(SinglePipelineCachingTests, ComputePipelineNoCache) {
+    mMockCache.Disable();
+
+    // First time should create and since cache is disabled, it should not write out to the
+    // cache.
+    {
+        wgpu::Device device = CreateDevice();
+        wgpu::ComputePipelineDescriptor desc;
+        desc.compute.module = utils::CreateShaderModule(device, kComputeShader.data());
+        desc.compute.entryPoint = "main";
+        EXPECT_CACHE_HIT(mMockCache, 0u, device.CreateComputePipeline(&desc));
+    }
+    EXPECT_EQ(mMockCache.GetNumEntries(), 0u);
+
+    // Second time should create fine with no cache hits since cache is disabled.
+    {
+        wgpu::Device device = CreateDevice();
+        wgpu::ComputePipelineDescriptor desc;
+        desc.compute.module = utils::CreateShaderModule(device, kComputeShader.data());
+        desc.compute.entryPoint = "main";
+        EXPECT_CACHE_HIT(mMockCache, 0u, device.CreateComputePipeline(&desc));
+    }
+    EXPECT_EQ(mMockCache.GetNumEntries(), 0u);
+}
+
+// Tests that pipeline creation on the same device uses frontend cache when possible.
+TEST_P(SinglePipelineCachingTests, ComputePipelineFrontedCache) {
+    wgpu::ComputePipelineDescriptor desc;
+    desc.compute.module = utils::CreateShaderModule(device, kComputeShader.data());
+    desc.compute.entryPoint = "main";
+
+    // First creation should create a cache entry.
+    wgpu::ComputePipeline pipeline;
+    EXPECT_CACHE_HIT(mMockCache, 0u, pipeline = device.CreateComputePipeline(&desc));
+    EXPECT_EQ(mMockCache.GetNumEntries(), 1u);
+
+    // Second creation on the same device should just return from frontend cache and should not
+    // call out to the blob cache.
+    EXPECT_CALL(mMockCache, LoadData).Times(0);
+    wgpu::ComputePipeline samePipeline;
+    EXPECT_CACHE_HIT(mMockCache, 0u, samePipeline = device.CreateComputePipeline(&desc));
+    EXPECT_EQ(pipeline.Get() == samePipeline.Get(), !UsesWire());
+}
+
+// Tests that pipeline creation hits the cache when it is enabled.
+// Note: This test needs to use more than 1 device since the frontend cache on each device
+//   will prevent going out to the blob cache.
+TEST_P(SinglePipelineCachingTests, ComputePipelineBlobCache) {
+    // First time should create and write out to the cache.
+    {
+        wgpu::Device device = CreateDevice();
+        wgpu::ComputePipelineDescriptor desc;
+        desc.compute.module = utils::CreateShaderModule(device, kComputeShader.data());
+        desc.compute.entryPoint = "main";
+        EXPECT_CACHE_HIT(mMockCache, 0u, device.CreateComputePipeline(&desc));
+    }
+    EXPECT_EQ(mMockCache.GetNumEntries(), 1u);
+
+    // Second time should create using the cache.
+    {
+        wgpu::Device device = CreateDevice();
+        wgpu::ComputePipelineDescriptor desc;
+        desc.compute.module = utils::CreateShaderModule(device, kComputeShader.data());
+        desc.compute.entryPoint = "main";
+        EXPECT_CACHE_HIT(mMockCache, 1u, device.CreateComputePipeline(&desc));
+    }
+    EXPECT_EQ(mMockCache.GetNumEntries(), 1u);
+}
+
+// Tests that pipeline creation does not hits the cache when it is enabled but we use different
+// isolation keys.
+TEST_P(SinglePipelineCachingTests, ComputePipelineBlobCacheIsolationKey) {
+    // First time should create and write out to the cache.
+    {
+        wgpu::Device device = CreateDevice("isolation key 1");
+        wgpu::ComputePipelineDescriptor desc;
+        desc.compute.module = utils::CreateShaderModule(device, kComputeShader.data());
+        desc.compute.entryPoint = "main";
+        EXPECT_CACHE_HIT(mMockCache, 0u, device.CreateComputePipeline(&desc));
+    }
+    EXPECT_EQ(mMockCache.GetNumEntries(), 1u);
+
+    // Second time should also create and write out to the cache.
+    {
+        wgpu::Device device = CreateDevice("isolation key 2");
+        wgpu::ComputePipelineDescriptor desc;
+        desc.compute.module = utils::CreateShaderModule(device, kComputeShader.data());
+        desc.compute.entryPoint = "main";
+        EXPECT_CACHE_HIT(mMockCache, 0u, device.CreateComputePipeline(&desc));
+    }
+    EXPECT_EQ(mMockCache.GetNumEntries(), 2u);
+}
+
+// Tests that pipeline creation works fine even if the cache is disabled.
+// Note: This tests needs to use more than 1 device since the frontend cache on each device
+//   will prevent going out to the blob cache.
+TEST_P(SinglePipelineCachingTests, RenderPipelineNoCache) {
+    mMockCache.Disable();
+
+    // First time should create and since cache is disabled, it should not write out to the
+    // cache.
+    {
+        wgpu::Device device = CreateDevice();
+        utils::ComboRenderPipelineDescriptor desc;
+        desc.cTargets[0].writeMask = wgpu::ColorWriteMask::None;
+        desc.vertex.module = utils::CreateShaderModule(device, kVertexShader.data());
+        desc.vertex.entryPoint = "main";
+        desc.cFragment.module = utils::CreateShaderModule(device, kFragmentShader.data());
+        desc.cFragment.entryPoint = "main";
+        EXPECT_CACHE_HIT(mMockCache, 0u, device.CreateRenderPipeline(&desc));
+    }
+    EXPECT_EQ(mMockCache.GetNumEntries(), 0u);
+
+    // Second time should create fine with no cache hits since cache is disabled.
+    {
+        wgpu::Device device = CreateDevice();
+        utils::ComboRenderPipelineDescriptor desc;
+        desc.cTargets[0].writeMask = wgpu::ColorWriteMask::None;
+        desc.vertex.module = utils::CreateShaderModule(device, kVertexShader.data());
+        desc.vertex.entryPoint = "main";
+        desc.cFragment.module = utils::CreateShaderModule(device, kFragmentShader.data());
+        desc.cFragment.entryPoint = "main";
+        EXPECT_CACHE_HIT(mMockCache, 0u, device.CreateRenderPipeline(&desc));
+    }
+    EXPECT_EQ(mMockCache.GetNumEntries(), 0u);
+}
+
+// Tests that pipeline creation on the same device uses frontend cache when possible.
+TEST_P(SinglePipelineCachingTests, RenderPipelineFrontedCache) {
+    utils::ComboRenderPipelineDescriptor desc;
+    desc.cTargets[0].writeMask = wgpu::ColorWriteMask::None;
+    desc.vertex.module = utils::CreateShaderModule(device, kVertexShader.data());
+    desc.vertex.entryPoint = "main";
+    desc.cFragment.module = utils::CreateShaderModule(device, kFragmentShader.data());
+    desc.cFragment.entryPoint = "main";
+
+    // First creation should create a cache entry.
+    wgpu::RenderPipeline pipeline;
+    EXPECT_CACHE_HIT(mMockCache, 0u, pipeline = device.CreateRenderPipeline(&desc));
+    EXPECT_EQ(mMockCache.GetNumEntries(), 1u);
+
+    // Second creation on the same device should just return from frontend cache and should not
+    // call out to the blob cache.
+    EXPECT_CALL(mMockCache, LoadData).Times(0);
+    wgpu::RenderPipeline samePipeline;
+    EXPECT_CACHE_HIT(mMockCache, 0u, samePipeline = device.CreateRenderPipeline(&desc));
+    EXPECT_EQ(pipeline.Get() == samePipeline.Get(), !UsesWire());
+}
+
+// Tests that pipeline creation hits the cache when it is enabled.
+// Note: This test needs to use more than 1 device since the frontend cache on each device
+//   will prevent going out to the blob cache.
+TEST_P(SinglePipelineCachingTests, RenderPipelineBlobCache) {
+    // First time should create and write out to the cache.
+    {
+        wgpu::Device device = CreateDevice();
+        utils::ComboRenderPipelineDescriptor desc;
+        desc.cTargets[0].writeMask = wgpu::ColorWriteMask::None;
+        desc.vertex.module = utils::CreateShaderModule(device, kVertexShader.data());
+        desc.vertex.entryPoint = "main";
+        desc.cFragment.module = utils::CreateShaderModule(device, kFragmentShader.data());
+        desc.cFragment.entryPoint = "main";
+        EXPECT_CACHE_HIT(mMockCache, 0u, device.CreateRenderPipeline(&desc));
+    }
+    EXPECT_EQ(mMockCache.GetNumEntries(), 1u);
+
+    // Second time should create using the cache.
+    {
+        wgpu::Device device = CreateDevice();
+        utils::ComboRenderPipelineDescriptor desc;
+        desc.cTargets[0].writeMask = wgpu::ColorWriteMask::None;
+        desc.vertex.module = utils::CreateShaderModule(device, kVertexShader.data());
+        desc.vertex.entryPoint = "main";
+        desc.cFragment.module = utils::CreateShaderModule(device, kFragmentShader.data());
+        desc.cFragment.entryPoint = "main";
+        EXPECT_CACHE_HIT(mMockCache, 1u, device.CreateRenderPipeline(&desc));
+    }
+    EXPECT_EQ(mMockCache.GetNumEntries(), 1u);
+}
+
+// Tests that pipeline creation does not hits the cache when it is enabled but we use different
+// isolation keys.
+TEST_P(SinglePipelineCachingTests, RenderPipelineBlobCacheIsolationKey) {
+    // First time should create and write out to the cache.
+    {
+        wgpu::Device device = CreateDevice("isolation key 1");
+        utils::ComboRenderPipelineDescriptor desc;
+        desc.cTargets[0].writeMask = wgpu::ColorWriteMask::None;
+        desc.vertex.module = utils::CreateShaderModule(device, kVertexShader.data());
+        desc.vertex.entryPoint = "main";
+        desc.cFragment.module = utils::CreateShaderModule(device, kFragmentShader.data());
+        desc.cFragment.entryPoint = "main";
+        EXPECT_CACHE_HIT(mMockCache, 0u, device.CreateRenderPipeline(&desc));
+    }
+    EXPECT_EQ(mMockCache.GetNumEntries(), 1u);
+
+    // Second time should also create and write out to the cache.
+    {
+        wgpu::Device device = CreateDevice("isolation key 2");
+        utils::ComboRenderPipelineDescriptor desc;
+        desc.cTargets[0].writeMask = wgpu::ColorWriteMask::None;
+        desc.vertex.module = utils::CreateShaderModule(device, kVertexShader.data());
+        desc.vertex.entryPoint = "main";
+        desc.cFragment.module = utils::CreateShaderModule(device, kFragmentShader.data());
+        desc.cFragment.entryPoint = "main";
+        EXPECT_CACHE_HIT(mMockCache, 0u, device.CreateRenderPipeline(&desc));
+    }
+    EXPECT_EQ(mMockCache.GetNumEntries(), 2u);
+}
+
+DAWN_INSTANTIATE_TEST(SinglePipelineCachingTests, VulkanBackend());
+
+}  // namespace

src/dawn/tests/white_box/VulkanErrorInjectorTests.cpp

@@ -21,8 +21,9 @@
 #include "dawn/native/vulkan/DeviceVk.h"
 #include "dawn/native/vulkan/VulkanError.h"
 
-namespace {
+namespace dawn::native::vulkan {
 
+namespace {
 class VulkanErrorInjectorTests : public DawnTest {
   public:
     void SetUp() override {
@@ -35,7 +36,6 @@ class VulkanErrorInjectorTests : public DawnTest {
   protected:
     dawn::native::vulkan::Device* mDeviceVk;
 };
-
 }  // anonymous namespace
 
 TEST_P(VulkanErrorInjectorTests, InjectErrorOnCreateBuffer) {
@@ -49,15 +49,15 @@ TEST_P(VulkanErrorInjectorTests, InjectErrorOnCreateBuffer) {
     {
         VkBuffer buffer = VK_NULL_HANDLE;
         EXPECT_EQ(
-            mDeviceVk->fn.CreateBuffer(mDeviceVk->GetVkDevice(), &createInfo, nullptr, &buffer),
+            mDeviceVk->fn.CreateBuffer(mDeviceVk->GetVkDevice(), &createInfo, nullptr, &*buffer),
             VK_SUCCESS);
         mDeviceVk->fn.DestroyBuffer(mDeviceVk->GetVkDevice(), buffer, nullptr);
     }
 
     auto CreateTestBuffer = [&]() -> bool {
         VkBuffer buffer = VK_NULL_HANDLE;
-        dawn::native::MaybeError err = CheckVkSuccess(
-            mDeviceVk->fn.CreateBuffer(mDeviceVk->GetVkDevice(), &createInfo, nullptr, &buffer),
+        MaybeError err = CheckVkSuccess(
+            mDeviceVk->fn.CreateBuffer(mDeviceVk->GetVkDevice(), &createInfo, nullptr, &*buffer),
             "vkCreateBuffer");
         if (err.IsError()) {
             // The handle should never be written to, even for mock failures.
@@ -78,47 +78,49 @@ TEST_P(VulkanErrorInjectorTests, InjectErrorOnCreateBuffer) {
         EXPECT_TRUE(CreateTestBuffer());
 
         // The error injector call count should be empty
-        EXPECT_EQ(dawn::native::AcquireErrorInjectorCallCount(), 0u);
+        EXPECT_EQ(AcquireErrorInjectorCallCount(), 0u);
     }
 
     // Test error injection works.
-    dawn::native::EnableErrorInjector();
+    EnableErrorInjector();
     {
         EXPECT_TRUE(CreateTestBuffer());
         EXPECT_TRUE(CreateTestBuffer());
 
         // The error injector call count should be two.
-        EXPECT_EQ(dawn::native::AcquireErrorInjectorCallCount(), 2u);
+        EXPECT_EQ(AcquireErrorInjectorCallCount(), 2u);
 
         // Inject an error at index 0. The first should fail, the second succeed.
         {
-            dawn::native::InjectErrorAt(0u);
+            InjectErrorAt(0u);
             EXPECT_FALSE(CreateTestBuffer());
             EXPECT_TRUE(CreateTestBuffer());
 
-            dawn::native::ClearErrorInjector();
+            ClearErrorInjector();
         }
 
         // Inject an error at index 1. The second should fail, the first succeed.
         {
-            dawn::native::InjectErrorAt(1u);
+            InjectErrorAt(1u);
             EXPECT_TRUE(CreateTestBuffer());
             EXPECT_FALSE(CreateTestBuffer());
 
-            dawn::native::ClearErrorInjector();
+            ClearErrorInjector();
         }
 
         // Inject an error and then clear the injector. Calls should be successful.
         {
-            dawn::native::InjectErrorAt(0u);
-            dawn::native::DisableErrorInjector();
+            InjectErrorAt(0u);
+            DisableErrorInjector();
 
             EXPECT_TRUE(CreateTestBuffer());
             EXPECT_TRUE(CreateTestBuffer());
 
-            dawn::native::ClearErrorInjector();
+            ClearErrorInjector();
         }
     }
 }
 
 DAWN_INSTANTIATE_TEST(VulkanErrorInjectorTests, VulkanBackend());
+
+}  // namespace dawn::native::vulkan