no message

cocos/renderer/gfx-metal/MTLBuffer.h

@@ -0,0 +1,84 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#include "gfx-base/GFXBuffer.h"
+
+#import <Metal/MTLBuffer.h>
+#import <Metal/MTLRenderCommandEncoder.h>
+#import <Metal/MTLStageInputOutputDescriptor.h>
+
+namespace cc {
+namespace gfx {
+
+class CCMTLCommandEncoder;
+struct CCMTLGPUBuffer;
+
+class CCMTLBuffer final : public Buffer {
+public:
+    explicit CCMTLBuffer();
+    ~CCMTLBuffer();
+    CCMTLBuffer(const CCMTLBuffer &) = delete;
+    CCMTLBuffer(CCMTLBuffer &&) = delete;
+    CCMTLBuffer &operator=(const CCMTLBuffer &) = delete;
+    CCMTLBuffer &operator=(CCMTLBuffer &&) = delete;
+
+    void update(const void *buffer, uint32_t offset) override;
+
+    void encodeBuffer(CCMTLCommandEncoder &encoder, uint32_t offset, uint32_t binding, ShaderStageFlags stages);
+
+    uint32_t currentOffset() const;
+    id<MTLBuffer> mtlBuffer() const;
+
+    inline CCMTLGPUBuffer *gpuBuffer() { return _gpuBuffer; }
+    inline MTLIndexType getIndexType() const { return _indexType; }
+    inline bool isDrawIndirectByIndex() const { return _isDrawIndirectByIndex; }
+    inline const DrawInfoList &getDrawInfos() const { return _drawInfos; }
+
+protected:
+    void doInit(const BufferInfo &info) override;
+    void doInit(const BufferViewInfo &info) override;
+    void doDestroy() override;
+    void doResize(uint32_t size, uint32_t count) override;
+
+    bool createMTLBuffer(uint32_t size, MemoryUsage usage);
+    void updateMTLBuffer(const void *buffer, uint32_t offset, uint32_t size);
+
+    MTLIndexType _indexType = MTLIndexTypeUInt16;
+    MTLResourceOptions _mtlResourceOptions = MTLResourceStorageModePrivate;
+    bool _isIndirectDrawSupported = false;
+    uint32_t _bufferViewOffset = 0;
+    uint8_t _lastUpdateCycle{0};
+
+    bool _isDrawIndirectByIndex = false;
+    ccstd::vector<MTLDrawIndexedPrimitivesIndirectArguments> _indexedPrimitivesIndirectArguments;
+    ccstd::vector<MTLDrawPrimitivesIndirectArguments> _primitiveIndirectArguments;
+    DrawInfoList _drawInfos;
+
+    CCMTLGPUBuffer *_gpuBuffer = nullptr;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLBuffer.mm

@@ -0,0 +1,308 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#import <Foundation/Foundation.h>
+#import <Metal/Metal.h>
+
+#include "MTLBuffer.h"
+#include "MTLCommandBuffer.h"
+#include "MTLDevice.h"
+#include "MTLQueue.h"
+#include "MTLRenderCommandEncoder.h"
+#include "MTLUtils.h"
+#include "MTLGPUObjects.h"
+#import "profiler/Profiler.h"
+#import "base/Log.h"
+
+namespace cc {
+namespace gfx {
+
+CCMTLBuffer::CCMTLBuffer() : Buffer() {
+    _typedID = generateObjectID<decltype(this)>();
+}
+
+CCMTLBuffer::~CCMTLBuffer() {
+    destroy();
+}
+
+void CCMTLBuffer::doInit(const BufferInfo &info) {
+    _gpuBuffer = ccnew CCMTLGPUBuffer;
+    _gpuBuffer->count = _count;
+    _gpuBuffer->mappedData = nullptr;
+    _gpuBuffer->instanceSize = _size;
+    _gpuBuffer->startOffset = _offset;
+    _gpuBuffer->stride = _stride;
+
+    _isIndirectDrawSupported = CCMTLDevice::getInstance()->isIndirectDrawSupported();
+    if (hasFlag(_usage, BufferUsage::INDEX)) {
+        switch (_stride) {
+            case 4: _indexType = MTLIndexTypeUInt32; break;
+            case 2: _indexType = MTLIndexTypeUInt16; break;
+            default:
+                CC_LOG_ERROR("CCMTLBuffer:: Illegal index buffer stride.");
+                break;
+        }
+    }
+
+    if (hasFlag(_usage, BufferUsageBit::VERTEX) ||
+        hasFlag(_usage, BufferUsageBit::UNIFORM) ||
+        hasFlag(_usage, BufferUsageBit::INDEX) ||
+        hasFlag(_usage, BufferUsageBit::STORAGE)) {
+        createMTLBuffer(_size, _memUsage);
+    } else if (hasFlag(_usage, BufferUsageBit::INDIRECT)) {
+        if (_isIndirectDrawSupported) {
+            createMTLBuffer(_size, _memUsage);
+            _primitiveIndirectArguments.resize(_count);
+            _indexedPrimitivesIndirectArguments.resize(_count);
+        } else {
+            _drawInfos.resize(_count);
+        }
+    }
+    CCMTLDevice::getInstance()->getMemoryStatus().bufferSize += _size;
+    CC_PROFILE_MEMORY_INC(Buffer, _size);
+}
+
+void CCMTLBuffer::doInit(const BufferViewInfo &info) {
+    auto *ccBuffer = static_cast<CCMTLBuffer *>(info.buffer);
+    _gpuBuffer = ccBuffer->gpuBuffer();
+    _indexType = ccBuffer->getIndexType();
+    _mtlResourceOptions = ccBuffer->_mtlResourceOptions;
+    _isIndirectDrawSupported = ccBuffer->_isIndirectDrawSupported;
+    _isDrawIndirectByIndex = ccBuffer->_isDrawIndirectByIndex;
+    _indexedPrimitivesIndirectArguments = ccBuffer->_indexedPrimitivesIndirectArguments;
+    _primitiveIndirectArguments = ccBuffer->_primitiveIndirectArguments;
+    _drawInfos = ccBuffer->_drawInfos;
+    _bufferViewOffset = info.offset;
+    _isBufferView = true;
+}
+
+bool CCMTLBuffer::createMTLBuffer(uint32_t size, MemoryUsage usage) {
+    if (!size) {
+        return false;
+    }
+
+    _mtlResourceOptions = mu::toMTLResourceOption(usage);
+
+    if (_gpuBuffer->mtlBuffer) {
+        id<MTLBuffer> mtlBuffer = _gpuBuffer->mtlBuffer;
+
+        std::function<void(void)> destroyFunc = [=]() {
+            if (mtlBuffer) {
+                //TODO_Zeqiang: [mac12 | ios15, ...) validate here
+                //                [mtlBuffer setPurgeableState:MTLPurgeableStateEmpty];
+                [mtlBuffer release];
+            }
+        };
+        //gpu object only
+        CCMTLGPUGarbageCollectionPool::getInstance()->collect(destroyFunc);
+    }
+
+    auto allocatedSize = size;
+    if(hasFlag(_memUsage, MemoryUsageBit::HOST)) {
+        constexpr uint8_t backBufferCount = MAX_FRAMES_IN_FLIGHT;
+        auto alignedSize = utils::alignTo(size, CCMTLDevice::getInstance()->getCapabilities().uboOffsetAlignment);
+        allocatedSize = alignedSize * backBufferCount;
+        _gpuBuffer->instanceSize = alignedSize;
+    }
+
+    id<MTLDevice> mtlDevice = id<MTLDevice>(CCMTLDevice::getInstance()->getMTLDevice());
+    _gpuBuffer->mtlBuffer = [mtlDevice newBufferWithLength:allocatedSize options:_mtlResourceOptions];
+    if (_gpuBuffer->mtlBuffer == nil) {
+        return false;
+    }
+    return true;
+}
+
+void CCMTLBuffer::doDestroy() {
+    if (_isBufferView) {
+        return;
+    }
+
+    CCMTLDevice::getInstance()->getMemoryStatus().bufferSize -= _size;
+    CC_PROFILE_MEMORY_DEC(Buffer, _size);
+
+    if (!_indexedPrimitivesIndirectArguments.empty()) {
+        _indexedPrimitivesIndirectArguments.clear();
+    }
+
+    if (!_primitiveIndirectArguments.empty()) {
+        _primitiveIndirectArguments.clear();
+    }
+
+    if (!_drawInfos.empty()) {
+        _drawInfos.clear();
+    }
+
+    if (_gpuBuffer) {
+        id<MTLBuffer> mtlBuffer = _gpuBuffer->mtlBuffer;
+        _gpuBuffer->mtlBuffer = nil;
+
+        std::function<void(void)> destroyFunc = [=]() {
+            if (mtlBuffer) {
+                //TODO_Zeqiang: [mac12 | ios15, ...) validate here
+                //                [mtlBuffer setPurgeableState:MTLPurgeableStateEmpty];
+                [mtlBuffer release];
+            }
+        };
+        //gpu object only
+        CCMTLGPUGarbageCollectionPool::getInstance()->collect(destroyFunc);
+    }
+
+    CC_SAFE_DELETE(_gpuBuffer);
+}
+
+void CCMTLBuffer::doResize(uint32_t size, uint32_t count) {
+    if (hasFlag(_usage, BufferUsageBit::VERTEX) ||
+        hasFlag(_usage, BufferUsageBit::INDEX) ||
+        hasFlag(_usage, BufferUsageBit::UNIFORM)) {
+        createMTLBuffer(size, _memUsage);
+    }
+
+    CCMTLDevice::getInstance()->getMemoryStatus().bufferSize -= _size;
+    CCMTLDevice::getInstance()->getMemoryStatus().bufferSize += size;
+    CC_PROFILE_MEMORY_DEC(Buffer, _size);
+    CC_PROFILE_MEMORY_INC(Buffer, size);
+
+    _size = size;
+    _count = count;
+    if (hasFlag(_usage, BufferUsageBit::INDIRECT)) {
+        if (_isIndirectDrawSupported) {
+            createMTLBuffer(size, _memUsage);
+            _primitiveIndirectArguments.resize(_count);
+            _indexedPrimitivesIndirectArguments.resize(_count);
+        } else {
+            _drawInfos.resize(_count);
+        }
+    }
+}
+
+void CCMTLBuffer::update(const void *buffer, uint32_t size) {
+    CC_PROFILE(CCMTLBufferUpdate);
+    if (_isBufferView) {
+        CC_LOG_WARNING("Cannot update a buffer view.");
+        return;
+    }
+
+    _isDrawIndirectByIndex = false;
+
+    if (hasFlag(_usage, BufferUsageBit::INDIRECT)) {
+        uint32_t drawInfoCount = size / _stride;
+        const auto *drawInfo = static_cast<const DrawInfo *>(buffer);
+        if (drawInfoCount > 0) {
+            if (drawInfo->indexCount) {
+                _isDrawIndirectByIndex = true;
+            }
+        }
+
+        if (_isIndirectDrawSupported) {
+            if (drawInfoCount > 0) {
+                if (_isDrawIndirectByIndex) {
+                    uint32_t stride = sizeof(MTLDrawIndexedPrimitivesIndirectArguments);
+
+                    for (uint32_t i = 0; i < drawInfoCount; ++i) {
+                        auto &arguments = _indexedPrimitivesIndirectArguments[i];
+                        arguments.indexCount = drawInfo->indexCount;
+                        arguments.instanceCount = std::max(drawInfo->instanceCount, 1U);
+                        arguments.indexStart = drawInfo->firstIndex;
+                        arguments.baseVertex = drawInfo->firstVertex;
+                        arguments.baseInstance = drawInfo->firstInstance;
+                        ++drawInfo;
+                    }
+                    updateMTLBuffer(_indexedPrimitivesIndirectArguments.data(), 0, drawInfoCount * stride);
+                } else {
+                    uint32_t stride = sizeof(MTLDrawPrimitivesIndirectArguments);
+
+                    for (uint32_t i = 0; i < drawInfoCount; ++i) {
+                        auto &arguments = _primitiveIndirectArguments[i];
+                        arguments.vertexCount = drawInfo->vertexCount;
+                        arguments.instanceCount = std::max(drawInfo->instanceCount, 1U);
+                        arguments.vertexStart = drawInfo->firstVertex;
+                        arguments.baseInstance = drawInfo->firstInstance;
+                        ++drawInfo;
+                    }
+                    updateMTLBuffer(_primitiveIndirectArguments.data(), 0, drawInfoCount * stride);
+                }
+            }
+        } else {
+            memcpy(_drawInfos.data(), buffer, size);
+        }
+    } else {
+        updateMTLBuffer(buffer, 0, size);
+    }
+}
+
+void CCMTLBuffer::updateMTLBuffer(const void *buffer, uint32_t /*offset*/, uint32_t size) {
+    id<MTLBuffer> mtlBuffer = _gpuBuffer->mtlBuffer;
+    auto* ccDevice = CCMTLDevice::getInstance();
+    if(mtlBuffer.storageMode != MTLStorageModePrivate) {
+        auto& lastUpdateCycle = _gpuBuffer->lastUpdateCycle;
+        lastUpdateCycle = ccDevice->currentFrameIndex();
+        bool backBuffer = hasFlag(_memUsage, MemoryUsageBit::HOST);
+        uint32_t offset = backBuffer ? lastUpdateCycle * _gpuBuffer->instanceSize : 0;
+        uint8_t* mappedData = static_cast<uint8_t*>(mtlBuffer.contents) + offset;
+        memcpy(mappedData, buffer, size);
+#if (CC_PLATFORM == CC_PLATFORM_MACOS)
+        if (mtlBuffer.storageMode == MTLStorageModeManaged) {
+            [mtlBuffer didModifyRange:NSMakeRange(offset, size)]; // Synchronize the managed buffer.
+        }
+#endif
+    } else {
+        auto* cmdBuffer = ccDevice->getCommandBuffer();
+        cmdBuffer->updateBuffer(this, buffer, size);
+    }
+}
+
+void CCMTLBuffer::encodeBuffer(CCMTLCommandEncoder &encoder, uint32_t offset, uint32_t binding, ShaderStageFlags stages) {
+    if (hasFlag(stages, ShaderStageFlagBit::VERTEX)) {
+        CCMTLRenderCommandEncoder *renderEncoder = static_cast<CCMTLRenderCommandEncoder *>(&encoder);
+        renderEncoder->setVertexBuffer(_gpuBuffer->mtlBuffer, offset + currentOffset(), binding);
+    }
+
+    if (hasFlag(stages, ShaderStageFlagBit::FRAGMENT)) {
+        CCMTLRenderCommandEncoder *renderEncoder = static_cast<CCMTLRenderCommandEncoder *>(&encoder);
+        renderEncoder->setFragmentBuffer(_gpuBuffer->mtlBuffer, offset + currentOffset(), binding);
+    }
+
+    if (hasFlag(stages, ShaderStageFlagBit::COMPUTE)) {
+        CCMTLComputeCommandEncoder *computeEncoder = static_cast<CCMTLComputeCommandEncoder *>(&encoder);
+        computeEncoder->setBuffer(_gpuBuffer->mtlBuffer, offset + currentOffset(), binding);
+    }
+}
+
+uint32_t CCMTLBuffer::currentOffset() const {
+    bool backBuffer = hasFlag(_memUsage, MemoryUsageBit::HOST);
+    uint32_t offset = backBuffer ? _gpuBuffer->lastUpdateCycle * _gpuBuffer->instanceSize : 0;
+    if(_isBufferView) {
+        offset += _offset; // buffer view offset
+    }
+    return offset;
+}
+
+id<MTLBuffer> CCMTLBuffer::mtlBuffer() const {
+    return _gpuBuffer->mtlBuffer;
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLCommandBuffer.h

@@ -0,0 +1,135 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#include "gfx-base/GFXCommandBuffer.h"
+
+#import <Metal/MTLCommandQueue.h>
+#import <MetalKit/MTKView.h>
+#include <bitset>
+#include "MTLComputeCommandEncoder.h"
+#include "MTLGPUObjects.h"
+#include "MTLRenderCommandEncoder.h"
+
+namespace cc {
+namespace gfx {
+
+struct CCMTLGPUPipelineState;
+class CCMTLInputAssembler;
+class CCMTLDevice;
+class CCMTLRenderPass;
+class CCMTLFence;
+class CCMTLFramebuffer;
+
+class CCMTLCommandBuffer final : public CommandBuffer {
+public:
+    explicit CCMTLCommandBuffer();
+    ~CCMTLCommandBuffer();
+    CCMTLCommandBuffer(const CCMTLCommandBuffer &) = delete;
+    CCMTLCommandBuffer(CCMTLCommandBuffer &&) = delete;
+    CCMTLCommandBuffer &operator=(const CCMTLCommandBuffer &) = delete;
+    CCMTLCommandBuffer &operator=(CCMTLCommandBuffer &&) = delete;
+
+    void begin(RenderPass *renderPass, uint32_t subpass, Framebuffer *frameBuffer) override;
+    void end() override;
+    void beginRenderPass(RenderPass *renderPass, Framebuffer *fbo, const Rect &renderArea, const Color *colors, float depth, uint32_t stencil, CommandBuffer *const *secondaryCBs, uint32_t secondaryCBCount) override;
+    void endRenderPass() override;
+    void insertMarker(const MarkerInfo &marker) override;
+    void beginMarker(const MarkerInfo &marker) override;
+    void endMarker() override;
+    void bindPipelineState(PipelineState *pso) override;
+    void bindDescriptorSet(uint32_t set, DescriptorSet *descriptorSet, uint32_t dynamicOffsetCount, const uint32_t *dynamicOffsets) override;
+    void bindInputAssembler(InputAssembler *ia) override;
+    void setViewport(const Viewport &vp) override;
+    void setScissor(const Rect &rect) override;
+    void setLineWidth(float width) override;
+    void setDepthBias(float constant, float clamp, float slope) override;
+    void setBlendConstants(const Color &constants) override;
+    void setDepthBound(float minBounds, float maxBounds) override;
+    void setStencilWriteMask(StencilFace face, uint32_t mask) override;
+    void setStencilCompareMask(StencilFace face, uint32_t ref, uint32_t mask) override;
+    void nextSubpass() override;
+    void draw(const DrawInfo &info) override;
+    void updateBuffer(Buffer *buff, const void *data, uint32_t size) override;
+    void copyBuffersToTexture(const uint8_t *const *buffers, Texture *texture, const BufferTextureCopy *regions, uint32_t count) override;
+    void blitTexture(Texture *srcTexture, Texture *dstTexture, const TextureBlit *regions, uint32_t count, Filter filter) override;
+    void copyTexture(Texture *srcTexture, Texture *dstTexture, const TextureCopy *regions, uint32_t count) override;
+    void resolveTexture(Texture *srcTexture, Texture *dstTexture, const TextureCopy *regions, uint32_t count) override;
+    void execute(CommandBuffer *const *cmdBuffs, uint32_t count) override;
+    void dispatch(const DispatchInfo &info) override;
+    void pipelineBarrier(const GeneralBarrier *barrier, const BufferBarrier *const *bufferBarriers, const Buffer *const *buffers, uint32_t bufferBarrierCount, const TextureBarrier *const *textureBarriers, const Texture *const *textures, uint32_t textureBarrierCount) override;
+    void copyTextureToBuffers(Texture *src, uint8_t *const *buffers, const BufferTextureCopy *regions, uint32_t count);
+    void beginQuery(QueryPool *queryPool, uint32_t id) override;
+    void endQuery(QueryPool *queryPool, uint32_t id) override;
+    void resetQueryPool(QueryPool *queryPool) override;
+    void completeQueryPool(QueryPool *queryPool) override;
+    inline bool isCommandBufferBegan() const { return _commandBufferBegan; }
+    inline CCMTLGPUCommandBufferObject *gpuCommandBufferObj() const { return _gpuCommandBufferObj; }
+
+    void afterCommit();
+
+    void signalFence();
+    void waitFence();
+
+protected:
+    friend class CCMTLQueue;
+
+    void doInit(const CommandBufferInfo &info) override;
+    void doDestroy() override;
+
+    void bindDescriptorSets();
+    void updateDepthStencilState(uint32_t subPassIndex, MTLRenderPassDescriptor *descriptor);
+    static bool isRenderingEntireDrawable(const Rect &rect, const CCMTLFramebuffer *renderPass);
+
+    id<MTLCommandBuffer> getMTLCommandBuffer();
+
+    ccstd::vector<CCMTLGPUDescriptorSet *> _GPUDescriptorSets; // NOLINT(bugprone-reserved-identifier)
+    ccstd::vector<ccstd::vector<uint32_t>> _dynamicOffsets;
+    uint32_t _firstDirtyDescriptorSet = UINT_MAX;
+
+    bool _indirectDrawSuppotred = false;
+    bool _commandBufferBegan = false;
+    bool _firstRenderPass = true;
+    CCMTLDevice *_mtlDevice = nullptr;
+    id<MTLCommandQueue> _mtlCommandQueue = nil;
+    CCMTLRenderCommandEncoder _renderEncoder;
+    CCMTLComputeCommandEncoder _computeEncoder;
+    id<MTLParallelRenderCommandEncoder> _parallelEncoder = nil;
+    MTLPrimitiveType _mtlPrimitiveType = MTLPrimitiveType::MTLPrimitiveTypeTriangle;
+
+    CCMTLGPUCommandBufferObject *_gpuCommandBufferObj = nullptr;
+
+    CCMTLSemaphore *_texCopySemaphore = nullptr;
+
+    std::bitset<MAX_COLORATTACHMENTS> _colorAppearedBefore;
+
+    CCMTLSemaphore *_inFlightSem{nullptr};
+
+    int32_t _currentFbWidth = 0;
+    int32_t _currentFbHeight = 0;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLCommandEncoder.h

@@ -0,0 +1,37 @@
+/****************************************************************************
+ Copyright (c) 2021-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+namespace cc {
+namespace gfx {
+
+class CCMTLCommandEncoder {
+public:
+    CCMTLCommandEncoder() = default;
+    virtual ~CCMTLCommandEncoder() = default;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLComputeCommandEncoder.h

@@ -0,0 +1,108 @@
+/****************************************************************************
+ Copyright (c) 2021-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#include <Metal/MTLComputeCommandEncoder.h>
+#include <Metal/MTLComputePipeline.h>
+#include "MTLCommandEncoder.h"
+#include "MTLUtils.h"
+#include "base/Macros.h"
+#include "math/Math.h"
+
+namespace cc {
+namespace gfx {
+class CCMTLComputeCommandEncoder final : public CCMTLCommandEncoder {
+public:
+    CCMTLComputeCommandEncoder() = default;
+    ~CCMTLComputeCommandEncoder() = default;
+    CCMTLComputeCommandEncoder(const CCMTLComputeCommandEncoder &) = delete;
+    CCMTLComputeCommandEncoder(CCMTLComputeCommandEncoder &&) = delete;
+    CCMTLComputeCommandEncoder &operator=(const CCMTLComputeCommandEncoder &) = delete;
+    CCMTLComputeCommandEncoder &operator=(CCMTLComputeCommandEncoder &&) = delete;
+
+    void initialize(id<MTLCommandBuffer> commandBuffer) {
+        _mtlEncoder = [[commandBuffer computeCommandEncoder] retain];
+        _initialized = true;
+    }
+
+    inline const bool isInitialized() {
+        return _initialized;
+    }
+
+    inline void setComputePipelineState(id<MTLComputePipelineState> pipelineState) {
+        if (_pipelineState == pipelineState)
+            return;
+        [_mtlEncoder setComputePipelineState:pipelineState];
+        _pipelineState = pipelineState;
+    }
+
+    inline void setBuffer(const id<MTLBuffer> buffer, uint32_t offset, uint32_t index) {
+        [_mtlEncoder setBuffer:buffer offset:offset atIndex:index];
+    }
+
+    inline void setTexture(const id<MTLTexture> texture, uint32_t index) {
+        [_mtlEncoder setTexture:texture atIndex:index];
+        _resourceSize = {texture.width, texture.height, texture.depth};
+    }
+
+    inline void dispatch(MTLSize groupsPerGrid) {
+        // GLSL -> SPIRV -> MSL
+        // GLSL shader request to specify the compute thread size,
+        // no such limit in Metal and have to set compute thread size explicity
+        NSUInteger w = _pipelineState.threadExecutionWidth;
+        NSUInteger h = _pipelineState.maxTotalThreadsPerThreadgroup / w;
+        MTLSize threadsPerThreadgroup = MTLSizeMake(w, h, 1);
+        [_mtlEncoder dispatchThreadgroups:groupsPerGrid threadsPerThreadgroup:threadsPerThreadgroup];
+    }
+
+    inline void dispatch(MTLSize groupsPerGrid, MTLSize workGroupSize) {
+        [_mtlEncoder dispatchThreadgroups:groupsPerGrid threadsPerThreadgroup:workGroupSize];
+    }
+
+    inline void dispatch(id<MTLBuffer> indirectBuffer, NSUInteger offset, MTLSize localSize) {
+        [_mtlEncoder dispatchThreadgroupsWithIndirectBuffer:indirectBuffer indirectBufferOffset:offset threadsPerThreadgroup:localSize];
+    }
+
+    inline void endEncoding() {
+        [_mtlEncoder endEncoding];
+        [_mtlEncoder release];
+        _mtlEncoder = nil;
+        _pipelineState = nil;
+        _initialized = false;
+    }
+
+    inline id<MTLComputeCommandEncoder> const getMTLEncoder() {
+        return _mtlEncoder;
+    }
+
+private:
+    bool _initialized = false;
+    MTLSize _resourceSize;
+    id<MTLComputeCommandEncoder> _mtlEncoder = nil;
+    id<MTLComputePipelineState> _pipelineState = nil;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLConfig.h

@@ -0,0 +1,28 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#define MAX_FRAMES_IN_FLIGHT     3
+#define MAX_COMMAND_BUFFER_COUNT 256

cocos/renderer/gfx-metal/MTLDescriptorSet.h

@@ -0,0 +1,55 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#include "gfx-base/GFXDescriptorSet.h"
+
+namespace cc {
+namespace gfx {
+
+class CCMTLGPUDescriptorSet;
+
+class CCMTLDescriptorSet final : public DescriptorSet {
+public:
+    explicit CCMTLDescriptorSet();
+    ~CCMTLDescriptorSet();
+    CCMTLDescriptorSet(const CCMTLDescriptorSet &) = delete;
+    CCMTLDescriptorSet(CCMTLDescriptorSet &&) = delete;
+    CCMTLDescriptorSet &operator=(const CCMTLDescriptorSet &) = delete;
+    CCMTLDescriptorSet &operator=(CCMTLDescriptorSet &&) = delete;
+
+    void update() override;
+    void forceUpdate() override;
+
+    inline CCMTLGPUDescriptorSet *gpuDescriptorSet() const { return _gpuDescriptorSet; }
+
+protected:
+    void doInit(const DescriptorSetInfo &info) override;
+    void doDestroy() override;
+
+    CCMTLGPUDescriptorSet *_gpuDescriptorSet = nullptr;
+};
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLDescriptorSet.mm

@@ -0,0 +1,96 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#include "MTLBuffer.h"
+#include "MTLDescriptorSet.h"
+#include "MTLDescriptorSetLayout.h"
+#include "MTLGPUObjects.h"
+#include "MTLSampler.h"
+#include "MTLTexture.h"
+
+namespace cc {
+namespace gfx {
+CCMTLDescriptorSet::CCMTLDescriptorSet() : DescriptorSet() {
+    _typedID = generateObjectID<decltype(this)>();
+}
+
+void CCMTLDescriptorSet::doInit(const DescriptorSetInfo &info) {
+    const auto gpuDescriptorSetLayout = static_cast<const CCMTLDescriptorSetLayout *>(_layout)->gpuDescriptorSetLayout();
+    const auto descriptorCount = gpuDescriptorSetLayout->descriptorCount;
+    const auto bindingCount = gpuDescriptorSetLayout->bindings.size();
+
+    _gpuDescriptorSet = ccnew CCMTLGPUDescriptorSet;
+    _gpuDescriptorSet->descriptorIndices = &gpuDescriptorSetLayout->descriptorIndices;
+    _gpuDescriptorSet->gpuDescriptors.resize(descriptorCount);
+    for (auto i = 0u, k = 0u; i < bindingCount; i++) {
+        const auto &binding = gpuDescriptorSetLayout->bindings[i];
+        for (auto j = 0u; j < binding.count; j++, k++) {
+            _gpuDescriptorSet->gpuDescriptors[k].type = binding.descriptorType;
+        }
+    }
+}
+
+CCMTLDescriptorSet::~CCMTLDescriptorSet() {
+    destroy();
+}
+
+void CCMTLDescriptorSet::doDestroy() {
+    CC_SAFE_DELETE(_gpuDescriptorSet);
+}
+
+void CCMTLDescriptorSet::update() {
+    if (_isDirty && _gpuDescriptorSet) {
+        const auto &descriptors = _gpuDescriptorSet->gpuDescriptors;
+        for (size_t i = 0; i < descriptors.size(); i++) {
+            if (hasAnyFlags(descriptors[i].type, DESCRIPTOR_BUFFER_TYPE)) {
+                if (_buffers[i].ptr) {
+                    _gpuDescriptorSet->gpuDescriptors[i].buffer = static_cast<CCMTLBuffer *>(_buffers[i].ptr);
+                }
+            } else if (hasAnyFlags(descriptors[i].type, DESCRIPTOR_TEXTURE_TYPE)) {
+                if (!_textures[i].ptr && !_samplers[i].ptr)
+                    continue;
+
+                Texture *tex = _textures[i].ptr;
+                if (!tex)
+                    tex = CCMTLTexture::getDefaultTexture();
+                _gpuDescriptorSet->gpuDescriptors[i].texture = static_cast<CCMTLTexture *>(tex);
+
+                Sampler *sampler = _samplers[i].ptr;
+                if (!sampler)
+                    sampler = CCMTLSampler::getDefaultSampler();
+                _gpuDescriptorSet->gpuDescriptors[i].sampler = static_cast<CCMTLSampler *>(sampler);
+            }
+        }
+        _isDirty = false;
+    }
+}
+
+void CCMTLDescriptorSet::forceUpdate() {
+    _isDirty = true;
+    update();
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLDescriptorSetLayout.h

@@ -0,0 +1,52 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#include "gfx-base/GFXDescriptorSetLayout.h"
+
+namespace cc {
+namespace gfx {
+
+class CCMTLGPUDescriptorSetLayout;
+class CCMTLDescriptorSetLayout final : public DescriptorSetLayout {
+public:
+    explicit CCMTLDescriptorSetLayout();
+    ~CCMTLDescriptorSetLayout();
+    CCMTLDescriptorSetLayout(const CCMTLDescriptorSetLayout &) = delete;
+    CCMTLDescriptorSetLayout(CCMTLDescriptorSetLayout &&) = delete;
+    CCMTLDescriptorSetLayout &operator=(const CCMTLDescriptorSetLayout &) = delete;
+    CCMTLDescriptorSetLayout &operator=(CCMTLDescriptorSetLayout &&) = delete;
+
+    inline CCMTLGPUDescriptorSetLayout *gpuDescriptorSetLayout() const { return _gpuDescriptorSetLayout; }
+
+protected:
+    void doInit(const DescriptorSetLayoutInfo &info) override;
+    void doDestroy() override;
+
+    CCMTLGPUDescriptorSetLayout *_gpuDescriptorSetLayout = nullptr;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLDescriptorSetLayout.mm

@@ -0,0 +1,63 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#include "MTLDescriptorSetLayout.h"
+#include "MTLGPUObjects.h"
+
+namespace cc {
+namespace gfx {
+
+CCMTLDescriptorSetLayout::CCMTLDescriptorSetLayout() : DescriptorSetLayout() {
+    _typedID = generateObjectID<decltype(this)>();
+}
+
+CCMTLDescriptorSetLayout::~CCMTLDescriptorSetLayout() {
+    destroy();
+}
+
+void CCMTLDescriptorSetLayout::doInit(const DescriptorSetLayoutInfo &info) {
+    _gpuDescriptorSetLayout = ccnew CCMTLGPUDescriptorSetLayout;
+
+    _gpuDescriptorSetLayout->descriptorCount = _descriptorCount;
+    _gpuDescriptorSetLayout->descriptorIndices = _descriptorIndices;
+    _gpuDescriptorSetLayout->bindingIndices = _bindingIndices;
+    _gpuDescriptorSetLayout->bindings = _bindings;
+
+    for (size_t i = 0; i < _bindings.size(); i++) {
+        const auto binding = _bindings[i];
+        if (hasAnyFlags(binding.descriptorType, DESCRIPTOR_DYNAMIC_TYPE)) {
+            for (uint32_t j = 0; j < binding.count; j++) {
+                _gpuDescriptorSetLayout->dynamicBindings.push_back(binding.binding);
+            }
+        }
+    }
+}
+
+void CCMTLDescriptorSetLayout::doDestroy() {
+    CC_SAFE_DELETE(_gpuDescriptorSetLayout);
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLDevice.h

@@ -0,0 +1,144 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#import "MTLConfig.h"
+#import "gfx-base/GFXDevice.h"
+
+namespace cc {
+namespace gfx {
+
+struct CCMTLGPUDeviceObject;
+
+class CCMTLGPUStagingBufferPool;
+class CCMTLSemaphore;
+class CCMTLSwapchain;
+
+class CCMTLDevice final : public Device {
+public:
+    static CCMTLDevice *getInstance();
+
+    CCMTLDevice();
+    ~CCMTLDevice();
+
+    CCMTLDevice(const CCMTLDevice &) = delete;
+    CCMTLDevice(CCMTLDevice &&) = delete;
+    CCMTLDevice &operator=(const CCMTLDevice &) = delete;
+    CCMTLDevice &operator=(CCMTLDevice &&) = delete;
+
+    using Device::copyBuffersToTexture;
+    using Device::createBuffer;
+    using Device::createCommandBuffer;
+    using Device::createDescriptorSet;
+    using Device::createDescriptorSetLayout;
+    using Device::createFramebuffer;
+    using Device::createInputAssembler;
+    using Device::createPipelineLayout;
+    using Device::createPipelineState;
+    using Device::createQueryPool;
+    using Device::createQueue;
+    using Device::createRenderPass;
+    using Device::createShader;
+    using Device::createTexture;
+
+    void frameSync() override;
+
+    void acquire(Swapchain *const *swapchains, uint32_t count) override;
+    void present() override;
+
+    void onPresentCompleted(uint32_t index);
+
+    inline void *getMTLCommandQueue() const { return _mtlCommandQueue; }
+    inline void *getMTLDevice() const { return _mtlDevice; }
+    inline uint32_t getMaximumSamplerUnits() const { return _maxSamplerUnits; }
+    inline uint32_t getMaximumColorRenderTargets() const { return _caps.maxColorRenderTargets; }
+    inline uint32_t getMaximumBufferBindingIndex() const { return _maxBufferBindingIndex; }
+    inline bool isIndirectCommandBufferSupported() const { return _icbSuppored; }
+    inline bool isIndirectDrawSupported() const { return _indirectDrawSupported; }
+    inline CCMTLGPUStagingBufferPool *gpuStagingBufferPool() const { return _gpuStagingBufferPools[_currentFrameIndex]; }
+    inline bool isSamplerDescriptorCompareFunctionSupported() const { return _isSamplerDescriptorCompareFunctionSupported; }
+    inline uint32_t currentFrameIndex() const { return _currentFrameIndex; }
+
+    inline void registerSwapchain(CCMTLSwapchain *swapchain) { _swapchains.push_back(swapchain); }
+    inline void unRegisterSwapchain(CCMTLSwapchain *swapchain) {
+        auto iter = std::find(_swapchains.begin(), _swapchains.end(), swapchain);
+        if (iter != _swapchains.end()) {
+            _swapchains.erase(iter);
+        }
+    }
+    
+    inline CCMTLGPUDeviceObject* gpuObject() const { return _gpuDeviceObj; }
+
+protected:
+    static CCMTLDevice *_instance;
+
+    friend class DeviceManager;
+
+    bool doInit(const DeviceInfo &info) override;
+    void doDestroy() override;
+    CommandBuffer *createCommandBuffer(const CommandBufferInfo &info, bool hasAgent) override;
+    Queue *createQueue() override;
+    QueryPool *createQueryPool() override;
+    Buffer *createBuffer() override;
+    Texture *createTexture() override;
+    Shader *createShader() override;
+    InputAssembler *createInputAssembler() override;
+    RenderPass *createRenderPass() override;
+    Framebuffer *createFramebuffer() override;
+    DescriptorSet *createDescriptorSet() override;
+    DescriptorSetLayout *createDescriptorSetLayout() override;
+    PipelineLayout *createPipelineLayout() override;
+    PipelineState *createPipelineState() override;
+    Sampler *createSampler(const SamplerInfo &info) override;
+    Swapchain *createSwapchain() override;
+    void copyBuffersToTexture(const uint8_t *const *buffers, Texture *dst, const BufferTextureCopy *regions, uint32_t count) override;
+    void copyTextureToBuffers(Texture *src, uint8_t *const *buffers, const BufferTextureCopy *region, uint32_t count) override;
+    void getQueryPoolResults(QueryPool *queryPool) override;
+    SampleCount getMaxSampleCount(Format format, TextureUsage usage, TextureFlags flags) const override;
+
+    void onMemoryWarning();
+    void initFormatFeatures(uint32_t family);
+
+    void *_mtlCommandQueue = nullptr;
+    void *_mtlDevice = nullptr;
+    void *_activeDrawable = nullptr;
+    unsigned long _mtlFeatureSet = 0;
+    uint32_t _maxSamplerUnits = 0;
+    uint32_t _maxBufferBindingIndex = 0;
+    bool _icbSuppored = false;
+    bool _indirectDrawSupported = false;
+    bool _isSamplerDescriptorCompareFunctionSupported = false;
+    CCMTLGPUStagingBufferPool *_gpuStagingBufferPools[MAX_FRAMES_IN_FLIGHT] = {nullptr};
+    uint32_t _currentBufferPoolId = 0;
+    uint32_t _currentFrameIndex = 0;
+    CC_UNUSED uint32_t _memoryAlarmListenerId = 0;
+
+    ccstd::vector<CCMTLSwapchain *> _swapchains;
+
+    CCMTLGPUDeviceObject *_gpuDeviceObj = nullptr;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLDevice.mm

@@ -0,0 +1,561 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#import "MTLDevice.h"
+#import "MTLBuffer.h"
+#import "MTLCommandBuffer.h"
+#import "MTLConfig.h"
+#import "MTLDescriptorSet.h"
+#import "MTLDescriptorSetLayout.h"
+#import "MTLFramebuffer.h"
+#import "MTLGPUObjects.h"
+#import "MTLInputAssembler.h"
+#import "MTLPipelineLayout.h"
+#import "MTLPipelineState.h"
+#import "MTLQueryPool.h"
+#import "MTLQueue.h"
+#import "MTLRenderPass.h"
+#import "MTLSampler.h"
+#import "MTLSemaphore.h"
+#import "MTLShader.h"
+#import "MTLSwapchain.h"
+#import "MTLTexture.h"
+#import "base/Log.h"
+#import "profiler/Profiler.h"
+#import <thread>
+
+namespace cc {
+namespace gfx {
+
+CCMTLDevice *CCMTLDevice::_instance = nullptr;
+
+CCMTLDevice *CCMTLDevice::getInstance() {
+    return CCMTLDevice::_instance;
+}
+
+CCMTLDevice::CCMTLDevice() {
+    _api = API::METAL;
+    _deviceName = "Metal";
+
+    _caps.supportQuery = true;
+    _caps.clipSpaceMinZ = 0.0f;
+    _caps.screenSpaceSignY = -1.0f;
+    _caps.clipSpaceSignY = 1.0f;
+
+    CCMTLDevice::_instance = this;
+}
+
+CCMTLDevice::~CCMTLDevice() {
+    CCMTLDevice::_instance = nullptr;
+}
+
+bool CCMTLDevice::doInit(const DeviceInfo &info) {
+    _gpuDeviceObj = ccnew CCMTLGPUDeviceObject;
+
+    _currentFrameIndex = 0;
+
+    id<MTLDevice> mtlDevice = MTLCreateSystemDefaultDevice();
+    _mtlDevice = mtlDevice;
+
+    NSString *deviceName = [mtlDevice name];
+    _renderer = [deviceName UTF8String];
+    NSArray* nameArr = [deviceName componentsSeparatedByString:@" "];
+    if ([nameArr count] > 0) {
+        _vendor = [nameArr[0] UTF8String];
+    }
+    _mtlFeatureSet = mu::highestSupportedFeatureSet(mtlDevice);
+    _version = std::to_string(_mtlFeatureSet);
+
+    const auto gpuFamily = mu::getGPUFamily(MTLFeatureSet(_mtlFeatureSet));
+    _indirectDrawSupported = mu::isIndirectDrawSupported(gpuFamily);
+    _caps.maxVertexAttributes = mu::getMaxVertexAttributes(gpuFamily);
+    _caps.maxTextureUnits = _caps.maxVertexTextureUnits = mu::getMaxEntriesInTextureArgumentTable(gpuFamily);
+    _maxSamplerUnits = mu::getMaxEntriesInSamplerStateArgumentTable(gpuFamily);
+    _caps.maxTextureSize = mu::getMaxTexture2DWidthHeight(gpuFamily);
+    _caps.maxCubeMapTextureSize = mu::getMaxCubeMapTextureWidthHeight(gpuFamily);
+    _caps.maxColorRenderTargets = mu::getMaxColorRenderTarget(gpuFamily);
+    _caps.uboOffsetAlignment = mu::getMinBufferOffsetAlignment(gpuFamily);
+    _caps.maxComputeWorkGroupInvocations = mu::getMaxThreadsPerGroup(gpuFamily);
+    _caps.maxFragmentUniformVectors = 256;
+    _caps.maxVertexUniformVectors = 256;
+    _caps.maxUniformBufferBindings = mu::getMaxUniformBufferBindings(gpuFamily);
+    _maxBufferBindingIndex = mu::getMaxEntriesInBufferArgumentTable(gpuFamily);
+    _icbSuppored = mu::isIndirectCommandBufferSupported(MTLFeatureSet(_mtlFeatureSet));
+    _isSamplerDescriptorCompareFunctionSupported = mu::isSamplerDescriptorCompareFunctionSupported(gpuFamily);
+
+    for (int i = 0; i < MAX_FRAMES_IN_FLIGHT; ++i) {
+        _gpuStagingBufferPools[i] = ccnew CCMTLGPUStagingBufferPool(mtlDevice);
+    }
+
+    initFormatFeatures(gpuFamily);
+
+    ccstd::string compressedFormats;
+
+    if (getFormatFeatures(Format::BC1_SRGB_ALPHA) != FormatFeature::NONE) {
+        compressedFormats += "dxt ";
+    }
+
+    if (getFormatFeatures(Format::ETC2_RGBA8) != FormatFeature::NONE) {
+        compressedFormats += "etc2 ";
+    }
+
+    if (getFormatFeatures(Format::ASTC_RGBA_4X4) != FormatFeature::NONE) {
+        compressedFormats += "astc ";
+    }
+
+    if (getFormatFeatures(Format::PVRTC_RGBA2) != FormatFeature::NONE) {
+        compressedFormats += "pvrtc ";
+    }
+
+    _features[toNumber(Feature::INSTANCED_ARRAYS)] = true;
+    _features[toNumber(Feature::MULTIPLE_RENDER_TARGETS)] = true;
+    _features[toNumber(Feature::BLEND_MINMAX)] = true;
+    _features[toNumber(Feature::ELEMENT_INDEX_UINT)] = true;
+    _features[toNumber(Feature::COMPUTE_SHADER)] = true;
+    _features[toNumber(Feature::INPUT_ATTACHMENT_BENEFIT)] = true;
+    _features[toNumber(Feature::SUBPASS_COLOR_INPUT)] = true;
+    _features[toNumber(Feature::SUBPASS_DEPTH_STENCIL_INPUT)] = false;
+    _features[toNumber(Feature::RASTERIZATION_ORDER_NOCOHERENT)] = true;
+
+    if (@available(iOS 13.0, macOS 10.15, *)) {
+        // detph resolve requires MTLGPUFamilyApple3 while stencil resolve requires MTLGPUFamilyApple5
+        _features[toNumber(Feature::MULTI_SAMPLE_RESOLVE_DEPTH_STENCIL)] = [mtlDevice supportsFamily:MTLGPUFamilyApple5];
+        _features[toNumber(Feature::MULTI_SAMPLE_RESOLVE_DEPTH_STENCIL)] |= [mtlDevice supportsFamily:MTLGPUFamilyMac2];
+    } else {
+#if CC_PLATFOTM == CC_PLATFORM_IOS
+        id<MTLDevice> device = static_cast<id<MTLDevice>>(_mtlDevice);
+        _features[toNumber(Feature::MULTI_SAMPLE_RESOLVE_DEPTH_STENCIL)] = [device supportsFeatureSet:MTLFeatureSet_iOS_GPUFamily2_v4];
+#elif CC_PLATFOTM == CC_PLATFORM_MACOS
+        _features[toNumber(Feature::MULTI_SAMPLE_RESOLVE_DEPTH_STENCIL)] = false;
+#endif
+    }
+
+    QueueInfo queueInfo;
+    queueInfo.type = QueueType::GRAPHICS;
+    _queue = createQueue(queueInfo);
+
+    QueryPoolInfo queryPoolInfo{QueryType::OCCLUSION, DEFAULT_MAX_QUERY_OBJECTS, true};
+    _queryPool = createQueryPool(queryPoolInfo);
+
+    CommandBufferInfo cmdBuffInfo;
+    cmdBuffInfo.type = CommandBufferType::PRIMARY;
+    cmdBuffInfo.queue = _queue;
+    _cmdBuff = createCommandBuffer(cmdBuffInfo);
+
+    CCMTLGPUGarbageCollectionPool::getInstance()->initialize(std::bind(&CCMTLDevice::currentFrameIndex, this));
+
+    CC_LOG_INFO("Metal Feature Set: %s", mu::featureSetToString(MTLFeatureSet(_mtlFeatureSet)).c_str());
+
+    return true;
+}
+
+void CCMTLDevice::doDestroy() {
+    CC_SAFE_DELETE(_gpuDeviceObj);
+
+    CC_SAFE_DESTROY_AND_DELETE(_queryPool)
+    CC_SAFE_DESTROY_AND_DELETE(_queue);
+    CC_SAFE_DESTROY_AND_DELETE(_cmdBuff);
+
+    CCMTLGPUGarbageCollectionPool::getInstance()->flush();
+
+    for (int i = 0; i < MAX_FRAMES_IN_FLIGHT; ++i) {
+        CC_SAFE_DELETE(_gpuStagingBufferPools[i]);
+        _gpuStagingBufferPools[i] = nullptr;
+    }
+
+    cc::gfx::mu::clearUtilResource();
+
+    CCMTLTexture::deleteDefaultTexture();
+    CCMTLSampler::deleteDefaultSampler();
+
+    CC_ASSERT(!_memoryStatus.bufferSize);  // Buffer memory leaked
+    CC_ASSERT(!_memoryStatus.textureSize); // Texture memory leaked
+}
+
+void CCMTLDevice::frameSync() {
+    CC_ASSERT(_cmdBuff);
+    auto* cmdBuff = static_cast<CCMTLCommandBuffer*>(_cmdBuff);
+    cmdBuff->waitFence();
+}
+
+void CCMTLDevice::acquire(Swapchain *const *swapchains, uint32_t count) {
+    if (_onAcquire) _onAcquire->execute();
+
+    for (CCMTLSwapchain *swapchain : _swapchains) {
+        swapchain->acquire();
+    }
+
+    // Clear queue stats
+    CCMTLQueue *queue = static_cast<CCMTLQueue *>(_queue);
+    queue->gpuQueueObj()->numDrawCalls = 0;
+    queue->gpuQueueObj()->numInstances = 0;
+    queue->gpuQueueObj()->numTriangles = 0;
+}
+
+void CCMTLDevice::present() {
+    CC_PROFILE(CCMTLDevicePresent);
+    CCMTLQueue *queue = (CCMTLQueue *)_queue;
+    _numDrawCalls = queue->gpuQueueObj()->numDrawCalls;
+    _numInstances = queue->gpuQueueObj()->numInstances;
+    _numTriangles = queue->gpuQueueObj()->numTriangles;
+
+    //hold this pointer before update _currentFrameIndex
+    auto tempIndex = _currentFrameIndex;
+    _currentBufferPoolId = _currentFrameIndex;
+    _currentFrameIndex = (_currentFrameIndex + 1) % MAX_FRAMES_IN_FLIGHT;
+
+    ccstd::vector<id<CAMetalDrawable>> releaseQ;
+    for (auto *swapchain : _swapchains) {
+        auto drawable = swapchain->currentDrawable();
+        if (drawable) {
+            releaseQ.push_back([drawable retain]);
+        }
+        swapchain->release();
+    }
+
+    // present drawable
+    {
+        id<MTLCommandBuffer> cmdBuffer = [queue->gpuQueueObj()->mtlCommandQueue commandBuffer];
+        [cmdBuffer enqueue];
+
+        for (auto drawable : releaseQ) {
+            [cmdBuffer presentDrawable:drawable];
+            [drawable release];
+        }
+        [cmdBuffer addCompletedHandler:^(id<MTLCommandBuffer> commandBuffer) {
+            onPresentCompleted(tempIndex);
+        }];
+        [cmdBuffer commit];
+    }
+}
+
+void CCMTLDevice::onPresentCompleted(uint32_t index) {
+    if (index >= 0 && index < MAX_FRAMES_IN_FLIGHT) {
+        CCMTLGPUStagingBufferPool *bufferPool = _gpuStagingBufferPools[index];
+        if (bufferPool) {
+            bufferPool->reset();
+            CCMTLGPUGarbageCollectionPool::getInstance()->clear(index);
+            static_cast<CCMTLCommandBuffer*>(_cmdBuff)->signalFence();
+        }
+    }
+}
+
+Queue *CCMTLDevice::createQueue() {
+    return ccnew CCMTLQueue;
+}
+
+QueryPool *CCMTLDevice::createQueryPool() {
+    return ccnew CCMTLQueryPool;
+}
+
+CommandBuffer *CCMTLDevice::createCommandBuffer(const CommandBufferInfo &info, bool /*hasAgent*/) {
+    return ccnew CCMTLCommandBuffer;
+}
+
+Buffer *CCMTLDevice::createBuffer() {
+    return ccnew CCMTLBuffer;
+}
+
+Texture *CCMTLDevice::createTexture() {
+    return ccnew CCMTLTexture;
+}
+
+Shader *CCMTLDevice::createShader() {
+    return ccnew CCMTLShader;
+}
+
+InputAssembler *CCMTLDevice::createInputAssembler() {
+    return ccnew CCMTLInputAssembler;
+}
+
+RenderPass *CCMTLDevice::createRenderPass() {
+    return ccnew CCMTLRenderPass;
+}
+
+Framebuffer *CCMTLDevice::createFramebuffer() {
+    return ccnew CCMTLFramebuffer;
+}
+
+DescriptorSet *CCMTLDevice::createDescriptorSet() {
+    return ccnew CCMTLDescriptorSet;
+}
+
+DescriptorSetLayout *CCMTLDevice::createDescriptorSetLayout() {
+    return ccnew CCMTLDescriptorSetLayout;
+}
+
+PipelineLayout *CCMTLDevice::createPipelineLayout() {
+    return ccnew CCMTLPipelineLayout;
+}
+
+PipelineState *CCMTLDevice::createPipelineState() {
+    return ccnew CCMTLPipelineState;
+}
+
+Sampler *CCMTLDevice::createSampler(const SamplerInfo &info) {
+    return ccnew CCMTLSampler(info);
+}
+
+Swapchain *CCMTLDevice::createSwapchain() {
+    return ccnew CCMTLSwapchain;
+}
+
+void CCMTLDevice::copyBuffersToTexture(const uint8_t *const *buffers, Texture *texture, const BufferTextureCopy *regions, uint32_t count) {
+    CC_PROFILE(CCMTLDeviceCopyBuffersToTexture);
+    // This assumes the default command buffer will get submitted every frame,
+    // which is true for now but may change in the future. This approach gives us
+    // the wiggle room to leverage immediate update vs. copy-upload strategies without
+    // breaking compatibilities. When we reached some conclusion on this subject,
+    // getting rid of this interface all together may become a better option.
+    static_cast<CCMTLCommandBuffer *>(_cmdBuff)->copyBuffersToTexture(buffers, texture, regions, count);
+}
+
+void CCMTLDevice::copyTextureToBuffers(Texture *src, uint8_t *const *buffers, const BufferTextureCopy *region, uint32_t count) {
+    CC_PROFILE(CCMTLDeviceCopyTextureToBuffers);
+    static_cast<CCMTLCommandBuffer *>(_cmdBuff)->copyTextureToBuffers(src, buffers, region, count);
+}
+
+void CCMTLDevice::getQueryPoolResults(QueryPool *queryPool) {
+    CC_PROFILE(CCMTLDeviceGetQueryPoolResults);
+    auto *mtlQueryPool = static_cast<CCMTLQueryPool *>(queryPool);
+    CCMTLGPUQueryPool *gpuQueryPool = mtlQueryPool->gpuQueryPool();
+    auto queryCount = static_cast<uint32_t>(mtlQueryPool->_ids.size());
+    CC_ASSERT(queryCount <= mtlQueryPool->getMaxQueryObjects()); // Too many query commands.
+
+    gpuQueryPool->semaphore->wait();
+    uint64_t *results = queryCount > 0U ? static_cast<uint64_t *>(gpuQueryPool->visibilityResultBuffer.contents) : nullptr;
+
+    ccstd::unordered_map<uint32_t, uint64_t> mapResults;
+    for (auto queryId = 0U; queryId < queryCount; queryId++) {
+        uint32_t id = mtlQueryPool->_ids[queryId];
+        auto iter = mapResults.find(id);
+        if (iter != mapResults.end()) {
+            iter->second += results[queryId];
+        } else {
+            mapResults[id] = results[queryId];
+        }
+    }
+
+    {
+        std::lock_guard<std::mutex> lock(mtlQueryPool->_mutex);
+        mtlQueryPool->_results = std::move(mapResults);
+    }
+}
+
+void CCMTLDevice::onMemoryWarning() {
+    for (int i = 0; i < MAX_FRAMES_IN_FLIGHT; ++i) {
+        _gpuStagingBufferPools[i]->shrinkSize();
+    }
+}
+void CCMTLDevice::initFormatFeatures(uint32_t gpuFamily) {
+    const FormatFeature completeFeature = FormatFeature::RENDER_TARGET | FormatFeature::SAMPLED_TEXTURE | FormatFeature::LINEAR_FILTER | FormatFeature::STORAGE_TEXTURE;
+
+    FormatFeature tempFeature = FormatFeature::RENDER_TARGET | FormatFeature::SAMPLED_TEXTURE | FormatFeature::STORAGE_TEXTURE;
+
+    _formatFeatures[toNumber(Format::R8UI)] = tempFeature;
+    _formatFeatures[toNumber(Format::RG8UI)] = tempFeature;
+    _formatFeatures[toNumber(Format::RGBA8UI)] = tempFeature;
+
+    _formatFeatures[toNumber(Format::R16UI)] = tempFeature;
+    _formatFeatures[toNumber(Format::RG16UI)] = tempFeature;
+    _formatFeatures[toNumber(Format::RGBA16UI)] = tempFeature;
+
+    _formatFeatures[toNumber(Format::R32F)] = tempFeature;
+    _formatFeatures[toNumber(Format::RG32F)] = tempFeature;
+    _formatFeatures[toNumber(Format::RGBA32F)] = tempFeature;
+
+    if (mu::isUISamplerSupported(gpuFamily)) {
+        tempFeature = FormatFeature::RENDER_TARGET | FormatFeature::STORAGE_TEXTURE;
+        _formatFeatures[toNumber(Format::R32UI)] = tempFeature;
+        _formatFeatures[toNumber(Format::RG32UI)] = tempFeature;
+        _formatFeatures[toNumber(Format::RGBA32UI)] = tempFeature;
+    } else {
+        tempFeature = FormatFeature::RENDER_TARGET | FormatFeature::SAMPLED_TEXTURE | FormatFeature::STORAGE_TEXTURE;
+        _formatFeatures[toNumber(Format::R32UI)] = tempFeature;
+        _formatFeatures[toNumber(Format::RG32UI)] = tempFeature;
+        _formatFeatures[toNumber(Format::RGBA32UI)] = tempFeature;
+    }
+
+    if (mu::isRGB10A2UIStorageSupported(gpuFamily)) {
+        _formatFeatures[toNumber(Format::RGB10A2UI)] = FormatFeature::RENDER_TARGET | FormatFeature::SAMPLED_TEXTURE;
+    } else {
+        _formatFeatures[toNumber(Format::RGB10A2UI)] = FormatFeature::RENDER_TARGET | FormatFeature::SAMPLED_TEXTURE | FormatFeature::STORAGE_TEXTURE;
+    }
+
+    tempFeature = FormatFeature::RENDER_TARGET | FormatFeature::SAMPLED_TEXTURE | FormatFeature::LINEAR_FILTER | FormatFeature::STORAGE_TEXTURE;
+    _formatFeatures[toNumber(Format::BGRA8)] = tempFeature;
+    _formatFeatures[toNumber(Format::R8SN)] = tempFeature;
+    _formatFeatures[toNumber(Format::RG8SN)] = tempFeature;
+    _formatFeatures[toNumber(Format::RGBA8SN)] = tempFeature;
+
+    _formatFeatures[toNumber(Format::R8)] = tempFeature;
+    _formatFeatures[toNumber(Format::RG8)] = tempFeature;
+    _formatFeatures[toNumber(Format::RGBA8)] = tempFeature;
+
+    _formatFeatures[toNumber(Format::R16F)] = tempFeature;
+    _formatFeatures[toNumber(Format::RG16F)] = tempFeature;
+    _formatFeatures[toNumber(Format::RGBA16F)] = tempFeature;
+
+    _formatFeatures[toNumber(Format::R11G11B10F)] = FormatFeature::RENDER_TARGET | FormatFeature::SAMPLED_TEXTURE;
+    _formatFeatures[toNumber(Format::RGB9E5)] = FormatFeature::RENDER_TARGET | FormatFeature::SAMPLED_TEXTURE;
+
+    if (mu::isDDepthStencilFilterSupported(gpuFamily)) {
+        _formatFeatures[toNumber(Format::DEPTH)] = FormatFeature::RENDER_TARGET | FormatFeature::SAMPLED_TEXTURE;
+        _formatFeatures[toNumber(Format::DEPTH_STENCIL)] = FormatFeature::RENDER_TARGET | FormatFeature::SAMPLED_TEXTURE;
+    } else {
+        _formatFeatures[toNumber(Format::DEPTH)] = FormatFeature::RENDER_TARGET | FormatFeature::SAMPLED_TEXTURE | FormatFeature::LINEAR_FILTER;
+        _formatFeatures[toNumber(Format::DEPTH_STENCIL)] = FormatFeature::RENDER_TARGET | FormatFeature::SAMPLED_TEXTURE | FormatFeature::LINEAR_FILTER;
+    }
+
+    const FormatFeature compressedFeature = FormatFeature::SAMPLED_TEXTURE | FormatFeature::LINEAR_FILTER;
+    if (mu::isPVRTCSuppported(gpuFamily)) {
+        _formatFeatures[toNumber(Format::PVRTC_RGB2)] = compressedFeature;
+        _formatFeatures[toNumber(Format::PVRTC_RGBA2)] = compressedFeature;
+        _formatFeatures[toNumber(Format::PVRTC_RGB4)] = compressedFeature;
+        _formatFeatures[toNumber(Format::PVRTC_RGBA4)] = compressedFeature;
+    }
+    if (mu::isEAC_ETCCSuppported(gpuFamily)) {
+        _formatFeatures[toNumber(Format::ETC2_RGB8)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ETC2_RGBA8)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ETC2_SRGB8)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ETC2_SRGB8_A8)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ETC2_RGB8_A1)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ETC2_SRGB8_A1)] = compressedFeature;
+    }
+    if (mu::isASTCSuppported(gpuFamily)) {
+        _formatFeatures[toNumber(Format::ASTC_RGBA_4X4)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_5X4)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_5X5)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_6X5)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_6X6)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_8X5)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_8X6)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_8X8)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_10X5)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_10X6)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_10X8)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_10X10)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_12X10)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_RGBA_12X12)] = compressedFeature;
+
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_4X4)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_5X4)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_5X5)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_6X5)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_6X6)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_8X5)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_8X6)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_8X8)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_10X5)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_10X6)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_10X8)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_10X10)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_12X10)] = compressedFeature;
+        _formatFeatures[toNumber(Format::ASTC_SRGBA_12X12)] = compressedFeature;
+    }
+
+    if (mu::isBCSupported(gpuFamily)) {
+        _formatFeatures[toNumber(Format::BC1)] = compressedFeature;
+        _formatFeatures[toNumber(Format::BC1_ALPHA)] = compressedFeature;
+        _formatFeatures[toNumber(Format::BC1_SRGB)] = compressedFeature;
+        _formatFeatures[toNumber(Format::BC1_SRGB_ALPHA)] = compressedFeature;
+        _formatFeatures[toNumber(Format::BC2)] = compressedFeature;
+        _formatFeatures[toNumber(Format::BC2_SRGB)] = compressedFeature;
+        _formatFeatures[toNumber(Format::BC3)] = compressedFeature;
+        _formatFeatures[toNumber(Format::BC3_SRGB)] = compressedFeature;
+    }
+
+    _formatFeatures[toNumber(Format::R8)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RG8)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGB8)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGBA8)] |= FormatFeature::VERTEX_ATTRIBUTE;
+
+    _formatFeatures[toNumber(Format::R8SN)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RG8SN)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGB8SN)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGBA8SN)] |= FormatFeature::VERTEX_ATTRIBUTE;
+
+    _formatFeatures[toNumber(Format::R8I)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RG8I)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGB8I)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGBA8I)] |= FormatFeature::VERTEX_ATTRIBUTE;
+
+    _formatFeatures[toNumber(Format::R8UI)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RG8UI)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGB8UI)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGBA8UI)] |= FormatFeature::VERTEX_ATTRIBUTE;
+
+    _formatFeatures[toNumber(Format::R16I)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RG16I)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGB16I)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGBA16I)] |= FormatFeature::VERTEX_ATTRIBUTE;
+
+    _formatFeatures[toNumber(Format::R16UI)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RG16UI)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGB16UI)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGBA16UI)] |= FormatFeature::VERTEX_ATTRIBUTE;
+
+    _formatFeatures[toNumber(Format::R16F)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RG16F)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGB16F)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGBA16F)] |= FormatFeature::VERTEX_ATTRIBUTE;
+
+    _formatFeatures[toNumber(Format::R32UI)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RG32UI)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGB32UI)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGBA32UI)] |= FormatFeature::VERTEX_ATTRIBUTE;
+
+    _formatFeatures[toNumber(Format::R32I)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RG32I)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGB32I)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGBA32I)] |= FormatFeature::VERTEX_ATTRIBUTE;
+
+    _formatFeatures[toNumber(Format::R32F)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RG32F)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGB32F)] |= FormatFeature::VERTEX_ATTRIBUTE;
+    _formatFeatures[toNumber(Format::RGBA32F)] |= FormatFeature::VERTEX_ATTRIBUTE;
+
+    _formatFeatures[toNumber(Format::RGB10A2)] |= FormatFeature::VERTEX_ATTRIBUTE;
+}
+
+SampleCount CCMTLDevice::getMaxSampleCount(Format format, TextureUsage usage, TextureFlags flags) const {
+    const SampleCount sampleCounts[] = {
+        SampleCount::X64,
+        SampleCount::X32,
+        SampleCount::X16,
+        SampleCount::X8,
+        SampleCount::X4,
+        SampleCount::X2,
+    };
+    for (auto sampleCount : sampleCounts) {
+        if  ([_mtlDevice supportsTextureSampleCount: static_cast<uint32_t>(sampleCount)]) {
+            return sampleCount;
+        }
+    }
+    return SampleCount::X1;
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLFramebuffer.h

@@ -0,0 +1,53 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#import "gfx-base/GFXFramebuffer.h"
+
+namespace cc {
+namespace gfx {
+class CCMTLSwapchain;
+class CCMTLFramebuffer final : public Framebuffer {
+public:
+    explicit CCMTLFramebuffer();
+    ~CCMTLFramebuffer();
+    CCMTLFramebuffer(const CCMTLFramebuffer &) = delete;
+    CCMTLFramebuffer(CCMTLFramebuffer &&) = delete;
+    CCMTLFramebuffer &operator=(const CCMTLFramebuffer &) = delete;
+    CCMTLFramebuffer &operator=(CCMTLFramebuffer &&) = delete;
+
+    inline bool isOffscreen() const { return _isOffscreen; }
+    inline CCMTLSwapchain *swapChain() const { return _swapChain; }
+
+protected:
+    void doInit(const FramebufferInfo &info) override;
+    void doDestroy() override;
+
+    bool _isOffscreen = false;
+    CCMTLSwapchain *_swapChain = nullptr;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLFramebuffer.mm

@@ -0,0 +1,66 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#import "MTLFramebuffer.h"
+#import "MTLRenderPass.h"
+#import "MTLTexture.h"
+
+namespace cc {
+namespace gfx {
+
+CCMTLFramebuffer::CCMTLFramebuffer() : Framebuffer() {
+    _typedID = generateObjectID<decltype(this)>();
+}
+
+CCMTLFramebuffer::~CCMTLFramebuffer() {
+    destroy();
+}
+
+void CCMTLFramebuffer::doInit(const FramebufferInfo&) {
+    _isOffscreen = true;
+    for (Texture* tex : _colorTextures) {
+        auto* ccTex = static_cast<CCMTLTexture*>(tex);
+        if (ccTex->swapChain()) {
+            _swapChain = ccTex->swapChain();
+            _isOffscreen = false;
+            break;
+        }
+    }
+
+    if (_depthStencilTexture) {
+        auto* ccTex = static_cast<CCMTLTexture*>(_depthStencilTexture);
+        if (ccTex->swapChain()) {
+            _swapChain = ccTex->swapChain();
+            _isOffscreen = false;
+        }
+    }
+}
+
+void CCMTLFramebuffer::doDestroy() {
+    _swapChain = nullptr;
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLGPUObjects.h

@@ -0,0 +1,340 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#import <Metal/MTLBuffer.h>
+#import <Metal/MTLCommandQueue.h>
+#import <Metal/MTLRenderCommandEncoder.h>
+#import <Metal/MTLSampler.h>
+#import <QuartzCore/CAMetalLayer.h>
+#include <array>
+#import "../../base/Utils.h"
+#import "MTLConfig.h"
+#import "MTLDevice.h"
+#import "MTLUtils.h"
+#include "base/std/container/queue.h"
+
+namespace cc {
+namespace gfx {
+class CCMTLBuffer;
+class CCMTLTexture;
+class CCMTLSampler;
+class CCMTLShader;
+class CCMTLQueue;
+class CCMTLRenderPass;
+class CCMTLFramebuffer;
+class CCMTLInputAssembler;
+class CCMTLPipelineState;
+class CCMTLSemaphore;
+class CCMTLCommandBuffer;
+
+namespace {
+constexpr size_t MegaBytesToBytes = 1024 * 1024;
+}
+
+constexpr size_t MAX_COLORATTACHMENTS = 16u;
+
+struct CCMTLGPUDescriptorSetLayout {
+    DescriptorSetLayoutBindingList bindings;
+    ccstd::vector<uint32_t> dynamicBindings;
+    ccstd::vector<uint32_t> descriptorIndices;
+    ccstd::vector<uint32_t> bindingIndices;
+    uint32_t descriptorCount = 0;
+};
+typedef ccstd::vector<CCMTLGPUDescriptorSetLayout *> MTLGPUDescriptorSetLayoutList;
+
+struct CCMTLGPUPipelineLayout {
+    MTLGPUDescriptorSetLayoutList setLayouts;
+    ccstd::vector<ccstd::vector<int>> dynamicOffsetIndices;
+};
+
+struct CCMTLGPUUniformBlock {
+    ccstd::string name;
+    uint32_t set = INVALID_BINDING;
+    uint32_t binding = INVALID_BINDING;
+    uint32_t mappedBinding = INVALID_BINDING;
+    ShaderStageFlags stages = ShaderStageFlagBit::NONE;
+    size_t size = 0;
+    uint32_t count = 0;
+};
+
+struct CCMTLGPUSamplerBlock {
+    ccstd::string name;
+    uint32_t set = INVALID_BINDING;
+    uint32_t binding = INVALID_BINDING;
+    uint32_t textureBinding = INVALID_BINDING;
+    uint32_t samplerBinding = INVALID_BINDING;
+    ShaderStageFlags stages = ShaderStageFlagBit::NONE;
+    Type type = Type::UNKNOWN;
+    uint32_t count = 0;
+};
+
+struct CCMTLGPUSubpassAttachment {
+    ccstd::string name;
+    uint32_t set = INVALID_BINDING;
+    uint32_t binding = INVALID_BINDING;
+};
+
+struct ResourceBinding {
+    uint32_t bufferBinding{0};
+    uint32_t textureBinding{0};
+    uint32_t samplerBinding{0};
+};
+
+struct CCMTLGPUShader {
+    ccstd::unordered_map<uint32_t, CCMTLGPUUniformBlock> blocks;
+    ccstd::unordered_map<uint32_t, CCMTLGPUSamplerBlock> samplers;
+
+    ccstd::unordered_map<uint32_t, ResourceBinding> resourceBinding;
+
+    ccstd::vector<CCMTLGPUSubpassAttachment> inputs;
+    ccstd::vector<CCMTLGPUSubpassAttachment> outputs;
+
+    std::array<uint32_t, 3> workGroupSize{0, 0, 0};
+
+    NSString *shaderSrc = nil;
+    bool specializeColor = true;
+
+    uint32_t bufferIndex = 0;
+    uint32_t samplerIndex = 0;
+
+    std::string name;
+};
+
+struct CCMTLGPUPipelineState {
+    MTLCullMode cullMode;
+    MTLWinding winding;
+    MTLTriangleFillMode fillMode;
+    MTLDepthClipMode depthClipMode;
+    MTLPrimitiveType primitiveType;
+    id<MTLRenderPipelineState> mtlRenderPipelineState = nil;
+    id<MTLDepthStencilState> mtlDepthStencilState = nil;
+    id<MTLComputePipelineState> mtlComputePipelineState = nil;
+    uint32_t stencilRefFront = 0;
+    uint32_t stencilRefBack = 0;
+    ccstd::vector<std::tuple<int /**vertexBufferBindingIndex*/, uint32_t /**stream*/>> vertexBufferBindingInfo;
+    const CCMTLGPUPipelineLayout *gpuPipelineLayout = nullptr;
+    const CCMTLGPUShader *gpuShader = nullptr;
+};
+
+struct CCMTLGPUBuffer {
+    uint32_t stride = 0;
+    uint32_t count = 0;
+    uint32_t instanceSize = 0;
+    uint32_t startOffset = 0;
+    id<MTLBuffer> mtlBuffer = nil;
+    uint8_t lastUpdateCycle = 0;
+    uint8_t *mappedData = nullptr;
+};
+
+struct CCMTLGPUTextureObject {
+    TextureInfo info;
+    id<MTLTexture> mtlTexture;
+};
+
+struct CCMTLGPUTextureViewObject {
+    TextureViewInfo viewInfo;
+    id<MTLTexture> mtlTextureView;
+};
+
+struct CCMTLGPUInputAssembler {
+    //
+};
+
+struct CCMTLGPUDescriptor {
+    DescriptorType type = DescriptorType::UNKNOWN;
+    CCMTLBuffer *buffer = nullptr;
+    CCMTLTexture *texture = nullptr;
+    CCMTLSampler *sampler = nullptr;
+};
+typedef ccstd::vector<CCMTLGPUDescriptor> MTLGPUDescriptorList;
+
+struct CCMTLGPUDescriptorSet {
+    MTLGPUDescriptorList gpuDescriptors;
+    const ccstd::vector<uint32_t> *descriptorIndices = nullptr;
+};
+
+class CCMTLGPUStagingBufferPool final {
+public:
+    CCMTLGPUStagingBufferPool(id<MTLDevice> device)
+    : _device(device) {}
+
+    ~CCMTLGPUStagingBufferPool() {
+        for (auto &buffer : _pool) {
+            [buffer.mtlBuffer release];
+            buffer.mtlBuffer = nil;
+        }
+
+        _pool.clear();
+    }
+
+    inline void alloc(CCMTLGPUBuffer *gpuBuffer) { alloc(gpuBuffer, 1); }
+    void alloc(CCMTLGPUBuffer *gpuBuffer, uint32_t alignment) {
+        size_t bufferCount = _pool.size();
+        Buffer *buffer = nullptr;
+        uint32_t offset = 0;
+        for (size_t idx = 0; idx < bufferCount; idx++) {
+            auto *cur = &_pool[idx];
+            offset = mu::alignUp(cur->curOffset, alignment);
+            if (gpuBuffer->instanceSize + offset <= [cur->mtlBuffer length]) {
+                buffer = cur;
+                break;
+            }
+        }
+        if (!buffer) {
+            uint32_t needs = mu::alignUp(gpuBuffer->instanceSize, MegaBytesToBytes);
+
+            _pool.resize(bufferCount + 1);
+            buffer = &_pool.back();
+
+            buffer->mtlBuffer = [_device newBufferWithLength:needs options:MTLResourceStorageModeShared];
+            buffer->mappedData = (uint8_t *)buffer->mtlBuffer.contents;
+            offset = 0;
+        }
+        gpuBuffer->mtlBuffer = buffer->mtlBuffer;
+        gpuBuffer->startOffset = offset;
+        gpuBuffer->mappedData = buffer->mappedData + offset;
+        buffer->curOffset = offset + gpuBuffer->instanceSize;
+    }
+
+    void reset() {
+        for (auto &buffer : _pool) {
+            buffer.curOffset = 0;
+        }
+    }
+
+    void shrinkSize() {
+        for (auto iter = _pool.begin(); iter != _pool.end() && _pool.size() > 1;) {
+            if (iter->curOffset == 0) {
+                [iter->mtlBuffer release];
+                iter = _pool.erase(iter);
+            } else {
+                ++iter;
+            }
+        }
+    }
+
+protected:
+    struct Buffer {
+        id<MTLBuffer> mtlBuffer = nil;
+        uint8_t *mappedData = nullptr;
+        uint32_t curOffset = 0;
+    };
+
+    id<MTLDevice> _device = nil;
+    ccstd::vector<Buffer> _pool;
+};
+
+struct CCMTLGPUBufferImageCopy {
+    NSUInteger sourceBytesPerRow = 0;
+    NSUInteger sourceBytesPerImage = 0;
+    MTLSize sourceSize = {0, 0, 0};
+    NSUInteger destinationSlice = 0;
+    NSUInteger destinationLevel = 0;
+    MTLOrigin destinationOrigin = {0, 0, 0};
+};
+
+//destroy GPU resource only, delete the owner object mannually.
+class CCMTLGPUGarbageCollectionPool final {
+    using GCFunc = std::function<void(void)>;
+
+    CCMTLGPUGarbageCollectionPool() = default;
+
+public:
+    static CCMTLGPUGarbageCollectionPool *getInstance() {
+        static CCMTLGPUGarbageCollectionPool gcPoolSingleton;
+        return &gcPoolSingleton;
+    }
+
+    void initialize(std::function<uint8_t(void)> getFrameIndex) {
+        CC_ASSERT(getFrameIndex);
+        _getFrameIndex = getFrameIndex;
+    }
+
+    void collect(std::function<void(void)> destroyFunc) {
+        uint8_t curFrameIndex = _getFrameIndex();
+        _releaseQueue[curFrameIndex].push(destroyFunc);
+    }
+
+    void clear(uint8_t currentFrameIndex) {
+        CC_ASSERT_LT(currentFrameIndex, MAX_FRAMES_IN_FLIGHT);
+        while (!_releaseQueue[currentFrameIndex].empty()) {
+            auto &&gcFunc = _releaseQueue[currentFrameIndex].front();
+            gcFunc();
+            _releaseQueue[currentFrameIndex].pop();
+        }
+    }
+
+    void flush() {
+        for (size_t i = 0; i < MAX_FRAMES_IN_FLIGHT; i++) {
+            while (!_releaseQueue[i].empty()) {
+                auto &&gcFunc = _releaseQueue[i].front();
+                gcFunc();
+                _releaseQueue[i].pop();
+            }
+        }
+    }
+
+protected:
+    //avoid cross-reference with CCMTLDevice
+    std::function<uint8_t(void)> _getFrameIndex;
+    ccstd::queue<GCFunc> _releaseQueue[MAX_FRAMES_IN_FLIGHT];
+};
+
+struct CCMTLGPUSwapChainObject {
+    id<CAMetalDrawable> currentDrawable = nil;
+    CAMetalLayer *mtlLayer = nullptr;
+};
+
+struct CCMTLGPUQueueObject {
+    id<MTLCommandQueue> mtlCommandQueue = nil;
+    uint32_t numDrawCalls = 0;
+    uint32_t numInstances = 0;
+    uint32_t numTriangles = 0;
+};
+
+struct CCMTLGPUCommandBufferObject {
+    CCMTLRenderPass *renderPass = nullptr;
+    CCMTLFramebuffer *fbo = nullptr;
+    CCMTLInputAssembler *inputAssembler = nullptr;
+    CCMTLPipelineState *pipelineState = nullptr;
+    id<MTLCommandBuffer> mtlCommandBuffer = nil;
+    bool isSecondary = false;
+};
+
+struct CCMTLGPUDeviceObject {
+    CCMTLCommandBuffer *_transferCmdBuffer{nullptr};
+};
+
+struct CCMTLGPUQueryPool {
+    QueryType type = QueryType::OCCLUSION;
+    uint32_t maxQueryObjects = 0;
+    bool forceWait = true;
+    id<MTLBuffer> visibilityResultBuffer = nil;
+    CCMTLSemaphore *semaphore = nullptr;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLInputAssembler.h

@@ -0,0 +1,53 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#include "gfx-base/GFXInputAssembler.h"
+
+namespace cc {
+namespace gfx {
+
+class CCMTLGPUInputAssembler;
+
+class CCMTLInputAssembler final : public InputAssembler {
+public:
+    explicit CCMTLInputAssembler();
+    ~CCMTLInputAssembler();
+    CCMTLInputAssembler(const CCMTLInputAssembler &) = delete;
+    CCMTLInputAssembler(CCMTLInputAssembler &&) = delete;
+    CCMTLInputAssembler &operator=(const CCMTLInputAssembler &) = delete;
+    CCMTLInputAssembler &operator=(CCMTLInputAssembler &&) = delete;
+
+protected:
+    friend class CCMTLQueue;
+
+    void doInit(const InputAssemblerInfo &info) override;
+    void doDestroy() override;
+
+    CCMTLGPUInputAssembler *_GPUInputAssembler = nullptr;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLInputAssembler.mm

@@ -0,0 +1,53 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#include "MTLBuffer.h"
+#include "MTLGPUObjects.h"
+#include "MTLInputAssembler.h"
+
+namespace cc {
+namespace gfx {
+
+CCMTLInputAssembler::CCMTLInputAssembler() : InputAssembler() {
+    _typedID = generateObjectID<decltype(this)>();
+}
+
+CCMTLInputAssembler::~CCMTLInputAssembler() {
+    destroy();
+}
+
+void CCMTLInputAssembler::doInit(const InputAssemblerInfo &info) {
+    _GPUInputAssembler = ccnew CCMTLGPUInputAssembler;
+    if (!_GPUInputAssembler) {
+        return;
+    }
+}
+
+void CCMTLInputAssembler::doDestroy() {
+    CC_SAFE_DELETE(_GPUInputAssembler);
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLPipelineLayout.h

@@ -0,0 +1,48 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#include "gfx-base/GFXPipelineLayout.h"
+
+namespace cc {
+namespace gfx {
+
+class CCMTLGPUPipelineLayout;
+
+class CCMTLPipelineLayout final : public PipelineLayout {
+public:
+    CCMTLPipelineLayout();
+    ~CCMTLPipelineLayout();
+
+    inline CCMTLGPUPipelineLayout *gpuPipelineLayout() const { return _gpuPipelineLayout; }
+
+protected:
+    void doInit(const PipelineLayoutInfo &info) override;
+    void doDestroy() override;
+
+    CCMTLGPUPipelineLayout *_gpuPipelineLayout = nullptr;
+};
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLPipelineLayout.mm

@@ -0,0 +1,66 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#include "MTLDescriptorSetLayout.h"
+#include "MTLGPUObjects.h"
+#include "MTLPipelineLayout.h"
+namespace cc {
+namespace gfx {
+
+CCMTLPipelineLayout::CCMTLPipelineLayout() : PipelineLayout() {
+    _typedID = generateObjectID<decltype(this)>();
+}
+
+CCMTLPipelineLayout::~CCMTLPipelineLayout() {
+    destroy();
+}
+
+void CCMTLPipelineLayout::doInit(const PipelineLayoutInfo &info) {
+    const auto setCount = _setLayouts.size();
+    _gpuPipelineLayout = ccnew CCMTLGPUPipelineLayout;
+    _gpuPipelineLayout->dynamicOffsetIndices.resize(setCount);
+
+    for (size_t i = 0; i < setCount; i++) {
+        const auto *setLayout = _setLayouts[i];
+        CC_ASSERT_NOT_NULL(setLayout);
+        auto gpuDescriptorSetLayout = static_cast<const CCMTLDescriptorSetLayout *>(setLayout)->gpuDescriptorSetLayout();
+        auto dynamicCount = gpuDescriptorSetLayout->dynamicBindings.size();
+        auto &indices = _gpuPipelineLayout->dynamicOffsetIndices[i];
+        indices.assign(setLayout->getBindingIndices().size(), -1);
+
+        for (int j = 0; j < dynamicCount; j++) {
+            auto binding = gpuDescriptorSetLayout->dynamicBindings[j];
+            if (indices[binding] < 0) indices[binding] = j;
+        }
+        _gpuPipelineLayout->setLayouts.emplace_back(gpuDescriptorSetLayout);
+    }
+}
+
+void CCMTLPipelineLayout::doDestroy() {
+    CC_SAFE_DELETE(_gpuPipelineLayout);
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLPipelineState.h

@@ -0,0 +1,74 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#include "gfx-base/GFXPipelineState.h"
+
+#include "MTLGPUObjects.h"
+
+#import <Metal/MTLDepthStencil.h>
+#import <Metal/MTLRenderPipeline.h>
+
+namespace cc {
+namespace gfx {
+
+class CCMTLRenderPass;
+
+class CCMTLPipelineState final : public PipelineState {
+public:
+    explicit CCMTLPipelineState();
+    ~CCMTLPipelineState();
+    CCMTLPipelineState(const CCMTLPipelineState &) = delete;
+    CCMTLPipelineState(CCMTLPipelineState &&) = delete;
+    CCMTLPipelineState &operator=(const CCMTLPipelineState &) = delete;
+    CCMTLPipelineState &operator=(CCMTLPipelineState &&) = delete;
+
+    inline CCMTLGPUPipelineState *getGPUPipelineState() const { return _GPUPipelineState; }
+
+    void check(CCMTLRenderPass *renderPass = nullptr);
+
+protected:
+    void doInit(const PipelineStateInfo &info) override;
+    void doDestroy() override;
+
+    bool initRenderPipeline();
+    bool createMTLDepthStencilState();
+    bool createGPUPipelineState();
+    bool createMTLComputePipelineState();
+    bool createMTLRenderPipelineState();
+    bool setVertexDescriptor(MTLRenderPipelineDescriptor *);
+    bool setMTLFunctionsAndFormats(MTLRenderPipelineDescriptor *);
+    bool setBlendStates(MTLRenderPipelineDescriptor *);
+    bool createMTLRenderPipeline(MTLRenderPipelineDescriptor *);
+
+    bool _renderPipelineReady = false;
+    id<MTLRenderPipelineState> _mtlRenderPipelineState = nil;
+    id<MTLDepthStencilState> _mtlDepthStencilState = nil;
+    id<MTLComputePipelineState> _mtlComputePipeline = nil;
+    CCMTLGPUPipelineState *_GPUPipelineState = nullptr;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLPipelineState.mm

@@ -0,0 +1,402 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#import "base/std/container/set.h"
+#import "base/Log.h"
+#import "MTLDevice.h"
+#import "MTLGPUObjects.h"
+#import "MTLPipelineLayout.h"
+#import "MTLPipelineState.h"
+#import "MTLRenderPass.h"
+#import "MTLSampler.h"
+#import "MTLShader.h"
+#import "MTLTexture.h"
+#import "MTLUtils.h"
+
+#import <Metal/MTLComputePipeline.h>
+#import <Metal/MTLDevice.h>
+#import <Metal/MTLVertexDescriptor.h>
+
+namespace cc {
+namespace gfx {
+
+CCMTLPipelineState::CCMTLPipelineState() : PipelineState() {
+    _typedID = generateObjectID<decltype(this)>();
+}
+
+CCMTLPipelineState::~CCMTLPipelineState() {
+    destroy();
+}
+
+void CCMTLPipelineState::doInit(const PipelineStateInfo &info) {
+    createGPUPipelineState();
+}
+
+void CCMTLPipelineState::doDestroy() {
+    CC_SAFE_DELETE(_GPUPipelineState);
+
+    id<MTLRenderPipelineState> renderPipelineState = _mtlRenderPipelineState;
+    _mtlRenderPipelineState = nil;
+    id<MTLDepthStencilState> depthStencilState = _mtlDepthStencilState;
+    _mtlDepthStencilState = nil;
+
+    std::function<void(void)> destroyFunc = [=]() {
+        if (renderPipelineState) {
+            [renderPipelineState release];
+        }
+        if (depthStencilState) {
+            [depthStencilState release];
+        }
+    };
+    CCMTLGPUGarbageCollectionPool::getInstance()->collect(destroyFunc);
+}
+
+bool CCMTLPipelineState::initRenderPipeline() {
+    if (!createMTLRenderPipelineState()) {
+        return false;
+    }
+    // Application can run with wrong depth/stencil state.
+    if (!createMTLDepthStencilState()) {
+        return false;
+    }
+
+    _GPUPipelineState->mtlDepthStencilState = _mtlDepthStencilState;
+    _GPUPipelineState->mtlRenderPipelineState = _mtlRenderPipelineState;
+    _GPUPipelineState->cullMode = static_cast<MTLCullMode>(mu::toMTLCullMode(_rasterizerState.cullMode));
+    _GPUPipelineState->fillMode = static_cast<MTLTriangleFillMode>(mu::toMTLTriangleFillMode(_rasterizerState.polygonMode));
+    _GPUPipelineState->depthClipMode = static_cast<MTLDepthClipMode>(mu::toMTLDepthClipMode(_rasterizerState.isDepthClip != 0));
+    _GPUPipelineState->winding = static_cast<MTLWinding>(mu::toMTLWinding(_rasterizerState.isFrontFaceCCW != 0));
+    _GPUPipelineState->stencilRefFront = _depthStencilState.stencilRefFront;
+    _GPUPipelineState->stencilRefBack = _depthStencilState.stencilRefBack;
+    _GPUPipelineState->primitiveType = mu::toMTLPrimitiveType(_primitive);
+    if (_pipelineLayout)
+        _GPUPipelineState->gpuPipelineLayout = static_cast<CCMTLPipelineLayout*>(_pipelineLayout)->gpuPipelineLayout();
+    _GPUPipelineState->gpuShader = static_cast<CCMTLShader *>(_shader)->gpuShader(
+        static_cast<CCMTLRenderPass*>(_renderPass), _subpass);
+
+    _renderPipelineReady = true;
+    return true;
+}
+
+void CCMTLPipelineState::check(CCMTLRenderPass *renderPass) {
+    if (renderPass)
+        _renderPass = renderPass;
+    if (!_renderPipelineReady) {
+        initRenderPipeline();
+    }
+}
+
+bool CCMTLPipelineState::createGPUPipelineState() {
+    _GPUPipelineState = ccnew CCMTLGPUPipelineState;
+    if (!_GPUPipelineState) {
+        CC_LOG_ERROR("CCMTLPipelineState: new CCMTLGPUPipelineState failed.");
+        return false;
+    }
+
+    if (_bindPoint == PipelineBindPoint::GRAPHICS) {
+        if (_renderPass->getSubpasses().empty()) {
+            initRenderPipeline();
+        }
+    } else if (_bindPoint == PipelineBindPoint::COMPUTE) {
+        if (!createMTLComputePipelineState()) {
+            return false;
+        }
+        _GPUPipelineState->mtlComputePipelineState = _mtlComputePipeline;
+        _GPUPipelineState->gpuShader = static_cast<CCMTLShader *>(_shader)->gpuShader(nullptr, 0);
+        if (_pipelineLayout)
+            _GPUPipelineState->gpuPipelineLayout = static_cast<CCMTLPipelineLayout *>(_pipelineLayout)->gpuPipelineLayout();
+    }
+
+    return true;
+}
+
+bool CCMTLPipelineState::createMTLComputePipelineState() {
+    //create with function
+    id<MTLDevice> mtlDevice = id<MTLDevice>(CCMTLDevice::getInstance()->getMTLDevice());
+    NSError *err;
+    _mtlComputePipeline = [mtlDevice newComputePipelineStateWithFunction:((CCMTLShader *)_shader)->getComputeMTLFunction()
+                                                                   error:&err];
+    if (!_mtlComputePipeline) {
+        CC_LOG_ERROR("Create compute pipeline failed: %s", [err.localizedDescription UTF8String]);
+        return false;
+    }
+    return true;
+}
+
+bool CCMTLPipelineState::createMTLDepthStencilState() {
+    const auto& colorAttachments = _renderPass->getColorAttachments();
+    bool hasDS = std::any_of(colorAttachments.begin(), colorAttachments.end(), [](const ColorAttachment& attachemnt){
+        return attachemnt.format == Format::DEPTH ||attachemnt.format == Format::DEPTH_STENCIL;
+    });
+    hasDS |= _renderPass->getDepthStencilAttachment().format != Format::UNKNOWN;
+
+    if(!hasDS) {
+        // default nil
+        return true;
+    }
+
+    MTLDepthStencilDescriptor *descriptor = [[MTLDepthStencilDescriptor alloc] init];
+    if (descriptor == nil) {
+        CC_LOG_ERROR("CCMTLPipelineState: MTLDepthStencilDescriptor could not be allocated.");
+        return false;
+    }
+
+    descriptor.depthWriteEnabled = _depthStencilState.depthWrite != 0;
+    if (!_depthStencilState.depthTest)
+        descriptor.depthCompareFunction = MTLCompareFunctionAlways;
+    else
+        descriptor.depthCompareFunction = mu::toMTLCompareFunction(_depthStencilState.depthFunc);
+
+    if (_depthStencilState.stencilTestFront) {
+        descriptor.frontFaceStencil.stencilCompareFunction = mu::toMTLCompareFunction(_depthStencilState.stencilFuncFront);
+        descriptor.frontFaceStencil.readMask = _depthStencilState.stencilReadMaskFront;
+        descriptor.frontFaceStencil.writeMask = _depthStencilState.stencilWriteMaskFront;
+        descriptor.frontFaceStencil.stencilFailureOperation = mu::toMTLStencilOperation(_depthStencilState.stencilFailOpFront);
+        descriptor.frontFaceStencil.depthFailureOperation = mu::toMTLStencilOperation(_depthStencilState.stencilZFailOpFront);
+        descriptor.frontFaceStencil.depthStencilPassOperation = mu::toMTLStencilOperation(_depthStencilState.stencilPassOpFront);
+    } else {
+        descriptor.frontFaceStencil = nil;
+    }
+
+    if (_depthStencilState.stencilTestBack) {
+        descriptor.backFaceStencil.stencilCompareFunction = mu::toMTLCompareFunction(_depthStencilState.stencilFuncBack);
+        descriptor.backFaceStencil.readMask = _depthStencilState.stencilReadMaskBack;
+        descriptor.backFaceStencil.writeMask = _depthStencilState.stencilWriteMaskBack;
+        descriptor.backFaceStencil.stencilFailureOperation = mu::toMTLStencilOperation(_depthStencilState.stencilFailOpBack);
+        descriptor.backFaceStencil.depthFailureOperation = mu::toMTLStencilOperation(_depthStencilState.stencilZFailOpBack);
+        descriptor.backFaceStencil.depthStencilPassOperation = mu::toMTLStencilOperation(_depthStencilState.stencilPassOpBack);
+    } else {
+        descriptor.backFaceStencil = nil;
+    }
+
+    id<MTLDevice> mtlDevice = id<MTLDevice>(CCMTLDevice::getInstance()->getMTLDevice());
+    _mtlDepthStencilState = [mtlDevice newDepthStencilStateWithDescriptor:descriptor];
+    [descriptor release];
+
+    if (!_mtlDepthStencilState) {
+        CC_LOG_ERROR("Failed to create MTLDepthStencilState.");
+        return false;
+    }
+    return true;
+}
+
+bool CCMTLPipelineState::createMTLRenderPipelineState() {
+    bool ret = true;
+    MTLRenderPipelineDescriptor *descriptor = [[MTLRenderPipelineDescriptor alloc] init];
+    if (descriptor == nil) {
+        CC_LOG_ERROR("CCMTLPipelineState: MTLRenderPipelineDescriptor could not be allocated.");
+        ret = false;
+    }
+    if(ret) ret = setMTLFunctionsAndFormats(descriptor);
+    if(ret) ret = setVertexDescriptor(descriptor);
+    if(ret) ret = setBlendStates(descriptor);
+    if(ret) ret = createMTLRenderPipeline(descriptor);
+    [descriptor release];
+
+    if(!ret) {
+        CC_LOG_ERROR("Failed to create pipeline state, please check if shader/pileinelayout match with each other!");
+    }
+
+    return ret;
+}
+
+//TODO: reconstruction
+bool CCMTLPipelineState::setVertexDescriptor(MTLRenderPipelineDescriptor *descriptor) {
+    bool res = true;
+    auto activeAttributes = static_cast<CCMTLShader *>(_shader)->getAttributes();
+    ccstd::vector<std::tuple<int /**vertexBufferBindingIndex*/, uint32_t /**stream*/>> layouts;
+    ccstd::unordered_map<int /**vertexBufferBindingIndex*/, std::tuple<uint32_t /**stride*/, bool /**isInstanced*/>> map;
+    ccstd::vector<uint32_t> streamOffsets(CCMTLDevice::getInstance()->getCapabilities().maxVertexAttributes, 0u);
+    ccstd::vector<bool> activeAttribIdx(activeAttributes.size(), false);
+    for (const auto &inputAttrib : _inputState.attributes) {
+        auto bufferIndex = static_cast<CCMTLShader *>(_shader)->getAvailableBufferBindingIndex(ShaderStageFlagBit::VERTEX, inputAttrib.stream);
+
+        for (auto i = 0; i < activeAttributes.size(); i++) {
+            const auto &activeAttribute = activeAttributes[i];
+            if (inputAttrib.name == activeAttribute.name) {
+                descriptor.vertexDescriptor.attributes[activeAttribute.location].format = mu::toMTLVertexFormat(inputAttrib.format, inputAttrib.isNormalized);
+                descriptor.vertexDescriptor.attributes[activeAttribute.location].offset = streamOffsets[inputAttrib.stream];
+                descriptor.vertexDescriptor.attributes[activeAttribute.location].bufferIndex = bufferIndex;
+                auto tuple = std::make_tuple(bufferIndex, inputAttrib.stream);
+                if (std::find(layouts.begin(), layouts.end(), tuple) == layouts.end())
+                    layouts.emplace_back(tuple);
+                activeAttribIdx[i] = true;
+                break;
+            }
+        }
+
+        uint32_t attributeSize = GFX_FORMAT_INFOS[(int)inputAttrib.format].size;
+        // NOTE: Metal requires 4 bytes alignment for attribute
+        if (attributeSize > 0) {
+            attributeSize = utils::alignTo(attributeSize, 4U);
+        }
+
+        streamOffsets[inputAttrib.stream] += attributeSize;
+        map[bufferIndex] = std::make_tuple(streamOffsets[inputAttrib.stream], inputAttrib.isInstanced);
+    }
+
+    for (auto i = 0; i < activeAttribIdx.size(); i++) {
+        if (activeAttribIdx[i]) continue;
+
+        const auto &dummy = activeAttributes[i];
+        descriptor.vertexDescriptor.attributes[dummy.location].format = MTLVertexFormatFloat;
+        descriptor.vertexDescriptor.attributes[dummy.location].offset = 0;
+        descriptor.vertexDescriptor.attributes[dummy.location].bufferIndex = static_cast<CCMTLShader *>(_shader)->getAvailableBufferBindingIndex(ShaderStageFlagBit::VERTEX, dummy.stream);
+        CC_LOG_WARNING("Attribute %s is missing, add a dummy data for it.", dummy.name.c_str());
+    }
+
+    if(layouts.empty()) {
+        res = false;
+    }
+
+    // layouts
+    for (const auto &layout : layouts) {
+        auto index = std::get<0>(layout);
+        descriptor.vertexDescriptor.layouts[index].stride = std::get<0>(map[index]);
+        descriptor.vertexDescriptor.layouts[index].stepFunction = std::get<1>(map[index]) ? MTLVertexStepFunctionPerInstance : MTLVertexStepFunctionPerVertex;
+        descriptor.vertexDescriptor.layouts[index].stepRate = 1;
+        //to improve performance: https://developer.apple.com/documentation/metal/mtlpipelinebufferdescriptor?language=objc
+        if (@available(iOS 11.0, macOS 10.13, *)) {
+            descriptor.vertexBuffers[index].mutability = MTLMutabilityImmutable;
+        }
+    }
+
+    _GPUPipelineState->vertexBufferBindingInfo = std::move(layouts);
+    return res;
+}
+
+bool CCMTLPipelineState::setMTLFunctionsAndFormats(MTLRenderPipelineDescriptor *descriptor) {
+    auto *mtlPass = static_cast<CCMTLRenderPass*>(_renderPass);
+    const SubpassInfoList &subpasses = _renderPass->getSubpasses();
+    const ColorAttachmentList &colorAttachments = _renderPass->getColorAttachments();
+    const auto &ccShader = static_cast<CCMTLShader *>(_shader);
+
+    const CCMTLGPUShader *gpuShader = ccShader->gpuShader(mtlPass, _subpass);
+
+    MTLPixelFormat mtlPixelFormat;
+    ccstd::set<uint32_t> inputs;
+    uint32_t depthStencilTexIndex = INVALID_BINDING;
+    if (!subpasses.empty()) {
+        for (size_t passIndex = 0; passIndex < subpasses.size(); ++passIndex) {
+            const SubpassInfo &subpass = subpasses[passIndex];
+            depthStencilTexIndex = subpass.depthStencil;
+            for (size_t i = 0; i < subpass.inputs.size(); ++i) {
+                uint32_t input = subpass.inputs[i];
+
+                if (inputs.find(input) == inputs.end()) {
+                    inputs.insert(input);
+                    if(_renderPass->getColorAttachments()[input].format == Format::DEPTH ||
+                       _renderPass->getColorAttachments()[input].format == Format::DEPTH_STENCIL) {
+                        continue;
+                    }
+                    mtlPixelFormat = mu::toMTLPixelFormat(colorAttachments[input].format);
+                    descriptor.colorAttachments[input].pixelFormat = mtlPixelFormat;
+                }
+            }
+
+            for (size_t i = 0; i < subpass.colors.size(); ++i) {
+                uint32_t output = subpass.colors[i];
+                if(output >= colorAttachments.size()) {
+                    depthStencilTexIndex = output;
+                    continue;
+                }
+                mtlPixelFormat = mu::toMTLPixelFormat(colorAttachments[output].format);
+                descriptor.colorAttachments[output].pixelFormat = mtlPixelFormat;
+            }
+        }
+    } else {
+        for (size_t i = 0; i < colorAttachments.size(); ++i) {
+            mtlPixelFormat = mu::toMTLPixelFormat(colorAttachments[i].format);
+            descriptor.colorAttachments[i].pixelFormat = mtlPixelFormat;
+            depthStencilTexIndex = INVALID_BINDING;
+        }
+    }
+
+    SampleCount sample = SampleCount::X1;
+    Format depthStencilFormat;
+    if (depthStencilTexIndex != INVALID_BINDING && depthStencilTexIndex < _renderPass->getColorAttachments().size()) {
+        sample = _renderPass->getColorAttachments()[depthStencilTexIndex].sampleCount;
+        depthStencilFormat = _renderPass->getColorAttachments()[depthStencilTexIndex].format;
+    } else {
+        sample = _renderPass->getDepthStencilAttachment().sampleCount;
+        depthStencilFormat = _renderPass->getDepthStencilAttachment().format;
+    }
+    descriptor.sampleCount = mu::toMTLSampleCount(sample);
+
+    auto *ccMTLShader = static_cast<CCMTLShader *>(_shader);
+    descriptor.vertexFunction = ccMTLShader->getVertMTLFunction();
+    descriptor.fragmentFunction = ccMTLShader->getFragmentMTLFunction();
+
+    mtlPixelFormat = mu::toMTLPixelFormat(depthStencilFormat);
+    if (mtlPixelFormat != MTLPixelFormatInvalid) {
+        descriptor.depthAttachmentPixelFormat = mtlPixelFormat;
+
+        if (depthStencilFormat == Format::DEPTH_STENCIL)
+            descriptor.stencilAttachmentPixelFormat = mtlPixelFormat;
+    }
+    return true;
+}
+
+bool CCMTLPipelineState::setBlendStates(MTLRenderPipelineDescriptor *descriptor) {
+    //FIXME: how to handle these two attributes?
+    //    BlendState::isIndepend
+    //    BlendState::blendColor;
+
+    descriptor.alphaToCoverageEnabled = _blendState.isA2C != 0;
+
+    int i = 0;
+    for (const auto blendTarget : _blendState.targets) {
+        MTLRenderPipelineColorAttachmentDescriptor *colorDescriptor = descriptor.colorAttachments[i];
+        colorDescriptor.writeMask = mu::toMTLColorWriteMask(blendTarget.blendColorMask);
+        colorDescriptor.blendingEnabled = blendTarget.blend != 0;
+        if (!blendTarget.blend)
+            continue;
+
+        colorDescriptor.sourceRGBBlendFactor = mu::toMTLBlendFactor(blendTarget.blendSrc);
+        colorDescriptor.destinationRGBBlendFactor = mu::toMTLBlendFactor(blendTarget.blendDst);
+        colorDescriptor.rgbBlendOperation = mu::toMTLBlendOperation(blendTarget.blendEq);
+        colorDescriptor.sourceAlphaBlendFactor = mu::toMTLBlendFactor(blendTarget.blendSrcAlpha);
+        colorDescriptor.destinationAlphaBlendFactor = mu::toMTLBlendFactor(blendTarget.blendDstAlpha);
+        colorDescriptor.alphaBlendOperation = mu::toMTLBlendOperation(blendTarget.blendAlphaEq);
+
+        ++i;
+    }
+    return true;
+}
+
+bool CCMTLPipelineState::createMTLRenderPipeline(MTLRenderPipelineDescriptor *descriptor) {
+    id<MTLDevice> mtlDevice = id<MTLDevice>(CCMTLDevice::getInstance()->getMTLDevice());
+    NSError *nsError = nil;
+    _mtlRenderPipelineState = [mtlDevice newRenderPipelineStateWithDescriptor:descriptor error:&nsError];
+    if (!_mtlRenderPipelineState) {
+        CC_LOG_ERROR("Failed to create MTLRenderPipelineState: %s", [nsError.localizedDescription UTF8String]);
+        return false;
+    }
+
+    return true;
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLQueryPool.h

@@ -0,0 +1,53 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#import "gfx-base/GFXQueryPool.h"
+
+namespace cc {
+namespace gfx {
+
+struct CCMTLGPUQueryPool;
+
+class CCMTLQueryPool final : public QueryPool {
+public:
+    CCMTLQueryPool();
+    ~CCMTLQueryPool() override;
+
+    inline CCMTLGPUQueryPool *gpuQueryPool() const { return _gpuQueryPool; }
+
+protected:
+    friend class CCMTLCommandBuffer;
+    friend class CCMTLDevice;
+
+    void doInit(const QueryPoolInfo &info) override;
+    void doDestroy() override;
+
+    CCMTLGPUQueryPool *_gpuQueryPool{nullptr};
+    ccstd::vector<uint32_t> _ids;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLQueryPool.mm

@@ -0,0 +1,77 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#import "MTLCommandBuffer.h"
+#import "MTLDevice.h"
+#import "MTLFramebuffer.h"
+#import "MTLGPUObjects.h"
+#import "MTLQueryPool.h"
+#import "MTLSemaphore.h"
+#import "MTLSwapchain.h"
+
+namespace cc {
+namespace gfx {
+
+CCMTLQueryPool::CCMTLQueryPool() {
+    _typedID = generateObjectID<decltype(this)>();
+}
+
+CCMTLQueryPool::~CCMTLQueryPool() {
+    destroy();
+}
+
+void CCMTLQueryPool::doInit(const QueryPoolInfo& info) {
+    id<MTLDevice> mtlDevice = id<MTLDevice>(CCMTLDevice::getInstance()->getMTLDevice());
+    _gpuQueryPool = ccnew CCMTLGPUQueryPool;
+    _gpuQueryPool->type = _type;
+    _gpuQueryPool->maxQueryObjects = _maxQueryObjects;
+    _gpuQueryPool->forceWait = _forceWait;
+    _gpuQueryPool->visibilityResultBuffer = [mtlDevice newBufferWithLength:_maxQueryObjects * sizeof(uint64_t) options:MTLResourceStorageModeShared];
+    _gpuQueryPool->semaphore = ccnew CCMTLSemaphore(1);
+}
+
+void CCMTLQueryPool::doDestroy() {
+    if (_gpuQueryPool) {
+        if (_gpuQueryPool->semaphore) {
+            _gpuQueryPool->semaphore->syncAll();
+            CC_SAFE_DELETE(_gpuQueryPool->semaphore);
+        }
+
+        id<MTLBuffer> mtlBuffer = _gpuQueryPool->visibilityResultBuffer;
+        _gpuQueryPool->visibilityResultBuffer = nil;
+
+        auto destroyFunc = [mtlBuffer]() {
+            if (mtlBuffer) {
+                [mtlBuffer release];
+            }
+        };
+        CCMTLGPUGarbageCollectionPool::getInstance()->collect(destroyFunc);
+
+        CC_SAFE_DELETE(_gpuQueryPool);
+    }
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLQueue.h

@@ -0,0 +1,55 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#import "gfx-base/GFXQueue.h"
+
+namespace cc {
+namespace gfx {
+struct CCMTLGPUQueueObject;
+class CCMTLQueue final : public Queue {
+    friend class CCMTLDevice;
+
+public:
+    explicit CCMTLQueue();
+    ~CCMTLQueue();
+    CCMTLQueue(const CCMTLQueue &) = delete;
+    CCMTLQueue(CCMTLQueue &&) = delete;
+    CCMTLQueue &operator=(const CCMTLQueue &) = delete;
+    CCMTLQueue &operator=(CCMTLQueue &&) = delete;
+
+    void submit(CommandBuffer *const *cmdBuffs, uint32_t count) override;
+
+    inline CCMTLGPUQueueObject *gpuQueueObj() { return _gpuQueueObj; }
+
+protected:
+    void doInit(const QueueInfo &info) override;
+    void doDestroy() override;
+
+    CCMTLGPUQueueObject *_gpuQueueObj = nullptr;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLQueue.mm

@@ -0,0 +1,79 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#import "MTLQueue.h"
+#import "MTLDevice.h"
+#import "MTLCommandBuffer.h"
+#import "MTLGPUObjects.h"
+#import "MTLFramebuffer.h"
+#import "MTLSwapchain.h"
+namespace cc {
+namespace gfx {
+
+CCMTLQueue::CCMTLQueue()
+: Queue() {
+    _typedID = generateObjectID<decltype(this)>();
+}
+
+CCMTLQueue::~CCMTLQueue() {
+    destroy();
+}
+
+void CCMTLQueue::doInit(const QueueInfo &info) {
+    _gpuQueueObj = ccnew CCMTLGPUQueueObject;
+    auto device = static_cast<id<MTLDevice>>(CCMTLDevice::getInstance()->getMTLDevice());
+    _gpuQueueObj->mtlCommandQueue = [device newCommandQueue];
+}
+
+void CCMTLQueue::doDestroy() {
+    if (_gpuQueueObj) {
+        id<MTLCommandQueue> mtlCmdQueue = _gpuQueueObj->mtlCommandQueue;
+        _gpuQueueObj->mtlCommandQueue = nil;
+
+        auto destroyFunc = [mtlCmdQueue]() {
+            if (mtlCmdQueue) {
+                [mtlCmdQueue release];
+            }
+        };
+        CCMTLGPUGarbageCollectionPool::getInstance()->collect(destroyFunc);
+
+        CC_SAFE_DELETE(_gpuQueueObj);
+    }
+}
+
+void CCMTLQueue::submit(CommandBuffer *const *cmdBuffs, uint32_t count) {
+    for (uint32_t i = 0u; i < count; ++i) {
+        CCMTLCommandBuffer *cmdBuffer = static_cast<CCMTLCommandBuffer *>(cmdBuffs[i]);
+        _gpuQueueObj->numDrawCalls += cmdBuffer->getNumDrawCalls();
+        _gpuQueueObj->numInstances += cmdBuffer->getNumInstances();
+        _gpuQueueObj->numTriangles += cmdBuffer->getNumTris();
+        id<MTLCommandBuffer> mtlCmdBuffer = cmdBuffer->gpuCommandBufferObj()->mtlCommandBuffer;
+        [mtlCmdBuffer commit];
+        cmdBuffer->afterCommit();
+    }
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLRenderCommandEncoder.h

@@ -0,0 +1,309 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#import <Metal/MTLCommandBuffer.h>
+#import <Metal/MTLRenderCommandEncoder.h>
+#import <Metal/MTLRenderPass.h>
+#include "MTLCommandEncoder.h"
+#include "MTLUtils.h"
+#include "base/Macros.h"
+#include "base/std/container/unordered_map.h"
+#include "math/Math.h"
+
+namespace cc {
+namespace gfx {
+
+struct Color;
+
+class CCMTLRenderCommandEncoder final : public CCMTLCommandEncoder {
+    struct BufferBinding final {
+        id<MTLBuffer> buffer;
+        uint32_t offset = 0;
+    };
+
+public:
+    CCMTLRenderCommandEncoder() = default;
+    ~CCMTLRenderCommandEncoder() = default;
+    CCMTLRenderCommandEncoder(const CCMTLRenderCommandEncoder &) = delete;
+    CCMTLRenderCommandEncoder(CCMTLRenderCommandEncoder &&) = delete;
+    CCMTLRenderCommandEncoder &operator=(const CCMTLRenderCommandEncoder &) = delete;
+    CCMTLRenderCommandEncoder &operator=(CCMTLRenderCommandEncoder &&) = delete;
+
+    void initialize(id<MTLCommandBuffer> commandBuffer, MTLRenderPassDescriptor *descriptor) {
+        _mtlEncoder = [[commandBuffer renderCommandEncoderWithDescriptor:descriptor] retain];
+        clearStates();
+    }
+    void initialize(id<MTLParallelRenderCommandEncoder> parallelEncoder) {
+        _mtlEncoder = [[parallelEncoder renderCommandEncoder] retain];
+        clearStates();
+    }
+
+    inline void clearStates() {
+        _isViewportSet = false;
+        _isScissorRectSet = false;
+        _isCullModeSet = false;
+        _isFrontFacingWinding = false;
+        _isTriangleFillModeSet = false;
+        _isDepthClipModeSet = false;
+        _isDepthBiasSet = false;
+        _isBlendColorSet = false;
+
+        _pipelineState = nil;
+        _depthStencilState = nil;
+
+        _frontReferenceValue = UINT_MAX;
+        _backReferenceValue = UINT_MAX;
+
+        _vertexBufferMap.clear();
+        _fragmentBufferMap.clear();
+        _vertexTextureMap.clear();
+        _fragmentTextureMap.clear();
+        _vertexSamplerMap.clear();
+        _fragmentSamplerMap.clear();
+    }
+
+    inline void setViewport(const Rect &rect) {
+        Viewport viewport = {rect.x, rect.y, rect.width, rect.height};
+        setViewport(viewport);
+    }
+
+    inline void setViewport(const Viewport &vp) {
+        if (_isViewportSet && _viewport == vp)
+            return;
+
+        _viewport = vp;
+        _isViewportSet = true;
+        [_mtlEncoder setViewport:mu::toMTLViewport(_viewport)];
+    }
+
+    inline void setScissor(const Rect &rect) {
+        if (_isScissorRectSet && _scissorRect == rect)
+            return;
+
+        _scissorRect = rect;
+        _isScissorRectSet = true;
+        [_mtlEncoder setScissorRect:mu::toMTLScissorRect(_scissorRect)];
+    }
+
+    inline void setCullMode(MTLCullMode mode) {
+        if (_isCullModeSet && (_cullMode == mode))
+            return;
+
+        _cullMode = mode;
+        _isCullModeSet = true;
+        [_mtlEncoder setCullMode:mode];
+    }
+
+    inline void setFrontFacingWinding(MTLWinding winding) {
+        if (_isFrontFacingWinding && (_frontFacingWinding == winding))
+            return;
+
+        _frontFacingWinding = winding;
+        _isFrontFacingWinding = true;
+        [_mtlEncoder setFrontFacingWinding:_frontFacingWinding];
+    }
+
+    inline void setDepthClipMode(MTLDepthClipMode mode) {
+#ifndef TARGET_OS_SIMULATOR
+        if (@available(iOS 11.0, macOS 10.11)) {
+            if (_isDepthClipModeSet && (_depthClipMode == mode))
+                return;
+
+            _depthClipMode = mode;
+            _isDepthClipModeSet = true;
+            [_mtlEncoder setDepthClipMode:_depthClipMode];
+        }
+#endif
+    }
+
+    inline void setTriangleFillMode(MTLTriangleFillMode mode) {
+        if (_isTriangleFillModeSet && (_triangleFillMode == mode))
+            return;
+
+        _triangleFillMode = mode;
+        _isTriangleFillModeSet = true;
+        [_mtlEncoder setTriangleFillMode:_triangleFillMode];
+    }
+
+    inline void setRenderPipelineState(id<MTLRenderPipelineState> pipelineState) {
+        if (_pipelineState == pipelineState)
+            return;
+
+        [_mtlEncoder setRenderPipelineState:pipelineState];
+        _pipelineState = pipelineState;
+    }
+
+    inline void setStencilFrontBackReferenceValue(uint32_t frontReferenceValue, uint32_t backReferenceValue) {
+        if (_frontReferenceValue == frontReferenceValue && _backReferenceValue == backReferenceValue)
+            return;
+
+        _frontReferenceValue = frontReferenceValue;
+        _backReferenceValue = backReferenceValue;
+        [_mtlEncoder setStencilFrontReferenceValue:_frontReferenceValue
+                                backReferenceValue:_backReferenceValue];
+    }
+
+    inline void setDepthStencilState(id<MTLDepthStencilState> depthStencilState) {
+        if (_depthStencilState == depthStencilState)
+            return;
+
+        [_mtlEncoder setDepthStencilState:depthStencilState];
+        _depthStencilState = depthStencilState;
+    }
+
+    inline void setDepthBias(float depthBias, float clamp, float slope) {
+        if (_isDepthBiasSet &&
+            math::isEqualF(_depthBias, depthBias) &&
+            math::isEqualF(_clamp, clamp) &&
+            math::isEqualF(_slope, slope)) {
+            return;
+        }
+
+        _depthBias = depthBias;
+        _clamp = clamp;
+        _slope = slope;
+        _isDepthBiasSet = true;
+        [_mtlEncoder setDepthBias:_depthBias
+                       slopeScale:_slope
+                            clamp:_clamp];
+    }
+
+    inline void setBlendColor(const Color &color) {
+        if (_isBlendColorSet && _blendColor == color)
+            return;
+
+        _blendColor = color;
+        _isBlendColorSet = true;
+        [_mtlEncoder setBlendColorRed:_blendColor.x
+                                green:_blendColor.y
+                                 blue:_blendColor.z
+                                alpha:_blendColor.w];
+    }
+
+    inline void setVertexBuffer(const id<MTLBuffer> buffer, uint32_t offset, uint32_t index) {
+        if (_vertexBufferMap.count(index) > 0) {
+            const auto &bufferBinding = _vertexBufferMap[index];
+            if (buffer == bufferBinding.buffer && offset == bufferBinding.offset) {
+                return;
+            }
+        }
+
+        _vertexBufferMap[index] = {buffer, offset};
+        [_mtlEncoder setVertexBuffer:buffer
+                              offset:offset
+                             atIndex:index];
+    }
+
+    inline void setFragmentBuffer(const id<MTLBuffer> buffer, uint32_t offset, uint32_t index) {
+        if (_fragmentBufferMap.count(index) > 0) {
+            const auto &bufferBinding = _fragmentBufferMap[index];
+            if (buffer == bufferBinding.buffer && offset == bufferBinding.offset) {
+                return;
+            }
+        }
+
+        _fragmentBufferMap[index] = {buffer, offset};
+        [_mtlEncoder setFragmentBuffer:buffer
+                                offset:offset
+                               atIndex:index];
+    }
+
+    void setVertexTexture(const id<MTLTexture> texture, uint32_t index) {
+        if (_vertexTextureMap.count(index) > 0 && (texture == _vertexTextureMap[index]))
+            return;
+
+        _vertexTextureMap[index] = texture;
+        [_mtlEncoder setVertexTexture:texture atIndex:index];
+    }
+
+    void setFragmentTexture(const id<MTLTexture> texture, uint32_t index) {
+        if (_fragmentTextureMap.count(index) > 0 && (texture == _fragmentTextureMap[index]))
+            return;
+
+        _fragmentTextureMap[index] = texture;
+        [_mtlEncoder setFragmentTexture:texture atIndex:index];
+    }
+
+    void setVertexSampler(const id<MTLSamplerState> sampler, uint32_t index) {
+        if (_vertexSamplerMap.count(index) > 0 && (sampler == _vertexSamplerMap[index]))
+            return;
+
+        _vertexSamplerMap[index] = sampler;
+        [_mtlEncoder setVertexSamplerState:sampler atIndex:index];
+    }
+
+    void setFragmentSampler(const id<MTLSamplerState> sampler, uint32_t index) {
+        if (_fragmentSamplerMap.count(index) > 0 && (sampler == _fragmentSamplerMap[index]))
+            return;
+
+        _fragmentSamplerMap[index] = sampler;
+        [_mtlEncoder setFragmentSamplerState:sampler atIndex:index];
+    }
+
+    inline void endEncoding() {
+        [_mtlEncoder endEncoding];
+        [_mtlEncoder release];
+        _mtlEncoder = nil;
+    }
+
+    inline id<MTLRenderCommandEncoder> const getMTLEncoder() {
+        return _mtlEncoder;
+    }
+
+protected:
+    bool _isViewportSet = false;
+    bool _isScissorRectSet = false;
+    bool _isCullModeSet = false;
+    bool _isFrontFacingWinding = false;
+    bool _isTriangleFillModeSet = false;
+    bool _isDepthClipModeSet = false;
+    bool _isDepthBiasSet = false;
+    bool _isBlendColorSet = false;
+    id<MTLRenderCommandEncoder> _mtlEncoder = nil;
+    id<MTLRenderPipelineState> _pipelineState = nil;
+    id<MTLDepthStencilState> _depthStencilState = nil;
+    uint32_t _frontReferenceValue = UINT_MAX;
+    uint32_t _backReferenceValue = UINT_MAX;
+    float _depthBias = 0.f;
+    float _clamp = 0.f;
+    float _slope = 0.f;
+    MTLCullMode _cullMode = MTLCullModeNone;
+    MTLWinding _frontFacingWinding = MTLWindingClockwise;
+    CC_UNUSED MTLDepthClipMode _depthClipMode = MTLDepthClipModeClip;
+    MTLTriangleFillMode _triangleFillMode = MTLTriangleFillModeFill;
+    Viewport _viewport;
+    Rect _scissorRect;
+    Color _blendColor;
+    ccstd::unordered_map<uint32_t, BufferBinding> _vertexBufferMap;
+    ccstd::unordered_map<uint32_t, BufferBinding> _fragmentBufferMap;
+    ccstd::unordered_map<uint32_t, id<MTLTexture>> _vertexTextureMap;
+    ccstd::unordered_map<uint32_t, id<MTLTexture>> _fragmentTextureMap;
+    ccstd::unordered_map<uint32_t, id<MTLSamplerState>> _vertexSamplerMap;
+    ccstd::unordered_map<uint32_t, id<MTLSamplerState>> _fragmentSamplerMap;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLRenderPass.h

@@ -0,0 +1,72 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#include "gfx-base/GFXRenderPass.h"
+
+#import <Metal/MTLPixelFormat.h>
+#import <Metal/MTLRenderPass.h>
+#import <Metal/MTLTexture.h>
+#include "math/Vec2.h"
+
+namespace cc {
+namespace gfx {
+
+class CCMTLTexture;
+
+class CCMTLRenderPass final : public RenderPass {
+public:
+    explicit CCMTLRenderPass();
+    ~CCMTLRenderPass();
+
+    using BufferList = ccstd::vector<uint32_t>; // offset to draw buffer, also means [[color(N)]] of input
+
+    void setColorAttachment(size_t slot, CCMTLTexture *texture, int level);
+    void setDepthStencilAttachment(CCMTLTexture *texture, int level);
+
+    inline MTLRenderPassDescriptor *getMTLRenderPassDescriptor() const { return _mtlRenderPassDescriptor; }
+    inline uint32_t getColorRenderTargetNums() const { return _colorRenderTargetNums; }
+    inline const ccstd::vector<Vec2> &getRenderTargetSizes() const { return _renderTargetSizes; }
+    inline void nextSubpass() { _currentSubpassIndex++; }
+    inline uint32_t getCurrentSubpassIndex() { return _currentSubpassIndex; }
+    inline void reset() { _currentSubpassIndex = 0; }
+    inline const BufferList &getDrawBuffer(uint32_t subPassIndex) { return _drawBuffers[subPassIndex]; }
+    inline const BufferList &getReadBuffer(uint32_t subPassIndex) { return _readBuffers[subPassIndex]; }
+
+protected:
+    void doInit(const RenderPassInfo &info) override;
+    void doDestroy() override;
+
+    uint32_t _currentSubpassIndex = 0;
+    MTLRenderPassDescriptor *_mtlRenderPassDescriptor = nil;
+    uint32_t _colorRenderTargetNums = 0;
+    ccstd::vector<Vec2> _renderTargetSizes;
+    ccstd::vector<BufferList> _drawBuffers;
+    ccstd::vector<BufferList> _readBuffers;
+    ccstd::vector<uint32_t> _colorIndices; // attachment index to draw buffer index
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLRenderPass.mm

@@ -0,0 +1,176 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#include <QuartzCore/CAMetalLayer.h>
+#include "MTLDevice.h"
+#include "MTLRenderPass.h"
+#include "MTLUtils.h"
+#include "MTLTexture.h"
+#include "MTLSwapchain.h"
+#include "base/Log.h"
+
+namespace cc {
+namespace gfx {
+
+CCMTLRenderPass::CCMTLRenderPass() : RenderPass() {
+    _typedID = generateObjectID<decltype(this)>();
+}
+
+CCMTLRenderPass::~CCMTLRenderPass() {
+    destroy();
+}
+
+void CCMTLRenderPass::doInit(const RenderPassInfo& info) {
+    _renderTargetSizes.resize(_colorAttachments.size());
+    _mtlRenderPassDescriptor = [[MTLRenderPassDescriptor alloc] init];
+
+    uint32_t i = 0;
+    for (const auto& colorAttachment : _colorAttachments) {
+        _mtlRenderPassDescriptor.colorAttachments[i].loadAction = mu::toMTLLoadAction(colorAttachment.loadOp);
+        _mtlRenderPassDescriptor.colorAttachments[i].storeAction = mu::toMTLStoreAction(colorAttachment.storeOp);
+
+        ++i;
+    }
+    _colorRenderTargetNums = i;
+    _mtlRenderPassDescriptor.depthAttachment.loadAction = mu::toMTLLoadAction(_depthStencilAttachment.depthLoadOp);
+    _mtlRenderPassDescriptor.depthAttachment.storeAction = mu::toMTLStoreAction(_depthStencilAttachment.depthStoreOp);
+    _mtlRenderPassDescriptor.stencilAttachment.loadAction = mu::toMTLLoadAction(_depthStencilAttachment.stencilLoadOp);
+    _mtlRenderPassDescriptor.stencilAttachment.storeAction = mu::toMTLStoreAction(_depthStencilAttachment.stencilStoreOp);
+    
+    uint32_t colorCount = utils::toUint(_colorAttachments.size());
+    if (_subpasses.empty()) {
+        _subpasses.emplace_back();
+        auto &subpass = _subpasses.back();
+        subpass.colors.resize(_colorAttachments.size());
+        for (uint32_t i = 0U; i < _colorAttachments.size(); ++i) {
+            subpass.colors[i] = i;
+        }
+        if (_depthStencilAttachment.format != Format::UNKNOWN) {
+            subpass.depthStencil = colorCount;
+        }
+    } else {
+        // unify depth stencil index
+        for (auto &subpass : _subpasses) {
+            if (subpass.depthStencil != INVALID_BINDING && subpass.depthStencil > colorCount) {
+                subpass.depthStencil = colorCount;
+            }
+            if (subpass.depthStencilResolve != INVALID_BINDING && subpass.depthStencilResolve > colorCount) {
+                subpass.depthStencilResolve = colorCount;
+            }
+        }
+    }
+
+    _colorIndices.resize(_colorAttachments.size(), INVALID_BINDING);
+    _drawBuffers.resize(_subpasses.size());
+    _readBuffers.resize(_subpasses.size());
+    std::vector<uint32_t> colors;
+
+    for (uint32_t i = 0; i < _subpasses.size(); ++i) {
+        auto &subPass = _subpasses[i];
+        auto &readBuffer = _readBuffers[i];
+        auto &drawBuffer = _drawBuffers[i];
+
+        for (auto &input : subPass.inputs) {
+            auto index = INVALID_BINDING;
+            if(input < _colorAttachments.size()) {
+                index = _colorIndices[input];
+            } else {
+                // ds input
+                index = input;
+            }
+            CC_ASSERT(index != INVALID_BINDING); // input should not appear before color or depthstencil.
+            readBuffer.emplace_back(index);
+        }
+
+        for (auto &color : subPass.colors) {
+            auto &index = _colorIndices[color];
+            if (index == INVALID_BINDING) {
+                index = static_cast<uint32_t>(colors.size());
+                colors.emplace_back(color);
+            }
+            drawBuffer.emplace_back(index);
+        }
+    }
+}
+
+void CCMTLRenderPass::doDestroy() {
+    if (_mtlRenderPassDescriptor) {
+        [_mtlRenderPassDescriptor release];
+        _mtlRenderPassDescriptor = nil;
+    }
+}
+
+void CCMTLRenderPass::setColorAttachment(size_t slot, CCMTLTexture* cctex, int level) {
+    if (!_mtlRenderPassDescriptor) {
+        CC_LOG_ERROR("CCMTLRenderPass: MTLRenderPassDescriptor should not be nullptr.");
+        return;
+    }
+
+    if (_colorRenderTargetNums < slot) {
+        CC_LOG_ERROR("CCMTLRenderPass: invalid color attachment slot %d.", slot);
+        return;
+    }
+
+    id<MTLTexture> texture = nil;
+    if (cctex->swapChain()) {
+        auto* swapchain = static_cast<CCMTLSwapchain*>(cctex->swapChain());
+        texture = swapchain->colorTexture()->getMTLTexture();
+    } else {
+        texture = cctex->getMTLTexture();
+    }
+
+    _mtlRenderPassDescriptor.colorAttachments[slot].texture = texture;
+    _mtlRenderPassDescriptor.colorAttachments[slot].level = level;
+    _renderTargetSizes[slot] = {static_cast<float>(texture.width), static_cast<float>(texture.height)};
+}
+
+void CCMTLRenderPass::setDepthStencilAttachment(CCMTLTexture* cctex, int level) {
+    if (!_mtlRenderPassDescriptor) {
+        CC_LOG_ERROR("CCMTLRenderPass: MTLRenderPassDescriptor should not be nullptr.");
+        return;
+    }
+
+    id<MTLTexture> texture = nil;
+    if (cctex->swapChain()) {
+        auto* swapchain = static_cast<CCMTLSwapchain*>(cctex->swapChain());
+        texture = swapchain->depthStencilTexture()->getMTLTexture();
+    } else {
+        texture = cctex->getMTLTexture();
+    }
+
+    _mtlRenderPassDescriptor.depthAttachment.texture = texture;
+    _mtlRenderPassDescriptor.depthAttachment.level = level;
+    if (cctex->getFormat() == Format::DEPTH_STENCIL) {
+        _mtlRenderPassDescriptor.stencilAttachment.texture = texture;
+        _mtlRenderPassDescriptor.stencilAttachment.level = level;
+    }
+
+    if(_renderTargetSizes.empty()) {
+        _renderTargetSizes.emplace_back(static_cast<float>(texture.width), static_cast<float>(texture.height));
+    }
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLSampler.h

@@ -0,0 +1,54 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#import "gfx-base/states/GFXSampler.h"
+
+#import <Metal/MTLSampler.h>
+
+namespace cc {
+namespace gfx {
+
+class CCMTLSampler final : public Sampler {
+public:
+    explicit CCMTLSampler(const SamplerInfo &info);
+    ~CCMTLSampler();
+    CCMTLSampler() = delete;
+    CCMTLSampler(const CCMTLSampler &) = delete;
+    CCMTLSampler(CCMTLSampler &&) = delete;
+    CCMTLSampler &operator=(const CCMTLSampler &) = delete;
+    CCMTLSampler &operator=(CCMTLSampler &&) = delete;
+
+    inline id<MTLSamplerState> getMTLSamplerState() const { return _mtlSamplerState; }
+
+    static CCMTLSampler *getDefaultSampler();
+    static void deleteDefaultSampler();
+
+protected:
+    id<MTLSamplerState> _mtlSamplerState = nil;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLSampler.mm

@@ -0,0 +1,87 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#import "MTLDevice.h"
+#import "MTLGPUObjects.h"
+#import "MTLSampler.h"
+#import "MTLUtils.h"
+#import <Metal/MTLDevice.h>
+
+namespace cc {
+namespace gfx {
+namespace {
+CCMTLSampler* defaultSampler = nullptr;
+};
+
+CCMTLSampler::CCMTLSampler(const SamplerInfo& info) : Sampler(info) {
+    _typedID = generateObjectID<decltype(this)>();
+    MTLSamplerDescriptor* descriptor = [[MTLSamplerDescriptor alloc] init];
+#if (CC_PLATFORM == CC_PLATFORM_MACOS)
+    descriptor.borderColor = MTLSamplerBorderColorTransparentBlack;
+#endif
+    descriptor.sAddressMode = mu::toMTLSamplerAddressMode(info.addressU);
+    descriptor.tAddressMode = mu::toMTLSamplerAddressMode(info.addressV);
+    descriptor.rAddressMode = mu::toMTLSamplerAddressMode(info.addressW);
+    descriptor.minFilter = mu::toMTLSamplerMinMagFilter(info.minFilter);
+    descriptor.magFilter = mu::toMTLSamplerMinMagFilter(info.magFilter);
+    descriptor.mipFilter = mu::toMTLSamplerMipFilter(info.mipFilter);
+    descriptor.maxAnisotropy = info.maxAnisotropy == 0 ? 1 : info.maxAnisotropy;
+    if (CCMTLDevice::getInstance()->isSamplerDescriptorCompareFunctionSupported()) {
+        descriptor.compareFunction = mu::toMTLCompareFunction(info.cmpFunc);
+    }
+
+    id<MTLDevice> mtlDevice = id<MTLDevice>(CCMTLDevice::getInstance()->getMTLDevice());
+    _mtlSamplerState = [mtlDevice newSamplerStateWithDescriptor:descriptor];
+
+    [descriptor release];
+}
+
+CCMTLSampler::~CCMTLSampler() {
+    id<MTLSamplerState> samplerState = _mtlSamplerState;
+    _mtlSamplerState = nil;
+
+    std::function<void(void)> destroyFunc = [samplerState]() {
+        if (samplerState) {
+            [samplerState release];
+        }
+    };
+    CCMTLGPUGarbageCollectionPool::getInstance()->collect(destroyFunc);
+}
+
+CCMTLSampler* CCMTLSampler::getDefaultSampler() {
+    if (!defaultSampler) {
+        SamplerInfo info;
+        defaultSampler = ccnew CCMTLSampler(info);
+    }
+    return defaultSampler;
+}
+void CCMTLSampler::deleteDefaultSampler() {
+    if (defaultSampler) {
+        delete defaultSampler;
+        defaultSampler = nullptr;
+    }
+}
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLSemaphore.h

@@ -0,0 +1,78 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#include <dispatch/dispatch.h>
+
+namespace cc {
+namespace gfx {
+
+class CCMTLSemaphore final {
+public:
+    explicit CCMTLSemaphore(uint32_t initialValue) : _semaphoreCount(initialValue) {
+        _semaphore = dispatch_semaphore_create(initialValue);
+    }
+
+    ~CCMTLSemaphore() {
+        if (_semaphore) {
+            for (size_t i = 0; i < _semaphoreCount; ++i) {
+                dispatch_semaphore_signal(_semaphore);
+            }
+            dispatch_release(_semaphore);
+            _semaphore = nullptr;
+        }
+    };
+    CCMTLSemaphore(const CCMTLSemaphore &) = delete;
+    CCMTLSemaphore(CCMTLSemaphore &&) = delete;
+    CCMTLSemaphore &operator=(const CCMTLSemaphore &) = delete;
+    CCMTLSemaphore &operator=(CCMTLSemaphore &&) = delete;
+
+    void signal() const {
+        dispatch_semaphore_signal(_semaphore);
+    }
+
+    void wait() const {
+        dispatch_semaphore_wait(_semaphore, DISPATCH_TIME_FOREVER);
+    }
+
+    void trySyncAll(uint64_t nanoSec) {
+        for (uint32_t i = 0; i < _semaphoreCount; i++) {
+            dispatch_semaphore_wait(_semaphore, nanoSec);
+        }
+    }
+
+    void syncAll() {
+        for (uint32_t i = 0; i < _semaphoreCount; i++) {
+            dispatch_semaphore_wait(_semaphore, DISPATCH_TIME_FOREVER);
+        }
+    }
+
+protected:
+    dispatch_semaphore_t _semaphore = nullptr;
+    uint32_t _semaphoreCount = 0;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLShader.h

@@ -0,0 +1,103 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#include "gfx-base/GFXShader.h"
+
+#import <Metal/MTLLibrary.h>
+
+namespace cc {
+namespace gfx {
+class CCMTLRenderPass;
+class CCMTLGPUShader;
+class SPIRVUtils;
+class CCMTLShader final : public Shader {
+public:
+    explicit CCMTLShader();
+    ~CCMTLShader();
+    CCMTLShader(const CCMTLShader &) = delete;
+    CCMTLShader(CCMTLShader &&) = delete;
+    CCMTLShader &operator=(const CCMTLShader &) = delete;
+    CCMTLShader &operator=(CCMTLShader &&) = delete;
+
+    inline id<MTLFunction> getVertMTLFunction() const { return _vertFunction; }
+    inline id<MTLFunction> getFragmentMTLFunction() const { return _fragFunction; }
+    inline id<MTLFunction> getComputeMTLFunction() const { return _cmptFunction; }
+    inline const ccstd::unordered_map<uint32_t, uint32_t> &getFragmentSamplerBindings() const { return _mtlFragmentSamplerBindings; }
+    const CCMTLGPUShader *gpuShader(CCMTLRenderPass *renderPass, uint32_t subPass);
+
+    uint32_t getAvailableBufferBindingIndex(ShaderStageFlagBit stage, uint32_t stream);
+
+    id<MTLFunction> getSpecializedFragFunction(uint32_t *index, int *val, uint32_t count);
+
+#ifdef DEBUG_SHADER
+    inline const ccstd::string &getVertGlslShader() const { return _vertGlslShader; }
+    inline const ccstd::string &getVertMtlSahder() const { return _vertMtlShader; }
+    inline const ccstd::string &getFragGlslShader() const { return _fragGlslShader; }
+    inline const ccstd::string &getFragMtlSahder() const { return _fragMtlShader; }
+    inline const ccstd::string &getcompGlslShader() const { return _cmptGlslShader; }
+    inline const ccstd::string &getcompMtlSahder() const { return _cmptMtlShader; }
+#endif
+
+protected:
+    void doInit(const ShaderInfo &info) override;
+    void doDestroy() override;
+
+    bool checkInputAttachment(const ShaderInfo& info) const;
+    bool createMTLFunction(const ShaderStage &, CCMTLRenderPass *renderPass, uint32_t subPass);
+    void setAvailableBufferBindingIndex();
+
+    id<MTLFunction> _vertFunction = nil;
+    id<MTLFunction> _fragFunction = nil;
+    id<MTLFunction> _cmptFunction = nil;
+
+    id<MTLLibrary> _vertLibrary = nil;
+    id<MTLLibrary> _fragLibrary = nil;
+    id<MTLLibrary> _cmptLibrary = nil;
+
+    // function constant hash , specialized MTLFunction
+    NSMutableDictionary<NSString *, id<MTLFunction>> *_specializedFragFuncs = nil;
+
+    ccstd::unordered_map<uint32_t, uint32_t> _mtlFragmentSamplerBindings;
+    ccstd::vector<uint32_t> _availableVertexBufferBindingIndex;
+    ccstd::vector<uint32_t> _availableFragmentBufferBindingIndex;
+
+    CCMTLGPUShader *_gpuShader = nullptr;
+
+    static SPIRVUtils *spirv;
+
+    // For debug
+#ifdef DEBUG_SHADER
+    ccstd::string _vertGlslShader;
+    ccstd::string _vertMtlShader;
+    ccstd::string _fragGlslShader;
+    ccstd::string _fragMtlShader;
+    ccstd::string _cmptGlslShader;
+    ccstd::string _cmptMtlShader;
+#endif
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLShader.mm

@@ -0,0 +1,353 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#import "MTLDevice.h"
+#import "MTLGPUObjects.h"
+#import "MTLShader.h"
+#import "MTLRenderPass.h"
+#import <Metal/MTLDevice.h>
+#import "gfx-base/SPIRVUtils.h"
+#include "base/Log.h"
+
+namespace cc {
+namespace gfx {
+
+SPIRVUtils* CCMTLShader::spirv = nullptr;
+
+CCMTLShader::CCMTLShader() : Shader() {
+    _typedID = generateObjectID<decltype(this)>();
+}
+
+CCMTLShader::~CCMTLShader() {
+    destroy();
+}
+
+const CCMTLGPUShader *CCMTLShader::gpuShader(CCMTLRenderPass *renderPass, uint32_t subPass)
+{
+    if (_gpuShader != nullptr) {
+        return _gpuShader;
+    }
+    _gpuShader = ccnew CCMTLGPUShader;
+    for (const auto& stage : _stages) {
+        if (!createMTLFunction(stage, renderPass, subPass)) {
+            destroy();
+            return nullptr;
+        }
+    }
+
+    setAvailableBufferBindingIndex();
+
+    CC_LOG_INFO("%s compile succeed.", _name.c_str());
+
+    // Clear shader source after they're uploaded to GPU
+    for (auto &stage : _stages) {
+        stage.source.clear();
+        stage.source.shrink_to_fit();
+    }
+    return _gpuShader;
+}
+
+void CCMTLShader::doInit(const ShaderInfo& info) {
+    _specializedFragFuncs = [[NSMutableDictionary alloc] init];
+    // spirv-cross for input attachment needs RenderPass to build [[color(index)]],
+    // build gpu shader only when there is no subPass input.
+//    if (!checkInputAttachment(info)) {
+//        gpuShader(nullptr, 0);
+//    }
+}
+
+void CCMTLShader::doDestroy() {
+    id<MTLLibrary> vertLib = _vertLibrary;
+    _vertLibrary = nil;
+    id<MTLLibrary> fragLib = _fragLibrary;
+    _fragLibrary = nil;
+    id<MTLLibrary> cmptLib = _cmptLibrary;
+    _cmptLibrary = nil;
+
+    id<MTLFunction> vertFunc = _vertFunction;
+    _vertFunction = nil;
+    id<MTLFunction> fragFunc = _fragFunction;
+    _fragFunction = nil;
+    id<MTLFunction> cmptFunc = _cmptFunction;
+    _cmptFunction = nil;
+
+    if (_gpuShader) {
+        [_gpuShader->shaderSrc release];
+        CC_SAFE_DELETE(_gpuShader);
+    }
+
+    // [_specializedFragFuncs release];
+    NSMutableDictionary<NSString*, id<MTLFunction>>* specFragFuncs = nil;
+    if (_specializedFragFuncs) {
+        specFragFuncs = _specializedFragFuncs;
+        _specializedFragFuncs = nil;
+    }
+
+    std::function<void(void)> destroyFunc = [=]() {
+        if ([specFragFuncs count] > 0) {
+            for (NSString* key in [specFragFuncs allKeys]) {
+                [[specFragFuncs valueForKey:key] release];
+            }
+        }
+        [specFragFuncs release];
+
+        if (vertFunc) {
+            [vertFunc release];
+        }
+        if (fragFunc) {
+            [fragFunc release];
+        }
+        if (cmptFunc) {
+            [cmptFunc release];
+        }
+
+        if (vertLib) {
+            [vertLib release];
+        }
+        if (fragLib) {
+            [fragLib release];
+        }
+        if (cmptLib) {
+            [cmptLib release];
+        }
+    };
+    CCMTLGPUGarbageCollectionPool::getInstance()->collect(destroyFunc);
+}
+
+bool CCMTLShader::checkInputAttachment(const ShaderInfo& info) const {
+    return !info.subpassInputs.empty();
+}
+
+bool CCMTLShader::createMTLFunction(const ShaderStage& stage, CCMTLRenderPass *renderPass, uint32_t subPass) {
+    bool isVertexShader = false;
+    bool isFragmentShader = false;
+    bool isComputeShader = false;
+
+    if (stage.stage == ShaderStageFlagBit::VERTEX) {
+        isVertexShader = true;
+    } else if (stage.stage == ShaderStageFlagBit::FRAGMENT) {
+        isFragmentShader = true;
+    } else if (stage.stage == ShaderStageFlagBit::COMPUTE) {
+        isComputeShader = true;
+    }
+
+    id<MTLDevice> mtlDevice = id<MTLDevice>(CCMTLDevice::getInstance()->getMTLDevice());
+    if (!spirv) {
+        spirv = SPIRVUtils::getInstance();
+        spirv->initialize(2); // vulkan >= 1.2  spirv >= 1.5
+    }
+
+    spirv->compileGLSL(stage.stage, "#version 450\n#define CC_USE_METAL 1\n" + stage.source);
+    if (stage.stage == ShaderStageFlagBit::VERTEX) spirv->compressInputLocations(_attributes);
+
+    auto* spvData = spirv->getOutputData();
+    size_t unitSize = sizeof(std::remove_pointer<decltype(spvData)>::type);
+
+    static const ccstd::vector<uint32_t> emptyBuffer;
+    const auto &drawBuffer = renderPass != nullptr ? renderPass->getDrawBuffer(subPass) : emptyBuffer;
+    const auto &readBuffer = renderPass != nullptr ? renderPass->getReadBuffer(subPass) : emptyBuffer;
+    ccstd::string mtlShaderSrc = mu::spirv2MSL(spirv->getOutputData(), spirv->getOutputSize() / unitSize, stage.stage,
+        _gpuShader, renderPass, subPass);
+
+    NSString* shader = [NSString stringWithUTF8String:mtlShaderSrc.c_str()];
+    NSError* error = nil;
+    MTLCompileOptions* opts = [[MTLCompileOptions alloc] init];
+    //opts.languageVersion = MTLLanguageVersion2_3;
+    id<MTLLibrary>& library = isVertexShader ? _vertLibrary : isFragmentShader ? _fragLibrary : _cmptLibrary;
+    ccstd::string shaderStage = isVertexShader ? "vertex" : isFragmentShader ? "fragment" : "compute";
+
+    if (isFragmentShader) {
+        if (@available(iOS 11.0, *)) {
+            library = [mtlDevice newLibraryWithSource:shader options:opts error:&error];
+            if (!library) {
+                CC_LOG_ERROR("Can not compile %s shader: %s", shaderStage.c_str(), [[error localizedDescription] UTF8String]);
+                CC_LOG_ERROR("%s", stage.source.c_str());
+                [opts release];
+                return false;
+            }
+        } else {
+            //delayed instance and pretend tobe specialized function.
+            _gpuShader->specializeColor = false;
+            _gpuShader->shaderSrc = [shader retain];
+            [opts release];
+            CC_ASSERT(_gpuShader->shaderSrc != nil);
+            return true;
+        }
+    } else {
+        library = [mtlDevice newLibraryWithSource:shader options:opts error:&error];
+        if (!library) {
+            CC_LOG_ERROR("Can not compile %s shader: %s", shaderStage.c_str(), [[error localizedDescription] UTF8String]);
+            CC_LOG_ERROR("%s", stage.source.c_str());
+            [opts release];
+            return false;
+        }
+    }
+
+    [opts release];
+
+    if (isVertexShader) {
+        _vertFunction = [library newFunctionWithName:@"main0"];
+        if (!_vertFunction) {
+            [library release];
+            CC_LOG_ERROR("Can not create vertex function: main0");
+            return false;
+        }
+    } else if (isFragmentShader) {
+        _fragFunction = [library newFunctionWithName:@"main0"];
+        if (!_fragFunction) {
+            [library release];
+            CC_LOG_ERROR("Can not create fragment function: main0");
+            return false;
+        }
+    } else if (isComputeShader) {
+        _cmptFunction = [library newFunctionWithName:@"main0"];
+        if (!_cmptFunction) {
+            [library release];
+            CC_LOG_ERROR("Can not create compute function: main0");
+            return false;
+        }
+    } else {
+        [library release];
+        CC_LOG_ERROR("Shader type not supported yet!");
+        return false;
+    }
+
+#ifdef DEBUG_SHADER
+    if (isVertexShader) {
+        _vertGlslShader = stage.source;
+        _vertMtlShader = mtlShader;
+    } else if (isFragmenShader) {
+        _fragGlslShader = stage.source;
+        _fragMtlShader = mtlShader;
+    } else if (isComputeShader) {
+        _cmptGlslShader = stage.source;
+        _cmptMtlShader = mtlShader;
+    }
+#endif
+    return true;
+}
+
+id<MTLFunction> CCMTLShader::getSpecializedFragFunction(uint32_t* index, int* val, uint32_t count) {
+    uint32_t notEvenHash = 0;
+    for (size_t i = 0; i < count; i++) {
+        notEvenHash += val[i] * std::pow(10, index[i]);
+    }
+    NSString* hashStr = [NSString stringWithFormat:@"%d", notEvenHash];
+    id<MTLFunction> specFunc = [_specializedFragFuncs objectForKey:hashStr];
+
+    if (!specFunc) {
+        if (_gpuShader->specializeColor) {
+            MTLFunctionConstantValues* constantValues = [MTLFunctionConstantValues new];
+            for (size_t i = 0; i < count; i++) {
+                [constantValues setConstantValue:&val[i] type:MTLDataTypeInt atIndex:index[i]];
+            }
+
+            NSError* error = nil;
+            id<MTLFunction> specFragFunc = [_fragLibrary newFunctionWithName:@"main0" constantValues:constantValues error:&error];
+            [constantValues release];
+            if (!specFragFunc) {
+                CC_LOG_ERROR("Can not specialize shader: %s", [[error localizedDescription] UTF8String]);
+            }
+            [_specializedFragFuncs setObject:specFragFunc forKey:hashStr];
+        } else {
+            NSString* res = nil;
+            for (size_t i = 0; i < count; i++) {
+                NSString* targetStr = [NSString stringWithFormat:@"(indexOffset%u)", static_cast<unsigned int>(i)];
+                NSString* index = [NSString stringWithFormat:@"(%u)", static_cast<unsigned int>(i)];
+                res = [_gpuShader->shaderSrc stringByReplacingOccurrencesOfString:targetStr withString:index];
+            }
+            id<MTLDevice> mtlDevice = id<MTLDevice>(CCMTLDevice::getInstance()->getMTLDevice());
+            NSError* error = nil;
+            MTLCompileOptions* opts = [[MTLCompileOptions alloc] init];
+            // always current
+            if (_fragLibrary) {
+                [_fragLibrary release];
+            }
+            _fragLibrary = [mtlDevice newLibraryWithSource:res options:opts error:&error];
+            if (!_fragLibrary) {
+                CC_LOG_ERROR("Can not compile frag shader: %s", [[error localizedDescription] UTF8String]);
+            }
+            [opts release];
+            _fragFunction = [_fragLibrary newFunctionWithName:@"main0"];
+            if (!_fragFunction) {
+                [_fragLibrary release];
+                CC_LOG_ERROR("Can not create fragment function: main0");
+            }
+
+            [_specializedFragFuncs setObject:_fragFunction forKey:hashStr];
+        }
+    }
+    return [_specializedFragFuncs valueForKey:hashStr];
+}
+
+uint32_t CCMTLShader::getAvailableBufferBindingIndex(ShaderStageFlagBit stage, uint32_t stream) {
+    if (hasFlag(stage, ShaderStageFlagBit::VERTEX)) {
+        return _availableVertexBufferBindingIndex.at(stream);
+    }
+
+    if (hasFlag(stage, ShaderStageFlagBit::FRAGMENT)) {
+        return _availableFragmentBufferBindingIndex.at(stream);
+    }
+
+    CC_LOG_ERROR("getAvailableBufferBindingIndex: invalid shader stage %d", stage);
+    return 0;
+}
+
+void CCMTLShader::setAvailableBufferBindingIndex() {
+    uint32_t usedVertexBufferBindingIndexes = 0;
+    uint32_t usedFragmentBufferBindingIndexes = 0;
+    uint32_t vertexBindingCount = 0;
+    uint32_t fragmentBindingCount = 0;
+    for (const auto& block : _gpuShader->blocks) {
+        if (hasFlag(block.second.stages, ShaderStageFlagBit::VERTEX)) {
+            vertexBindingCount++;
+            usedVertexBufferBindingIndexes |= 1 << block.second.mappedBinding;
+        }
+        if (hasFlag(block.second.stages, ShaderStageFlagBit::FRAGMENT)) {
+            fragmentBindingCount++;
+            usedFragmentBufferBindingIndexes |= 1 << block.second.mappedBinding;
+        }
+    }
+
+    auto maxBufferBindingIndex = CCMTLDevice::getInstance()->getMaximumBufferBindingIndex();
+    _availableVertexBufferBindingIndex.resize(maxBufferBindingIndex - vertexBindingCount);
+    _availableFragmentBufferBindingIndex.resize(maxBufferBindingIndex - fragmentBindingCount);
+    uint32_t availableVertexBufferBit = ~usedVertexBufferBindingIndexes;
+    uint32_t availableFragmentBufferBit = ~usedFragmentBufferBindingIndexes;
+    int theBit = maxBufferBindingIndex - 1;
+    uint32_t i = 0, j = 0;
+    for (; theBit >= 0; theBit--) {
+        if ((availableVertexBufferBit & (1 << theBit))) {
+            _availableVertexBufferBindingIndex[i++] = theBit;
+        }
+
+        if ((availableFragmentBufferBit & (1 << theBit))) {
+            _availableFragmentBufferBindingIndex[j++] = theBit;
+        }
+    }
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLSwapchain.h

@@ -0,0 +1,63 @@
+/****************************************************************************
+ Copyright (c) 2021-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+#import <QuartzCore/CAMetalLayer.h>
+#import "MTLTexture.h"
+#import "gfx-base/GFXSwapchain.h"
+
+namespace cc {
+namespace gfx {
+
+struct CCMTLGPUSwapChainObject;
+class CCMTLTexture;
+
+class CCMTLSwapchain final : public Swapchain {
+public:
+    CCMTLSwapchain();
+    ~CCMTLSwapchain();
+
+    inline CCMTLGPUSwapChainObject *gpuSwapChainObj() { return _gpuSwapchainObj; }
+    CCMTLTexture *colorTexture();
+    CCMTLTexture *depthStencilTexture();
+
+    id<CAMetalDrawable> currentDrawable();
+
+    void acquire();
+
+    void release();
+
+protected:
+    void doInit(const SwapchainInfo &info) override;
+    void doDestroy() override;
+    void doResize(uint32_t width, uint32_t height, SurfaceTransform transform) override;
+    void doDestroySurface() override;
+    void doCreateSurface(void *windowHandle) override;
+
+private:
+    CCMTLGPUSwapChainObject *_gpuSwapchainObj = nullptr;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLSwapchain.mm

@@ -0,0 +1,178 @@
+/****************************************************************************
+ Copyright (c) 2021-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#import "../gfx-base/GFXDef-common.h"
+#import "MTLSwapchain.h"
+#if CC_PLATFORM == CC_PLATFORM_MACOS
+    #import <AppKit/NSView.h>
+#else
+    #import <UIKit/UIView.h>
+#endif
+
+#import "MTLGPUObjects.h"
+#import "MTLDevice.h"
+#import "MTLGPUObjects.h"
+namespace cc {
+namespace gfx {
+
+namespace {
+#if CC_PLATFORM == CC_PLATFORM_MACOS
+using CCView = NSView;
+#else
+using CCView = UIView;
+#endif
+}; // namespace
+
+CCMTLSwapchain::CCMTLSwapchain() {
+}
+
+CCMTLSwapchain::~CCMTLSwapchain() {
+    destroy();
+}
+
+void CCMTLSwapchain::doInit(const SwapchainInfo& info) {
+    _gpuSwapchainObj = ccnew CCMTLGPUSwapChainObject;
+
+    //----------------------acquire layer-----------------------------------
+#if CC_EDITOR
+    CAMetalLayer* layer = (CAMetalLayer*)info.windowHandle;
+    if (!layer.device) {
+        layer.device = MTLCreateSystemDefaultDevice();
+    }
+#else
+    auto *view = (CCView *)info.windowHandle;
+    CAMetalLayer *layer = static_cast<CAMetalLayer *>(view.layer);
+#endif
+
+    if (layer.pixelFormat == MTLPixelFormatInvalid) {
+        layer.pixelFormat = MTLPixelFormatBGRA8Unorm;
+    }
+    layer.framebufferOnly = NO;
+    //setDisplaySyncEnabled : physic device refresh rate.
+    //setPresentsWithTransaction : Core Animation transactions update rate.
+    auto syncModeFunc = [&](BOOL sync, BOOL transaction) {
+#if CC_PLATFORM == CC_PLATFORM_MACOS
+        [layer setDisplaySyncEnabled:sync];
+#endif
+        [layer setPresentsWithTransaction:transaction];
+    };
+    switch (_vsyncMode) {
+        case VsyncMode::OFF:
+            syncModeFunc(NO, NO);
+            break;
+        case VsyncMode::ON:
+            syncModeFunc(YES, YES);
+        case VsyncMode::RELAXED:
+        case VsyncMode::MAILBOX:
+        case VsyncMode::HALF:
+            syncModeFunc(YES, NO);
+        default:
+            break;
+    }
+    _gpuSwapchainObj->mtlLayer = layer;
+
+    //    MTLPixelFormatBGRA8Unorm
+    //    MTLPixelFormatBGRA8Unorm_sRGB
+    //    MTLPixelFormatRGBA16Float
+    //    MTLPixelFormatRGB10A2Unorm (macOS only)
+    //    MTLPixelFormatBGR10A2Unorm (macOS only)
+    //    MTLPixelFormatBGRA10_XR
+    //    MTLPixelFormatBGRA10_XR_sRGB
+    //    MTLPixelFormatBGR10_XR
+    //    MTLPixelFormatBGR10_XR_sRGB
+    Format colorFmt = Format::BGRA8;
+    Format depthStencilFmt = Format::DEPTH_STENCIL;
+
+    _colorTexture = ccnew CCMTLTexture;
+    _depthStencilTexture = ccnew CCMTLTexture;
+
+    SwapchainTextureInfo textureInfo;
+    textureInfo.swapchain = this;
+    textureInfo.format = colorFmt;
+    textureInfo.width = info.width;
+    textureInfo.height = info.height;
+    initTexture(textureInfo, _colorTexture);
+
+    textureInfo.format = depthStencilFmt;
+    initTexture(textureInfo, _depthStencilTexture);
+
+    CCMTLDevice::getInstance()->registerSwapchain(this);
+}
+
+void CCMTLSwapchain::doDestroy() {
+    CCMTLDevice::getInstance()->unRegisterSwapchain(this);
+    if (_gpuSwapchainObj) {
+        _gpuSwapchainObj->currentDrawable = nil;
+        _gpuSwapchainObj->mtlLayer = nil;
+
+        CC_SAFE_DELETE(_gpuSwapchainObj);
+    }
+
+    CC_SAFE_DESTROY_NULL(_colorTexture);
+    CC_SAFE_DESTROY_NULL(_depthStencilTexture);
+}
+
+void CCMTLSwapchain::doDestroySurface() {
+    if (_gpuSwapchainObj) {
+        _gpuSwapchainObj->currentDrawable = nil;
+        _gpuSwapchainObj->mtlLayer = nil;
+    }
+}
+
+void CCMTLSwapchain::doResize(uint32_t width, uint32_t height, SurfaceTransform /*transform*/) {
+    _colorTexture->resize(width, height);
+    _depthStencilTexture->resize(width, height);
+}
+
+CCMTLTexture* CCMTLSwapchain::colorTexture() {
+    return static_cast<CCMTLTexture*>(_colorTexture.get());
+}
+
+CCMTLTexture* CCMTLSwapchain::depthStencilTexture() {
+    return static_cast<CCMTLTexture*>(_depthStencilTexture.get());
+}
+
+id<CAMetalDrawable> CCMTLSwapchain::currentDrawable() {
+    return _gpuSwapchainObj->currentDrawable;
+}
+
+void CCMTLSwapchain::release() {
+    _gpuSwapchainObj->currentDrawable = nil;
+    static_cast<CCMTLTexture*>(_colorTexture.get())->update();
+}
+
+void CCMTLSwapchain::acquire() {
+    // hang on here if next drawable not available
+    while (!_gpuSwapchainObj->currentDrawable) {
+        _gpuSwapchainObj->currentDrawable = [_gpuSwapchainObj->mtlLayer nextDrawable];
+        static_cast<CCMTLTexture*>(_colorTexture.get())->update();
+    }
+}
+
+void CCMTLSwapchain::doCreateSurface(void* windowHandle) {
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLTexture.h

@@ -0,0 +1,82 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#import "gfx-base/GFXTexture.h"
+
+#import <Metal/MTLTexture.h>
+
+namespace cc {
+namespace gfx {
+
+struct CCMTLGPUTextureObject;
+class CCMTLSwapchain;
+
+class CCMTLTexture final : public Texture {
+public:
+    explicit CCMTLTexture();
+    ~CCMTLTexture();
+    CCMTLTexture(const CCMTLTexture &) = delete;
+    CCMTLTexture(CCMTLTexture &&) = delete;
+    CCMTLTexture &operator=(const CCMTLTexture &) = delete;
+    CCMTLTexture &operator=(CCMTLTexture &&) = delete;
+
+    inline id<MTLTexture> getMTLTexture() const {
+        return _isTextureView ? _mtlTextureView : _mtlTexture;
+    }
+    inline Format getConvertedFormat() const { return _convertedFormat; }
+    inline bool isArray() const { return _isArray; }
+    inline bool isPVRTC() const { return _isPVRTC; }
+
+    //update drawable from swapchain.
+    void update();
+
+    const TextureInfo &textureInfo();
+    CCMTLSwapchain *swapChain();
+
+    static CCMTLTexture *getDefaultTexture();
+    static void deleteDefaultTexture();
+
+protected:
+    void doInit(const TextureInfo &info) override;
+    void doInit(const TextureViewInfo &info) override;
+    void doDestroy() override;
+    void doResize(uint32_t width, uint32_t height, uint32_t size) override;
+    void doInit(const SwapchainTextureInfo &info) override;
+
+    bool createMTLTexture();
+
+    Format _convertedFormat = Format::UNKNOWN;
+    bool _isArray = false;
+    bool _isPVRTC = false;
+
+    bool _allocateMemory = true; // allocate device memory by metal driver.
+
+    id<MTLTexture> _mtlTexture = nil;
+    id<MTLTexture> _mtlTextureView = nil;
+};
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLTexture.mm

@@ -0,0 +1,368 @@
+/****************************************************************************
+ Copyright (c) 2019-2022 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated engine source code (the "Software"), a limited,
+ worldwide, royalty-free, non-assignable, revocable and non-exclusive license
+ to use Cocos Creator solely to develop games on your target platforms. You shall
+ not use Cocos Creator software for developing other software or tools that's
+ used for developing games. You are not granted to publish, distribute,
+ sublicense, and/or sell copies of Cocos Creator.
+
+ The software or tools in this License Agreement are licensed, not sold.
+ Xiamen Yaji Software Co., Ltd. reserves all rights not expressly granted to you.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#import "MTLDevice.h"
+#import "MTLGPUObjects.h"
+#import "MTLTexture.h"
+#import "MTLUtils.h"
+#import "MTLSwapchain.h"
+#import "profiler/Profiler.h"
+#include "base/Log.h"
+#import <CoreVideo/CVPixelBuffer.h>
+#import <CoreVideo/CVMetalTexture.h>
+#import <CoreVideo/CVMetalTextureCache.h>
+
+// deferred testcase 'camera'
+#define MEMLESS_ON 0
+
+namespace cc {
+namespace gfx {
+
+namespace {
+CCMTLTexture *defaultTexture = nullptr;
+}
+
+CCMTLTexture::CCMTLTexture() : Texture() {
+    _typedID = generateObjectID<decltype(this)>();
+}
+
+CCMTLTexture::~CCMTLTexture() {
+    destroy();
+}
+
+void CCMTLTexture::doInit(const TextureInfo &info) {
+    _isArray = _info.type == TextureType::TEX1D_ARRAY || info.type == TextureType::TEX2D_ARRAY;
+    if (_info.format == Format::PVRTC_RGB2 ||
+        _info.format == Format::PVRTC_RGBA2 ||
+        _info.format == Format::PVRTC_RGB4 ||
+        _info.format == Format::PVRTC_RGBA4 ||
+        _info.format == Format::PVRTC2_2BPP ||
+        _info.format == Format::PVRTC2_4BPP) {
+        _isPVRTC = true;
+    }
+
+    if (_info.externalRes) {
+        auto pixelBuffer = static_cast<CVPixelBufferRef>(_info.externalRes);
+
+// for separating y tex and cbcr tex from arkit background pixelbuffer
+#if CC_USE_AR_MODULE
+        size_t width = CVPixelBufferGetWidthOfPlane(pixelBuffer, _info.layerCount);
+        size_t height = CVPixelBufferGetHeightOfPlane(pixelBuffer, _info.layerCount);
+        CVReturn cvret;
+        CVMetalTextureCacheRef CVMTLTextureCache;
+        cvret = CVMetalTextureCacheCreate(
+            kCFAllocatorDefault,
+            nil,
+            (id<MTLDevice>)CCMTLDevice::getInstance()->getMTLDevice(),
+            nil,
+            &CVMTLTextureCache);
+
+        CC_ASSERT_EQ(cvret, kCVReturnSuccess); // Failed to create Metal texture cache.
+
+        _convertedFormat = mu::convertGFXPixelFormat(_info.format);
+        MTLPixelFormat mtlFormat = mu::toMTLPixelFormat(_convertedFormat);
+        CVMetalTextureRef CVMTLTexture;
+        cvret = CVMetalTextureCacheCreateTextureFromImage(
+            kCFAllocatorDefault,
+            CVMTLTextureCache,
+            pixelBuffer, nil,
+            mtlFormat,
+            width, height,
+            _info.layerCount,
+            &CVMTLTexture);
+
+        CC_ASSERT_EQ(cvret, kCVReturnSuccess); // Failed to create CoreVideo Metal texture from image.
+
+        _mtlTexture = CVMetalTextureGetTexture(CVMTLTexture);
+
+        CC_ASSERT(_mtlTexture); // Failed to create Metal texture CoreVideo Metal Texture
+
+        return;
+#else
+
+        size_t width = CVPixelBufferGetWidth(pixelBuffer);
+        size_t height = CVPixelBufferGetHeight(pixelBuffer);
+
+        CVReturn cvret;
+        CVMetalTextureCacheRef CVMTLTextureCache;
+        cvret = CVMetalTextureCacheCreate(
+            kCFAllocatorDefault,
+            nil,
+            (id<MTLDevice>)CCMTLDevice::getInstance()->getMTLDevice(),
+            nil,
+            &CVMTLTextureCache);
+
+        CC_ASSERT_EQ(cvret, kCVReturnSuccess); // Failed to create Metal texture cache.
+
+        _convertedFormat = mu::convertGFXPixelFormat(_info.format);
+        MTLPixelFormat mtlFormat = mu::toMTLPixelFormat(_convertedFormat);
+        CVMetalTextureRef CVMTLTexture;
+        cvret = CVMetalTextureCacheCreateTextureFromImage(
+            kCFAllocatorDefault,
+            CVMTLTextureCache,
+            pixelBuffer, nil,
+            mtlFormat,
+            width, height,
+            0,
+            &CVMTLTexture);
+
+        CC_ASSERT_EQ(cvret, kCVReturnSuccess); // Failed to create CoreVideo Metal texture from image.
+
+        _mtlTexture = CVMetalTextureGetTexture(CVMTLTexture);
+
+        CFRelease(CVMTLTexture);
+        CFRelease(CVMTLTextureCache);
+
+        CC_ASSERT(_mtlTexture); // Failed to create Metal texture CoreVideo Metal Texture
+#endif
+    }
+
+    if (!createMTLTexture()) {
+        CC_LOG_ERROR("CCMTLTexture: create MTLTexture failed.");
+        return;
+    }
+
+    if (_allocateMemory) {
+        CCMTLDevice::getInstance()->getMemoryStatus().textureSize += _size;
+        CC_PROFILE_MEMORY_INC(Texture, _size);
+    }
+}
+
+void CCMTLTexture::doInit(const TextureViewInfo &info) {
+    _isArray = info.type == TextureType::TEX1D_ARRAY || _viewInfo.type == TextureType::TEX2D_ARRAY;
+    if (_viewInfo.format == Format::PVRTC_RGB2 ||
+        _viewInfo.format == Format::PVRTC_RGBA2 ||
+        _viewInfo.format == Format::PVRTC_RGB4 ||
+        _viewInfo.format == Format::PVRTC_RGBA4 ||
+        _viewInfo.format == Format::PVRTC2_2BPP ||
+        _viewInfo.format == Format::PVRTC2_4BPP) {
+        _isPVRTC = true;
+    }
+    _convertedFormat = mu::convertGFXPixelFormat(_viewInfo.format);
+    auto mtlTextureType = mu::toMTLTextureType(_viewInfo.type);
+
+    MTLPixelFormat format = mu::toMTLPixelFormat(_convertedFormat);
+    if(_viewInfo.format == Format::DEPTH_STENCIL) {
+        format = _viewInfo.basePlane == 0 ? mu::toMTLPixelFormat(_viewInfo.texture->getFormat()) : MTLPixelFormatX32_Stencil8;
+    }
+
+    _mtlTextureView = [static_cast<CCMTLTexture *>(_viewInfo.texture)->_mtlTexture
+        newTextureViewWithPixelFormat:format
+                          textureType:mtlTextureType
+                               levels:NSMakeRange(_viewInfo.baseLevel, _viewInfo.levelCount)
+                               slices:NSMakeRange(_viewInfo.baseLayer, _viewInfo.layerCount)];
+}
+
+void CCMTLTexture::doInit(const SwapchainTextureInfo &info) {
+    _swapchain = info.swapchain;
+    if (info.format == Format::DEPTH_STENCIL) {
+        createMTLTexture();
+    } else {
+        _mtlTexture = [static_cast<CCMTLSwapchain *>(_swapchain)->currentDrawable() texture];
+    }
+}
+
+void CCMTLTexture::update() {
+    if (_swapchain) {
+        id<CAMetalDrawable> drawable = static_cast<CCMTLSwapchain *>(_swapchain)->currentDrawable();
+        _mtlTexture = drawable ? [drawable texture] : nil;
+    }
+}
+
+bool CCMTLTexture::createMTLTexture() {
+    if (_info.width == 0 || _info.height == 0) {
+        CC_LOG_ERROR("CCMTLTexture: width or height should not be zero.");
+        return false;
+    }
+    _convertedFormat = mu::convertGFXPixelFormat(_info.format);
+    MTLPixelFormat mtlFormat = mu::toMTLPixelFormat(_convertedFormat);
+    if (mtlFormat == MTLPixelFormatInvalid)
+        return false;
+
+    MTLTextureDescriptor *descriptor = nullptr;
+    auto mtlTextureType = mu::toMTLTextureType(_info.type);
+    switch (mtlTextureType) {
+        case MTLTextureType2D:
+        case MTLTextureType2DArray:
+            // No need to set mipmapped flag since mipmapLevelCount was explicty set via `_levelCount`.
+            descriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:mtlFormat
+                                                                            width:_info.width
+                                                                           height:_info.height
+                                                                        mipmapped:NO];
+            break;
+        case MTLTextureTypeCube:
+            descriptor = [MTLTextureDescriptor textureCubeDescriptorWithPixelFormat:mtlFormat
+                                                                               size:_info.width
+                                                                          mipmapped:NO];
+            break;
+        default:
+            CC_ABORT();
+            break;
+    }
+
+    if (descriptor == nullptr)
+        return false;
+
+    descriptor.usage = mu::toMTLTextureUsage(_info.usage);
+    if(hasFlag(_info.flags, TextureFlags::MUTABLE_VIEW_FORMAT)) {
+        descriptor.usage |= MTLTextureUsagePixelFormatView;
+    }
+    descriptor.sampleCount = mu::toMTLSampleCount(_info.samples);
+    descriptor.textureType = descriptor.sampleCount > 1 ? MTLTextureType2DMultisample : mu::toMTLTextureType(_info.type);
+    descriptor.mipmapLevelCount = _info.levelCount;
+    descriptor.arrayLength = _info.type == TextureType::CUBE ? 1 : _info.layerCount;
+
+    bool memoryless = false;
+#if !(CC_PLATFORM == CC_PLATFORM_MACOS && __MAC_OS_X_VERSION_MIN_REQUIRED < 110000)
+    if (@available(macos 11.0, ios 10.0, *)) {
+        memoryless = hasFlag(_info.flags, TextureFlagBit::LAZILY_ALLOCATED) &&
+            hasAllFlags(TextureUsageBit::COLOR_ATTACHMENT | TextureUsageBit::DEPTH_STENCIL_ATTACHMENT | TextureUsageBit::INPUT_ATTACHMENT, _info.usage);
+        if (memoryless) {
+            descriptor.storageMode = MTLStorageModeMemoryless;
+            _allocateMemory = false;
+        }
+    }
+#endif
+    if (!memoryless && !_isPVRTC) {
+        // pvrtc can not use blit encoder to upload data.
+        descriptor.storageMode = MTLStorageModePrivate;
+    }
+
+    id<MTLDevice> mtlDevice = id<MTLDevice>(CCMTLDevice::getInstance()->getMTLDevice());
+    _mtlTexture = [mtlDevice newTextureWithDescriptor:descriptor];
+
+    return _mtlTexture != nil;
+}
+
+CCMTLSwapchain *CCMTLTexture::swapChain() {
+    return static_cast<CCMTLSwapchain *>(_swapchain);
+}
+
+const TextureInfo &CCMTLTexture::textureInfo() {
+    return _isTextureView ? static_cast<CCMTLTexture *>(_viewInfo.texture)->_info : _info;
+}
+
+void CCMTLTexture::doDestroy() {
+    //decrease only non-swapchain tex and have had been inited.
+    if (!_swapchain && _mtlTexture && _allocateMemory) {
+        CCMTLDevice::getInstance()->getMemoryStatus().textureSize -= _size;
+        CC_PROFILE_MEMORY_DEC(Texture, _size);
+    }
+
+    if (_swapchain) {
+        _swapchain = nullptr;
+        _mtlTexture = nil;
+        return;
+    }
+
+    id<MTLTexture> mtlTexure = nil;
+    if (_isTextureView) {
+        mtlTexure = _mtlTextureView;
+        _mtlTextureView = nil;
+    } else {
+        mtlTexure = _mtlTexture;
+        _mtlTexture = nil;
+    }
+
+    std::function<void(void)> destroyFunc = [mtlTexure]() {
+        if (mtlTexure) {
+            //TODO_Zeqiang: [mac12 | ios15, ...) validate here
+            //            [mtlTexure setPurgeableState:MTLPurgeableStateEmpty];
+            [mtlTexure release];
+        }
+    };
+    //gpu object only
+    CCMTLGPUGarbageCollectionPool::getInstance()->collect(destroyFunc);
+}
+
+void CCMTLTexture::doResize(uint32_t width, uint32_t height, uint32_t size) {
+    if (_isTextureView) {
+        CC_LOG_ERROR("TextureView does not support resize! at %p", this);
+        return;
+    }
+    auto oldSize = _size;
+    auto oldWidth = _info.width;
+    auto oldHeight = _info.height;
+
+    id<MTLTexture> oldMTLTexture = _mtlTexture;
+
+    _info.width = width;
+    _info.height = height;
+    _size = size;
+    if (!createMTLTexture()) {
+        _info.width = oldWidth;
+        _info.height = oldHeight;
+        _size = oldSize;
+        _mtlTexture = oldMTLTexture;
+        CC_LOG_ERROR("CCMTLTexture: create MTLTexture failed when try to resize the texture.");
+        return;
+    }
+
+    // texture is a wrapper of drawable when _swapchain is active, drawable is not a resource alloc by gfx,
+    // but the system so skip here.
+    if (!_swapchain && _allocateMemory) {
+        CCMTLDevice::getInstance()->getMemoryStatus().textureSize -= oldSize;
+        CCMTLDevice::getInstance()->getMemoryStatus().textureSize += size;
+        CC_PROFILE_MEMORY_DEC(Texture, oldSize);
+        CC_PROFILE_MEMORY_INC(Texture, size);
+    }
+
+    if (oldMTLTexture) {
+        std::function<void(void)> destroyFunc = [=]() {
+            if (oldMTLTexture) {
+                //TODO_Zeqiang: [mac12 | ios15, ...) validate here
+                //                [oldMTLTexture setPurgeableState:MTLPurgeableStateEmpty];
+                [oldMTLTexture release];
+            }
+        };
+        //gpu object only
+        CCMTLGPUGarbageCollectionPool::getInstance()->collect(destroyFunc);
+    }
+}
+
+CCMTLTexture *CCMTLTexture::getDefaultTexture() {
+    if (!defaultTexture) {
+        TextureInfo info;
+        info.type = TextureType::TEX2D;
+        info.usage = TextureUsage::SAMPLED;
+        info.format = Format::BGRA8;
+        info.width = 2;
+        info.height = 2;
+
+        defaultTexture = ccnew CCMTLTexture();
+        defaultTexture->initialize(info);
+    }
+    return defaultTexture;
+}
+
+void CCMTLTexture::deleteDefaultTexture() {
+    if (defaultTexture) {
+        delete defaultTexture;
+        defaultTexture = nullptr;
+    }
+}
+
+} // namespace gfx
+} // namespace cc

cocos/renderer/gfx-metal/MTLUtils.h

@@ -0,0 +1,118 @@
+/****************************************************************************
+ Copyright (c) 2019-2023 Xiamen Yaji Software Co., Ltd.
+
+ http://www.cocos.com
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights to
+ use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+ of the Software, and to permit persons to whom the Software is furnished to do so,
+ subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+****************************************************************************/
+
+#pragma once
+
+#import <Metal/MTLDepthStencil.h>
+#import <Metal/MTLRenderPass.h>
+#import <Metal/MTLRenderPipeline.h>
+#import <Metal/MTLSampler.h>
+#import <Metal/MTLTexture.h>
+#import <Metal/MTLVertexDescriptor.h>
+#include "base/std/container/unordered_map.h"
+#include "gfx-base/GFXDef.h"
+
+namespace cc {
+namespace gfx {
+class CCMTLGPUShader;
+class CCMTLDevice;
+namespace mu {
+
+API_AVAILABLE(ios(12.0))
+MTLMultisampleStencilResolveFilter toMTLStencilResolveMode(ResolveMode mode);
+
+MTLResourceOptions toMTLResourceOption(MemoryUsage usage);
+MTLLoadAction toMTLLoadAction(LoadOp op);
+MTLStoreAction toMTLStoreAction(StoreOp op);
+MTLStoreAction toMTLMSAAStoreAction(StoreOp op);
+MTLClearColor toMTLClearColor(const Color &clearColor);
+MTLVertexFormat toMTLVertexFormat(Format, bool);
+MTLPixelFormat toMTLPixelFormat(Format);
+MTLMultisampleDepthResolveFilter toMTLDepthResolveMode(ResolveMode mode);
+// Because some pixel format is not supported on metal, so need to convert to supported pixel format.
+Format convertGFXPixelFormat(Format);
+MTLColorWriteMask toMTLColorWriteMask(ColorMask);
+MTLBlendFactor toMTLBlendFactor(BlendFactor);
+MTLBlendOperation toMTLBlendOperation(BlendOp);
+MTLCullMode toMTLCullMode(CullMode);
+MTLWinding toMTLWinding(bool isFrontFaceCCW);
+MTLViewport toMTLViewport(const Viewport &);
+MTLScissorRect toMTLScissorRect(const Rect &);
+MTLTriangleFillMode toMTLTriangleFillMode(PolygonMode);
+MTLDepthClipMode toMTLDepthClipMode(bool isClip);
+MTLCompareFunction toMTLCompareFunction(ComparisonFunc);
+MTLStencilOperation toMTLStencilOperation(StencilOp);
+MTLPrimitiveType toMTLPrimitiveType(PrimitiveMode);
+MTLTextureUsage toMTLTextureUsage(TextureUsage);
+MTLTextureType toMTLTextureType(TextureType type);
+NSUInteger toMTLSampleCount(SampleCount);
+MTLSamplerAddressMode toMTLSamplerAddressMode(Address);
+int toMTLSamplerBorderColor(const Color &);
+MTLSamplerMinMagFilter toMTLSamplerMinMagFilter(Filter);
+MTLSamplerMipFilter toMTLSamplerMipFilter(Filter);
+ccstd::string spirv2MSL(const uint32_t *ir, size_t word_count, ShaderStageFlagBit shaderType, CCMTLGPUShader *gpuShader, RenderPass* renderpass, uint32_t subpassIndex);
+const uint8_t *convertRGB8ToRGBA8(const uint8_t *source, uint32_t length);
+const uint8_t *convertRGB32FToRGBA32F(const uint8_t *source, uint32_t length);
+NSUInteger highestSupportedFeatureSet(id<MTLDevice> device);
+uint32_t getGPUFamily(MTLFeatureSet featureSet);
+uint32_t getMaxVertexAttributes(uint32_t family);
+uint32_t getMaxUniformBufferBindings(uint32_t family);
+uint32_t getMaxEntriesInBufferArgumentTable(uint32_t family);
+uint32_t getMaxEntriesInTextureArgumentTable(uint32_t family);
+uint32_t getMaxEntriesInSamplerStateArgumentTable(uint32_t family);
+uint32_t getMaxTexture2DWidthHeight(uint32_t family);
+uint32_t getMaxCubeMapTextureWidthHeight(uint32_t family);
+uint32_t getMaxColorRenderTarget(uint32_t family);
+uint32_t getMinBufferOffsetAlignment(uint32_t family);
+uint32_t getMaxThreadsPerGroup(uint32_t family);
+bool isPVRTCSuppported(uint32_t family);
+bool isEAC_ETCCSuppported(uint32_t family);
+bool isASTCSuppported(uint32_t family);
+bool isBCSupported(uint32_t family);
+bool isColorBufferFloatSupported(uint32_t family);
+bool isColorBufferHalfFloatSupported(uint32_t family);
+bool isLinearTextureSupported(uint32_t family);
+bool isUISamplerSupported(uint32_t family);
+bool isRGB10A2UIStorageSupported(uint32_t family);
+bool isDDepthStencilFilterSupported(uint32_t family);
+bool isIndirectCommandBufferSupported(MTLFeatureSet featureSet);
+bool isDepthStencilFormatSupported(id<MTLDevice> device, Format format, uint32_t family);
+MTLPixelFormat getSupportedDepthStencilFormat(id<MTLDevice> device, uint32_t family, uint32_t &depthBits);
+bool isIndirectDrawSupported(uint32_t family);
+bool isImageBlockSupported();
+bool isFramebufferFetchSupported();
+ccstd::string featureSetToString(MTLFeatureSet featureSet);
+const uint8_t *const convertData(const uint8_t *source, uint32_t length, Format type);
+uint32_t getBlockSize(Format format);
+uint32_t getBytesPerRow(Format format, uint32_t width);
+bool pixelFormatIsColorRenderable(Format format);
+bool isSamplerDescriptorCompareFunctionSupported(uint32_t family);
+void clearRenderArea(CCMTLDevice *device, id<MTLRenderCommandEncoder> renderEncoder, RenderPass *renderPass, const Rect &renderArea, const Color *colors, float depth, uint32_t stencil);
+inline uint32_t alignUp(uint32_t inSize, uint32_t align) { return ((inSize + align - 1) / align) * align; }
+void clearUtilResource();
+inline uint32_t roundUp(uint32_t dividend, uint32_t divisor) { return (dividend - 1) / divisor + 1; }
+} // namespace mu
+
+} // namespace gfx
+} // namespace cc