Split Pipeline in Render and Compute, Metal part

src/backend/CMakeLists.txt

@@ -108,16 +108,18 @@ if (NXT_ENABLE_METAL)
         ${METAL_DIR}/BufferMTL.h
         ${METAL_DIR}/CommandBufferMTL.mm
         ${METAL_DIR}/CommandBufferMTL.h
+        ${METAL_DIR}/ComputePipelineMTL.mm
+        ${METAL_DIR}/ComputePipelineMTL.h
         ${METAL_DIR}/DepthStencilStateMTL.mm
         ${METAL_DIR}/DepthStencilStateMTL.h
         ${METAL_DIR}/InputStateMTL.mm
         ${METAL_DIR}/InputStateMTL.h
         ${METAL_DIR}/MetalBackend.mm
         ${METAL_DIR}/MetalBackend.h
-        ${METAL_DIR}/PipelineMTL.mm
-        ${METAL_DIR}/PipelineMTL.h
         ${METAL_DIR}/PipelineLayoutMTL.mm
         ${METAL_DIR}/PipelineLayoutMTL.h
+        ${METAL_DIR}/RenderPipelineMTL.mm
+        ${METAL_DIR}/RenderPipelineMTL.h
         ${METAL_DIR}/ResourceUploader.mm
         ${METAL_DIR}/ResourceUploader.h
         ${METAL_DIR}/SamplerMTL.mm

src/backend/metal/CommandBufferMTL.mm

@@ -16,11 +16,12 @@
 
 #include "backend/Commands.h"
 #include "backend/metal/BufferMTL.h"
+#include "backend/metal/ComputePipelineMTL.h"
 #include "backend/metal/DepthStencilStateMTL.h"
 #include "backend/metal/InputStateMTL.h"
 #include "backend/metal/MetalBackend.h"
-#include "backend/metal/PipelineMTL.h"
 #include "backend/metal/PipelineLayoutMTL.h"
+#include "backend/metal/RenderPipelineMTL.h"
 #include "backend/metal/SamplerMTL.h"
 #include "backend/metal/TextureMTL.h"
 
@@ -146,7 +147,8 @@ namespace metal {
 
     void CommandBuffer::FillCommands(id<MTLCommandBuffer> commandBuffer) {
         Command type;
-        Pipeline* lastPipeline = nullptr;
+        ComputePipeline* lastComputePipeline = nullptr;
+        RenderPipeline* lastRenderPipeline = nullptr;
         id<MTLBuffer> indexBuffer = nil;
         uint32_t indexBufferOffset = 0;
         MTLIndexType indexType = MTLIndexTypeUInt32;
@@ -267,10 +269,9 @@ namespace metal {
                     {
                         DispatchCmd* dispatch = commands.NextCommand<DispatchCmd>();
                         ASSERT(encoders.compute);
-                        ASSERT(lastPipeline->IsCompute());
 
                         [encoders.compute dispatchThreadgroups:MTLSizeMake(dispatch->x, dispatch->y, dispatch->z)
-                            threadsPerThreadgroup: lastPipeline->GetLocalWorkGroupSize()];
+                            threadsPerThreadgroup: lastComputePipeline->GetLocalWorkGroupSize()];
                     }
                     break;
 
@@ -326,20 +327,25 @@ namespace metal {
                     }
                     break;
 
-                case Command::SetPipeline:
+                case Command::SetComputePipeline:
                     {
-                        SetPipelineCmd* cmd = commands.NextCommand<SetPipelineCmd>();
+                        SetComputePipelineCmd* cmd = commands.NextCommand<SetComputePipelineCmd>();
-                        lastPipeline = ToBackend(cmd->pipeline).Get();
+                        lastComputePipeline = ToBackend(cmd->pipeline).Get();
 
-                        if (lastPipeline->IsCompute()) {
+                        ASSERT(encoders.compute);
-                            ASSERT(encoders.compute);
+                        lastComputePipeline->Encode(encoders.compute);
-                            lastPipeline->Encode(encoders.compute);
+                    }
-                        } else {
+                    break;
-                            ASSERT(encoders.render);
+
-                            DepthStencilState* depthStencilState = ToBackend(lastPipeline->GetDepthStencilState());
+                case Command::SetRenderPipeline:
-                            [encoders.render setDepthStencilState:depthStencilState->GetMTLDepthStencilState()];
+                    {
-                            lastPipeline->Encode(encoders.render);
+                        SetRenderPipelineCmd* cmd = commands.NextCommand<SetRenderPipelineCmd>();
-                        }
+                        lastRenderPipeline = ToBackend(cmd->pipeline).Get();
+
+                        ASSERT(encoders.render);
+                        DepthStencilState* depthStencilState = ToBackend(lastRenderPipeline->GetDepthStencilState());
+                        [encoders.render setDepthStencilState:depthStencilState->GetMTLDepthStencilState()];
+                        lastRenderPipeline->Encode(encoders.render);
                     }
                     break;
 
@@ -369,12 +375,6 @@ namespace metal {
 
                         const auto& layout = group->GetLayout()->GetBindingInfo();
 
-                        if (lastPipeline->IsCompute()) {
-                            ASSERT(encoders.compute);
-                        } else {
-                            ASSERT(encoders.render);
-                        }
-
                         // TODO(kainino@chromium.org): Maintain buffers and offsets arrays in BindGroup so that we
                         // only have to do one setVertexBuffers and one setFragmentBuffers call here.
                         for (size_t binding = 0; binding < layout.mask.size(); ++binding) {
@@ -390,15 +390,18 @@ namespace metal {
                             uint32_t fragIndex = 0;
                             uint32_t computeIndex = 0;
                             if (vertStage) {
-                                vertIndex = ToBackend(lastPipeline->GetLayout())->
+                                ASSERT(lastRenderPipeline != nullptr);
+                                vertIndex = ToBackend(lastRenderPipeline->GetLayout())->
                                     GetBindingIndexInfo(nxt::ShaderStage::Vertex)[groupIndex][binding];
                             }
                             if (fragStage) {
-                                fragIndex = ToBackend(lastPipeline->GetLayout())->
+                                ASSERT(lastRenderPipeline != nullptr);
+                                fragIndex = ToBackend(lastRenderPipeline->GetLayout())->
                                     GetBindingIndexInfo(nxt::ShaderStage::Fragment)[groupIndex][binding];
                             }
                             if (computeStage) {
-                                computeIndex = ToBackend(lastPipeline->GetLayout())->
+                                ASSERT(lastComputePipeline != nullptr);
+                                computeIndex = ToBackend(lastComputePipeline->GetLayout())->
                                     GetBindingIndexInfo(nxt::ShaderStage::Compute)[groupIndex][binding];
                             }
 

src/backend/metal/ComputePipelineMTL.h

@@ -12,27 +12,25 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#ifndef BACKEND_METAL_PIPELINEMTL_H_
+#ifndef BACKEND_METAL_COMPUTEPIPELINEMTL_H_
-#define BACKEND_METAL_PIPELINEMTL_H_
+#define BACKEND_METAL_COMPUTEPIPELINEMTL_H_
 
-#include "backend/Pipeline.h"
+#include "backend/ComputePipeline.h"
 
 #import <Metal/Metal.h>
 
 namespace backend {
 namespace metal {
 
-    class Pipeline : public PipelineBase {
+    class ComputePipeline : public ComputePipelineBase {
         public:
-            Pipeline(PipelineBuilder* builder);
+            ComputePipeline(ComputePipelineBuilder* builder);
-            ~Pipeline();
+            ~ComputePipeline();
 
-            void Encode(id<MTLRenderCommandEncoder> encoder);
             void Encode(id<MTLComputeCommandEncoder> encoder);
             MTLSize GetLocalWorkGroupSize() const;
 
         private:
-            id<MTLRenderPipelineState> mtlRenderPipelineState = nil;
             id<MTLComputePipelineState> mtlComputePipelineState = nil;
             MTLSize localWorkgroupSize;
     };
@@ -40,4 +38,4 @@ namespace metal {
 }
 }
 
-#endif // BACKEND_METAL_PIPELINEMTL_H_
+#endif // BACKEND_METAL_COMPUTEPIPELINEMTL_H_

src/backend/metal/ComputePipelineMTL.mm

@@ -0,0 +1,59 @@
+// Copyright 2017 The NXT Authors
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "backend/metal/ComputePipelineMTL.h"
+
+#include "backend/metal/MetalBackend.h"
+#include "backend/metal/ShaderModuleMTL.h"
+
+namespace backend {
+namespace metal {
+
+    ComputePipeline::ComputePipeline(ComputePipelineBuilder* builder)
+        : ComputePipelineBase(builder) {
+
+        auto mtlDevice = ToBackend(builder->GetDevice())->GetMTLDevice();
+
+        const auto& module = ToBackend(builder->GetStageInfo(nxt::ShaderStage::Compute).module);
+        const auto& entryPoint = builder->GetStageInfo(nxt::ShaderStage::Compute).entryPoint;
+
+        id<MTLFunction> function = module->GetFunction(entryPoint.c_str());
+
+        NSError *error = nil;
+        mtlComputePipelineState = [mtlDevice
+            newComputePipelineStateWithFunction:function error:&error];
+        if (error != nil) {
+            NSLog(@" error => %@", error);
+            builder->HandleError("Error creating pipeline state");
+            return;
+        }
+
+        // Copy over the local workgroup size as it is passed to dispatch explicitly in Metal
+        localWorkgroupSize = module->GetLocalWorkGroupSize(entryPoint);
+    }
+
+    ComputePipeline::~ComputePipeline() {
+        [mtlComputePipelineState release];
+    }
+
+    void ComputePipeline::Encode(id<MTLComputeCommandEncoder> encoder) {
+        [encoder setComputePipelineState:mtlComputePipelineState];
+    }
+
+    MTLSize ComputePipeline::GetLocalWorkGroupSize() const {
+        return localWorkgroupSize;
+    }
+
+}
+}

src/backend/metal/GeneratedCodeIncludes.h

@@ -15,10 +15,11 @@
 #include "backend/metal/MetalBackend.h"
 #include "backend/metal/BufferMTL.h"
 #include "backend/metal/CommandBufferMTL.h"
+#include "backend/metal/ComputePipelineMTL.h"
 #include "backend/metal/DepthStencilStateMTL.h"
 #include "backend/metal/InputStateMTL.h"
-#include "backend/metal/PipelineMTL.h"
 #include "backend/metal/PipelineLayoutMTL.h"
+#include "backend/metal/RenderPipelineMTL.h"
 #include "backend/metal/SamplerMTL.h"
 #include "backend/metal/ShaderModuleMTL.h"
 #include "backend/metal/TextureMTL.h"

src/backend/metal/MetalBackend.h

@@ -38,14 +38,15 @@ namespace metal {
     class Buffer;
     class BufferView;
     class CommandBuffer;
+    class ComputePipeline;
     class DepthStencilState;
     class Device;
-    class InputState;
     class Framebuffer;
-    class Pipeline;
+    class InputState;
     class PipelineLayout;
     class Queue;
     class RenderPass;
+    class RenderPipeline;
     class Sampler;
     class ShaderModule;
     class Texture;
@@ -57,14 +58,15 @@ namespace metal {
         using BufferType = Buffer;
         using BufferViewType = BufferView;
         using CommandBufferType = CommandBuffer;
+        using ComputePipelineType = ComputePipeline;
         using DepthStencilStateType = DepthStencilState;
         using DeviceType = Device;
-        using InputStateType = InputState;
         using FramebufferType = Framebuffer;
-        using PipelineType = Pipeline;
+        using InputStateType = InputState;
         using PipelineLayoutType = PipelineLayout;
         using QueueType = Queue;
         using RenderPassType = RenderPass;
+        using RenderPipelineType = RenderPipeline;
         using SamplerType = Sampler;
         using ShaderModuleType = ShaderModule;
         using TextureType = Texture;
@@ -89,13 +91,14 @@ namespace metal {
             BufferBase* CreateBuffer(BufferBuilder* builder) override;
             BufferViewBase* CreateBufferView(BufferViewBuilder* builder) override;
             CommandBufferBase* CreateCommandBuffer(CommandBufferBuilder* builder) override;
+            ComputePipelineBase* CreateComputePipeline(ComputePipelineBuilder* builder) override;
             DepthStencilStateBase* CreateDepthStencilState(DepthStencilStateBuilder* builder) override;
             InputStateBase* CreateInputState(InputStateBuilder* builder) override;
             FramebufferBase* CreateFramebuffer(FramebufferBuilder* builder) override;
-            PipelineBase* CreatePipeline(PipelineBuilder* builder) override;
             PipelineLayoutBase* CreatePipelineLayout(PipelineLayoutBuilder* builder) override;
             QueueBase* CreateQueue(QueueBuilder* builder) override;
             RenderPassBase* CreateRenderPass(RenderPassBuilder* builder) override;
+            RenderPipelineBase* CreateRenderPipeline(RenderPipelineBuilder* builder) override;
             SamplerBase* CreateSampler(SamplerBuilder* builder) override;
             ShaderModuleBase* CreateShaderModule(ShaderModuleBuilder* builder) override;
             TextureBase* CreateTexture(TextureBuilder* builder) override;

src/backend/metal/MetalBackend.mm

@@ -16,9 +16,10 @@
 
 #include "backend/metal/BufferMTL.h"
 #include "backend/metal/CommandBufferMTL.h"
+#include "backend/metal/ComputePipelineMTL.h"
 #include "backend/metal/DepthStencilStateMTL.h"
 #include "backend/metal/InputStateMTL.h"
-#include "backend/metal/PipelineMTL.h"
+#include "backend/metal/RenderPipelineMTL.h"
 #include "backend/metal/PipelineLayoutMTL.h"
 #include "backend/metal/ResourceUploader.h"
 #include "backend/metal/SamplerMTL.h"
@@ -91,17 +92,17 @@ namespace metal {
     CommandBufferBase* Device::CreateCommandBuffer(CommandBufferBuilder* builder) {
         return new CommandBuffer(builder);
     }
+    ComputePipelineBase* Device::CreateComputePipeline(ComputePipelineBuilder* builder) {
+        return new ComputePipeline(builder);
+    }
     DepthStencilStateBase* Device::CreateDepthStencilState(DepthStencilStateBuilder* builder) {
         return new DepthStencilState(builder);
     }
-    InputStateBase* Device::CreateInputState(InputStateBuilder* builder) {
-        return new InputState(builder);
-    }
     FramebufferBase* Device::CreateFramebuffer(FramebufferBuilder* builder) {
         return new Framebuffer(builder);
     }
-    PipelineBase* Device::CreatePipeline(PipelineBuilder* builder) {
+    InputStateBase* Device::CreateInputState(InputStateBuilder* builder) {
-        return new Pipeline(builder);
+        return new InputState(builder);
     }
     PipelineLayoutBase* Device::CreatePipelineLayout(PipelineLayoutBuilder* builder) {
         return new PipelineLayout(builder);
@@ -112,6 +113,9 @@ namespace metal {
     RenderPassBase* Device::CreateRenderPass(RenderPassBuilder* builder) {
         return new RenderPass(builder);
     }
+    RenderPipelineBase* Device::CreateRenderPipeline(RenderPipelineBuilder* builder) {
+        return new RenderPipeline(builder);
+    }
     SamplerBase* Device::CreateSampler(SamplerBuilder* builder) {
         return new Sampler(builder);
     }

src/backend/metal/PipelineMTL.mm

@@ -1,112 +0,0 @@
-// Copyright 2017 The NXT Authors
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "backend/metal/PipelineMTL.h"
-
-#include "backend/metal/DepthStencilStateMTL.h"
-#include "backend/metal/InputStateMTL.h"
-#include "backend/metal/MetalBackend.h"
-#include "backend/metal/PipelineLayoutMTL.h"
-#include "backend/metal/ShaderModuleMTL.h"
-
-namespace backend {
-namespace metal {
-
-    Pipeline::Pipeline(PipelineBuilder* builder)
-        : PipelineBase(builder) {
-
-        auto mtlDevice = ToBackend(builder->GetDevice())->GetMTLDevice();
-
-        if (IsCompute()) {
-            const auto& module = ToBackend(builder->GetStageInfo(nxt::ShaderStage::Compute).module);
-            const auto& entryPoint = builder->GetStageInfo(nxt::ShaderStage::Compute).entryPoint;
-
-            id<MTLFunction> function = module->GetFunction(entryPoint.c_str());
-
-            NSError *error = nil;
-            mtlComputePipelineState = [mtlDevice
-                newComputePipelineStateWithFunction:function error:&error];
-            if (error != nil) {
-                NSLog(@" error => %@", error);
-                builder->HandleError("Error creating pipeline state");
-                return;
-            }
-
-            // Copy over the local workgroup size as it is passed to dispatch explicitly in Metal
-            localWorkgroupSize = module->GetLocalWorkGroupSize(entryPoint);
-
-        } else {
-            MTLRenderPipelineDescriptor* descriptor = [MTLRenderPipelineDescriptor new];
-
-            for (auto stage : IterateStages(GetStageMask())) {
-                const auto& module = ToBackend(builder->GetStageInfo(stage).module);
-
-                const auto& entryPoint = builder->GetStageInfo(stage).entryPoint;
-                id<MTLFunction> function = module->GetFunction(entryPoint.c_str());
-
-                switch (stage) {
-                    case nxt::ShaderStage::Vertex:
-                        descriptor.vertexFunction = function;
-                        break;
-                    case nxt::ShaderStage::Fragment:
-                        descriptor.fragmentFunction = function;
-                        break;
-                    case nxt::ShaderStage::Compute:
-                        UNREACHABLE();
-                }
-            }
-
-            // TODO(cwallez@chromium.org): get the attachment formats from the subpass
-            descriptor.colorAttachments[0].pixelFormat = MTLPixelFormatRGBA8Unorm;
-            descriptor.depthAttachmentPixelFormat = MTLPixelFormatDepth32Float;
-
-            InputState* inputState = ToBackend(GetInputState());
-            descriptor.vertexDescriptor = inputState->GetMTLVertexDescriptor();
-
-            // TODO(kainino@chromium.org): push constants, textures, samplers
-
-            NSError *error = nil;
-            mtlRenderPipelineState = [mtlDevice
-                newRenderPipelineStateWithDescriptor:descriptor error:&error];
-            if (error != nil) {
-                NSLog(@" error => %@", error);
-                builder->HandleError("Error creating pipeline state");
-                return;
-            }
-
-            [descriptor release];
-        }
-    }
-
-    Pipeline::~Pipeline() {
-        [mtlRenderPipelineState release];
-        [mtlComputePipelineState release];
-    }
-
-    void Pipeline::Encode(id<MTLRenderCommandEncoder> encoder) {
-        ASSERT(!IsCompute());
-        [encoder setRenderPipelineState:mtlRenderPipelineState];
-    }
-
-    void Pipeline::Encode(id<MTLComputeCommandEncoder> encoder) {
-        ASSERT(IsCompute());
-        [encoder setComputePipelineState:mtlComputePipelineState];
-    }
-
-    MTLSize Pipeline::GetLocalWorkGroupSize() const {
-        return localWorkgroupSize;
-    }
-
-}
-}

src/backend/metal/RenderPipelineMTL.h

@@ -0,0 +1,39 @@
+// Copyright 2017 The NXT Authors
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef BACKEND_METAL_RENDERPIPELINEMTL_H_
+#define BACKEND_METAL_RENDERPIPELINEMTL_H_
+
+#include "backend/RenderPipeline.h"
+
+#import <Metal/Metal.h>
+
+namespace backend {
+namespace metal {
+
+    class RenderPipeline : public RenderPipelineBase {
+        public:
+            RenderPipeline(RenderPipelineBuilder* builder);
+            ~RenderPipeline();
+
+            void Encode(id<MTLRenderCommandEncoder> encoder);
+
+        private:
+            id<MTLRenderPipelineState> mtlRenderPipelineState = nil;
+    };
+
+}
+}
+
+#endif // BACKEND_METAL_RENDERPIPELINEMTL_H_

src/backend/metal/RenderPipelineMTL.mm

@@ -0,0 +1,81 @@
+// Copyright 2017 The NXT Authors
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "backend/metal/RenderPipelineMTL.h"
+
+#include "backend/metal/DepthStencilStateMTL.h"
+#include "backend/metal/InputStateMTL.h"
+#include "backend/metal/MetalBackend.h"
+#include "backend/metal/PipelineLayoutMTL.h"
+#include "backend/metal/ShaderModuleMTL.h"
+
+namespace backend {
+namespace metal {
+
+    RenderPipeline::RenderPipeline(RenderPipelineBuilder* builder)
+        : RenderPipelineBase(builder) {
+
+        auto mtlDevice = ToBackend(builder->GetDevice())->GetMTLDevice();
+
+        MTLRenderPipelineDescriptor* descriptor = [MTLRenderPipelineDescriptor new];
+
+        for (auto stage : IterateStages(GetStageMask())) {
+            const auto& module = ToBackend(builder->GetStageInfo(stage).module);
+
+            const auto& entryPoint = builder->GetStageInfo(stage).entryPoint;
+            id<MTLFunction> function = module->GetFunction(entryPoint.c_str());
+
+            switch (stage) {
+                case nxt::ShaderStage::Vertex:
+                    descriptor.vertexFunction = function;
+                    break;
+                case nxt::ShaderStage::Fragment:
+                    descriptor.fragmentFunction = function;
+                    break;
+                case nxt::ShaderStage::Compute:
+                    UNREACHABLE();
+            }
+        }
+
+        // TODO(cwallez@chromium.org): get the attachment formats from the subpass
+        descriptor.colorAttachments[0].pixelFormat = MTLPixelFormatRGBA8Unorm;
+        descriptor.depthAttachmentPixelFormat = MTLPixelFormatDepth32Float;
+
+        InputState* inputState = ToBackend(GetInputState());
+        descriptor.vertexDescriptor = inputState->GetMTLVertexDescriptor();
+
+        // TODO(kainino@chromium.org): push constants, textures, samplers
+
+        NSError *error = nil;
+        mtlRenderPipelineState = [mtlDevice
+            newRenderPipelineStateWithDescriptor:descriptor error:&error];
+        if (error != nil) {
+            NSLog(@" error => %@", error);
+            builder->HandleError("Error creating pipeline state");
+            return;
+        }
+
+        [descriptor release];
+    }
+
+    RenderPipeline::~RenderPipeline() {
+        [mtlRenderPipelineState release];
+    }
+
+    void RenderPipeline::Encode(id<MTLRenderCommandEncoder> encoder) {
+        [encoder setRenderPipelineState:mtlRenderPipelineState];
+    }
+
+}
+}