Dynamitos/Seele
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Trying to fix metal shader compilation

Browse Source
This commit is contained in:
Dynamitos
2024-04-19 18:23:36 +02:00
parent 21636f2460
commit 194a0e8b91
43 changed files with 353 additions and 183 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/Engine/Graphics/Metal/Command.h
+4 -1
View File
@@ -2,6 +2,7 @@
#include "Buffer.h"
#include "Graphics/Command.h"
#include "Metal/MTLBlitCommandEncoder.hpp"
#include "Metal/MTLCaptureManager.hpp"
#include "Metal/MTLCommandBuffer.hpp"
#include "Metal/MTLComputeCommandEncoder.hpp"
#include "Metal/MTLRenderCommandEncoder.hpp"
@@ -16,6 +17,7 @@ DECLARE_REF(RenderCommand)
DECLARE_REF(Graphics)
DECLARE_REF(IndexBuffer)
DECLARE_REF(GraphicsPipeline)
DECLARE_REF(ComputePipeline)
class Command {
public:
Command(PGraphics graphics, MTL::CommandBuffer* cmdBuffer);
@@ -85,6 +87,7 @@ public:
virtual void dispatch(uint32 threadX, uint32 threadY, uint32 threadZ) override;
private:
PComputePipeline boundPipeline;
MTL::CommandBuffer* commandBuffer;
MTL::ComputeCommandEncoder* encoder;
std::string name;
@@ -123,4 +126,4 @@ private:
};
DEFINE_REF(IOCommandQueue)
} // namespace Metal
} // namespace Seele
} // namespace Seele
src/Engine/Graphics/Metal/Command.mm
+40 -7
View File
@@ -7,6 +7,7 @@
#include "Graphics/Enums.h"
#include "Graphics/Graphics.h"
#include "Graphics/Resources.h"
#include "Metal/MTLCaptureManager.hpp"
#include "Pipeline.h"
#include "Resources.h"
#include "Window.h"
@@ -77,10 +78,11 @@ void RenderCommand::bindPipeline(Gfx::PGraphicsPipeline pipeline) {
}
void RenderCommand::bindDescriptor(Gfx::PDescriptorSet descriptorSet) {
encoder->setVertexBuffer(descriptorSet.cast<DescriptorSet>()->getBuffer(), 0, descriptorSet->getSetIndex());
encoder->setFragmentBuffer(descriptorSet.cast<DescriptorSet>()->getBuffer(), 0, descriptorSet->getSetIndex());
encoder->setMeshBuffer(descriptorSet.cast<DescriptorSet>()->getBuffer(), 0, descriptorSet->getSetIndex());
encoder->setObjectBuffer(descriptorSet.cast<DescriptorSet>()->getBuffer(), 0, descriptorSet->getSetIndex());
uint32 parameterIndex = boundPipeline->getPipelineLayout()->findParameter(descriptorSet->getLayout()->getName());
encoder->setVertexBuffer(descriptorSet.cast<DescriptorSet>()->getBuffer(), 0, parameterIndex);
encoder->setFragmentBuffer(descriptorSet.cast<DescriptorSet>()->getBuffer(), 0, parameterIndex);
encoder->setMeshBuffer(descriptorSet.cast<DescriptorSet>()->getBuffer(), 0, parameterIndex);
encoder->setObjectBuffer(descriptorSet.cast<DescriptorSet>()->getBuffer(), 0, parameterIndex);
}
void RenderCommand::bindDescriptor(const Array<Gfx::PDescriptorSet>& descriptorSets) {
@@ -135,13 +137,44 @@ void ComputeCommand::end() {
}
void ComputeCommand::bindPipeline(Gfx::PComputePipeline pipeline) {
encoder->setComputePipelineState(pipeline.cast<ComputePipeline>()->getHandle());
boundPipeline = pipeline.cast<ComputePipeline>();
encoder->setComputePipelineState(boundPipeline->getHandle());
}
void ComputeCommand::bindDescriptor(Gfx::PDescriptorSet set) {
auto metalSet = set.cast<DescriptorSet>();
metalSet->bind();
encoder->setBuffer(metalSet->getBuffer(), 0, set->getSetIndex());
uint32 parameterIndex = boundPipeline->getPipelineLayout()->findParameter(set->getLayout()->getName());
encoder->setBuffer(metalSet->getBuffer(), 0, parameterIndex);
auto bindings =metalSet->getLayout()->getBindings();
for(size_t i = 0; i < bindings.size(); ++i)
{
auto binding = bindings[i];
if(binding.descriptorType == Gfx::SE_DESCRIPTOR_TYPE_SAMPLER)
{
continue;
}
MTL::ResourceUsage usage;
switch(binding.access) {
case Gfx::SE_DESCRIPTOR_ACCESS_READ_ONLY_BIT:
if(binding.descriptorType == Gfx::SE_DESCRIPTOR_TYPE_SAMPLED_IMAGE)
{
usage = MTL::ResourceUsageSample;
break;
}else
{
usage = MTL::ResourceUsageRead;
break;
}
case Gfx::SE_DESCRIPTOR_ACCESS_READ_WRITE_BIT:
usage = MTL::ResourceUsageRead | MTL::ResourceUsageWrite;
break;
case Gfx::SE_DESCRIPTOR_ACCESS_WRITE_ONLY_BIT:
usage = MTL::ResourceUsageWrite;
break;
}
encoder->useResource(metalSet->getBoundResources()[i], usage);
}
}
void ComputeCommand::bindDescriptor(const Array<Gfx::PDescriptorSet>& sets) {
@@ -157,7 +190,7 @@ void ComputeCommand::pushConstants(Gfx::PPipelineLayout, Gfx::SeShaderStageFlags
void ComputeCommand::dispatch(uint32 threadX, uint32 threadY, uint32 threadZ) {
// TODO
encoder->dispatchThreadgroups(MTL::Size(threadX, threadY, threadZ), MTL::Size(32, 32, 32));
encoder->dispatchThreadgroups(MTL::Size(threadX, threadY, threadZ), MTL::Size(32, 32, 1));
}
CommandQueue::CommandQueue(PGraphics graphics) : graphics(graphics) {
src/Engine/Graphics/Metal/Descriptor.h
+3 -2
View File
@@ -48,7 +48,6 @@ public:
virtual void updateSampler(uint32_t binding, Gfx::PSampler samplerState);
virtual void updateTexture(uint32_t binding, Gfx::PTexture texture, Gfx::PSampler sampler = nullptr);
virtual void updateTextureArray(uint32_t binding, Array<Gfx::PTexture> texture);
virtual bool operator<(Gfx::PDescriptorSet other);
constexpr bool isCurrentlyBound() const { return bindCount > 0; }
constexpr bool isCurrentlyInUse() const { return currentlyInUse; }
@@ -58,12 +57,14 @@ public:
constexpr void free() { currentlyInUse = false; }
constexpr MTL::Buffer* getBuffer() const { return buffer; }
constexpr const Array<MTL::Resource*>& getBoundResources() const { return boundResources; }
private:
PGraphics graphics;
PDescriptorPool owner;
MTL::Buffer* buffer;
MTL::ArgumentEncoder* encoder;
Array<MTL::Resource*> boundResources;
uint32 bindCount;
bool currentlyInUse;
};
@@ -80,4 +81,4 @@ private:
};
DEFINE_REF(PipelineLayout)
} // namespace Metal
} // namespace Seele
} // namespace Seele
src/Engine/Graphics/Metal/Descriptor.mm
+8 -9
View File
@@ -122,6 +122,7 @@ DescriptorSet::DescriptorSet(PGraphics graphics, PDescriptorPool owner)
if (owner->getArguments()->count() == 0) {
buffer = graphics->getDevice()->newBuffer(0, MTL::ResourceOptionCPUCacheModeDefault);
} else {
boundResources.resize(owner->getArguments()->count());
encoder = graphics->getDevice()->newArgumentEncoder(owner->getArguments());
buffer = graphics->getDevice()->newBuffer(encoder->encodedLength(), MTL::ResourceOptionCPUCacheModeDefault);
encoder->setArgumentBuffer(buffer, 0);
@@ -135,10 +136,12 @@ void DescriptorSet::writeChanges() {}
void DescriptorSet::updateBuffer(uint32_t binding, Gfx::PUniformBuffer uniformBuffer) {
PUniformBuffer metalBuffer = uniformBuffer.cast<UniformBuffer>();
encoder->setBuffer(metalBuffer->getHandle(), 0, binding);
boundResources[binding] = metalBuffer->getHandle();
}
void DescriptorSet::updateBuffer(uint32_t binding, Gfx::PShaderBuffer uniformBuffer) {
PShaderBuffer metalBuffer = uniformBuffer.cast<ShaderBuffer>();
encoder->setBuffer(metalBuffer->getHandle(), 0, binding);
boundResources[binding] = metalBuffer->getHandle();
}
void DescriptorSet::updateSampler(uint32_t binding, Gfx::PSampler samplerState) {
PSampler sampler = samplerState.cast<Sampler>();
@@ -147,30 +150,26 @@ void DescriptorSet::updateSampler(uint32_t binding, Gfx::PSampler samplerState)
void DescriptorSet::updateTexture(uint32_t binding, Gfx::PTexture texture, Gfx::PSampler samplerState) {
PTextureBase base = texture.cast<TextureBase>();
encoder->setTexture(base->getTexture(), binding);
if (samplerState != nullptr) {
PSampler sampler = samplerState.cast<Sampler>();
encoder->setSamplerState(sampler->getHandle(), binding);
}
boundResources[binding] = base->getTexture();
}
void DescriptorSet::updateTextureArray(uint32_t binding, Array<Gfx::PTexture> array) {
for (auto& t : array) {
PTextureBase metalTexture = t.cast<TextureBase>();
encoder->setTexture(metalTexture->getTexture(), binding++);
encoder->setTexture(metalTexture->getTexture(), binding);
boundResources[binding++] = metalTexture->getTexture();
}
}
bool DescriptorSet::operator<(Gfx::PDescriptorSet other) { return this < other.getHandle(); }
PipelineLayout::PipelineLayout(PGraphics graphics, Gfx::PPipelineLayout baseLayout)
: Gfx::PipelineLayout(baseLayout), graphics(graphics) {}
PipelineLayout::~PipelineLayout() {}
void PipelineLayout::create() {
for (auto& set : descriptorSetLayouts) {
for (auto& [_, set] : descriptorSetLayouts) {
set->create();
uint32 setHash = set->getHash();
layoutHash = CRC::Calculate(&setHash, sizeof(uint32), CRC::CRC_32(), layoutHash);
}
}
}
src/Engine/Graphics/Metal/Pipeline.h
+3 -3
View File
@@ -15,7 +15,7 @@ DECLARE_REF(PipelineLayout)
DECLARE_REF(Graphics)
class GraphicsPipeline : public Gfx::GraphicsPipeline {
public:
GraphicsPipeline(PGraphics graphics, MTL::PrimitiveType primitive, MTL::RenderPipelineState* pipeline, Gfx::OPipelineLayout createInfo);
GraphicsPipeline(PGraphics graphics, MTL::PrimitiveType primitive, MTL::RenderPipelineState* pipeline, Gfx::PPipelineLayout createInfo);
virtual ~GraphicsPipeline();
constexpr MTL::RenderPipelineState* getHandle() const { return state; }
constexpr MTL::PrimitiveType getPrimitive() const { return primitiveType; }
@@ -27,7 +27,7 @@ private:
DEFINE_REF(GraphicsPipeline)
class ComputePipeline : public Gfx::ComputePipeline {
public:
ComputePipeline(PGraphics graphics, MTL::ComputePipelineState* pipeline, Gfx::OPipelineLayout);
ComputePipeline(PGraphics graphics, MTL::ComputePipelineState* pipeline, Gfx::PPipelineLayout);
virtual ~ComputePipeline();
constexpr MTL::ComputePipelineState* getHandle() const { return state; }
@@ -37,4 +37,4 @@ private:
};
DEFINE_REF(ComputePipeline)
} // namespace Metal
} // namespace Seele
} // namespace Seele
src/Engine/Graphics/Metal/Pipeline.mm
+4 -4
View File
@@ -12,13 +12,13 @@ VertexInput::VertexInput(VertexInputStateCreateInfo createInfo) : Gfx::VertexInp
VertexInput::~VertexInput() {}
GraphicsPipeline::GraphicsPipeline(PGraphics graphics, MTL::PrimitiveType primitive, MTL::RenderPipelineState* pipeline,
Gfx::OPipelineLayout createInfo)
: Gfx::GraphicsPipeline(std::move(createInfo)), graphics(graphics), state(pipeline), primitiveType(primitive) {}
Gfx::PPipelineLayout layout)
: Gfx::GraphicsPipeline(layout), graphics(graphics), state(pipeline), primitiveType(primitive) {}
GraphicsPipeline::~GraphicsPipeline() {}
ComputePipeline::ComputePipeline(PGraphics graphics, MTL::ComputePipelineState* pipeline,
Gfx::OPipelineLayout createInfo)
: Gfx::ComputePipeline(std::move(createInfo)), graphics(graphics), state(pipeline) {}
Gfx::PPipelineLayout layout)
: Gfx::ComputePipeline(layout), graphics(graphics), state(pipeline) {}
ComputePipeline::~ComputePipeline() {}
src/Engine/Graphics/Metal/Shader.mm
+72 -6
View File
@@ -6,6 +6,7 @@
#include "Graphics/slang-compile.h"
#include "Metal/MTLDevice.hpp"
#include "Metal/MTLLibrary.hpp"
#include "Descriptor.h"
#include <fstream>
#include <iostream>
#include <metal_irconverter/metal_irconverter.h>
@@ -24,16 +25,75 @@ Shader::~Shader() {
}
void Shader::create(const ShaderCreateInfo& createInfo) {
Slang::ComPtr<slang::IBlob> kernelBlob = generateShader(createInfo, SLANG_DXIL);
Map<std::string, uint32> test;
Slang::ComPtr<slang::IBlob> kernelBlob = generateShader(createInfo, SLANG_DXIL, test);
hash = CRC::Calculate(kernelBlob->getBufferPointer(), kernelBlob->getBufferSize(), CRC::CRC_32());
IRCompiler* pCompiler = IRCompilerCreate();
IRCompilerSetMinimumGPUFamily(pCompiler, IRGPUFamilyMetal3);
IRCompilerIgnoreRootSignature(pCompiler, true);
IRCompilerSetEntryPointName(pCompiler, "main");
IRObject* pDXIL = IRObjectCreateFromDXIL((const uint8*)kernelBlob->getBufferPointer(), kernelBlob->getBufferSize(),
IRBytecodeOwnershipNone);
// Compile DXIL to Metal IR:
IRVersionedRootSignatureDescriptor descriptor;
descriptor.version = IRRootSignatureVersion_1_1;
Map<std::string, Gfx::PDescriptorLayout> layouts =createInfo.rootSignature->getLayouts();
descriptor.desc_1_1.NumParameters = layouts.size();
descriptor.desc_1_1.NumStaticSamplers = 0;
descriptor.desc_1_1.pStaticSamplers = nullptr;
Array<IRRootParameter1> parameters;
// each layout has an array for its bindings
Array<Array<IRDescriptorRange1>> ranges;
for(const auto& [name, layout] : layouts)
{
uint32 textureReg = 0; // SRV
uint32 samplerReg = 0; // Sampler
uint32 uavReg = 0; // UAV
uint32 constantReg = 0; // CBV
auto& bindingRanges = ranges.add();
for(auto binding : layout->getBindings())
{
IRDescriptorRangeType type;
uint32 reg = 0;
switch(binding.descriptorType)
{
case Gfx::SE_DESCRIPTOR_TYPE_SAMPLED_IMAGE:
case Gfx::SE_DESCRIPTOR_TYPE_STORAGE_BUFFER:
type = IRDescriptorRangeTypeSRV;
reg = textureReg++;
break;
case Gfx::SE_DESCRIPTOR_TYPE_UNIFORM_BUFFER:
case Gfx::SE_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER:
type = IRDescriptorRangeTypeCBV;
reg = constantReg++;
break;
case Gfx::SE_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER:
case Gfx::SE_DESCRIPTOR_TYPE_STORAGE_IMAGE:
type = IRDescriptorRangeTypeUAV;
reg = uavReg++;
break;
default: throw std::logic_error("Not implemented");
};
bindingRanges.add() = IRDescriptorRange1{
.BaseShaderRegister = reg,
.NumDescriptors = binding.descriptorCount,
.RangeType = type,
.RegisterSpace = createInfo.rootSignature->findParameter(name),
};
}
parameters.add() = IRRootParameter1{
.ParameterType = IRRootParameterTypeDescriptorTable,
.DescriptorTable = IRRootDescriptorTable1{ .NumDescriptorRanges = (uint32)(bindingRanges.size()), bindingRanges.data() },
.ShaderVisibility = IRShaderVisibilityAll,
};
}
IRError* signatureError = nullptr;
descriptor.desc_1_1.NumParameters = parameters.size();
descriptor.desc_1_1.pParameters = parameters.data();
IRRootSignature* rootSignature = IRRootSignatureCreateFromDescriptor(&descriptor, &signatureError);
assert(rootSignature);
IRCompilerSetGlobalRootSignature(pCompiler, rootSignature);
// Compile DXIL to Metal IR:
IRError* pError = nullptr;
IRObject* pOutIR = IRCompilerAllocCompileAndLink(pCompiler, NULL, pDXIL, &pError);
@@ -63,7 +123,13 @@ void Shader::create(const ShaderCreateInfo& createInfo) {
IRMetalLibBinary* pMetallib = IRMetalLibBinaryCreate();
IRObjectGetMetalLibBinary(pOutIR, irStage, pMetallib);
dispatch_data_t data = IRMetalLibGetBytecodeData(pMetallib);
IRShaderReflection* reflection = IRShaderReflectionCreate();
IRObjectGetReflection(pOutIR, irStage, reflection);
std::cout << " NumRes: " << IRShaderReflectionGetResourceCount(reflection) << std::endl;
std::cout << IRShaderReflectionAllocStringAndSerialize(reflection) << std::endl;
IRShaderReflectionDestroy(reflection);
// Store the metallib to custom format or disk, or use to create a MTLLibrary.
NS::Error* error;
library = graphics->getDevice()->newLibrary(data, &error);
@@ -78,4 +144,4 @@ void Shader::create(const ShaderCreateInfo& createInfo) {
IRCompilerDestroy(pCompiler);
}
uint32 Shader::getShaderHash() const { return hash; }
uint32 Shader::getShaderHash() const { return hash; }
src/Engine/Graphics/Metal/Texture.mm
+5 -5
View File
@@ -45,11 +45,11 @@ TextureBase::TextureBase(PGraphics graphics, MTL::TextureType type,
descriptor->release();
}
if(createInfo.sourceData.data != nullptr)
{
MTL::Region region(0, 0, 0, width, height, depth);
texture->replaceRegion(region, 0, createInfo.sourceData.data, createInfo.sourceData.size / (depth / height));
}
// if(createInfo.sourceData.data != nullptr)
// {
// MTL::Region region(0, 0, 0, width, height, depth);
// texture->replaceRegion(region, 0, createInfo.sourceData.data, createInfo.sourceData.size / (depth * height));
// }
}
TextureBase::~TextureBase() {
src/Engine/Graphics/Metal/Window.h
+1
View File
@@ -51,6 +51,7 @@ private:
CAMetalLayer* metalLayer;
CA::MetalDrawable* drawable;
OTexture2D backBuffer;
MTL::CaptureDescriptor* capture = nullptr;
std::function<void(KeyCode, InputAction, KeyModifier)> keyCallback;
std::function<void(double, double)> mouseMoveCallback;
src/Engine/Graphics/Metal/Window.mm
+19
View File
@@ -1,11 +1,15 @@
#include "Window.h"
#include "Command.h"
#include "Foundation/NSAutoreleasePool.hpp"
#include "Foundation/NSString.hpp"
#include "Graphics/Enums.h"
#include "Graphics/Initializer.h"
#include "Graphics/Metal/Enums.h"
#include "Graphics/Texture.h"
#include "Metal/MTLCaptureManager.hpp"
#include "Metal/MTLTexture.hpp"
#include <GLFW/glfw3.h>
#include <fmt/core.h>
#include <iostream>
using namespace Seele;
@@ -13,6 +17,8 @@ using namespace Seele::Metal;
double currentFrameDelta = 0;
double Gfx::getCurrentFrameDelta() { return currentFrameDelta; }
uint32 currentFrameIndex = 0;
uint32 Gfx::getCurrentFrameIndex() { return currentFrameIndex; }
void glfwKeyCallback(GLFWwindow* handle, int key, int, int action, int modifier) {
if (key == -1) {
@@ -88,6 +94,17 @@ Window::~Window() { glfwDestroyWindow(static_cast<GLFWwindow*>(windowHandle)); }
void Window::pollInput() { glfwPollEvents(); }
void Window::beginFrame() {
auto captureManager = MTL::CaptureManager::sharedCaptureManager();
capture = MTL::CaptureDescriptor::alloc()->init();
capture->setCaptureObject((__bridge id<MTLDevice>)graphics->getDevice());
capture->setDestination(MTL::CaptureDestinationDeveloperTools);
capture->setOutputURL(NS::URL::fileURLWithPath(NS::String::string(fmt::format("frame{0}.trace", Gfx::getCurrentFrameIndex()).c_str(), NS::ASCIIStringEncoding)));
NS::Error* error;
captureManager->startCapture(capture, &error);
if(error)
{
std::cout << error->localizedDescription()->cString(NS::ASCIIStringEncoding) << std::endl;
}
drawable = (__bridge CA::MetalDrawable*)[metalLayer nextDrawable];
createBackBuffer();
}
@@ -95,6 +112,8 @@ void Window::beginFrame() {
void Window::endFrame() {
graphics->getQueue()->getCommands()->present(drawable);
graphics->getQueue()->submitCommands();
MTL::CaptureManager::sharedCaptureManager()->stopCapture();
currentFrameIndex++;
}
void Window::onWindowCloseEvent() {}