Skip to content

Commit 6c31866

Browse files
azhirnovazhirnov
committed
Metal: added tests for PRS with TLAS, added test for threadgroup memory
1 parent b357896 commit 6c31866

File tree

6 files changed

+912
-1
lines changed

6 files changed

+912
-1
lines changed

‎Tests/DiligentCoreAPITest/include/InlineShaders/RayTracingTestMSL.h‎

Lines changed: 75 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -37,7 +37,6 @@ const std::string RayTracingTest8_CS{R"msl(
3737
#include <metal_stdlib>
3838
#include <simd/simd.h>
3939
#include <metal_raytracing>
40-
#include <metal_visible_function_table>
4140
4241
using namespace metal;
4342
using namespace raytracing;
@@ -88,6 +87,81 @@ void CSMain(uint2 DTid [[thread_position_in_grid]],
8887
}
8988
)msl"};
9089

90+
91+
const std::string RayTracingTest9_CS{R"msl(
92+
#include <metal_stdlib>
93+
#include <simd/simd.h>
94+
#include <metal_raytracing>
95+
96+
using namespace metal;
97+
using namespace raytracing;
98+
99+
#ifndef CONST_BUFFER_1
100+
# define CONST_BUFFER_1 0
101+
#endif
102+
#ifndef CONST_BUFFER_2
103+
# define CONST_BUFFER_2 1
104+
#endif
105+
#ifndef CONST_BUFFER_3
106+
# define CONST_BUFFER_3 2
107+
#endif
108+
#ifndef TLAS_1
109+
# define TLAS_1 3
110+
#endif
111+
#ifndef TLAS_2
112+
# define TLAS_2 4
113+
#endif
114+
115+
kernel
116+
void CSMain(uint2 DTid [[thread_position_in_grid]],
117+
constant float4& g_Constant1 [[buffer(CONST_BUFFER_1)]],
118+
constant float4& g_Constant2 [[buffer(CONST_BUFFER_2)]],
119+
const device float4& g_Constant3 [[buffer(CONST_BUFFER_3)]],
120+
texture2d<float, access::write> g_ColorBuffer [[texture(0)]],
121+
instance_acceleration_structure g_TLAS1 [[buffer(TLAS_1)]],
122+
instance_acceleration_structure g_TLAS2 [[buffer(TLAS_2)]] )
123+
{
124+
if (DTid.x >= g_ColorBuffer.get_width() || DTid.y >= g_ColorBuffer.get_height())
125+
return;
126+
127+
ray Ray;
128+
Ray.origin = float3((float(DTid.x) + 0.5) / float(g_ColorBuffer.get_width()), 1.0 - (float(DTid.y) + 0.5) / float(g_ColorBuffer.get_height()), -1.0);
129+
Ray.direction = float3(0.0, 0.0, 1.0);
130+
Ray.min_distance = 0.01;
131+
Ray.max_distance = 10.0;
132+
133+
intersector<triangle_data, instancing> Intersector;
134+
Intersector.assume_geometry_type( geometry_type::triangle );
135+
Intersector.force_opacity( forced_opacity::opaque );
136+
Intersector.accept_any_intersection( false );
137+
138+
intersection_result<triangle_data, instancing> Intersection = Intersector.intersect(Ray, g_TLAS1, 0x0F);
139+
140+
float4 color;
141+
if (Intersection.type != intersection_type::none)
142+
{
143+
float3 barycentrics;
144+
barycentrics.yz = Intersection.triangle_barycentric_coord;
145+
barycentrics.x = 1.0 - barycentrics.x - barycentrics.y;
146+
147+
color = float4(barycentrics, 1.0) * g_Constant1;
148+
}
149+
else
150+
{
151+
Ray.origin.z = 1.0;
152+
Ray.direction = float3(0.0, 0.0, 1.0);
153+
intersection_result<triangle_data, instancing> Intersection2 = Intersector.intersect(Ray, g_TLAS2, 0xF0);
154+
155+
if (Intersection2.type != intersection_type::none)
156+
color = g_Constant2;
157+
else
158+
color = g_Constant3;
159+
}
160+
161+
g_ColorBuffer.write(color, DTid.xy);
162+
}
163+
)msl"};
164+
91165
// clang-format on
92166

93167
} // namespace MSL

‎Tests/DiligentCoreAPITest/include/InlineShaders/ThreadgroupMemoryMtl.h‎

Lines changed: 87 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,87 @@
1+
/*
2+
* Copyright 2019-2021 Diligent Graphics LLC
3+
*
4+
* Licensed under the Apache License, Version 2.0 (the "License");
5+
* you may not use this file except in compliance with the License.
6+
* You may obtain a copy of the License at
7+
*
8+
* http://www.apache.org/licenses/LICENSE-2.0
9+
*
10+
* Unless required by applicable law or agreed to in writing, software
11+
* distributed under the License is distributed on an "AS IS" BASIS,
12+
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+
* See the License for the specific language governing permissions and
14+
* limitations under the License.
15+
*
16+
* In no event and under no legal theory, whether in tort (including negligence),
17+
* contract, or otherwise, unless required by applicable law (such as deliberate
18+
* and grossly negligent acts) or agreed to in writing, shall any Contributor be
19+
* liable for any damages, including any direct, indirect, special, incidental,
20+
* or consequential damages of any character arising as a result of this License or
21+
* out of the use or inability to use the software (including but not limited to damages
22+
* for loss of goodwill, work stoppage, computer failure or malfunction, or any and
23+
* all other commercial damages or losses), even if such Contributor has been advised
24+
* of the possibility of such damages.
25+
*/
26+
27+
#include <string>
28+
29+
namespace
30+
{
31+
32+
namespace MSL
33+
{
34+
// clang-format off
35+
36+
const std::string SetComputeThreadgroupMemoryLength_CS{R"msl(
37+
#include <metal_stdlib>
38+
#include <simd/simd.h>
39+
using namespace metal;
40+
41+
kernel void CSmain(threadgroup float4* groupPixels [[threadgroup(0)]],
42+
texture2d<float, access::write> g_OutImage [[texture(0)]],
43+
uint2 globalId [[thread_position_in_grid]],
44+
uint2 globalSize [[threads_per_grid]],
45+
uint2 localId [[thread_position_in_threadgroup]],
46+
uint2 localSize [[dispatch_threads_per_threadgroup]] )
47+
{
48+
if (globalId.x >= g_OutImage.get_width() || globalId.y >= g_OutImage.get_height())
49+
return;
50+
51+
// pass 1
52+
{
53+
float2 uv = float2(globalId) / float2(globalSize) * 10.0;
54+
float4 col = float4(1.0);
55+
float4 t = float4(1.2f, 0.25f, 1.1f, 0.14f);
56+
57+
col.r = sin(uv.x + t.x) * cos(uv.y + t.y);
58+
col.g = fract(uv.x + t.z) * fract(uv.y + t.w);
59+
60+
uint idx = localId.x + localId.y * localSize.x;
61+
groupPixels[idx] = col;
62+
}
63+
64+
threadgroup_barrier(mem_flags::mem_threadgroup);
65+
66+
// pass 2
67+
{
68+
float4 col = float4(0.0);
69+
col += groupPixels[ ((localId.x + 0) % localSize.x) + ((localId.y + 0) % localSize.y) * localSize.x ];
70+
col += groupPixels[ ((localId.x - 1) % localSize.x) + ((localId.y + 1) % localSize.y) * localSize.x ];
71+
col += groupPixels[ ((localId.x - 1) % localSize.x) + ((localId.y - 1) % localSize.y) * localSize.x ];
72+
col += groupPixels[ ((localId.x + 1) % localSize.x) + ((localId.y - 1) % localSize.y) * localSize.x ];
73+
col += groupPixels[ ((localId.x + 1) % localSize.x) + ((localId.y + 1) % localSize.y) * localSize.x ];
74+
col /= 5.0;
75+
76+
g_OutImage.write(col, globalId);
77+
}
78+
}
79+
)msl"};
80+
81+
82+
83+
// clang-format on
84+
85+
} // namespace MSL
86+
87+
} // namespace

‎Tests/DiligentCoreAPITest/src/Metal/RayTracingReferenceMtl.mm‎

Lines changed: 151 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -169,6 +169,157 @@ void InlineRayTracingInComputePplnReferenceMtl(ISwapChain* pSwapChain)
169169
}
170170
}
171171

172+
void RayTracingPRSReferenceMtl(ISwapChain* pSwapChain)
173+
{
174+
auto* const pEnv = TestingEnvironmentMtl::GetInstance();
175+
auto const mtlDevice = pEnv->GetMtlDevice();
176+
177+
if (@available(macos 11.0, ios 14.0, *))
178+
{
179+
@autoreleasepool
180+
{
181+
auto* progSrc = [NSString stringWithUTF8String: MSL::RayTracingTest9_CS.c_str()];
182+
NSError *errors = nil; // Autoreleased
183+
id <MTLLibrary> library = [mtlDevice newLibraryWithSource: progSrc
184+
options: nil
185+
error: &errors];
186+
ASSERT_TRUE(library != nil);
187+
[library autorelease];
188+
189+
id <MTLFunction> computeFunc = [library newFunctionWithName: @"CSMain"];
190+
ASSERT_TRUE(computeFunc != nil);
191+
[computeFunc autorelease];
192+
193+
auto* computePipeline = [mtlDevice newComputePipelineStateWithFunction: computeFunc error: &errors];
194+
ASSERT_TRUE(computePipeline != nil);
195+
[computePipeline autorelease];
196+
197+
auto* pTestingSwapChainMtl = ValidatedCast<TestingSwapChainMtl>(pSwapChain);
198+
auto* pUAV = pTestingSwapChainMtl->GetCurrentBackBufferUAV();
199+
auto* mtlTexture = ValidatedCast<ITextureViewMtl>(pUAV)->GetMtlTexture();
200+
const auto& SCDesc = pTestingSwapChainMtl->GetDesc();
201+
202+
const auto& Vertices = TestingConstants::TriangleClosestHit::Vertices;
203+
auto* mtlVertexBuf = [mtlDevice newBufferWithLength: sizeof(Vertices)
204+
options: (MTLResourceCPUCacheModeDefaultCache | MTLResourceStorageModeShared | MTLResourceHazardTrackingModeDefault)];
205+
ASSERT_TRUE(mtlVertexBuf != nil);
206+
[mtlVertexBuf autorelease];
207+
memcpy([mtlVertexBuf contents], Vertices, sizeof(Vertices));
208+
209+
auto* mtlBLASDesc = [MTLPrimitiveAccelerationStructureDescriptor descriptor];
210+
auto* mtlTriangleDesc = [MTLAccelerationStructureTriangleGeometryDescriptor descriptor];
211+
auto* mtlTriangleArr = [[[NSMutableArray<MTLAccelerationStructureTriangleGeometryDescriptor*> alloc] initWithCapacity: 1] autorelease];
212+
213+
mtlTriangleDesc.intersectionFunctionTableOffset = 0;
214+
mtlTriangleDesc.allowDuplicateIntersectionFunctionInvocation = false;
215+
mtlTriangleDesc.opaque = true;
216+
mtlTriangleDesc.triangleCount = _countof(Vertices) / 3;
217+
mtlTriangleDesc.vertexBuffer = mtlVertexBuf;
218+
mtlTriangleDesc.vertexBufferOffset = 0;
219+
mtlTriangleDesc.vertexStride = sizeof(Vertices[0]);
220+
mtlTriangleDesc.indexType = MTLIndexTypeUInt32;
221+
[mtlTriangleArr addObject: mtlTriangleDesc];
222+
223+
[mtlBLASDesc setUsage: MTLAccelerationStructureUsageNone];
224+
[mtlBLASDesc setGeometryDescriptors: mtlTriangleArr];
225+
226+
const auto mtlBLASSizes = [mtlDevice accelerationStructureSizesWithDescriptor: mtlBLASDesc];
227+
auto* mtlBLAS = [mtlDevice newAccelerationStructureWithSize: mtlBLASSizes.accelerationStructureSize];
228+
ASSERT_TRUE(mtlBLAS != nil);
229+
[mtlBLAS autorelease];
230+
231+
auto* mtlInstanceBuf = [mtlDevice newBufferWithLength: sizeof(MTLAccelerationStructureInstanceDescriptor)
232+
options: (MTLResourceCPUCacheModeDefaultCache | MTLResourceStorageModeShared | MTLResourceHazardTrackingModeDefault)];
233+
ASSERT_TRUE(mtlInstanceBuf != nil);
234+
[mtlInstanceBuf autorelease];
235+
auto* Instances = static_cast<MTLAccelerationStructureInstanceDescriptor*>([mtlInstanceBuf contents]);
236+
237+
Instances->transformationMatrix.columns[0] = MTLPackedFloat3{1.0f, 0.0f, 0.0f};
238+
Instances->transformationMatrix.columns[1] = MTLPackedFloat3{0.0f, 1.0f, 0.0f};
239+
Instances->transformationMatrix.columns[2] = MTLPackedFloat3{0.0f, 0.0f, 1.0f};
240+
Instances->transformationMatrix.columns[3] = MTLPackedFloat3{0.0f, 0.0f, 0.0f};
241+
Instances->accelerationStructureIndex = 0;
242+
Instances->intersectionFunctionTableOffset = 0;
243+
Instances->mask = ~0u;
244+
Instances->options = MTLAccelerationStructureInstanceOptionOpaque;
245+
246+
auto* mtlTLASDesc = [MTLInstanceAccelerationStructureDescriptor descriptor];
247+
auto* mtlAccelStrArr = [[[NSMutableArray<id<MTLAccelerationStructure>> alloc] initWithCapacity: 1] autorelease];
248+
[mtlAccelStrArr addObject: mtlBLAS];
249+
[mtlTLASDesc setUsage: MTLAccelerationStructureUsageNone];
250+
[mtlTLASDesc setInstanceCount: 1];
251+
[mtlTLASDesc setInstanceDescriptorBuffer: mtlInstanceBuf];
252+
[mtlTLASDesc setInstancedAccelerationStructures: mtlAccelStrArr];
253+
254+
const auto mtlTLASSizes = [mtlDevice accelerationStructureSizesWithDescriptor: mtlTLASDesc];
255+
auto* mtlTLAS = [mtlDevice newAccelerationStructureWithSize: mtlTLASSizes.accelerationStructureSize];
256+
ASSERT_TRUE(mtlTLAS != nil);
257+
[mtlTLAS autorelease];
258+
259+
const auto ScratchBuffSize = std::max(mtlBLASSizes.buildScratchBufferSize, mtlTLASSizes.buildScratchBufferSize);
260+
auto* ScratchBuf = [mtlDevice newBufferWithLength: ScratchBuffSize
261+
options: (MTLResourceCPUCacheModeDefaultCache | MTLResourceStorageModePrivate | MTLResourceHazardTrackingModeDefault)];
262+
ASSERT_TRUE(ScratchBuf != nil);
263+
[ScratchBuf autorelease];
264+
265+
auto* ConstBuf = [mtlDevice newBufferWithLength: 256 * 3
266+
options: (MTLResourceCPUCacheModeDefaultCache | MTLResourceStorageModeShared | MTLResourceHazardTrackingModeDefault)];
267+
ASSERT_TRUE(ConstBuf != nil);
268+
[ConstBuf autorelease];
269+
270+
uint8_t* pMapped = (uint8_t*)ConstBuf.contents;
271+
const float Const1[] = {0.5f, 0.9f, 0.75f, 1.0f};
272+
const float Const2[] = {0.2f, 0.0f, 1.0f, 0.5f};
273+
const float Const3[] = {0.9f, 0.1f, 0.2f, 1.0f};
274+
memcpy(pMapped + 0, Const1, sizeof(Const1));
275+
memcpy(pMapped + 256, Const2, sizeof(Const2));
276+
memcpy(pMapped + 512, Const3, sizeof(Const3));
277+
278+
auto* mtlCommandQueue = pEnv->GetMtlCommandQueue();
279+
auto* mtlCommandBuffer = [mtlCommandQueue commandBuffer]; // Autoreleased
280+
auto* asEncoder = [mtlCommandBuffer accelerationStructureCommandEncoder]; // Autoreleased
281+
ASSERT_TRUE(asEncoder != nil);
282+
283+
[asEncoder buildAccelerationStructure: mtlBLAS
284+
descriptor: mtlBLASDesc
285+
scratchBuffer: ScratchBuf
286+
scratchBufferOffset: 0];
287+
[asEncoder buildAccelerationStructure: mtlTLAS
288+
descriptor: mtlTLASDesc
289+
scratchBuffer: ScratchBuf
290+
scratchBufferOffset: 0];
291+
[asEncoder endEncoding];
292+
293+
auto* cmdEncoder = [mtlCommandBuffer computeCommandEncoder]; // Autoreleased
294+
ASSERT_TRUE(cmdEncoder != nil);
295+
296+
[cmdEncoder setComputePipelineState: computePipeline];
297+
[cmdEncoder setTexture: mtlTexture
298+
atIndex: 0];
299+
[cmdEncoder setBuffer: ConstBuf
300+
offset: 0
301+
atIndex: 0];
302+
[cmdEncoder setBuffer: ConstBuf
303+
offset: 256
304+
atIndex: 1];
305+
[cmdEncoder setBuffer: ConstBuf
306+
offset: 512
307+
atIndex: 2];
308+
[cmdEncoder setAccelerationStructure: mtlTLAS
309+
atBufferIndex: 3];
310+
[cmdEncoder setAccelerationStructure: mtlTLAS
311+
atBufferIndex: 4];
312+
[cmdEncoder dispatchThreadgroups: MTLSizeMake((SCDesc.Width + 15) / 16, (SCDesc.Height + 15) / 16, 1)
313+
threadsPerThreadgroup: MTLSizeMake(16, 16, 1)];
314+
315+
[cmdEncoder endEncoding];
316+
[mtlCommandBuffer commit];
317+
[mtlCommandBuffer waitUntilCompleted];
318+
319+
} // @autoreleasepool
320+
} // @available
321+
}
322+
172323
} // namespace Testing
173324

174325
} // namespace Diligent

0 commit comments

Comments
 (0)