xref: /aosp_15_r20/external/mesa3d/src/intel/vulkan/grl/gpu/traversal_shader.cl (revision 6104692788411f58d303aa86923a9ff6ecaded22) 
1//
2// Copyright (C) 2009-2021 Intel Corporation
3//
4// SPDX-License-Identifier: MIT
5//
6//
7
8#include "instance.h"
9#include "api_interface.h"
10
11#include "bvh_build_primref.h"
12#include "bvh_build_refit.h"
13
14/*
15  Create primrefs from array of instance descriptors.
16 */
17 GRL_ANNOTATE_IGC_DO_NOT_SPILL
18__attribute__((reqd_work_group_size(MAX_HW_SIMD_WIDTH, 1, 1)))
19__attribute__((intel_reqd_sub_group_size(MAX_HW_SIMD_WIDTH))) void kernel
20TS_primrefs_from_instances(
21    global struct Globals* globals,
22    global struct BVHBase* bvh,
23    global __const struct GRL_RAYTRACING_INSTANCE_DESC* instances,
24    uint numInstances,
25    global struct AABB* primrefs,
26    global uchar* pAABBs,
27    global uchar* pIsProcedural,
28    dword aabb_stride,
29    uint allowUpdate
30    )
31{
32    const uint instanceIndex = get_sub_group_local_id() + get_group_id(0) * MAX_HW_SIMD_WIDTH;
33    if (instanceIndex < numInstances)
34    {
35        global __const struct GRL_RAYTRACING_INSTANCE_DESC* instance = instances + instanceIndex;
36
37        global struct GRL_RAYTRACING_AABB* procedural_bb = 0;
38        if ( pIsProcedural[instanceIndex] )
39        {
40            procedural_bb = (global struct GRL_RAYTRACING_AABB*)(pAABBs + aabb_stride * instanceIndex);
41        }
42
43        primrefs_from_instances(
44            globals,
45            bvh,
46            instance,
47            instanceIndex,
48            primrefs,
49            procedural_bb,
50            allowUpdate);
51    }
52}
53
54/*
55  Create primrefs from array of instance descriptors.
56 */
57 GRL_ANNOTATE_IGC_DO_NOT_SPILL
58__attribute__((reqd_work_group_size(MAX_HW_SIMD_WIDTH, 1, 1)))
59void kernel
60TS_primrefs_from_instances_indirect(
61    global struct Globals* globals,
62    global struct BVHBase* bvh,
63    global __const struct GRL_RAYTRACING_INSTANCE_DESC* instances,
64    uint numInstances,
65    global struct AABB* primrefs,
66    global uchar* pAABBs,
67    global uchar* pIsProcedural,
68    dword aabb_stride,
69    uint allowUpdate,
70    global struct IndirectBuildRangeInfo* indirect_data
71    )
72{
73    const uint instanceIndex = get_local_id(0) + get_group_id(0) * MAX_HW_SIMD_WIDTH;
74    if (instanceIndex < indirect_data->primitiveCount)
75    {
76        instances = (global __const struct GRL_RAYTRACING_INSTANCE_DESC*)
77            (((global char*)instances) + indirect_data->primitiveOffset);
78        global __const struct GRL_RAYTRACING_INSTANCE_DESC* instance = instances + instanceIndex;
79
80        global struct GRL_RAYTRACING_AABB* procedural_bb = 0;
81        if ( pIsProcedural[instanceIndex] )
82        {
83            procedural_bb = (global struct GRL_RAYTRACING_AABB*)(pAABBs + aabb_stride * instanceIndex);
84        }
85
86        primrefs_from_instances(
87            globals,
88            bvh,
89            instance,
90            instanceIndex,
91            primrefs,
92            procedural_bb,
93            allowUpdate);
94    }
95}
96
97/*
98  Create primrefs from array of pointers to instance descriptors.
99 */
100 GRL_ANNOTATE_IGC_DO_NOT_SPILL
101__attribute__((reqd_work_group_size(MAX_HW_SIMD_WIDTH, 1, 1)))
102__attribute__((intel_reqd_sub_group_size(MAX_HW_SIMD_WIDTH))) void kernel
103TS_primrefs_from_instances_pointers(global struct Globals* globals,
104    global struct BVHBase* bvh,
105    global void* instances_in,
106    uint numInstances,
107    global struct AABB* primrefs,
108    global uchar* pAABBs,
109    global uchar* pIsProcedural,
110    dword aabb_stride,
111    uint allowUpdate
112    )
113{
114    global const struct GRL_RAYTRACING_INSTANCE_DESC** instances =
115        (global const struct GRL_RAYTRACING_INSTANCE_DESC**)instances_in;
116
117    const uint instanceIndex = get_sub_group_local_id() + get_group_id(0) * MAX_HW_SIMD_WIDTH;
118    if (instanceIndex < numInstances)
119    {
120        global __const struct GRL_RAYTRACING_INSTANCE_DESC* instance = instances[instanceIndex];
121
122        global struct GRL_RAYTRACING_AABB* procedural_bb = 0;
123        if (pIsProcedural[instanceIndex])
124        {
125            procedural_bb = (global struct GRL_RAYTRACING_AABB*)(pAABBs + aabb_stride * instanceIndex);
126        }
127
128        primrefs_from_instances(
129            globals,
130            bvh,
131            instance,
132            instanceIndex,
133            primrefs,
134            procedural_bb,
135            allowUpdate);
136    }
137}
138
139/*
140  Create primrefs from array of pointers to instance descriptors.
141 */
142 GRL_ANNOTATE_IGC_DO_NOT_SPILL
143__attribute__((reqd_work_group_size(MAX_HW_SIMD_WIDTH, 1, 1)))
144void kernel
145TS_primrefs_from_instances_pointers_indirect(global struct Globals* globals,
146    global struct BVHBase* bvh,
147    global void* instances_in,
148    global struct AABB* primrefs,
149    global uchar* pAABBs,
150    global uchar* pIsProcedural,
151    dword aabb_stride,
152    uint allowUpdate,
153    global struct IndirectBuildRangeInfo* indirect_data
154    )
155{
156    const uint instanceIndex = get_local_id(0) + get_group_id(0) * MAX_HW_SIMD_WIDTH;
157    if (instanceIndex < indirect_data->primitiveCount)
158    {
159        instances_in = ((global char*)instances_in) + indirect_data->primitiveOffset;
160        global const struct GRL_RAYTRACING_INSTANCE_DESC** instances =
161            (global const struct GRL_RAYTRACING_INSTANCE_DESC**)instances_in;
162        global __const struct GRL_RAYTRACING_INSTANCE_DESC* instance = instances[instanceIndex];
163
164        global struct GRL_RAYTRACING_AABB* procedural_bb = 0;
165        if (pIsProcedural[instanceIndex])
166        {
167            procedural_bb = (global struct GRL_RAYTRACING_AABB*)(pAABBs + aabb_stride * instanceIndex);
168        }
169
170        primrefs_from_instances(
171            globals,
172            bvh,
173            instance,
174            instanceIndex,
175            primrefs,
176            procedural_bb,
177            allowUpdate);
178    }
179}
180
181
182
183GRL_ANNOTATE_IGC_DO_NOT_SPILL
184__attribute__((reqd_work_group_size(16, 1, 1)))
185void kernel
186TS_update_instance_leaves(global struct BVHBase* bvh,
187    uint64_t dxrInstancesArray,
188    uint64_t dxrInstancesPtr,
189    global struct AABB3f* instance_aabb_scratch,
190    global uchar* aabbs,
191    global uchar* is_procedural,
192    dword aabb_stride
193)
194{
195    uint num_leaves = BVHBase_GetNumHWInstanceLeaves(bvh);
196    uint id = get_local_id(0) + get_local_size(0) * get_group_id(0);
197    if (id >= num_leaves)
198        return;
199
200    struct HwInstanceLeaf* leaves = BVHBase_GetHWInstanceLeaves(bvh);
201    uint idx = HwInstanceLeaf_GetInstanceIndex(&leaves[id]);
202
203    global GRL_RAYTRACING_AABB* procedural_box = 0;
204    if (is_procedural[idx])
205    {
206        procedural_box = (global GRL_RAYTRACING_AABB*)(aabbs + (aabb_stride * idx));
207    }
208
209    DO_update_instance_leaves(
210        bvh,
211        dxrInstancesArray,
212        dxrInstancesPtr,
213        instance_aabb_scratch,
214        id,
215        procedural_box);
216}
217
218
219GRL_ANNOTATE_IGC_DO_NOT_SPILL
220__attribute__((reqd_work_group_size(16, 1, 1)))
221void kernel
222TS_fixup_leaves( global struct BVHBase* bvh,
223                 global uchar* primref_index,
224                 global PrimRef* primrefs,
225                 uint stride )
226
227{
228    uint num_inners = BVHBase_GetNumInternalNodes(bvh);
229    uint id       = get_local_id(0) + get_local_size(0) * get_group_id(0);
230
231    // assign 8 lanes to each inner node, 6 of which will do useful work
232    uint node_id  = id / 8;
233    uint child_id = id % 8;
234
235    bool node_valid = (node_id < num_inners);
236
237    if (node_valid )
238    {
239        global InternalNode* nodes = (global InternalNode*) BVHBase_GetInternalNodes(bvh);
240        global InternalNode* my_node = nodes + node_id;
241
242        if (my_node->nodeType == BVH_INSTANCE_NODE)
243        {
244            bool child_valid = (child_id < 6) && InternalNode_IsChildValid(my_node, child_id);
245            if (child_valid)
246            {
247                global HwInstanceLeaf* leaves = (global HwInstanceLeaf*)InternalNode_GetChildren(my_node);
248                uint leafIndex = (leaves - BVHBase_GetHWInstanceLeaves(bvh)) + child_id;
249
250                const uint primrefID = *(uint*)(primref_index + leafIndex * stride);
251
252                uint type = PRIMREF_isProceduralInstance(&primrefs[primrefID]) ?
253                                BVH_PROCEDURAL_NODE : BVH_INSTANCE_NODE;
254
255                InternalNode_SetChildType(my_node, child_id, type);
256            }
257
258            if (child_id == 0)
259                my_node->nodeType = BVH_INTERNAL_NODE;
260        }
261    }
262}
263
264
265
266
267
268GRL_ANNOTATE_IGC_DO_NOT_SPILL
269__attribute__((reqd_work_group_size(SG_REFIT_WG_SIZE, 1, 1))) void kernel
270TS_Refit_per_one_startpoint_sg(
271    global struct BVHBase* bvh,
272    global struct AABB3f* instance_leaf_aabbs,
273    global uchar* procedural_instance_enable_buffer )
274{
275    DO_Refit_per_one_startpoint_sg(bvh, (global GRL_RAYTRACING_GEOMETRY_DESC*) bvh, instance_leaf_aabbs, procedural_instance_enable_buffer );
276
277}
278