Implemented compute shader test Metal reference

2 files changed, 94 insertions, 4 deletions

diff --git a/Tests/DiligentCoreAPITest/include/InlineShaders/ComputeShaderTestMSL.h b/Tests/DiligentCoreAPITest/include/InlineShaders/ComputeShaderTestMSL.h
new file mode 100644
index 00000000..f2751400
--- /dev/null
+++ b/Tests/DiligentCoreAPITest/include/InlineShaders/ComputeShaderTestMSL.h
@@ -0,0 +1,61 @@
+/*
+ *  Copyright 2019-2020 Diligent Graphics LLC
+ *  Copyright 2015-2019 Egor Yusov
+ *  
+ *  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.
+ *
+ *  In no event and under no legal theory, whether in tort (including negligence), 
+ *  contract, or otherwise, unless required by applicable law (such as deliberate 
+ *  and grossly negligent acts) or agreed to in writing, shall any Contributor be
+ *  liable for any damages, including any direct, indirect, special, incidental, 
+ *  or consequential damages of any character arising as a result of this License or 
+ *  out of the use or inability to use the software (including but not limited to damages 
+ *  for loss of goodwill, work stoppage, computer failure or malfunction, or any and 
+ *  all other commercial damages or losses), even if such Contributor has been advised 
+ *  of the possibility of such damages.
+ */
+
+#include <string>
+
+namespace
+{
+
+namespace MSL
+{
+
+// clang-format off
+const std::string FillTextureCS{
+R"(
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+kernel void CSMain(texture2d<float, access::write> g_tex2DUAV            [[texture(0)]],
+                   uint3                           gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    if (gl_GlobalInvocationID.x < g_tex2DUAV.get_width() &&
+        gl_GlobalInvocationID.y < g_tex2DUAV.get_height())
+    {
+        g_tex2DUAV.write(float4(float2(gl_GlobalInvocationID.xy % uint2(256u)) / 256.0, 0.0, 1.0),
+                         uint2(gl_GlobalInvocationID.xy));
+
+    }
+}
+)"
+};
+// clang-format on
+
+} // namespace MSL
+
+} // namespace
diff --git a/Tests/DiligentCoreAPITest/src/Metal/ComputeShaderReferenceMtl.mm b/Tests/DiligentCoreAPITest/src/Metal/ComputeShaderReferenceMtl.mm
index 6341ecdb..489f36f2 100644
--- a/Tests/DiligentCoreAPITest/src/Metal/ComputeShaderReferenceMtl.mm
+++ b/Tests/DiligentCoreAPITest/src/Metal/ComputeShaderReferenceMtl.mm
@@ -29,8 +29,9 @@
 #include "Metal/TestingSwapChainMtl.hpp"
 
 #include "DeviceContextMtl.h"
+#include "TextureViewMtl.h"
 
-//#include "InlineShaders/ComputeShaderTestGLSL.h"
+#include "InlineShaders/ComputeShaderTestMSL.h"
 
 namespace Diligent
 {
@@ -40,9 +41,37 @@ namespace Testing
 
 void ComputeShaderReferenceMtl(ISwapChain* pSwapChain)
 {
-    //auto* pEnv     = TestingEnvironmentMtl::GetInstance();
-    //auto  vkDevice = pEnv->GetVkDevice();
-    //auto* pContext = pEnv->GetDeviceContext();
+    auto* const pEnv      = TestingEnvironmentMtl::GetInstance();
+    auto const  mtlDevice = pEnv->GetMtlDevice();
+
+    auto* progSrc = [NSString stringWithUTF8String:MSL::FillTextureCS.c_str()];
+    NSError *errors = nil;
+    id <MTLLibrary> library = [mtlDevice newLibraryWithSource:progSrc
+                               options:nil
+                               error:&errors];
+    ASSERT_TRUE(library != nil);
+    id <MTLFunction> computeFunc = [library newFunctionWithName:@"CSMain"];
+    ASSERT_TRUE(computeFunc != nil);
+    auto* computePipeline = [mtlDevice newComputePipelineStateWithFunction:computeFunc error:&errors];
+    ASSERT_TRUE(computePipeline != nil);
+
+    auto* pTestingSwapChainMtl = ValidatedCast<TestingSwapChainMtl>(pSwapChain);
+    auto* pUAV = pTestingSwapChainMtl->GetCurrentBackBufferUAV();
+    auto* mtlTexture = ValidatedCast<ITextureViewMtl>(pUAV)->GetMtlTexture();
+    const auto& SCDesc = pTestingSwapChainMtl->GetDesc();
+
+    auto* mtlCommandQueue = pEnv->GetMtlCommandQueue();
+    id <MTLCommandBuffer> mtlCommandBuffer = [mtlCommandQueue commandBuffer];
+    auto* cmdEncoder = [mtlCommandBuffer computeCommandEncoder];
+    ASSERT_TRUE(cmdEncoder != nil);
+
+    [cmdEncoder setComputePipelineState:computePipeline];
+    [cmdEncoder setTexture:mtlTexture atIndex:0];
+    [cmdEncoder dispatchThreadgroups:MTLSizeMake((SCDesc.Width + 15) / 16, (SCDesc.Height + 15) / 16, 1)
+               threadsPerThreadgroup:MTLSizeMake(16, 16, 1)];
+
+    [cmdEncoder endEncoding];
+    [mtlCommandBuffer commit];
 }
 
 } // namespace Testing