[Metal] Ported BarrierReadAfterWrite test to Metal backend.

- Provide Metal shader for unit test. - Fixed memory issue in MTBuiltinPSOFactory. - Fixed MTLTexture usage flags for textures with Storage binding flags but without Sampled binding flags.
LukasBanana · Jan 25, 2025 · c68053f · c68053f
1 parent e8d86d7
commit c68053f
Show file tree

Hide file tree

Showing 13 changed files with 95 additions and 32 deletions.
diff --git a/sources/Core/Exception.cpp b/sources/Core/Exception.cpp
@@ -48,10 +48,10 @@ LLGL_EXPORT void Trap(Exception exception, const char* origin, const char* forma
     /* Throw exception with report and optional origin */
     switch (exception)
     {
-        case RuntimeError:      throw std::runtime_error(report);
-        case OutOfRange:        throw std::out_of_range(report);
-        case BadCast:           throw std::bad_cast();
-        case InvalidArgument:   throw std::invalid_argument(report);
+        case Exception::RuntimeError:       throw std::runtime_error(report);
+        case Exception::OutOfRange:         throw std::out_of_range(report);
+        case Exception::BadCast:            throw std::bad_cast();
+        case Exception::InvalidArgument:    throw std::invalid_argument(report);
     }
 
     #endif // /LLGL_EXCEPTIONS_SUPPORTED

diff --git a/sources/Renderer/Metal/Command/MTDirectCommandBuffer.mm b/sources/Renderer/Metal/Command/MTDirectCommandBuffer.mm
@@ -486,7 +486,7 @@
     std::uint32_t       /*numTextures*/,
     Texture* const *    /*textures*/)
 {
-    //TODO
+    // dummy
 }
 
 /* ----- Render Passes ----- */

diff --git a/sources/Renderer/Metal/Command/MTMultiSubmitCommandBuffer.mm b/sources/Renderer/Metal/Command/MTMultiSubmitCommandBuffer.mm
@@ -451,7 +451,7 @@
     std::uint32_t       /*numTextures*/,
     Texture* const *    /*textures*/)
 {
-    //TODO
+    // dummy
 }
 
 /* ----- Render Passes ----- */

diff --git a/sources/Renderer/Metal/RenderState/MTBuiltinPSOFactory.h b/sources/Renderer/Metal/RenderState/MTBuiltinPSOFactory.h
@@ -11,6 +11,9 @@
 
 #import <Metal/Metal.h>
 
+#include "../Shader/MTShader.h"
+#include <memory>
+
 
 namespace LLGL
 {
@@ -54,9 +57,10 @@ class MTBuiltinPSOFactory
 
     private:
 
-        static const std::size_t g_numComputePSOs = static_cast<std::size_t>(MTBuiltinComputePSO::Num);
+        static constexpr std::size_t k_numComputePSOs = static_cast<std::size_t>(MTBuiltinComputePSO::Num);
 
-        id<MTLComputePipelineState> builtinComputePSOs_[g_numComputePSOs];
+        id<MTLComputePipelineState> builtinComputePSOs_[k_numComputePSOs];
+        std::unique_ptr<MTShader>   builtinComputeShaders_[k_numComputePSOs];
 
 };
 

diff --git a/sources/Renderer/Metal/RenderState/MTBuiltinPSOFactory.mm b/sources/Renderer/Metal/RenderState/MTBuiltinPSOFactory.mm
@@ -6,9 +6,9 @@
  */
 
 #include "MTBuiltinPSOFactory.h"
-#include "../Shader/MTShader.h"
 #include "../Shader/Builtin/MTBuiltin.h"
 #include "../MTCore.h"
+#include "../../../Core/CoreUtils.h"
 #include "../../../Core/Exception.h"
 #include <LLGL/Report.h>
 
@@ -31,7 +31,7 @@
 id<MTLComputePipelineState> MTBuiltinPSOFactory::GetComputePSO(const MTBuiltinComputePSO builtin) const
 {
     const auto idx = static_cast<std::size_t>(builtin);
-    if (idx < g_numComputePSOs)
+    if (idx < k_numComputePSOs)
         return builtinComputePSOs_[idx];
     else
         return nil;
@@ -58,10 +58,10 @@
         shaderDesc.entryPoint   = "CS";
         shaderDesc.profile      = "1.1";
     }
-    MTShader cs{ device, shaderDesc };
+    std::unique_ptr<MTShader> cs = MakeUnique<MTShader>(device, shaderDesc);
 
     /* We cannot recover from a faulty built-in shader */
-    if (const Report* report = cs.GetReport())
+    if (const Report* report = cs->GetReport())
     {
         if (report->HasErrors())
             LLGL_TRAP("%s", report->GetText());
@@ -70,9 +70,12 @@
     /* Create native compute pipeline state */
     const std::size_t idx = static_cast<std::size_t>(builtin);
     NSError* error = nullptr;
-    builtinComputePSOs_[idx] = [device newComputePipelineStateWithFunction:cs.GetNative() error:&error];
+    builtinComputePSOs_[idx] = [device newComputePipelineStateWithFunction:cs->GetNative() error:&error];
     if (!builtinComputePSOs_[idx])
         MTThrowIfCreateFailed(error, "MTLComputePipelineState");
+
+    /* Keep shader in memory as we still need the MTLLibrary and MTLFunction objects */
+    builtinComputeShaders_[idx] = std::move(cs);
 }
 
 

diff --git a/sources/Renderer/Metal/Shader/MTShader.mm b/sources/Renderer/Metal/Shader/MTShader.mm
@@ -375,6 +375,9 @@ static void ReflectShaderArgument(MTLArgument* arg, ShaderReflection& reflection
 
 bool MTShader::ReflectComputePipeline(ShaderReflection& reflection) const
 {
+    if (GetNative() == nil)
+        return false;
+
     /* Create temporary compute pipeline state to retrieve shader reflection */
     MTLAutoreleasedComputePipelineReflection psoReflect = nullptr;
     MTLPipelineOption opt = (MTLPipelineOptionArgumentInfo | MTLPipelineOptionBufferTypeInfo);

diff --git a/sources/Renderer/Metal/Texture/MTTexture.mm b/sources/Renderer/Metal/Texture/MTTexture.mm
@@ -37,7 +37,7 @@ static MTLTextureUsage GetTextureUsage(const TextureDescriptor& desc)
     if ((desc.bindFlags & BindFlags::Sampled) != 0)
         usage |= MTLTextureUsageShaderRead | MTLTextureUsagePixelFormatView;
     if ((desc.bindFlags & BindFlags::Storage) != 0)
-        usage |= MTLTextureUsageShaderWrite;
+        usage |= MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite;
     if ((desc.bindFlags & (BindFlags::ColorAttachment | BindFlags::DepthStencilAttachment)) != 0)
         usage |= MTLTextureUsageRenderTarget;
 

diff --git a/tests/Testbed/Shaders/ReadAfterWrite/ReadAfterWrite.450core.comp b/tests/Testbed/Shaders/ReadAfterWrite/ReadAfterWrite.450core.comp
@@ -14,7 +14,7 @@ struct Entry
     uint b;
 };
 
-layout(binding = 1, std430) buffer type_RWStructuredBuffer_Entry
+layout(binding = 2, std430) buffer type_RWStructuredBuffer_Entry
 {
     Entry _m0[];
 } buf2;
@@ -30,9 +30,9 @@ uniform uint readPos;
 uniform uint writePos;
 #endif
 
-layout(binding = 0, r32ui) uniform uimageBuffer buf1;
-layout(binding = 2, r32ui) uniform uimage1D tex1;
-layout(binding = 3, rg32ui) uniform uimage2D tex2;
+layout(binding = 1, r32ui) uniform uimageBuffer buf1;
+layout(binding = 3, r32ui) uniform uimage1D tex1;
+layout(binding = 4, rg32ui) uniform uimage2D tex2;
 
 void main()
 {

diff --git a/tests/Testbed/Shaders/ReadAfterWrite/ReadAfterWrite.hlsl b/tests/Testbed/Shaders/ReadAfterWrite/ReadAfterWrite.hlsl
@@ -5,17 +5,17 @@
  * Licensed under the terms of the BSD 3-Clause license (see LICENSE.txt).
  */
 
-RWBuffer<uint> buf1 : register(u0);
+RWBuffer<uint> buf1 : register(u1);
 
 struct Entry
 {
     uint a, b;
 };
 
-RWStructuredBuffer<Entry> buf2 : register(u1);
+RWStructuredBuffer<Entry> buf2 : register(u2);
 
-RWTexture1D<uint> tex1 : register(u2);
-RWTexture2D<uint2> tex2 : register(u3);
+RWTexture1D<uint> tex1 : register(u3);
+RWTexture2D<uint2> tex2 : register(u4);
 
 uint readPos;
 uint writePos;

diff --git a/tests/Testbed/Shaders/ReadAfterWrite/ReadAfterWrite.metal b/tests/Testbed/Shaders/ReadAfterWrite/ReadAfterWrite.metal
@@ -0,0 +1,42 @@
+// Cross compiled with SPIRV-Cross from ReadAfterWrite.hlsl
+// DO NOT EDIT
+
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct Entry
+{
+    uint a;
+    uint b;
+};
+
+struct type_RWStructuredBuffer_Entry
+{
+    Entry _m0[1];
+};
+
+struct type_Globals
+{
+    uint readPos;
+    uint writePos;
+};
+
+kernel void CSMain(
+    constant type_Globals& _Globals [[buffer(0)]],
+    device uint* buf1 [[buffer(1)]],
+    device type_RWStructuredBuffer_Entry& buf2 [[buffer(2)]],
+    texture1d<uint, access::read_write> tex1 [[texture(3)]],
+    texture2d<uint, access::read_write> tex2 [[texture(4)]],
+    uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    uint _38 = _Globals.readPos + gl_GlobalInvocationID.x;
+    uint _44 = _Globals.writePos + gl_GlobalInvocationID.x;
+    buf1[_44] = buf1[_38];
+    buf2._m0[_44] = buf2._m0[_38];
+    tex1.write(uint4(tex1.read(uint(_38)).x), uint(_44));
+    tex2.write(tex2.read(uint2(uint2(_38, 0u))).xy.xyyy, uint2(uint2(_44, 0u)));
+}
+
diff --git a/tests/Testbed/Shaders/ReadAfterWrite/default.metallib b/tests/Testbed/Shaders/ReadAfterWrite/default.metallib
diff --git a/tests/Testbed/TestbedContext.cpp b/tests/Testbed/TestbedContext.cpp
@@ -995,6 +995,10 @@ bool TestbedContext::LoadShaders()
 //      shaders[CSResourceBinding]  = LoadShaderFromFile("ResourceBinding.metal",      ShaderType::Compute,  "CSMain",  "1.1");
         shaders[VSClear]            = LoadShaderFromFile("ClearScreen.metal",          ShaderType::Vertex,   "VSMain",  "1.1", nullptr, VertFmtEmpty);
         shaders[PSClear]            = LoadShaderFromFile("ClearScreen.metal",          ShaderType::Fragment, "PSMain",  "1.1");
+        if (IsShadingLanguageSupported(ShadingLanguage::Metal_1_2))
+        {
+            shaders[CSReadAfterWrite]   = LoadShaderFromFile("ReadAfterWrite.metal",   ShaderType::Compute,  "CSMain",  "1.2"); // access::read_write requires Metal 1.2
+        }
     }
     else if (IsShadingLanguageSupported(ShadingLanguage::SPIRV))
     {

diff --git a/tests/Testbed/UnitTests/TestBarrierReadAfterWrite.cpp b/tests/Testbed/UnitTests/TestBarrierReadAfterWrite.cpp
@@ -16,18 +16,25 @@ The test must validate that the correct memory barriers are inserted between the
 */
 DEF_TEST( BarrierReadAfterWrite )
 {
-    //TODO: not supported for Vulkan and Metal yet
-    if (renderer->GetRendererID() != RendererID::OpenGL &&
-        renderer->GetRendererID() != RendererID::Direct3D11 &&
-        renderer->GetRendererID() != RendererID::Direct3D12)
+    //TODO: not supported for Vulkan yet
+    if (renderer->GetRendererID() == RendererID::Vulkan)
     {
         return TestResult::Skipped;
     }
 
     if (shaders[CSReadAfterWrite] == nullptr)
     {
-        Log::Errorf("Missing shaders for backend\n");
-        return TestResult::FailedErrors;
+        if (renderer->GetRendererID() == RendererID::Metal)
+        {
+            if (opt.verbose)
+                Log::Printf("Read/write texture access not supported for this Metal device\n");
+            return TestResult::Skipped;
+        }
+        else
+        {
+            Log::Errorf("Missing shaders for backend\n");
+            return TestResult::FailedErrors;
+        }
     }
 
     constexpr unsigned      numFrames       = 2;
@@ -87,10 +94,10 @@ DEF_TEST( BarrierReadAfterWrite )
 
     // Create compute PSO
     PipelineLayoutDescriptor psoLayoutDesc = Parse(
-        "rwbuffer(buf1@0):comp,"
-        "rwbuffer(buf2@1):comp,"
-        "rwtexture(tex1@2):comp,"
-        "rwtexture(tex2@3):comp,"
+        "rwbuffer(buf1@1):comp,"
+        "rwbuffer(buf2@2):comp,"
+        "rwtexture(tex1@3):comp,"
+        "rwtexture(tex2@4):comp,"
         "uint(readPos),"
         "uint(writePos),"
     );