[WebGPU] Index buffer can read outside the range of a vertex buffer

https://bugs.webkit.org/show_bug.cgi?id=273828
<radar://127672770>

Reviewed by Tadeu Zagallo.

Prevent out of bounds accesses into vertex buffers via drawIndexed,
drawIndirect, and drawIndexedIndirect by running a non-rasterizing
vertex shader immedietly before the draw call.

ICB path rewrites the ICB render commands to avoid out of bounds
accesses.

* Source/WebGPU/WebGPU/BindableResource.h:
* Source/WebGPU/WebGPU/Buffer.h:
* Source/WebGPU/WebGPU/Buffer.mm:
(WebGPU::Buffer::Buffer):
(WebGPU::Buffer::maxIndex const):
(WebGPU::Buffer::indirectBuffer const):
(WebGPU::Buffer::indirectIndexedBuffer const):
(WebGPU::Buffer::indirectBufferRequiresRecomputation const):
(WebGPU::Buffer::indirectBufferRecomputed):
(WebGPU::Buffer::indirectBufferInvalidated):
(WebGPU::Buffer::recomputeMaxIndexValues const): Deleted.
* Source/WebGPU/WebGPU/CommandEncoder.h:
* Source/WebGPU/WebGPU/CommandEncoder.mm:
(WebGPU::CommandEncoder::copyBufferToBuffer):
(WebGPU::CommandEncoder::copyTextureToBuffer):
(WebGPU::CommandEncoder::clearBuffer):
* Source/WebGPU/WebGPU/ComputePassEncoder.mm:
(WebGPU::addResourceToActiveResources):
(WebGPU::ComputePassEncoder::runPredispatchIndirectCallValidation):
(WebGPU::setCommandEncoder):
* Source/WebGPU/WebGPU/Device.h:
* Source/WebGPU/WebGPU/Device.mm:
(WebGPU::GPUFrameCapture::captureFrame):
(WebGPU::Device::dispatchCallBuffer):
(WebGPU::Device::dispatchCallPipelineState):
(WebGPU::Device::copyIndexIndirectArgsPipeline):
(wgpuDeviceReference): Deleted.
(wgpuDeviceRelease): Deleted.
(wgpuDeviceCreateBindGroup): Deleted.
(wgpuDeviceCreateBindGroupLayout): Deleted.
(wgpuDeviceCreateBuffer): Deleted.
(wgpuDeviceCreateCommandEncoder): Deleted.
(wgpuDeviceCreateComputePipeline): Deleted.
(wgpuDeviceCreateComputePipelineAsync): Deleted.
(wgpuDeviceCreateComputePipelineAsyncWithBlock): Deleted.
(wgpuDeviceCreatePipelineLayout): Deleted.
(wgpuDeviceCreateQuerySet): Deleted.
(wgpuDeviceCreateRenderBundleEncoder): Deleted.
(wgpuDeviceCreateRenderPipeline): Deleted.
(wgpuDeviceCreateRenderPipelineAsync): Deleted.
(wgpuDeviceCreateRenderPipelineAsyncWithBlock): Deleted.
(wgpuDeviceCreateSampler): Deleted.
(wgpuDeviceImportExternalTexture): Deleted.
(wgpuDeviceCreateShaderModule): Deleted.
(wgpuDeviceCreateSwapChain): Deleted.
(wgpuDeviceCreateTexture): Deleted.
(wgpuDeviceDestroy): Deleted.
(wgpuDeviceEnumerateFeatures): Deleted.
(wgpuDeviceGetLimits): Deleted.
(wgpuDeviceGetQueue): Deleted.
(wgpuDeviceHasFeature): Deleted.
(wgpuDevicePopErrorScope): Deleted.
(wgpuDevicePopErrorScopeWithBlock): Deleted.
(wgpuDevicePushErrorScope): Deleted.
(wgpuDeviceSetDeviceLostCallback): Deleted.
(wgpuDeviceSetDeviceLostCallbackWithBlock): Deleted.
(wgpuDeviceSetUncapturedErrorCallback): Deleted.
(wgpuDeviceSetUncapturedErrorCallbackWithBlock): Deleted.
(wgpuDeviceSetLabel): Deleted.
* Source/WebGPU/WebGPU/Pipeline.mm:
(WebGPU::validateBindGroup):
* Source/WebGPU/WebGPU/RenderBundleEncoder.h:
* Source/WebGPU/WebGPU/RenderBundleEncoder.mm:
(-[RenderBundleICBWithResources initWithICB:containerBuffer:pipelineState:depthStencilState:cullMode:frontFace:depthClipMode:depthBias:depthBiasSlopeScale:depthBiasClamp:fragmentDynamicOffsetsBuffer:pipeline:]):
(-[RenderBundleICBWithResources minVertexCountForDrawCommand]):
(WebGPU::makeRenderBundleICBWithResources):
(WebGPU::RenderBundleEncoder::addResource):
(WebGPU::RenderBundleEncoder::computeMininumVertexCount const):
(WebGPU::RenderBundleEncoder::storeVertexBufferCountsForValidation):
(WebGPU::RenderBundleEncoder::drawIndexed):
(WebGPU::RenderBundleEncoder::drawIndexedIndirect):
(WebGPU::RenderBundleEncoder::drawIndirect):
(WebGPU::RenderBundleEncoder::endCurrentICB):
(WebGPU::RenderBundleEncoder::setIndexBuffer):
(WebGPU::RenderBundleEncoder::setVertexBuffer):
(-[RenderBundleICBWithResources initWithICB:pipelineState:depthStencilState:cullMode:frontFace:depthClipMode:depthBias:depthBiasSlopeScale:depthBiasClamp:fragmentDynamicOffsetsBuffer:pipeline:]): Deleted.
* Source/WebGPU/WebGPU/RenderPassEncoder.h:
* Source/WebGPU/WebGPU/RenderPassEncoder.mm:
(WebGPU::m_maxDrawCount):
(WebGPU::RenderPassEncoder::addResourceToActiveResources):
(WebGPU::RenderPassEncoder::computeMininumVertexCount const):
(WebGPU::RenderPassEncoder::clampIndexBufferToValidValues):
(WebGPU::RenderPassEncoder::clampIndirectIndexBufferToValidValues):
(WebGPU::RenderPassEncoder::clampIndirectBufferToValidValues):
(WebGPU::RenderPassEncoder::drawIndexed):
(WebGPU::RenderPassEncoder::drawIndexedIndirect):
(WebGPU::RenderPassEncoder::drawIndirect):
(WebGPU::RenderPassEncoder::setCommandEncoder):
(WebGPU::RenderPassEncoder::executeBundles):

Canonical link: https://commits.webkit.org/279182@main

Source/WebGPU/WebGPU/BindableResource.h

@@ -29,6 +29,7 @@
 #import <wtf/OptionSet.h>
 #import <wtf/RefPtr.h>
 #import <wtf/Vector.h>
+#import <wtf/WeakPtr.h>
 
 namespace WebGPU {
 
@@ -58,7 +59,7 @@ static constexpr auto isTextureBindGroupEntryUsage(OptionSet<BindGroupEntryUsage
 struct BindGroupEntryUsageData {
     OptionSet<BindGroupEntryUsage> usage { BindGroupEntryUsage::Undefined };
     uint32_t binding { 0 };
-    using Resource = std::variant<RefPtr<const Buffer>, RefPtr<const TextureView>, RefPtr<const ExternalTexture>>;
+    using Resource = std::variant<RefPtr<Buffer>, RefPtr<const TextureView>, RefPtr<const ExternalTexture>>;
     Resource resource;
     static constexpr uint32_t invalidBindingIndex = INT_MAX;
     static constexpr BindGroupEntryUsage invalidBindGroupUsage = static_cast<BindGroupEntryUsage>(std::numeric_limits<std::underlying_type<BindGroupEntryUsage>::type>::max());
@@ -71,4 +72,23 @@ struct BindableResources {
     MTLRenderStages renderStages;
 };
 
+struct IndexData {
+    uint64_t renderCommand { 0 };
+    uint32_t minVertexCount { UINT32_MAX };
+    uint64_t bufferGpuAddress { 0 };
+    uint32_t indexCount { 0 };
+    uint32_t instanceCount { 0 };
+    uint32_t firstIndex { 0 };
+    int32_t baseVertex { 0 };
+    uint32_t firstInstance { 0 };
+    MTLPrimitiveType primitiveType { MTLPrimitiveTypeTriangle };
+};
+
+struct IndexBufferAndIndexData {
+    WeakPtr<Buffer> indexBuffer;
+    MTLIndexType indexType { MTLIndexTypeUInt16 };
+    NSUInteger indexBufferOffsetInBytes { 0 };
+    IndexData indexData;
+};
+
 } // namespace WebGPU

Source/WebGPU/WebGPU/Buffer.h

@@ -82,6 +82,8 @@ class Buffer : public WGPUBufferImpl, public RefCounted<Buffer>, public CanMakeW
     };
 
     id<MTLBuffer> buffer() const { return m_buffer; }
+    id<MTLBuffer> indirectBuffer() const;
+    id<MTLBuffer> indirectIndexedBuffer() const;
     uint64_t initialSize() const;
     uint64_t currentSize() const;
     WGPUBufferUsageFlags usage() const { return m_usage; }
@@ -92,18 +94,22 @@ class Buffer : public WGPUBufferImpl, public RefCounted<Buffer>, public CanMakeW
     void setCommandEncoder(CommandEncoder&, bool mayModifyBuffer = false) const;
     uint32_t maxIndex(MTLIndexType) const;
     uint8_t* getBufferContents();
+    bool indirectBufferRequiresRecomputation(uint32_t baseIndex, uint32_t indexCount, uint32_t minVertexCount, MTLIndexType) const;
+    void indirectBufferRecomputed(uint32_t baseIndex, uint32_t indexCount, uint32_t minVertexCount, MTLIndexType);
+    void indirectBufferInvalidated();
 
 private:
     Buffer(id<MTLBuffer>, uint64_t initialSize, WGPUBufferUsageFlags, State initialState, MappingRange initialMappingRange, Device&);
     Buffer(Device&);
-    void recomputeMaxIndexValues() const;
 
     bool validateGetMappedRange(size_t offset, size_t rangeSize) const;
     NSString* errorValidatingMapAsync(WGPUMapModeFlags, size_t offset, size_t rangeSize) const;
     bool validateUnmap() const;
     void setState(State);
 
     id<MTLBuffer> m_buffer { nil };
+    id<MTLBuffer> m_indirectBuffer { nil };
+    id<MTLBuffer> m_indirectIndexedBuffer { nil };
 
     // https://gpuweb.github.io/gpuweb/#buffer-interface
 
@@ -115,6 +121,12 @@ class Buffer : public WGPUBufferImpl, public RefCounted<Buffer>, public CanMakeW
     using MappedRanges = RangeSet<Range<size_t>>;
     MappedRanges m_mappedRanges;
     WGPUMapModeFlags m_mapMode { WGPUMapMode_None };
+    struct IndirectArgsCache {
+        uint32_t lastBaseIndex { 0 };
+        uint32_t indexCount { 0 };
+        uint32_t minVertexCount { 0 };
+        MTLIndexType indexType { MTLIndexTypeUInt16 };
+    } m_indirectCache;
 
     const Ref<Device> m_device;
     mutable WeakPtr<CommandEncoder> m_commandEncoder;

Source/WebGPU/WebGPU/Buffer.mm

@@ -143,6 +143,10 @@ static MTLStorageMode storageMode(bool deviceHasUnifiedMemory, WGPUBufferUsageFl
     , m_mappingRange(initialMappingRange)
     , m_device(device)
 {
+    if (m_usage & WGPUBufferUsage_Indirect)
+        m_indirectBuffer = device.safeCreateBuffer(sizeof(MTLDrawPrimitivesIndirectArguments), MTLStorageModePrivate);
+    if (m_usage & (WGPUBufferUsage_Indirect | WGPUBufferUsage_Index))
+        m_indirectIndexedBuffer = device.safeCreateBuffer(sizeof(MTLDrawIndexedPrimitivesIndirectArguments), MTLStorageModePrivate);
 }
 
 Buffer::Buffer(Device& device)
@@ -386,29 +390,41 @@ static size_t computeRangeSize(uint64_t size, size_t offset)
     return state() == State::Destroyed;
 }
 
-void Buffer::recomputeMaxIndexValues() const
+uint32_t Buffer::maxIndex(MTLIndexType indexType) const
 {
-    if (!(m_usage & WGPUBufferUsage_Index))
-        return;
+    ASSERT(m_usage & WGPUBufferUsage_Index);
+    return indexType == MTLIndexTypeUInt16 ? m_max16BitIndex : m_max32BitIndex;
+}
 
-    NSUInteger lengthInBytes = m_buffer.length;
-    auto bufferPtr = static_cast<uint8_t*>(m_buffer.contents);
-    RELEASE_ASSERT(bufferPtr);
-    m_max16BitIndex = 0;
-    m_max32BitIndex = 0;
-    uint8_t* bufferEnd = bufferPtr + lengthInBytes;
-    for (; (bufferPtr += sizeof(uint32_t)) <= bufferEnd; bufferPtr += sizeof(uint32_t)) {
-        m_max32BitIndex = std::max(*reinterpret_cast<uint32_t*>(bufferPtr), m_max32BitIndex);
-        m_max16BitIndex = std::max(*(reinterpret_cast<uint16_t*>(bufferPtr) + 1), std::max(*reinterpret_cast<uint16_t*>(bufferPtr), m_max16BitIndex));
-    }
-    if (bufferPtr + sizeof(uint16_t) <= bufferEnd)
-        m_max16BitIndex = std::max(*reinterpret_cast<uint16_t*>(bufferPtr), m_max16BitIndex);
+id<MTLBuffer> Buffer::indirectBuffer() const
+{
+    return m_indirectBuffer;
 }
 
-uint32_t Buffer::maxIndex(MTLIndexType indexType) const
+id<MTLBuffer> Buffer::indirectIndexedBuffer() const
 {
-    ASSERT(m_usage & WGPUBufferUsage_Index);
-    return indexType == MTLIndexTypeUInt16 ? m_max16BitIndex : m_max32BitIndex;
+    return m_indirectIndexedBuffer;
+}
+
+bool Buffer::indirectBufferRequiresRecomputation(uint32_t baseIndex, uint32_t indexCount, uint32_t minVertexCount, MTLIndexType indexType) const
+{
+    auto rangeBegin = m_indirectCache.lastBaseIndex;
+    auto rangeEnd = m_indirectCache.lastBaseIndex + m_indirectCache.indexCount;
+    auto newRangeEnd = baseIndex + indexCount;
+    return baseIndex < rangeBegin || newRangeEnd > rangeEnd || minVertexCount > m_indirectCache.minVertexCount || m_indirectCache.indexType != indexType;
+}
+
+void Buffer::indirectBufferRecomputed(uint32_t baseIndex, uint32_t indexCount, uint32_t minVertexCount, MTLIndexType indexType)
+{
+    m_indirectCache.lastBaseIndex = baseIndex;
+    m_indirectCache.indexCount = indexCount;
+    m_indirectCache.minVertexCount = minVertexCount;
+    m_indirectCache.indexType = indexType;
+}
+
+void Buffer::indirectBufferInvalidated()
+{
+    indirectBufferRecomputed(0, 0, 0, MTLIndexTypeUInt16);
 }
 
 } // namespace WebGPU

Source/WebGPU/WebGPU/CommandEncoder.h

@@ -70,11 +70,11 @@ class CommandEncoder : public WGPUCommandEncoderImpl, public RefCounted<CommandE
 
     Ref<ComputePassEncoder> beginComputePass(const WGPUComputePassDescriptor&);
     Ref<RenderPassEncoder> beginRenderPass(const WGPURenderPassDescriptor&);
-    void copyBufferToBuffer(const Buffer& source, uint64_t sourceOffset, const Buffer& destination, uint64_t destinationOffset, uint64_t size);
+    void copyBufferToBuffer(const Buffer& source, uint64_t sourceOffset, Buffer& destination, uint64_t destinationOffset, uint64_t size);
     void copyBufferToTexture(const WGPUImageCopyBuffer& source, const WGPUImageCopyTexture& destination, const WGPUExtent3D& copySize);
     void copyTextureToBuffer(const WGPUImageCopyTexture& source, const WGPUImageCopyBuffer& destination, const WGPUExtent3D& copySize);
     void copyTextureToTexture(const WGPUImageCopyTexture& source, const WGPUImageCopyTexture& destination, const WGPUExtent3D& copySize);
-    void clearBuffer(const Buffer&, uint64_t offset, uint64_t size);
+    void clearBuffer(Buffer&, uint64_t offset, uint64_t size);
     Ref<CommandBuffer> finish(const WGPUCommandBufferDescriptor&);
     void insertDebugMarker(String&& markerLabel);
     void popDebugGroup();

Source/WebGPU/WebGPU/CommandEncoder.mm

@@ -711,7 +711,7 @@ static bool isRenderableTextureView(const TextureView& texture)
         m_cachedCommandBuffer->setBufferMapCount(m_bufferMapCount);
 }
 
-void CommandEncoder::copyBufferToBuffer(const Buffer& source, uint64_t sourceOffset, const Buffer& destination, uint64_t destinationOffset, uint64_t size)
+void CommandEncoder::copyBufferToBuffer(const Buffer& source, uint64_t sourceOffset, Buffer& destination, uint64_t destinationOffset, uint64_t size)
 {
     // https://gpuweb.github.io/gpuweb/#dom-gpucommandencoder-copybuffertobuffer
     if (!prepareTheEncoderState()) {
@@ -726,6 +726,7 @@ static bool isRenderableTextureView(const TextureView& texture)
 
     source.setCommandEncoder(*this);
     destination.setCommandEncoder(*this);
+    destination.indirectBufferInvalidated();
     if (!size || source.isDestroyed() || destination.isDestroyed())
         return;
 
@@ -1261,6 +1262,7 @@ static bool hasValidDimensions(WGPUTextureDimension dimension, NSUInteger width,
     auto& apiDestinationBuffer = fromAPI(destination.buffer);
     sourceTexture.setCommandEncoder(*this);
     apiDestinationBuffer.setCommandEncoder(*this);
+    apiDestinationBuffer.indirectBufferInvalidated();
     if (sourceTexture.isDestroyed() || apiDestinationBuffer.isDestroyed())
         return;
 
@@ -1673,7 +1675,7 @@ static bool areCopyCompatible(WGPUTextureFormat format1, WGPUTextureFormat forma
     return true;
 }
 
-void CommandEncoder::clearBuffer(const Buffer& buffer, uint64_t offset, uint64_t size)
+void CommandEncoder::clearBuffer(Buffer& buffer, uint64_t offset, uint64_t size)
 {
     // https://gpuweb.github.io/gpuweb/#dom-gpucommandencoder-clearbuffer
 
@@ -1697,6 +1699,7 @@ static bool areCopyCompatible(WGPUTextureFormat format1, WGPUTextureFormat forma
     }
 
     buffer.setCommandEncoder(*this);
+    buffer.indirectBufferInvalidated();
     auto range = NSMakeRange(static_cast<NSUInteger>(offset), static_cast<NSUInteger>(size));
     if (buffer.isDestroyed() || !size || NSMaxRange(range) > buffer.buffer().length)
         return;

Source/WebGPU/WebGPU/ComputePassEncoder.mm

@@ -137,9 +137,12 @@ static bool addResourceToActiveResources(const TextureView& texture, OptionSet<B
 
 static bool addResourceToActiveResources(const BindGroupEntryUsageData::Resource& resource, id<MTLResource> mtlResource, OptionSet<BindGroupEntryUsage> resourceUsage, BindGroupId bindGroup, EntryMapContainer& usagesForResource)
 {
-    return WTF::switchOn(resource, [&](const RefPtr<const Buffer>& buffer) {
-        if (buffer.get())
+    return WTF::switchOn(resource, [&](const RefPtr<Buffer>& buffer) {
+        if (buffer.get()) {
+            if (resourceUsage.contains(BindGroupEntryUsage::Storage))
+                buffer->indirectBufferInvalidated();
             return addResourceToActiveResources(buffer.get(), buffer->buffer(), resourceUsage, usagesForResource, bindGroup);
+        }
         return true;
     }, [&](const RefPtr<const TextureView>& textureView) {
         if (textureView.get())
@@ -251,27 +254,26 @@ static bool addResourceToActiveResources(const BindGroupEntryUsageData::Resource
 
 id<MTLBuffer> ComputePassEncoder::runPredispatchIndirectCallValidation(const Buffer& indirectBuffer, uint64_t indirectOffset)
 {
-    static id<MTLComputePipelineState> computePipelineState = nil;
+    static id<MTLFunction> function = nil;
     id<MTLDevice> device = m_device->device();
-    if (!computePipelineState) {
+    if (!function) {
         auto dimensionMax = m_device->limits().maxComputeWorkgroupsPerDimension;
-        NSError *error = nil;
         MTLCompileOptions* options = [MTLCompileOptions new];
         ALLOW_DEPRECATED_DECLARATIONS_BEGIN
         options.fastMathEnabled = YES;
         ALLOW_DEPRECATED_DECLARATIONS_END
+        NSError *error = nil;
         id<MTLLibrary> library = [device newLibraryWithSource:[NSString stringWithFormat:@"[[kernel]] void cs(device const uint* indirectBuffer, device uint* dispatchCallBuffer, uint index [[thread_position_in_grid]]) { dispatchCallBuffer[index] = metal::select(indirectBuffer[index], 0u, indirectBuffer[index] > %u); }", dimensionMax] options:options error:&error];
         if (error)
             return nil;
 
-        id<MTLFunction> function = [library newFunctionWithName:@"cs"];
-        computePipelineState = [device newComputePipelineStateWithFunction:function error:&error];
-
+        function = [library newFunctionWithName:@"cs"];
         if (error)
             return nil;
     }
 
-    static id<MTLBuffer> dispatchCallBuffer = [device newBufferWithLength:sizeof(MTLDispatchThreadgroupsIndirectArguments) options:MTLResourceStorageModePrivate];
+    id<MTLComputePipelineState> computePipelineState = m_device->dispatchCallPipelineState(function);
+    id<MTLBuffer> dispatchCallBuffer = m_device->dispatchCallBuffer();
     [computeCommandEncoder() setComputePipelineState:computePipelineState];
     [computeCommandEncoder() setBuffer:indirectBuffer.buffer() offset:indirectOffset atIndex:0];
     [computeCommandEncoder() setBuffer:dispatchCallBuffer offset:0 atIndex:1];
@@ -393,7 +395,7 @@ static bool addResourceToActiveResources(const BindGroupEntryUsageData::Resource
 
 static void setCommandEncoder(const BindGroupEntryUsageData::Resource& resource, CommandEncoder& parentEncoder)
 {
-    WTF::switchOn(resource, [&](const RefPtr<const Buffer>& buffer) {
+    WTF::switchOn(resource, [&](const RefPtr<Buffer>& buffer) {
         if (buffer)
             buffer->setCommandEncoder(parentEncoder);
         }, [&](const RefPtr<const TextureView>& textureView) {

Source/WebGPU/WebGPU/Device.h

@@ -135,6 +135,14 @@ class Device : public WGPUDeviceImpl, public ThreadSafeRefCountedAndCanMakeThrea
     id<MTLTexture> placeholderTexture(WGPUTextureFormat) const;
     bool isDestroyed() const;
     NSString *errorValidatingTextureCreation(const WGPUTextureDescriptor&, const Vector<WGPUTextureFormat>& viewFormats);
+    id<MTLBuffer> dispatchCallBuffer();
+    id<MTLComputePipelineState> dispatchCallPipelineState(id<MTLFunction>);
+    id<MTLRenderPipelineState> indexBufferClampPipeline(MTLIndexType, NSUInteger rasterSampleCount);
+    id<MTLRenderPipelineState> indexedIndirectBufferClampPipeline(NSUInteger rasterSampleCount);
+    id<MTLRenderPipelineState> indirectBufferClampPipeline(NSUInteger rasterSampleCount);
+    id<MTLRenderPipelineState> icbCommandClampPipeline(MTLIndexType, NSUInteger rasterSampleCount);
+    id<MTLRenderPipelineState> copyIndexIndirectArgsPipeline(NSUInteger rasterSampleCount);
+    id<MTLBuffer> safeCreateBuffer(NSUInteger length, MTLStorageMode, MTLCPUCacheMode = MTLCPUCacheModeDefaultCache, MTLHazardTrackingMode = MTLHazardTrackingModeDefault) const;
 
 private:
     Device(id<MTLDevice>, id<MTLCommandQueue> defaultQueue, HardwareCapabilities&&, Adapter&);
@@ -143,7 +151,6 @@ class Device : public WGPUDeviceImpl, public ThreadSafeRefCountedAndCanMakeThrea
     struct ErrorScope;
     ErrorScope* currentErrorScope(WGPUErrorFilter);
     std::optional<WGPUErrorType> validatePopErrorScope() const;
-    id<MTLBuffer> safeCreateBuffer(NSUInteger length, MTLStorageMode, MTLCPUCacheMode = MTLCPUCacheModeDefaultCache, MTLHazardTrackingMode = MTLHazardTrackingModeDefault) const;
     bool validateCreateIOSurfaceBackedTexture(const WGPUTextureDescriptor&, const Vector<WGPUTextureFormat>& viewFormats, IOSurfaceRef backing);
 
     bool validateRenderPipeline(const WGPURenderPipelineDescriptor&);
@@ -188,6 +195,27 @@ class Device : public WGPUDeviceImpl, public ThreadSafeRefCountedAndCanMakeThrea
     id<MTLBuffer> m_placeholderBuffer { nil };
     id<MTLTexture> m_placeholderTexture { nil };
     id<MTLTexture> m_placeholderDepthStencilTexture { nil };
+    id<MTLBuffer> m_dispatchCallBuffer { nil };
+    id<MTLComputePipelineState> m_dispatchCallPipelineState { nil };
+
+    id<MTLRenderPipelineState> m_indexBufferClampUintPSO { nil };
+    id<MTLRenderPipelineState> m_indexBufferClampUshortPSO { nil };
+    id<MTLRenderPipelineState> m_indexBufferClampUintPSOMS { nil };
+    id<MTLRenderPipelineState> m_indexBufferClampUshortPSOMS { nil };
+
+    id<MTLRenderPipelineState> m_indexedIndirectBufferClampPSO { nil };
+    id<MTLRenderPipelineState> m_indexedIndirectBufferClampPSOMS { nil };
+
+    id<MTLRenderPipelineState> m_indirectBufferClampPSO { nil };
+    id<MTLRenderPipelineState> m_indirectBufferClampPSOMS { nil };
+
+    id<MTLRenderPipelineState> m_icbCommandClampUintPSO { nil };
+    id<MTLRenderPipelineState> m_icbCommandClampUshortPSO { nil };
+    id<MTLRenderPipelineState> m_icbCommandClampUintPSOMS { nil };
+    id<MTLRenderPipelineState> m_icbCommandClampUshortPSOMS { nil };
+
+    id<MTLRenderPipelineState> m_copyIndexedIndirectArgsPSO { nil };
+    id<MTLRenderPipelineState> m_copyIndexedIndirectArgsPSOMS { nil };
 
     const Ref<Adapter> m_adapter;
 #if HAVE(COREVIDEO_METAL_SUPPORT)