diff --git a/src/backend/CMakeLists.txt b/src/backend/CMakeLists.txt index 5dd0e06f1b..49ef85f746 100644 --- a/src/backend/CMakeLists.txt +++ b/src/backend/CMakeLists.txt @@ -143,8 +143,26 @@ if (APPLE) SetPIC(metal_autogen) list(APPEND BACKEND_SOURCES + ${METAL_DIR}/BufferMTL.mm + ${METAL_DIR}/BufferMTL.h + ${METAL_DIR}/CommandBufferMTL.mm + ${METAL_DIR}/CommandBufferMTL.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}/SamplerMTL.mm + ${METAL_DIR}/SamplerMTL.h + ${METAL_DIR}/ShaderModuleMTL.mm + ${METAL_DIR}/ShaderModuleMTL.h + ${METAL_DIR}/TextureMTL.mm + ${METAL_DIR}/TextureMTL.h ) endif() diff --git a/src/backend/metal/BufferMTL.h b/src/backend/metal/BufferMTL.h new file mode 100644 index 0000000000..f519dfb3d3 --- /dev/null +++ b/src/backend/metal/BufferMTL.h @@ -0,0 +1,53 @@ +// 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_BUFFERMTL_H_ +#define BACKEND_METAL_BUFFERMTL_H_ + +#include "common/Buffer.h" + +#import + +#include + +namespace backend { +namespace metal { + + class Buffer : public BufferBase { + public: + Buffer(BufferBuilder* builder); + ~Buffer(); + + id GetMTLBuffer(); + std::mutex& GetMutex(); + + private: + void SetSubDataImpl(uint32_t start, uint32_t count, const uint32_t* data) override; + void MapReadAsyncImpl(uint32_t serial, uint32_t start, uint32_t count) override; + void UnmapImpl() override; + void TransitionUsageImpl(nxt::BufferUsageBit currentUsage, nxt::BufferUsageBit targetUsage) override; + + std::mutex mutex; + id mtlBuffer = nil; + }; + + class BufferView : public BufferViewBase { + public: + BufferView(BufferViewBuilder* builder); + }; + +} +} + +#endif // BACKEND_METAL_BUFFERMTL_H_ diff --git a/src/backend/metal/BufferMTL.mm b/src/backend/metal/BufferMTL.mm new file mode 100644 index 0000000000..425770e6c2 --- /dev/null +++ b/src/backend/metal/BufferMTL.mm @@ -0,0 +1,67 @@ +// 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 "BufferMTL.h" + +#include "MetalBackend.h" + +namespace backend { +namespace metal { + + Buffer::Buffer(BufferBuilder* builder) + : BufferBase(builder) { + mtlBuffer = [ToBackend(GetDevice())->GetMTLDevice() newBufferWithLength:GetSize() + options:MTLResourceStorageModeManaged]; + } + + Buffer::~Buffer() { + std::lock_guard lock(mutex); + [mtlBuffer release]; + mtlBuffer = nil; + } + + id Buffer::GetMTLBuffer() { + return mtlBuffer; + } + + std::mutex& Buffer::GetMutex() { + return mutex; + } + + void Buffer::SetSubDataImpl(uint32_t start, uint32_t count, const uint32_t* data) { + uint32_t* dest = reinterpret_cast([mtlBuffer contents]); + { + std::lock_guard lock(mutex); + memcpy(&dest[start], data, count * sizeof(uint32_t)); + } + [mtlBuffer didModifyRange:NSMakeRange(start * sizeof(uint32_t), count * sizeof(uint32_t))]; + } + + void Buffer::MapReadAsyncImpl(uint32_t serial, uint32_t start, uint32_t count) { + // TODO(cwallez@chromium.org): Implement Map Read for the metal backend + } + + void Buffer::UnmapImpl() { + // TODO(cwallez@chromium.org): Implement Map Read for the metal backend + } + + void Buffer::TransitionUsageImpl(nxt::BufferUsageBit currentUsage, nxt::BufferUsageBit targetUsage) { + } + + BufferView::BufferView(BufferViewBuilder* builder) + : BufferViewBase(builder) { + } + +} +} diff --git a/src/backend/metal/CommandBufferMTL.h b/src/backend/metal/CommandBufferMTL.h new file mode 100644 index 0000000000..67c3a0722d --- /dev/null +++ b/src/backend/metal/CommandBufferMTL.h @@ -0,0 +1,45 @@ +// 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_COMMANDBUFFERMTL_H_ +#define BACKEND_METAL_COMMANDBUFFERMTL_H_ + +#include "common/CommandBuffer.h" + +#import + +#include +#include + +namespace backend { +namespace metal { + + class Device; + + class CommandBuffer : public CommandBufferBase { + public: + CommandBuffer(Device* device, CommandBufferBuilder* builder); + ~CommandBuffer(); + + void FillCommands(id commandBuffer, std::unordered_set* mutexes); + + private: + Device* device; + CommandIterator commands; + }; + +} +} + +#endif // BACKEND_METAL_COMMANDBUFFERMTL_H_ diff --git a/src/backend/metal/CommandBufferMTL.mm b/src/backend/metal/CommandBufferMTL.mm new file mode 100644 index 0000000000..3d49dbde20 --- /dev/null +++ b/src/backend/metal/CommandBufferMTL.mm @@ -0,0 +1,474 @@ +// 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 "CommandBufferMTL.h" + +#include "common/Commands.h" +#include "BufferMTL.h" +#include "DepthStencilStateMTL.h" +#include "InputStateMTL.h" +#include "MetalBackend.h" +#include "PipelineMTL.h" +#include "PipelineLayoutMTL.h" +#include "SamplerMTL.h" +#include "TextureMTL.h" + +namespace backend { +namespace metal { + + namespace { + MTLIndexType IndexFormatType(nxt::IndexFormat format) { + switch (format) { + case nxt::IndexFormat::Uint16: + return MTLIndexTypeUInt16; + case nxt::IndexFormat::Uint32: + return MTLIndexTypeUInt32; + } + } + + struct CurrentEncoders { + Device* device; + + id blit = nil; + id compute = nil; + id render = nil; + + RenderPass* currentRenderPass = nullptr; + Framebuffer* currentFramebuffer = nullptr; + + void FinishEncoders() { + ASSERT(render == nil); + if (blit != nil) { + [blit endEncoding]; + blit = nil; + } + if (compute != nil) { + [compute endEncoding]; + compute = nil; + } + } + + void EnsureBlit(id commandBuffer) { + if (blit == nil) { + FinishEncoders(); + blit = [commandBuffer blitCommandEncoder]; + } + } + void EnsureCompute(id commandBuffer) { + if (compute == nil) { + FinishEncoders(); + compute = [commandBuffer computeCommandEncoder]; + // TODO(cwallez@chromium.org): does any state need to be reset? + } + } + void BeginSubpass(id commandBuffer, uint32_t subpass) { + ASSERT(currentRenderPass); + if (render != nil) { + [render endEncoding]; + render = nil; + } + + const auto& info = currentRenderPass->GetSubpassInfo(subpass); + + MTLRenderPassDescriptor* descriptor = [MTLRenderPassDescriptor renderPassDescriptor]; + bool usingBackbuffer = false; // HACK(kainino@chromium.org): workaround for not having depth attachments + for (uint32_t index = 0; index < info.colorAttachments.size(); ++index) { + uint32_t attachment = info.colorAttachments[index]; + + // TODO(kainino@chromium.org): currently a 'null' texture view + // falls back to the 'back buffer' but this should go away + // when we have WSI. + id texture = nil; + if (auto textureView = currentFramebuffer->GetTextureView(attachment)) { + texture = ToBackend(textureView->GetTexture())->GetMTLTexture(); + } else { + texture = device->GetCurrentTexture(); + usingBackbuffer = true; + } + descriptor.colorAttachments[index].texture = texture; + descriptor.colorAttachments[index].loadAction = MTLLoadActionLoad; + descriptor.colorAttachments[index].storeAction = MTLStoreActionStore; + } + // TODO(kainino@chromium.org): load depth attachment from subpass + if (usingBackbuffer) { + descriptor.depthAttachment.texture = device->GetCurrentDepthTexture(); + descriptor.depthAttachment.loadAction = MTLLoadActionLoad; + descriptor.depthAttachment.storeAction = MTLStoreActionStore; + } + + render = [commandBuffer renderCommandEncoderWithDescriptor:descriptor]; + // TODO(cwallez@chromium.org): does any state need to be reset? + } + void EndRenderPass() { + ASSERT(render != nil); + [render endEncoding]; + render = nil; + } + }; + } + + CommandBuffer::CommandBuffer(Device* device, CommandBufferBuilder* builder) + : CommandBufferBase(builder), device(device), commands(builder->AcquireCommands()) { + } + + CommandBuffer::~CommandBuffer() { + FreeCommands(&commands); + } + + void CommandBuffer::FillCommands(id commandBuffer, std::unordered_set* mutexes) { + Command type; + Pipeline* lastPipeline = nullptr; + id indexBuffer = nil; + uint32_t indexBufferOffset = 0; + MTLIndexType indexType = MTLIndexTypeUInt32; + + CurrentEncoders encoders; + encoders.device = device; + + uint32_t currentSubpass = 0; + id renderEncoder = nil; + + while (commands.NextCommandId(&type)) { + switch (type) { + case Command::AdvanceSubpass: + { + commands.NextCommand(); + currentSubpass += 1; + encoders.BeginSubpass(commandBuffer, currentSubpass); + } + break; + + case Command::BeginRenderPass: + { + BeginRenderPassCmd* beginRenderPassCmd = commands.NextCommand(); + encoders.currentRenderPass = ToBackend(beginRenderPassCmd->renderPass.Get()); + encoders.currentFramebuffer = ToBackend(beginRenderPassCmd->framebuffer.Get()); + encoders.FinishEncoders(); + currentSubpass = 0; + encoders.BeginSubpass(commandBuffer, currentSubpass); + } + break; + + case Command::CopyBufferToBuffer: + { + CopyBufferToBufferCmd* copy = commands.NextCommand(); + + encoders.EnsureBlit(commandBuffer); + [encoders.blit + copyFromBuffer:ToBackend(copy->source)->GetMTLBuffer() + sourceOffset:copy->sourceOffset + toBuffer:ToBackend(copy->destination)->GetMTLBuffer() + destinationOffset:copy->destinationOffset + size:copy->size]; + } + break; + + case Command::CopyBufferToTexture: + { + CopyBufferToTextureCmd* copy = commands.NextCommand(); + Buffer* buffer = ToBackend(copy->buffer.Get()); + Texture* texture = ToBackend(copy->texture.Get()); + + unsigned rowSize = copy->width * TextureFormatPixelSize(texture->GetFormat()); + MTLOrigin origin; + origin.x = copy->x; + origin.y = copy->y; + origin.z = copy->z; + + MTLSize size; + size.width = copy->width; + size.height = copy->height; + size.depth = copy->depth; + + encoders.EnsureBlit(commandBuffer); + [encoders.blit + copyFromBuffer:buffer->GetMTLBuffer() + sourceOffset:copy->bufferOffset + sourceBytesPerRow:rowSize + sourceBytesPerImage:(rowSize * copy->height) + sourceSize:size + toTexture:texture->GetMTLTexture() + destinationSlice:0 + destinationLevel:copy->level + destinationOrigin:origin]; + } + break; + + case Command::Dispatch: + { + DispatchCmd* dispatch = commands.NextCommand(); + encoders.EnsureCompute(commandBuffer); + ASSERT(lastPipeline->IsCompute()); + + [encoders.compute dispatchThreadgroups:MTLSizeMake(dispatch->x, dispatch->y, dispatch->z) + threadsPerThreadgroup: lastPipeline->GetLocalWorkGroupSize()]; + } + break; + + case Command::DrawArrays: + { + DrawArraysCmd* draw = commands.NextCommand(); + + ASSERT(encoders.render); + [encoders.render + drawPrimitives:MTLPrimitiveTypeTriangle + vertexStart:draw->firstVertex + vertexCount:draw->vertexCount + instanceCount:draw->instanceCount + baseInstance:draw->firstInstance]; + } + break; + + case Command::DrawElements: + { + DrawElementsCmd* draw = commands.NextCommand(); + + ASSERT(encoders.render); + [encoders.render + drawIndexedPrimitives:MTLPrimitiveTypeTriangle + indexCount:draw->indexCount + indexType:indexType + indexBuffer:indexBuffer + indexBufferOffset:indexBufferOffset + instanceCount:draw->instanceCount + baseVertex:0 + baseInstance:draw->firstInstance]; + } + break; + + case Command::EndRenderPass: + { + commands.NextCommand(); + encoders.EndRenderPass(); + } + break; + + case Command::SetPipeline: + { + SetPipelineCmd* cmd = commands.NextCommand(); + lastPipeline = ToBackend(cmd->pipeline).Get(); + + if (lastPipeline->IsCompute()) { + encoders.EnsureCompute(commandBuffer); + lastPipeline->Encode(encoders.compute); + } else { + ASSERT(encoders.render); + DepthStencilState* depthStencilState = ToBackend(lastPipeline->GetDepthStencilState()); + [encoders.render setDepthStencilState:depthStencilState->GetMTLDepthStencilState()]; + lastPipeline->Encode(encoders.render); + } + } + break; + + case Command::SetPushConstants: + { + SetPushConstantsCmd* cmd = commands.NextCommand(); + uint32_t* valuesUInt = commands.NextData(cmd->count); + int32_t* valuesInt = reinterpret_cast(valuesUInt); + float* valuesFloat = reinterpret_cast(valuesUInt); + + // TODO(kainino@chromium.org): implement SetPushConstants + } + break; + + case Command::SetStencilReference: + { + SetStencilReferenceCmd* cmd = commands.NextCommand(); + + ASSERT(encoders.render); + + [encoders.render setStencilReferenceValue:cmd->reference]; + } + break; + + case Command::SetBindGroup: + { + SetBindGroupCmd* cmd = commands.NextCommand(); + BindGroup* group = ToBackend(cmd->group.Get()); + uint32_t groupIndex = cmd->index; + + const auto& layout = group->GetLayout()->GetBindingInfo(); + + if (lastPipeline->IsCompute()) { + encoders.EnsureCompute(commandBuffer); + } 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) { + if (!layout.mask[binding]) { + continue; + } + + auto stage = layout.visibilities[binding]; + bool vertStage = stage & nxt::ShaderStageBit::Vertex; + bool fragStage = stage & nxt::ShaderStageBit::Fragment; + bool computeStage = stage & nxt::ShaderStageBit::Compute; + uint32_t vertIndex = 0; + uint32_t fragIndex = 0; + uint32_t computeIndex = 0; + if (vertStage) { + vertIndex = ToBackend(lastPipeline->GetLayout())-> + GetBindingIndexInfo(nxt::ShaderStage::Vertex)[groupIndex][binding]; + } + if (fragStage) { + fragIndex = ToBackend(lastPipeline->GetLayout())-> + GetBindingIndexInfo(nxt::ShaderStage::Fragment)[groupIndex][binding]; + } + if (computeStage) { + computeIndex = ToBackend(lastPipeline->GetLayout())-> + GetBindingIndexInfo(nxt::ShaderStage::Compute)[groupIndex][binding]; + } + + switch (layout.types[binding]) { + case nxt::BindingType::UniformBuffer: + case nxt::BindingType::StorageBuffer: + { + BufferView* view = ToBackend(group->GetBindingAsBufferView(binding)); + auto b = ToBackend(view->GetBuffer()); + mutexes->insert(&b->GetMutex()); + const id buffer = b->GetMTLBuffer(); + const NSUInteger offset = view->GetOffset(); + if (vertStage) { + [encoders.render + setVertexBuffers:&buffer + offsets:&offset + withRange:NSMakeRange(vertIndex, 1)]; + } + if (fragStage) { + [encoders.render + setFragmentBuffers:&buffer + offsets:&offset + withRange:NSMakeRange(fragIndex, 1)]; + } + if (computeStage) { + [encoders.compute + setBuffers:&buffer + offsets:&offset + withRange:NSMakeRange(computeIndex, 1)]; + } + + } + break; + + case nxt::BindingType::Sampler: + { + auto sampler = ToBackend(group->GetBindingAsSampler(binding)); + if (vertStage) { + [encoders.render + setVertexSamplerState:sampler->GetMTLSamplerState() + atIndex:vertIndex]; + } + if (fragStage) { + [encoders.render + setFragmentSamplerState:sampler->GetMTLSamplerState() + atIndex:fragIndex]; + } + if (computeStage) { + [encoders.compute + setSamplerState:sampler->GetMTLSamplerState() + atIndex:computeIndex]; + } + } + break; + + case nxt::BindingType::SampledTexture: + { + auto texture = ToBackend(group->GetBindingAsTextureView(binding)->GetTexture()); + if (vertStage) { + [encoders.render + setVertexTexture:texture->GetMTLTexture() + atIndex:vertIndex]; + } + if (fragStage) { + [encoders.render + setFragmentTexture:texture->GetMTLTexture() + atIndex:fragIndex]; + } + if (computeStage) { + [encoders.compute + setTexture:texture->GetMTLTexture() + atIndex:computeIndex]; + } + } + break; + } + } + } + break; + + case Command::SetIndexBuffer: + { + SetIndexBufferCmd* cmd = commands.NextCommand(); + auto b = ToBackend(cmd->buffer.Get()); + mutexes->insert(&b->GetMutex()); + indexBuffer = b->GetMTLBuffer(); + indexBufferOffset = cmd->offset; + indexType = IndexFormatType(cmd->format); + } + break; + + case Command::SetVertexBuffers: + { + SetVertexBuffersCmd* cmd = commands.NextCommand(); + auto buffers = commands.NextData>(cmd->count); + auto offsets = commands.NextData(cmd->count); + + auto inputState = lastPipeline->GetInputState(); + + std::array, kMaxVertexInputs> mtlBuffers; + std::array mtlOffsets; + + // Perhaps an "array of vertex buffers(+offsets?)" should be + // a NXT API primitive to avoid reconstructing this array? + for (uint32_t i = 0; i < cmd->count; ++i) { + Buffer* buffer = ToBackend(buffers[i].Get()); + mutexes->insert(&buffer->GetMutex()); + mtlBuffers[i] = buffer->GetMTLBuffer(); + mtlOffsets[i] = offsets[i]; + } + + ASSERT(encoders.render); + [encoders.render + setVertexBuffers:mtlBuffers.data() + offsets:mtlOffsets.data() + withRange:NSMakeRange(kMaxBindingsPerGroup + cmd->startSlot, cmd->count)]; + } + break; + + case Command::TransitionBufferUsage: + { + TransitionBufferUsageCmd* cmd = commands.NextCommand(); + + cmd->buffer->UpdateUsageInternal(cmd->usage); + } + break; + + case Command::TransitionTextureUsage: + { + TransitionTextureUsageCmd* cmd = commands.NextCommand(); + + cmd->texture->UpdateUsageInternal(cmd->usage); + } + break; + } + } + + encoders.FinishEncoders(); + } + +} +} diff --git a/src/backend/metal/DepthStencilStateMTL.h b/src/backend/metal/DepthStencilStateMTL.h new file mode 100644 index 0000000000..e0dc97fc4c --- /dev/null +++ b/src/backend/metal/DepthStencilStateMTL.h @@ -0,0 +1,41 @@ +// 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_DEPTHSTENCILSTATEMTL_H_ +#define BACKEND_METAL_DEPTHSTENCILSTATEMTL_H_ + +#include "common/DepthStencilState.h" + +#import + +namespace backend { +namespace metal { + + class Device; + + class DepthStencilState : public DepthStencilStateBase { + public: + DepthStencilState(DepthStencilStateBuilder* builder); + ~DepthStencilState(); + + id GetMTLDepthStencilState(); + + private: + id mtlDepthStencilState = nil; + }; + +} +} + +#endif // BACKEND_METAL_DEPTHSTENCILSTATEMTL_H_ diff --git a/src/backend/metal/DepthStencilStateMTL.mm b/src/backend/metal/DepthStencilStateMTL.mm new file mode 100644 index 0000000000..3c727f90a8 --- /dev/null +++ b/src/backend/metal/DepthStencilStateMTL.mm @@ -0,0 +1,117 @@ +// 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 "DepthStencilStateMTL.h" + +#include "MetalBackend.h" + +namespace backend { +namespace metal { + + namespace { + MTLCompareFunction MetalDepthStencilCompareFunction(nxt::CompareFunction compareFunction) { + switch (compareFunction) { + case nxt::CompareFunction::Never: + return MTLCompareFunctionNever; + case nxt::CompareFunction::Less: + return MTLCompareFunctionLess; + case nxt::CompareFunction::LessEqual: + return MTLCompareFunctionLessEqual; + case nxt::CompareFunction::Greater: + return MTLCompareFunctionGreater; + case nxt::CompareFunction::GreaterEqual: + return MTLCompareFunctionGreaterEqual; + case nxt::CompareFunction::NotEqual: + return MTLCompareFunctionNotEqual; + case nxt::CompareFunction::Equal: + return MTLCompareFunctionEqual; + case nxt::CompareFunction::Always: + return MTLCompareFunctionAlways; + } + } + + MTLStencilOperation MetalStencilOperation(nxt::StencilOperation stencilOperation) { + switch (stencilOperation) { + case nxt::StencilOperation::Keep: + return MTLStencilOperationKeep; + case nxt::StencilOperation::Zero: + return MTLStencilOperationZero; + case nxt::StencilOperation::Replace: + return MTLStencilOperationReplace; + case nxt::StencilOperation::Invert: + return MTLStencilOperationInvert; + case nxt::StencilOperation::IncrementClamp: + return MTLStencilOperationIncrementClamp; + case nxt::StencilOperation::DecrementClamp: + return MTLStencilOperationDecrementClamp; + case nxt::StencilOperation::IncrementWrap: + return MTLStencilOperationIncrementWrap; + case nxt::StencilOperation::DecrementWrap: + return MTLStencilOperationDecrementWrap; + } + } + } + + DepthStencilState::DepthStencilState(DepthStencilStateBuilder* builder) + : DepthStencilStateBase(builder) { + MTLDepthStencilDescriptor* mtlDepthStencilDescriptor = [MTLDepthStencilDescriptor new]; + + if (DepthTestEnabled()) { + auto& depth = GetDepth(); + mtlDepthStencilDescriptor.depthCompareFunction = MetalDepthStencilCompareFunction(depth.compareFunction); + mtlDepthStencilDescriptor.depthWriteEnabled = depth.depthWriteEnabled; + } + + auto& stencil = GetStencil(); + + if (StencilTestEnabled()) { + MTLStencilDescriptor* backFaceStencil = [MTLStencilDescriptor new]; + MTLStencilDescriptor* frontFaceStencil = [MTLStencilDescriptor new]; + + backFaceStencil.stencilCompareFunction = MetalDepthStencilCompareFunction(stencil.back.compareFunction); + backFaceStencil.stencilFailureOperation = MetalStencilOperation(stencil.back.stencilFail); + backFaceStencil.depthFailureOperation = MetalStencilOperation(stencil.back.depthFail); + backFaceStencil.depthStencilPassOperation = MetalStencilOperation(stencil.back.depthStencilPass); + backFaceStencil.readMask = stencil.readMask; + backFaceStencil.writeMask = stencil.writeMask; + + frontFaceStencil.stencilCompareFunction = MetalDepthStencilCompareFunction(stencil.front.compareFunction); + frontFaceStencil.stencilFailureOperation = MetalStencilOperation(stencil.front.stencilFail); + frontFaceStencil.depthFailureOperation = MetalStencilOperation(stencil.front.depthFail); + frontFaceStencil.depthStencilPassOperation = MetalStencilOperation(stencil.front.depthStencilPass); + frontFaceStencil.readMask = stencil.readMask; + frontFaceStencil.writeMask = stencil.writeMask; + + mtlDepthStencilDescriptor.backFaceStencil = backFaceStencil; + mtlDepthStencilDescriptor.frontFaceStencil = frontFaceStencil; + [backFaceStencil release]; + [frontFaceStencil release]; + } + + auto mtlDevice = ToBackend(builder->GetDevice())->GetMTLDevice(); + mtlDepthStencilState = [mtlDevice newDepthStencilStateWithDescriptor:mtlDepthStencilDescriptor]; + [mtlDepthStencilDescriptor release]; + } + + DepthStencilState::~DepthStencilState() { + [mtlDepthStencilState release]; + mtlDepthStencilState = nil; + } + + id DepthStencilState::GetMTLDepthStencilState() { + return mtlDepthStencilState; + } + +} +} diff --git a/src/backend/metal/GeneratedCodeIncludes.h b/src/backend/metal/GeneratedCodeIncludes.h index 16d9bfe4b5..808462e434 100644 --- a/src/backend/metal/GeneratedCodeIncludes.h +++ b/src/backend/metal/GeneratedCodeIncludes.h @@ -13,6 +13,12 @@ // limitations under the License. #include "MetalBackend.h" - -#include "common/Device.h" -#include "common/CommandBuffer.h" +#include "BufferMTL.h" +#include "CommandBufferMTL.h" +#include "DepthStencilStateMTL.h" +#include "InputStateMTL.h" +#include "PipelineMTL.h" +#include "PipelineLayoutMTL.h" +#include "SamplerMTL.h" +#include "ShaderModuleMTL.h" +#include "TextureMTL.h" diff --git a/src/backend/metal/InputStateMTL.h b/src/backend/metal/InputStateMTL.h new file mode 100644 index 0000000000..8011f40f32 --- /dev/null +++ b/src/backend/metal/InputStateMTL.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_INPUTSTATEMTL_H_ +#define BACKEND_METAL_INPUTSTATEMTL_H_ + +#include "common/InputState.h" + +#import + +namespace backend { +namespace metal { + + class InputState : public InputStateBase { + public: + InputState(InputStateBuilder* builder); + ~InputState(); + + MTLVertexDescriptor* GetMTLVertexDescriptor(); + + private: + MTLVertexDescriptor* mtlVertexDescriptor = nil; + }; + +} +} + +#endif // BACKEND_METAL_COMMANDINPUTSTATEMTL_H_ diff --git a/src/backend/metal/InputStateMTL.mm b/src/backend/metal/InputStateMTL.mm new file mode 100644 index 0000000000..852ddb5a55 --- /dev/null +++ b/src/backend/metal/InputStateMTL.mm @@ -0,0 +1,97 @@ +// 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 "InputStateMTL.h" + +#include "MetalBackend.h" + +namespace backend { +namespace metal { + + namespace { + MTLVertexFormat VertexFormatType(nxt::VertexFormat format) { + switch (format) { + case nxt::VertexFormat::FloatR32G32B32A32: + return MTLVertexFormatFloat4; + case nxt::VertexFormat::FloatR32G32B32: + return MTLVertexFormatFloat3; + case nxt::VertexFormat::FloatR32G32: + return MTLVertexFormatFloat2; + } + } + + MTLVertexStepFunction InputStepModeFunction(nxt::InputStepMode mode) { + switch (mode) { + case nxt::InputStepMode::Vertex: + return MTLVertexStepFunctionPerVertex; + case nxt::InputStepMode::Instance: + return MTLVertexStepFunctionPerInstance; + } + } + } + + InputState::InputState(InputStateBuilder* builder) + : InputStateBase(builder) { + mtlVertexDescriptor = [MTLVertexDescriptor new]; + + const auto& attributesSetMask = GetAttributesSetMask(); + for (size_t i = 0; i < attributesSetMask.size(); ++i) { + if (!attributesSetMask[i]) { + continue; + } + const AttributeInfo& info = GetAttribute(i); + + auto attribDesc = [MTLVertexAttributeDescriptor new]; + attribDesc.format = VertexFormatType(info.format); + attribDesc.offset = info.offset; + attribDesc.bufferIndex = kMaxBindingsPerGroup + info.bindingSlot; + mtlVertexDescriptor.attributes[i] = attribDesc; + [attribDesc release]; + } + + const auto& inputsSetMask = GetInputsSetMask(); + for (size_t i = 0; i < inputsSetMask.size(); ++i) { + if (!inputsSetMask[i]) { + continue; + } + const InputInfo& info = GetInput(i); + + auto layoutDesc = [MTLVertexBufferLayoutDescriptor new]; + if (info.stride == 0) { + // For MTLVertexStepFunctionConstant, the stepRate must be 0, + // but the stride must NOT be 0, so I made up a value (256). + layoutDesc.stepFunction = MTLVertexStepFunctionConstant; + layoutDesc.stepRate = 0; + layoutDesc.stride = 256; + } else { + layoutDesc.stepFunction = InputStepModeFunction(info.stepMode); + layoutDesc.stepRate = 1; + layoutDesc.stride = info.stride; + } + mtlVertexDescriptor.layouts[kMaxBindingsPerGroup + i] = layoutDesc; + [layoutDesc release]; + } + } + + InputState::~InputState() { + [mtlVertexDescriptor release]; + mtlVertexDescriptor = nil; + } + + MTLVertexDescriptor* InputState::GetMTLVertexDescriptor() { + return mtlVertexDescriptor; + } + +} +} diff --git a/src/backend/metal/MetalBackend.h b/src/backend/metal/MetalBackend.h index f0f582dc8a..968dedf434 100644 --- a/src/backend/metal/MetalBackend.h +++ b/src/backend/metal/MetalBackend.h @@ -17,35 +17,18 @@ #include "nxt/nxtcpp.h" -#include -#include -#include - -#include "common/Buffer.h" #include "common/BindGroup.h" #include "common/BindGroupLayout.h" #include "common/Device.h" -#include "common/CommandBuffer.h" -#include "common/DepthStencilState.h" -#include "common/InputState.h" #include "common/Framebuffer.h" -#include "common/Pipeline.h" -#include "common/PipelineLayout.h" #include "common/Queue.h" #include "common/RenderPass.h" -#include "common/Sampler.h" -#include "common/ShaderModule.h" -#include "common/Texture.h" #include "common/ToBackend.h" #include #import #import -namespace spirv_cross { - class CompilerMSL; -} - namespace backend { namespace metal { @@ -152,70 +135,6 @@ namespace metal { Device* device; }; - class Buffer : public BufferBase { - public: - Buffer(Device* device, BufferBuilder* builder); - ~Buffer(); - - id GetMTLBuffer(); - std::mutex& GetMutex(); - - private: - void SetSubDataImpl(uint32_t start, uint32_t count, const uint32_t* data) override; - void MapReadAsyncImpl(uint32_t serial, uint32_t start, uint32_t count) override; - void UnmapImpl() override; - void TransitionUsageImpl(nxt::BufferUsageBit currentUsage, nxt::BufferUsageBit targetUsage) override; - - Device* device; - std::mutex mutex; - id mtlBuffer = nil; - }; - - class BufferView : public BufferViewBase { - public: - BufferView(Device* device, BufferViewBuilder* builder); - - private: - Device* device; - }; - - class CommandBuffer : public CommandBufferBase { - public: - CommandBuffer(Device* device, CommandBufferBuilder* builder); - ~CommandBuffer(); - - void FillCommands(id commandBuffer, std::unordered_set* mutexes); - - private: - Device* device; - CommandIterator commands; - }; - - class DepthStencilState : public DepthStencilStateBase { - public: - DepthStencilState(Device* device, DepthStencilStateBuilder* builder); - ~DepthStencilState(); - - id GetMTLDepthStencilState(); - - private: - Device* device; - - id mtlDepthStencilState = nil; - }; - - class InputState : public InputStateBase { - public: - InputState(Device* device, InputStateBuilder* builder); - ~InputState(); - - MTLVertexDescriptor* GetMTLVertexDescriptor(); - - private: - Device* device; - MTLVertexDescriptor* mtlVertexDescriptor = nil; - }; - class Framebuffer : public FramebufferBase { public: Framebuffer(Device* device, FramebufferBuilder* builder); @@ -225,35 +144,6 @@ namespace metal { Device* device; }; - class Pipeline : public PipelineBase { - public: - Pipeline(Device* device, PipelineBuilder* builder); - ~Pipeline(); - - void Encode(id encoder); - void Encode(id encoder); - MTLSize GetLocalWorkGroupSize() const; - - private: - Device* device; - - id mtlRenderPipelineState = nil; - id mtlComputePipelineState = nil; - MTLSize localWorkgroupSize; - }; - - class PipelineLayout : public PipelineLayoutBase { - public: - PipelineLayout(Device* device, PipelineLayoutBuilder* builder); - - using BindingIndexInfo = std::array, kMaxBindGroups>; - const BindingIndexInfo& GetBindingIndexInfo(nxt::ShaderStage stage) const; - - private: - Device* device; - PerStage indexInfo; - }; - class Queue : public QueueBase { public: Queue(Device* device, QueueBuilder* builder); @@ -278,54 +168,6 @@ namespace metal { Device* device; }; - class Sampler : public SamplerBase { - public: - Sampler(Device* device, SamplerBuilder* builder); - ~Sampler(); - - id GetMTLSamplerState(); - - private: - Device* device; - id mtlSamplerState = nil; - }; - - class ShaderModule : public ShaderModuleBase { - public: - ShaderModule(Device* device, ShaderModuleBuilder* builder); - ~ShaderModule(); - - id GetFunction(const char* functionName) const; - MTLSize GetLocalWorkGroupSize(const std::string& entryPoint) const; - - private: - Device* device; - id mtlLibrary = nil; - spirv_cross::CompilerMSL* compiler = nullptr; - }; - - class Texture : public TextureBase { - public: - Texture(Device* device, TextureBuilder* builder); - ~Texture(); - - id GetMTLTexture(); - - private: - void TransitionUsageImpl(nxt::TextureUsageBit currentUsage, nxt::TextureUsageBit targetUsage) override; - - Device* device; - id mtlTexture = nil; - }; - - class TextureView : public TextureViewBase { - public: - TextureView(Device* device, TextureViewBuilder* builder); - - private: - Device* device; - }; - } } diff --git a/src/backend/metal/MetalBackend.mm b/src/backend/metal/MetalBackend.mm index e406bbfd1e..5251603c43 100644 --- a/src/backend/metal/MetalBackend.mm +++ b/src/backend/metal/MetalBackend.mm @@ -16,11 +16,15 @@ #include "MetalBackend.h" -#include - -#include - -#include "common/Commands.h" +#include "BufferMTL.h" +#include "CommandBufferMTL.h" +#include "DepthStencilStateMTL.h" +#include "InputStateMTL.h" +#include "PipelineMTL.h" +#include "PipelineLayoutMTL.h" +#include "SamplerMTL.h" +#include "ShaderModuleMTL.h" +#include "TextureMTL.h" namespace backend { namespace metal { @@ -72,28 +76,28 @@ namespace metal { return new BindGroupLayout(this, builder); } BufferBase* Device::CreateBuffer(BufferBuilder* builder) { - return new Buffer(this, builder); + return new Buffer(builder); } BufferViewBase* Device::CreateBufferView(BufferViewBuilder* builder) { - return new BufferView(this, builder); + return new BufferView(builder); } CommandBufferBase* Device::CreateCommandBuffer(CommandBufferBuilder* builder) { return new CommandBuffer(this, builder); } DepthStencilStateBase* Device::CreateDepthStencilState(DepthStencilStateBuilder* builder) { - return new DepthStencilState(this, builder); + return new DepthStencilState(builder); } InputStateBase* Device::CreateInputState(InputStateBuilder* builder) { - return new InputState(this, builder); + return new InputState(builder); } FramebufferBase* Device::CreateFramebuffer(FramebufferBuilder* builder) { return new Framebuffer(this, builder); } PipelineBase* Device::CreatePipeline(PipelineBuilder* builder) { - return new Pipeline(this, builder); + return new Pipeline(builder); } PipelineLayoutBase* Device::CreatePipelineLayout(PipelineLayoutBuilder* builder) { - return new PipelineLayout(this, builder); + return new PipelineLayout(builder); } QueueBase* Device::CreateQueue(QueueBuilder* builder) { return new Queue(this, builder); @@ -102,16 +106,16 @@ namespace metal { return new RenderPass(this, builder); } SamplerBase* Device::CreateSampler(SamplerBuilder* builder) { - return new Sampler(this, builder); + return new Sampler(builder); } ShaderModuleBase* Device::CreateShaderModule(ShaderModuleBuilder* builder) { - return new ShaderModule(this, builder); + return new ShaderModule(builder); } TextureBase* Device::CreateTexture(TextureBuilder* builder) { - return new Texture(this, builder); + return new Texture(builder); } TextureViewBase* Device::CreateTextureView(TextureViewBuilder* builder) { - return new TextureView(this, builder); + return new TextureView(builder); } void Device::TickImpl() { @@ -197,669 +201,6 @@ namespace metal { : BindGroupLayoutBase(builder), device(device) { } - // Buffer - - Buffer::Buffer(Device* device, BufferBuilder* builder) - : BufferBase(builder), device(device) { - mtlBuffer = [device->GetMTLDevice() newBufferWithLength:GetSize() - options:MTLResourceStorageModeManaged]; - } - - Buffer::~Buffer() { - std::lock_guard lock(mutex); - [mtlBuffer release]; - mtlBuffer = nil; - } - - id Buffer::GetMTLBuffer() { - return mtlBuffer; - } - - std::mutex& Buffer::GetMutex() { - return mutex; - } - - void Buffer::SetSubDataImpl(uint32_t start, uint32_t count, const uint32_t* data) { - uint32_t* dest = reinterpret_cast([mtlBuffer contents]); - { - std::lock_guard lock(mutex); - memcpy(&dest[start], data, count * sizeof(uint32_t)); - } - [mtlBuffer didModifyRange:NSMakeRange(start * sizeof(uint32_t), count * sizeof(uint32_t))]; - } - - void Buffer::MapReadAsyncImpl(uint32_t serial, uint32_t start, uint32_t count) { - // TODO(cwallez@chromium.org): Implement Map Read for the metal backend - } - - void Buffer::UnmapImpl() { - // TODO(cwallez@chromium.org): Implement Map Read for the metal backend - } - - void Buffer::TransitionUsageImpl(nxt::BufferUsageBit currentUsage, nxt::BufferUsageBit targetUsage) { - } - - // BufferView - - BufferView::BufferView(Device* device, BufferViewBuilder* builder) - : BufferViewBase(builder), device(device) { - } - - // CommandBuffer - - static MTLIndexType IndexFormatType(nxt::IndexFormat format) { - switch (format) { - case nxt::IndexFormat::Uint16: - return MTLIndexTypeUInt16; - case nxt::IndexFormat::Uint32: - return MTLIndexTypeUInt32; - } - } - - CommandBuffer::CommandBuffer(Device* device, CommandBufferBuilder* builder) - : CommandBufferBase(builder), device(device), commands(builder->AcquireCommands()) { - } - - CommandBuffer::~CommandBuffer() { - FreeCommands(&commands); - } - - namespace { - - struct CurrentEncoders { - Device* device; - - id blit = nil; - id compute = nil; - id render = nil; - - RenderPass* currentRenderPass = nullptr; - Framebuffer* currentFramebuffer = nullptr; - - void FinishEncoders() { - ASSERT(render == nil); - if (blit != nil) { - [blit endEncoding]; - blit = nil; - } - if (compute != nil) { - [compute endEncoding]; - compute = nil; - } - } - - void EnsureBlit(id commandBuffer) { - if (blit == nil) { - FinishEncoders(); - blit = [commandBuffer blitCommandEncoder]; - } - } - void EnsureCompute(id commandBuffer) { - if (compute == nil) { - FinishEncoders(); - compute = [commandBuffer computeCommandEncoder]; - // TODO(cwallez@chromium.org): does any state need to be reset? - } - } - void BeginSubpass(id commandBuffer, uint32_t subpass) { - ASSERT(currentRenderPass); - if (render != nil) { - [render endEncoding]; - render = nil; - } - - const auto& info = currentRenderPass->GetSubpassInfo(subpass); - - MTLRenderPassDescriptor* descriptor = [MTLRenderPassDescriptor renderPassDescriptor]; - bool usingBackbuffer = false; // HACK(kainino@chromium.org): workaround for not having depth attachments - for (uint32_t index = 0; index < info.colorAttachments.size(); ++index) { - uint32_t attachment = info.colorAttachments[index]; - - // TODO(kainino@chromium.org): currently a 'null' texture view - // falls back to the 'back buffer' but this should go away - // when we have WSI. - id texture = nil; - if (auto textureView = currentFramebuffer->GetTextureView(attachment)) { - texture = ToBackend(textureView->GetTexture())->GetMTLTexture(); - } else { - texture = device->GetCurrentTexture(); - usingBackbuffer = true; - } - descriptor.colorAttachments[index].texture = texture; - descriptor.colorAttachments[index].loadAction = MTLLoadActionLoad; - descriptor.colorAttachments[index].storeAction = MTLStoreActionStore; - } - // TODO(kainino@chromium.org): load depth attachment from subpass - if (usingBackbuffer) { - descriptor.depthAttachment.texture = device->GetCurrentDepthTexture(); - descriptor.depthAttachment.loadAction = MTLLoadActionLoad; - descriptor.depthAttachment.storeAction = MTLStoreActionStore; - } - - render = [commandBuffer renderCommandEncoderWithDescriptor:descriptor]; - // TODO(cwallez@chromium.org): does any state need to be reset? - } - void EndRenderPass() { - ASSERT(render != nil); - [render endEncoding]; - render = nil; - } - }; - - } - - void CommandBuffer::FillCommands(id commandBuffer, std::unordered_set* mutexes) { - Command type; - Pipeline* lastPipeline = nullptr; - id indexBuffer = nil; - uint32_t indexBufferOffset = 0; - MTLIndexType indexType = MTLIndexTypeUInt32; - - CurrentEncoders encoders; - encoders.device = device; - - uint32_t currentSubpass = 0; - id renderEncoder = nil; - - while (commands.NextCommandId(&type)) { - switch (type) { - case Command::AdvanceSubpass: - { - commands.NextCommand(); - currentSubpass += 1; - encoders.BeginSubpass(commandBuffer, currentSubpass); - } - break; - - case Command::BeginRenderPass: - { - BeginRenderPassCmd* beginRenderPassCmd = commands.NextCommand(); - encoders.currentRenderPass = ToBackend(beginRenderPassCmd->renderPass.Get()); - encoders.currentFramebuffer = ToBackend(beginRenderPassCmd->framebuffer.Get()); - encoders.FinishEncoders(); - currentSubpass = 0; - encoders.BeginSubpass(commandBuffer, currentSubpass); - } - break; - - case Command::CopyBufferToBuffer: - { - CopyBufferToBufferCmd* copy = commands.NextCommand(); - - encoders.EnsureBlit(commandBuffer); - [encoders.blit - copyFromBuffer:ToBackend(copy->source)->GetMTLBuffer() - sourceOffset:copy->sourceOffset - toBuffer:ToBackend(copy->destination)->GetMTLBuffer() - destinationOffset:copy->destinationOffset - size:copy->size]; - } - break; - - case Command::CopyBufferToTexture: - { - CopyBufferToTextureCmd* copy = commands.NextCommand(); - Buffer* buffer = ToBackend(copy->buffer.Get()); - Texture* texture = ToBackend(copy->texture.Get()); - - unsigned rowSize = copy->width * TextureFormatPixelSize(texture->GetFormat()); - MTLOrigin origin; - origin.x = copy->x; - origin.y = copy->y; - origin.z = copy->z; - - MTLSize size; - size.width = copy->width; - size.height = copy->height; - size.depth = copy->depth; - - encoders.EnsureBlit(commandBuffer); - [encoders.blit - copyFromBuffer:buffer->GetMTLBuffer() - sourceOffset:copy->bufferOffset - sourceBytesPerRow:rowSize - sourceBytesPerImage:(rowSize * copy->height) - sourceSize:size - toTexture:texture->GetMTLTexture() - destinationSlice:0 - destinationLevel:copy->level - destinationOrigin:origin]; - } - break; - - case Command::Dispatch: - { - DispatchCmd* dispatch = commands.NextCommand(); - encoders.EnsureCompute(commandBuffer); - ASSERT(lastPipeline->IsCompute()); - - [encoders.compute dispatchThreadgroups:MTLSizeMake(dispatch->x, dispatch->y, dispatch->z) - threadsPerThreadgroup: lastPipeline->GetLocalWorkGroupSize()]; - } - break; - - case Command::DrawArrays: - { - DrawArraysCmd* draw = commands.NextCommand(); - - ASSERT(encoders.render); - [encoders.render - drawPrimitives:MTLPrimitiveTypeTriangle - vertexStart:draw->firstVertex - vertexCount:draw->vertexCount - instanceCount:draw->instanceCount - baseInstance:draw->firstInstance]; - } - break; - - case Command::DrawElements: - { - DrawElementsCmd* draw = commands.NextCommand(); - - ASSERT(encoders.render); - [encoders.render - drawIndexedPrimitives:MTLPrimitiveTypeTriangle - indexCount:draw->indexCount - indexType:indexType - indexBuffer:indexBuffer - indexBufferOffset:indexBufferOffset - instanceCount:draw->instanceCount - baseVertex:0 - baseInstance:draw->firstInstance]; - } - break; - - case Command::EndRenderPass: - { - commands.NextCommand(); - encoders.EndRenderPass(); - } - break; - - case Command::SetPipeline: - { - SetPipelineCmd* cmd = commands.NextCommand(); - lastPipeline = ToBackend(cmd->pipeline).Get(); - - if (lastPipeline->IsCompute()) { - encoders.EnsureCompute(commandBuffer); - lastPipeline->Encode(encoders.compute); - } else { - ASSERT(encoders.render); - DepthStencilState* depthStencilState = ToBackend(lastPipeline->GetDepthStencilState()); - [encoders.render setDepthStencilState:depthStencilState->GetMTLDepthStencilState()]; - lastPipeline->Encode(encoders.render); - } - } - break; - - case Command::SetPushConstants: - { - SetPushConstantsCmd* cmd = commands.NextCommand(); - uint32_t* valuesUInt = commands.NextData(cmd->count); - int32_t* valuesInt = reinterpret_cast(valuesUInt); - float* valuesFloat = reinterpret_cast(valuesUInt); - - // TODO(kainino@chromium.org): implement SetPushConstants - } - break; - - case Command::SetStencilReference: - { - SetStencilReferenceCmd* cmd = commands.NextCommand(); - - ASSERT(encoders.render); - - [encoders.render setStencilReferenceValue:cmd->reference]; - } - break; - - case Command::SetBindGroup: - { - SetBindGroupCmd* cmd = commands.NextCommand(); - BindGroup* group = ToBackend(cmd->group.Get()); - uint32_t groupIndex = cmd->index; - - const auto& layout = group->GetLayout()->GetBindingInfo(); - - if (lastPipeline->IsCompute()) { - encoders.EnsureCompute(commandBuffer); - } 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) { - if (!layout.mask[binding]) { - continue; - } - - auto stage = layout.visibilities[binding]; - bool vertStage = stage & nxt::ShaderStageBit::Vertex; - bool fragStage = stage & nxt::ShaderStageBit::Fragment; - bool computeStage = stage & nxt::ShaderStageBit::Compute; - uint32_t vertIndex = 0; - uint32_t fragIndex = 0; - uint32_t computeIndex = 0; - if (vertStage) { - vertIndex = ToBackend(lastPipeline->GetLayout())-> - GetBindingIndexInfo(nxt::ShaderStage::Vertex)[groupIndex][binding]; - } - if (fragStage) { - fragIndex = ToBackend(lastPipeline->GetLayout())-> - GetBindingIndexInfo(nxt::ShaderStage::Fragment)[groupIndex][binding]; - } - if (computeStage) { - computeIndex = ToBackend(lastPipeline->GetLayout())-> - GetBindingIndexInfo(nxt::ShaderStage::Compute)[groupIndex][binding]; - } - - switch (layout.types[binding]) { - case nxt::BindingType::UniformBuffer: - case nxt::BindingType::StorageBuffer: - { - BufferView* view = ToBackend(group->GetBindingAsBufferView(binding)); - auto b = ToBackend(view->GetBuffer()); - mutexes->insert(&b->GetMutex()); - const id buffer = b->GetMTLBuffer(); - const NSUInteger offset = view->GetOffset(); - if (vertStage) { - [encoders.render - setVertexBuffers:&buffer - offsets:&offset - withRange:NSMakeRange(vertIndex, 1)]; - } - if (fragStage) { - [encoders.render - setFragmentBuffers:&buffer - offsets:&offset - withRange:NSMakeRange(fragIndex, 1)]; - } - if (computeStage) { - [encoders.compute - setBuffers:&buffer - offsets:&offset - withRange:NSMakeRange(computeIndex, 1)]; - } - - } - break; - - case nxt::BindingType::Sampler: - { - auto sampler = ToBackend(group->GetBindingAsSampler(binding)); - if (vertStage) { - [encoders.render - setVertexSamplerState:sampler->GetMTLSamplerState() - atIndex:vertIndex]; - } - if (fragStage) { - [encoders.render - setFragmentSamplerState:sampler->GetMTLSamplerState() - atIndex:fragIndex]; - } - if (computeStage) { - [encoders.compute - setSamplerState:sampler->GetMTLSamplerState() - atIndex:computeIndex]; - } - } - break; - - case nxt::BindingType::SampledTexture: - { - auto texture = ToBackend(group->GetBindingAsTextureView(binding)->GetTexture()); - if (vertStage) { - [encoders.render - setVertexTexture:texture->GetMTLTexture() - atIndex:vertIndex]; - } - if (fragStage) { - [encoders.render - setFragmentTexture:texture->GetMTLTexture() - atIndex:fragIndex]; - } - if (computeStage) { - [encoders.compute - setTexture:texture->GetMTLTexture() - atIndex:computeIndex]; - } - } - break; - } - } - } - break; - - case Command::SetIndexBuffer: - { - SetIndexBufferCmd* cmd = commands.NextCommand(); - auto b = ToBackend(cmd->buffer.Get()); - mutexes->insert(&b->GetMutex()); - indexBuffer = b->GetMTLBuffer(); - indexBufferOffset = cmd->offset; - indexType = IndexFormatType(cmd->format); - } - break; - - case Command::SetVertexBuffers: - { - SetVertexBuffersCmd* cmd = commands.NextCommand(); - auto buffers = commands.NextData>(cmd->count); - auto offsets = commands.NextData(cmd->count); - - auto inputState = lastPipeline->GetInputState(); - - std::array, kMaxVertexInputs> mtlBuffers; - std::array mtlOffsets; - - // Perhaps an "array of vertex buffers(+offsets?)" should be - // a NXT API primitive to avoid reconstructing this array? - for (uint32_t i = 0; i < cmd->count; ++i) { - Buffer* buffer = ToBackend(buffers[i].Get()); - mutexes->insert(&buffer->GetMutex()); - mtlBuffers[i] = buffer->GetMTLBuffer(); - mtlOffsets[i] = offsets[i]; - } - - ASSERT(encoders.render); - [encoders.render - setVertexBuffers:mtlBuffers.data() - offsets:mtlOffsets.data() - withRange:NSMakeRange(kMaxBindingsPerGroup + cmd->startSlot, cmd->count)]; - } - break; - - case Command::TransitionBufferUsage: - { - TransitionBufferUsageCmd* cmd = commands.NextCommand(); - - cmd->buffer->UpdateUsageInternal(cmd->usage); - } - break; - - case Command::TransitionTextureUsage: - { - TransitionTextureUsageCmd* cmd = commands.NextCommand(); - - cmd->texture->UpdateUsageInternal(cmd->usage); - } - break; -; - } - } - - encoders.FinishEncoders(); - } - - // DepthStencilState - - static MTLCompareFunction MetalDepthStencilCompareFunction(nxt::CompareFunction compareFunction) { - switch (compareFunction) { - case nxt::CompareFunction::Never: - return MTLCompareFunctionNever; - case nxt::CompareFunction::Less: - return MTLCompareFunctionLess; - case nxt::CompareFunction::LessEqual: - return MTLCompareFunctionLessEqual; - case nxt::CompareFunction::Greater: - return MTLCompareFunctionGreater; - case nxt::CompareFunction::GreaterEqual: - return MTLCompareFunctionGreaterEqual; - case nxt::CompareFunction::NotEqual: - return MTLCompareFunctionNotEqual; - case nxt::CompareFunction::Equal: - return MTLCompareFunctionEqual; - case nxt::CompareFunction::Always: - return MTLCompareFunctionAlways; - } - } - - static MTLStencilOperation MetalStencilOperation(nxt::StencilOperation stencilOperation) { - switch (stencilOperation) { - case nxt::StencilOperation::Keep: - return MTLStencilOperationKeep; - case nxt::StencilOperation::Zero: - return MTLStencilOperationZero; - case nxt::StencilOperation::Replace: - return MTLStencilOperationReplace; - case nxt::StencilOperation::Invert: - return MTLStencilOperationInvert; - case nxt::StencilOperation::IncrementClamp: - return MTLStencilOperationIncrementClamp; - case nxt::StencilOperation::DecrementClamp: - return MTLStencilOperationDecrementClamp; - case nxt::StencilOperation::IncrementWrap: - return MTLStencilOperationIncrementWrap; - case nxt::StencilOperation::DecrementWrap: - return MTLStencilOperationDecrementWrap; - } - } - - DepthStencilState::DepthStencilState(Device* device, DepthStencilStateBuilder* builder) - : DepthStencilStateBase(builder), device(device) { - MTLDepthStencilDescriptor* mtlDepthStencilDescriptor = [MTLDepthStencilDescriptor new]; - - if (DepthTestEnabled()) { - auto& depth = GetDepth(); - mtlDepthStencilDescriptor.depthCompareFunction = MetalDepthStencilCompareFunction(depth.compareFunction); - mtlDepthStencilDescriptor.depthWriteEnabled = depth.depthWriteEnabled; - } - - auto& stencil = GetStencil(); - - if (StencilTestEnabled()) { - MTLStencilDescriptor* backFaceStencil = [MTLStencilDescriptor new]; - MTLStencilDescriptor* frontFaceStencil = [MTLStencilDescriptor new]; - - backFaceStencil.stencilCompareFunction = MetalDepthStencilCompareFunction(stencil.back.compareFunction); - backFaceStencil.stencilFailureOperation = MetalStencilOperation(stencil.back.stencilFail); - backFaceStencil.depthFailureOperation = MetalStencilOperation(stencil.back.depthFail); - backFaceStencil.depthStencilPassOperation = MetalStencilOperation(stencil.back.depthStencilPass); - backFaceStencil.readMask = stencil.readMask; - backFaceStencil.writeMask = stencil.writeMask; - - frontFaceStencil.stencilCompareFunction = MetalDepthStencilCompareFunction(stencil.front.compareFunction); - frontFaceStencil.stencilFailureOperation = MetalStencilOperation(stencil.front.stencilFail); - frontFaceStencil.depthFailureOperation = MetalStencilOperation(stencil.front.depthFail); - frontFaceStencil.depthStencilPassOperation = MetalStencilOperation(stencil.front.depthStencilPass); - frontFaceStencil.readMask = stencil.readMask; - frontFaceStencil.writeMask = stencil.writeMask; - - mtlDepthStencilDescriptor.backFaceStencil = backFaceStencil; - mtlDepthStencilDescriptor.frontFaceStencil = frontFaceStencil; - [backFaceStencil release]; - [frontFaceStencil release]; - } - - mtlDepthStencilState = [device->GetMTLDevice() newDepthStencilStateWithDescriptor:mtlDepthStencilDescriptor]; - [mtlDepthStencilDescriptor release]; - } - - DepthStencilState::~DepthStencilState() { - [mtlDepthStencilState release]; - mtlDepthStencilState = nil; - } - - id DepthStencilState::GetMTLDepthStencilState() { - return mtlDepthStencilState; - } - - // InputState - - static MTLVertexFormat VertexFormatType(nxt::VertexFormat format) { - switch (format) { - case nxt::VertexFormat::FloatR32G32B32A32: - return MTLVertexFormatFloat4; - case nxt::VertexFormat::FloatR32G32B32: - return MTLVertexFormatFloat3; - case nxt::VertexFormat::FloatR32G32: - return MTLVertexFormatFloat2; - } - } - - static MTLVertexStepFunction InputStepModeFunction(nxt::InputStepMode mode) { - switch (mode) { - case nxt::InputStepMode::Vertex: - return MTLVertexStepFunctionPerVertex; - case nxt::InputStepMode::Instance: - return MTLVertexStepFunctionPerInstance; - } - } - - InputState::InputState(Device* device, InputStateBuilder* builder) - : InputStateBase(builder), device(device) { - mtlVertexDescriptor = [MTLVertexDescriptor new]; - - const auto& attributesSetMask = GetAttributesSetMask(); - for (size_t i = 0; i < attributesSetMask.size(); ++i) { - if (!attributesSetMask[i]) { - continue; - } - const AttributeInfo& info = GetAttribute(i); - - auto attribDesc = [MTLVertexAttributeDescriptor new]; - attribDesc.format = VertexFormatType(info.format); - attribDesc.offset = info.offset; - attribDesc.bufferIndex = kMaxBindingsPerGroup + info.bindingSlot; - mtlVertexDescriptor.attributes[i] = attribDesc; - [attribDesc release]; - } - - const auto& inputsSetMask = GetInputsSetMask(); - for (size_t i = 0; i < inputsSetMask.size(); ++i) { - if (!inputsSetMask[i]) { - continue; - } - const InputInfo& info = GetInput(i); - - auto layoutDesc = [MTLVertexBufferLayoutDescriptor new]; - if (info.stride == 0) { - // For MTLVertexStepFunctionConstant, the stepRate must be 0, - // but the stride must NOT be 0, so I made up a value (256). - layoutDesc.stepFunction = MTLVertexStepFunctionConstant; - layoutDesc.stepRate = 0; - layoutDesc.stride = 256; - } else { - layoutDesc.stepFunction = InputStepModeFunction(info.stepMode); - layoutDesc.stepRate = 1; - layoutDesc.stride = info.stride; - } - mtlVertexDescriptor.layouts[kMaxBindingsPerGroup + i] = layoutDesc; - [layoutDesc release]; - } - } - - InputState::~InputState() { - [mtlVertexDescriptor release]; - mtlVertexDescriptor = nil; - } - - MTLVertexDescriptor* InputState::GetMTLVertexDescriptor() { - return mtlVertexDescriptor; - } - // Framebuffer Framebuffer::Framebuffer(Device* device, FramebufferBuilder* builder) @@ -869,135 +210,6 @@ namespace metal { Framebuffer::~Framebuffer() { } - // Pipeline - - Pipeline::Pipeline(Device* device, PipelineBuilder* builder) - : PipelineBase(builder), device(device) { - - if (IsCompute()) { - const auto& module = ToBackend(builder->GetStageInfo(nxt::ShaderStage::Compute).module); - const auto& entryPoint = builder->GetStageInfo(nxt::ShaderStage::Compute).entryPoint; - - id function = module->GetFunction(entryPoint.c_str()); - - NSError *error = nil; - mtlComputePipelineState = [device->GetMTLDevice() - 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 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: - ASSERT(false); - break; - } - } - - descriptor.colorAttachments[0].pixelFormat = MTLPixelFormatBGRA8Unorm; - descriptor.depthAttachmentPixelFormat = MTLPixelFormatDepth32Float; - - InputState* inputState = ToBackend(GetInputState()); - descriptor.vertexDescriptor = inputState->GetMTLVertexDescriptor(); - - // TODO(kainino@chromium.org): push constants, textures, samplers - - NSError *error = nil; - mtlRenderPipelineState = [device->GetMTLDevice() - 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 encoder) { - ASSERT(!IsCompute()); - [encoder setRenderPipelineState:mtlRenderPipelineState]; - } - - void Pipeline::Encode(id encoder) { - ASSERT(IsCompute()); - [encoder setComputePipelineState:mtlComputePipelineState]; - } - - MTLSize Pipeline::GetLocalWorkGroupSize() const { - return localWorkgroupSize; - } - - // PipelineLayout - - PipelineLayout::PipelineLayout(Device* device, PipelineLayoutBuilder* builder) - : PipelineLayoutBase(builder), device(device) { - // Each stage has its own numbering namespace in CompilerMSL. - for (auto stage : IterateStages(kAllStages)) { - uint32_t bufferIndex = 0; - uint32_t samplerIndex = 0; - uint32_t textureIndex = 0; - - for (size_t group = 0; group < kMaxBindGroups; ++group) { - const auto& groupInfo = GetBindGroupLayout(group)->GetBindingInfo(); - for (size_t binding = 0; binding < kMaxBindingsPerGroup; ++binding) { - if (!(groupInfo.visibilities[binding] & StageBit(stage))) { - continue; - } - if (!groupInfo.mask[binding]) { - continue; - } - - switch (groupInfo.types[binding]) { - case nxt::BindingType::UniformBuffer: - case nxt::BindingType::StorageBuffer: - indexInfo[stage][group][binding] = bufferIndex; - bufferIndex++; - break; - case nxt::BindingType::Sampler: - indexInfo[stage][group][binding] = samplerIndex; - samplerIndex++; - break; - case nxt::BindingType::SampledTexture: - indexInfo[stage][group][binding] = textureIndex; - textureIndex++; - break; - } - } - } - } - } - - const PipelineLayout::BindingIndexInfo& PipelineLayout::GetBindingIndexInfo(nxt::ShaderStage stage) const { - return indexInfo[stage]; - } - // Queue Queue::Queue(Device* device, QueueBuilder* builder) @@ -1049,127 +261,5 @@ namespace metal { RenderPass::~RenderPass() { } - // Sampler - - MTLSamplerMinMagFilter FilterModeToMinMagFilter(nxt::FilterMode mode) { - switch (mode) { - case nxt::FilterMode::Nearest: - return MTLSamplerMinMagFilterNearest; - case nxt::FilterMode::Linear: - return MTLSamplerMinMagFilterLinear; - } - } - - MTLSamplerMipFilter FilterModeToMipFilter(nxt::FilterMode mode) { - switch (mode) { - case nxt::FilterMode::Nearest: - return MTLSamplerMipFilterNearest; - case nxt::FilterMode::Linear: - return MTLSamplerMipFilterLinear; - } - } - - Sampler::Sampler(Device* device, SamplerBuilder* builder) - : SamplerBase(builder), device(device) { - auto desc = [MTLSamplerDescriptor new]; - [desc autorelease]; - desc.minFilter = FilterModeToMinMagFilter(builder->GetMinFilter()); - desc.magFilter = FilterModeToMinMagFilter(builder->GetMagFilter()); - desc.mipFilter = FilterModeToMipFilter(builder->GetMipMapFilter()); - // TODO(kainino@chromium.org): wrap modes - mtlSamplerState = [device->GetMTLDevice() newSamplerStateWithDescriptor:desc]; - } - - Sampler::~Sampler() { - [mtlSamplerState release]; - } - - id Sampler::GetMTLSamplerState() { - return mtlSamplerState; - } - - // ShaderModule - - ShaderModule::ShaderModule(Device* device, ShaderModuleBuilder* builder) - : ShaderModuleBase(builder), device(device) { - compiler = new spirv_cross::CompilerMSL(builder->AcquireSpirv()); - ExtractSpirvInfo(*compiler); - - std::string msl = compiler->compile(); - - NSString* mslSource = [NSString stringWithFormat:@"%s", msl.c_str()]; - NSError *error = nil; - mtlLibrary = [device->GetMTLDevice() newLibraryWithSource:mslSource options:nil error:&error]; - if (error != nil) { - NSLog(@"MTLDevice newLibraryWithSource => %@", error); - builder->HandleError("Error creating MTLLibrary from MSL source"); - } - } - - ShaderModule::~ShaderModule() { - delete compiler; - } - - id ShaderModule::GetFunction(const char* functionName) const { - // TODO(kainino@chromium.org): make this somehow more robust; it needs to behave like clean_func_name: - // https://github.com/KhronosGroup/SPIRV-Cross/blob/4e915e8c483e319d0dd7a1fa22318bef28f8cca3/spirv_msl.cpp#L1213 - if (strcmp(functionName, "main") == 0) { - functionName = "main0"; - } - NSString* name = [NSString stringWithFormat:@"%s", functionName]; - return [mtlLibrary newFunctionWithName:name]; - } - - MTLSize ShaderModule::GetLocalWorkGroupSize(const std::string& entryPoint) const { - auto size = compiler->get_entry_point(entryPoint).workgroup_size; - return MTLSizeMake(size.x, size.y, size.z); - } - - // Texture - - MTLPixelFormat TextureFormatPixelFormat(nxt::TextureFormat format) { - switch (format) { - case nxt::TextureFormat::R8G8B8A8Unorm: - return MTLPixelFormatRGBA8Unorm; - } - } - - Texture::Texture(Device* device, TextureBuilder* builder) - : TextureBase(builder), device(device) { - auto desc = [MTLTextureDescriptor new]; - [desc autorelease]; - switch (GetDimension()) { - case nxt::TextureDimension::e2D: - desc.textureType = MTLTextureType2D; - break; - } - desc.usage = MTLTextureUsageShaderRead; - desc.pixelFormat = TextureFormatPixelFormat(GetFormat()); - desc.width = GetWidth(); - desc.height = GetHeight(); - desc.depth = GetDepth(); - desc.mipmapLevelCount = GetNumMipLevels(); - desc.arrayLength = 1; - - mtlTexture = [device->GetMTLDevice() newTextureWithDescriptor:desc]; - } - - Texture::~Texture() { - [mtlTexture release]; - } - - id Texture::GetMTLTexture() { - return mtlTexture; - } - - void Texture::TransitionUsageImpl(nxt::TextureUsageBit currentUsage, nxt::TextureUsageBit targetUsage) { - } - - // TextureView - - TextureView::TextureView(Device* device, TextureViewBuilder* builder) - : TextureViewBase(builder), device(device) { - } - } } diff --git a/src/backend/metal/PipelineLayoutMTL.h b/src/backend/metal/PipelineLayoutMTL.h new file mode 100644 index 0000000000..e944e33fa4 --- /dev/null +++ b/src/backend/metal/PipelineLayoutMTL.h @@ -0,0 +1,45 @@ +// 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_PIPELINELAYOUTMTL_H_ +#define BACKEND_METAL_PIPELINELAYOUTMTL_H_ + +#include "common/PipelineLayout.h" + +#include "common/PerStage.h" + +#import + +namespace spirv_cross { + class CompilerMSL; +} + +namespace backend { +namespace metal { + + class PipelineLayout : public PipelineLayoutBase { + public: + PipelineLayout(PipelineLayoutBuilder* builder); + + using BindingIndexInfo = std::array, kMaxBindGroups>; + const BindingIndexInfo& GetBindingIndexInfo(nxt::ShaderStage stage) const; + + private: + PerStage indexInfo; + }; + +} +} + +#endif // BACKEND_METAL_PIPELINELAYOUTMTL_H_ diff --git a/src/backend/metal/PipelineLayoutMTL.mm b/src/backend/metal/PipelineLayoutMTL.mm new file mode 100644 index 0000000000..5fcadb9a4d --- /dev/null +++ b/src/backend/metal/PipelineLayoutMTL.mm @@ -0,0 +1,65 @@ +// 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 "PipelineLayoutMTL.h" + +#include "MetalBackend.h" + +namespace backend { +namespace metal { + + PipelineLayout::PipelineLayout(PipelineLayoutBuilder* builder) + : PipelineLayoutBase(builder) { + // Each stage has its own numbering namespace in CompilerMSL. + for (auto stage : IterateStages(kAllStages)) { + uint32_t bufferIndex = 0; + uint32_t samplerIndex = 0; + uint32_t textureIndex = 0; + + for (size_t group = 0; group < kMaxBindGroups; ++group) { + const auto& groupInfo = GetBindGroupLayout(group)->GetBindingInfo(); + for (size_t binding = 0; binding < kMaxBindingsPerGroup; ++binding) { + if (!(groupInfo.visibilities[binding] & StageBit(stage))) { + continue; + } + if (!groupInfo.mask[binding]) { + continue; + } + + switch (groupInfo.types[binding]) { + case nxt::BindingType::UniformBuffer: + case nxt::BindingType::StorageBuffer: + indexInfo[stage][group][binding] = bufferIndex; + bufferIndex++; + break; + case nxt::BindingType::Sampler: + indexInfo[stage][group][binding] = samplerIndex; + samplerIndex++; + break; + case nxt::BindingType::SampledTexture: + indexInfo[stage][group][binding] = textureIndex; + textureIndex++; + break; + } + } + } + } + } + + const PipelineLayout::BindingIndexInfo& PipelineLayout::GetBindingIndexInfo(nxt::ShaderStage stage) const { + return indexInfo[stage]; + } + +} +} diff --git a/src/backend/metal/PipelineMTL.h b/src/backend/metal/PipelineMTL.h new file mode 100644 index 0000000000..d3e97f23a2 --- /dev/null +++ b/src/backend/metal/PipelineMTL.h @@ -0,0 +1,43 @@ +// 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_PIPELINEMTL_H_ +#define BACKEND_METAL_PIPELINEMTL_H_ + +#include "common/Pipeline.h" + +#import + +namespace backend { +namespace metal { + + class Pipeline : public PipelineBase { + public: + Pipeline(PipelineBuilder* builder); + ~Pipeline(); + + void Encode(id encoder); + void Encode(id encoder); + MTLSize GetLocalWorkGroupSize() const; + + private: + id mtlRenderPipelineState = nil; + id mtlComputePipelineState = nil; + MTLSize localWorkgroupSize; + }; + +} +} + +#endif // BACKEND_METAL_PIPELINEMTL_H_ diff --git a/src/backend/metal/PipelineMTL.mm b/src/backend/metal/PipelineMTL.mm new file mode 100644 index 0000000000..3732c9d928 --- /dev/null +++ b/src/backend/metal/PipelineMTL.mm @@ -0,0 +1,112 @@ +// 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 "PipelineMTL.h" + +#include "DepthStencilStateMTL.h" +#include "InputStateMTL.h" +#include "MetalBackend.h" +#include "PipelineLayoutMTL.h" +#include "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 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 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: + ASSERT(false); + break; + } + } + + descriptor.colorAttachments[0].pixelFormat = MTLPixelFormatBGRA8Unorm; + 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 encoder) { + ASSERT(!IsCompute()); + [encoder setRenderPipelineState:mtlRenderPipelineState]; + } + + void Pipeline::Encode(id encoder) { + ASSERT(IsCompute()); + [encoder setComputePipelineState:mtlComputePipelineState]; + } + + MTLSize Pipeline::GetLocalWorkGroupSize() const { + return localWorkgroupSize; + } + +} +} diff --git a/src/backend/metal/SamplerMTL.h b/src/backend/metal/SamplerMTL.h new file mode 100644 index 0000000000..5990e55307 --- /dev/null +++ b/src/backend/metal/SamplerMTL.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_SAMPLERMTL_H_ +#define BACKEND_METAL_SAMPLERMTL_H_ + +#include "common/Sampler.h" + +#import + +namespace backend { +namespace metal { + + class Sampler : public SamplerBase { + public: + Sampler(SamplerBuilder* builder); + ~Sampler(); + + id GetMTLSamplerState(); + + private: + id mtlSamplerState = nil; + }; + +} +} + +#endif // BACKEND_METAL_SAMPLERMTL_H_ diff --git a/src/backend/metal/SamplerMTL.mm b/src/backend/metal/SamplerMTL.mm new file mode 100644 index 0000000000..fbaa8b2718 --- /dev/null +++ b/src/backend/metal/SamplerMTL.mm @@ -0,0 +1,64 @@ +// 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 "SamplerMTL.h" + +#include "MetalBackend.h" + +namespace backend { +namespace metal { + + namespace { + MTLSamplerMinMagFilter FilterModeToMinMagFilter(nxt::FilterMode mode) { + switch (mode) { + case nxt::FilterMode::Nearest: + return MTLSamplerMinMagFilterNearest; + case nxt::FilterMode::Linear: + return MTLSamplerMinMagFilterLinear; + } + } + + MTLSamplerMipFilter FilterModeToMipFilter(nxt::FilterMode mode) { + switch (mode) { + case nxt::FilterMode::Nearest: + return MTLSamplerMipFilterNearest; + case nxt::FilterMode::Linear: + return MTLSamplerMipFilterLinear; + } + } + } + + Sampler::Sampler(SamplerBuilder* builder) + : SamplerBase(builder) { + auto desc = [MTLSamplerDescriptor new]; + [desc autorelease]; + desc.minFilter = FilterModeToMinMagFilter(builder->GetMinFilter()); + desc.magFilter = FilterModeToMinMagFilter(builder->GetMagFilter()); + desc.mipFilter = FilterModeToMipFilter(builder->GetMipMapFilter()); + + // TODO(kainino@chromium.org): wrap modes + auto mtlDevice = ToBackend(builder->GetDevice())->GetMTLDevice(); + mtlSamplerState = [mtlDevice newSamplerStateWithDescriptor:desc]; + } + + Sampler::~Sampler() { + [mtlSamplerState release]; + } + + id Sampler::GetMTLSamplerState() { + return mtlSamplerState; + } + +} +} diff --git a/src/backend/metal/ShaderModuleMTL.h b/src/backend/metal/ShaderModuleMTL.h new file mode 100644 index 0000000000..6b7c3f7923 --- /dev/null +++ b/src/backend/metal/ShaderModuleMTL.h @@ -0,0 +1,45 @@ +// 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_SHADERMODULEMTL_H_ +#define BACKEND_METAL_SHADERMODULEMTL_H_ + +#include "common/ShaderModule.h" + +#import + +namespace spirv_cross { + class CompilerMSL; +} + +namespace backend { +namespace metal { + + class ShaderModule : public ShaderModuleBase { + public: + ShaderModule(ShaderModuleBuilder* builder); + ~ShaderModule(); + + id GetFunction(const char* functionName) const; + MTLSize GetLocalWorkGroupSize(const std::string& entryPoint) const; + + private: + id mtlLibrary = nil; + spirv_cross::CompilerMSL* compiler = nullptr; + }; + +} +} + +#endif // BACKEND_METAL_SHADERMODULEMTL_H_ diff --git a/src/backend/metal/ShaderModuleMTL.mm b/src/backend/metal/ShaderModuleMTL.mm new file mode 100644 index 0000000000..5247eaf70e --- /dev/null +++ b/src/backend/metal/ShaderModuleMTL.mm @@ -0,0 +1,63 @@ +// 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 "ShaderModuleMTL.h" + +#include "MetalBackend.h" + +#include + +#include + +namespace backend { +namespace metal { + + ShaderModule::ShaderModule(ShaderModuleBuilder* builder) + : ShaderModuleBase(builder) { + compiler = new spirv_cross::CompilerMSL(builder->AcquireSpirv()); + ExtractSpirvInfo(*compiler); + + std::string msl = compiler->compile(); + NSString* mslSource = [NSString stringWithFormat:@"%s", msl.c_str()]; + + auto mtlDevice = ToBackend(builder->GetDevice())->GetMTLDevice(); + NSError *error = nil; + mtlLibrary = [mtlDevice newLibraryWithSource:mslSource options:nil error:&error]; + if (error != nil) { + NSLog(@"MTLDevice newLibraryWithSource => %@", error); + builder->HandleError("Error creating MTLLibrary from MSL source"); + } + } + + ShaderModule::~ShaderModule() { + delete compiler; + } + + id ShaderModule::GetFunction(const char* functionName) const { + // TODO(kainino@chromium.org): make this somehow more robust; it needs to behave like clean_func_name: + // https://github.com/KhronosGroup/SPIRV-Cross/blob/4e915e8c483e319d0dd7a1fa22318bef28f8cca3/spirv_msl.cpp#L1213 + if (strcmp(functionName, "main") == 0) { + functionName = "main0"; + } + NSString* name = [NSString stringWithFormat:@"%s", functionName]; + return [mtlLibrary newFunctionWithName:name]; + } + + MTLSize ShaderModule::GetLocalWorkGroupSize(const std::string& entryPoint) const { + auto size = compiler->get_entry_point(entryPoint).workgroup_size; + return MTLSizeMake(size.x, size.y, size.z); + } + +} +} diff --git a/src/backend/metal/TextureMTL.h b/src/backend/metal/TextureMTL.h new file mode 100644 index 0000000000..bc72cd1629 --- /dev/null +++ b/src/backend/metal/TextureMTL.h @@ -0,0 +1,46 @@ +// 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_TEXTUREMTL_H_ +#define BACKEND_METAL_TEXTUREMTL_H_ + +#include "common/Texture.h" + +#import + +namespace backend { +namespace metal { + + class Texture : public TextureBase { + public: + Texture(TextureBuilder* builder); + ~Texture(); + + id GetMTLTexture(); + + private: + void TransitionUsageImpl(nxt::TextureUsageBit currentUsage, nxt::TextureUsageBit targetUsage) override; + + id mtlTexture = nil; + }; + + class TextureView : public TextureViewBase { + public: + TextureView(TextureViewBuilder* builder); + }; + +} +} + +#endif // BACKEND_METAL_TEXTUREMTL_H_ diff --git a/src/backend/metal/TextureMTL.mm b/src/backend/metal/TextureMTL.mm new file mode 100644 index 0000000000..36b8a45f74 --- /dev/null +++ b/src/backend/metal/TextureMTL.mm @@ -0,0 +1,67 @@ +// 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 "TextureMTL.h" + +#include "MetalBackend.h" + +namespace backend { +namespace metal { + + namespace { + MTLPixelFormat TextureFormatPixelFormat(nxt::TextureFormat format) { + switch (format) { + case nxt::TextureFormat::R8G8B8A8Unorm: + return MTLPixelFormatRGBA8Unorm; + } + } + } + + Texture::Texture(TextureBuilder* builder) + : TextureBase(builder) { + auto desc = [MTLTextureDescriptor new]; + [desc autorelease]; + switch (GetDimension()) { + case nxt::TextureDimension::e2D: + desc.textureType = MTLTextureType2D; + break; + } + desc.usage = MTLTextureUsageShaderRead; + desc.pixelFormat = TextureFormatPixelFormat(GetFormat()); + desc.width = GetWidth(); + desc.height = GetHeight(); + desc.depth = GetDepth(); + desc.mipmapLevelCount = GetNumMipLevels(); + desc.arrayLength = 1; + + auto mtlDevice = ToBackend(builder->GetDevice())->GetMTLDevice(); + mtlTexture = [mtlDevice newTextureWithDescriptor:desc]; + } + + Texture::~Texture() { + [mtlTexture release]; + } + + id Texture::GetMTLTexture() { + return mtlTexture; + } + + void Texture::TransitionUsageImpl(nxt::TextureUsageBit currentUsage, nxt::TextureUsageBit targetUsage) { + } + + TextureView::TextureView(TextureViewBuilder* builder) + : TextureViewBase(builder) { + } +} +} diff --git a/src/backend/opengl/CommandBufferGL.h b/src/backend/opengl/CommandBufferGL.h index 1492551ed3..d380f42fe6 100644 --- a/src/backend/opengl/CommandBufferGL.h +++ b/src/backend/opengl/CommandBufferGL.h @@ -12,20 +12,16 @@ // See the License for the specific language governing permissions and // limitations under the License. -#ifndef BACKEND_OPENGL_COMMANDBUFFER_H_ -#define BACKEND_OPENGL_COMMANDBUFFER_H_ +#ifndef BACKEND_OPENGL_COMMANDBUFFERGL_H_ +#define BACKEND_OPENGL_COMMANDBUFFERGL_H_ #include "common/CommandAllocator.h" #include "common/CommandBuffer.h" -namespace backend { - class CommandBufferBuilder; -} - namespace backend { namespace opengl { - class Device; + class Device; class CommandBuffer : public CommandBufferBase { public: @@ -42,4 +38,4 @@ namespace opengl { } } -#endif // BACKEND_OPENGL_COMMANDBUFFER_H_ +#endif // BACKEND_OPENGL_COMMANDBUFFERGL_H_