1// 2// Copyright (C) 2009-2021 Intel Corporation 3// 4// SPDX-License-Identifier: MIT 5// 6// 7 8// @file bvh_debug.cl 9// 10// @brief routines to do basic integrity checks 11// 12// Notes: 13// 14 15#include "GRLGen12.h" 16#include "intrinsics.h" 17#include "libs/lsc_intrinsics.h" 18#include "GRLGen12IntegrityChecks.h" 19#include "api_interface.h" 20 21#define ERROR_PRINTF 0 22GRL_INLINE bool commit_err( 23 global uint* some_null, 24 global BVHBase* bvh, 25 global ERROR_INFO* err_info_slot, 26 ERROR_INFO err) 27{ 28 if (err.type != error_t_no_error) { 29 uint expected = error_t_no_error; 30 atomic_compare_exchange_global(&err_info_slot->type, &expected, err.type); 31 if (expected == error_t_no_error) 32 { 33 err_info_slot->offset_in_BVH = err.offset_in_BVH; 34 err_info_slot->when = err.when; 35 err_info_slot->reserved = 0xAAACCAAA; 36 mem_fence_evict_to_memory(); 37#if ERROR_PRINTF 38 printf("bvh = 0x%llX, err.type = %X, err.offset_in_BVH = %d\n", bvh, err.type, err.offset_in_BVH); 39#else 40 // This is to trigger PF. Note we have to write directly to memory. 41 // If write would stay in L3 it won't give a PF untill this will get evicted to mem. 42 store_uint_L1UC_L3UC(some_null, 0, 0x0EEE0000 + err.type); 43#endif 44 return true; 45 } 46 } 47 return false; 48} 49 50__attribute__((reqd_work_group_size(16, 1, 1))) 51void kernel check_tree_topology( 52 global uint* some_null, 53 global BVHBase* bvh, 54 global ERROR_INFO* err, 55 uint phase) 56{ 57 uint globalID = get_local_id(0) + get_group_id(0) * get_local_size(0); 58 59 if (err->type != error_t_no_error) return; 60 61 uint dummy1, dummy2, dummy3; 62 ERROR_INFO reterr = check_tree_topology_helper(bvh, globalID, &dummy1, &dummy2, &dummy3, false); 63 if (reterr.type == error_t_no_error) 64 { 65 reterr = check_backpointers(bvh, globalID); 66 } 67 if (reterr.type == error_t_no_error) 68 { 69 reterr = validate_atomic_update_structs(bvh, globalID); 70 } 71 reterr.when = phase; 72 commit_err(some_null, bvh, err, reterr); 73} 74 75GRL_INLINE bool IsValid48bPtr(qword ptr) 76{ 77 qword CANONIZED_BITS = 0xFFFFul << 48ul; 78 qword canonized_part = ptr & CANONIZED_BITS; 79 bool isIt = ptr != 0 && ( 80 canonized_part == 0 || canonized_part == CANONIZED_BITS); 81 return isIt; 82} 83 84__attribute__((reqd_work_group_size(16, 1, 1))) 85void kernel check_geos_before_quad_update( 86 global BVHBase* bvh, //dest bvh 87 global GRL_RAYTRACING_GEOMETRY_DESC* geomDesc, 88 global uint* some_null, 89 global ERROR_INFO* err, 90 uint phase, 91 uint numGeos, 92 uint numThreads) 93{ 94 uint globalID = get_local_id(0) + get_group_id(0) * get_local_size(0); 95 96 if (err->type != error_t_no_error) return; 97 98 // first check sanity of geos 99 ERROR_INFO geo_insanity_error = { error_t_input_geo_insane, 0 }; 100 101 for (uint ID = globalID; ID < numGeos; ID += numThreads * get_sub_group_size()) 102 { 103 bool IsSane = IsValid48bPtr((qword)(qword)geomDesc); 104 105 if (IsSane) { 106 GRL_RAYTRACING_GEOMETRY_DESC geo = geomDesc[globalID]; 107 IsSane = geo.Type < NUM_GEOMETRY_TYPES; 108 if (IsSane) { 109 if (geo.Type == GEOMETRY_TYPE_TRIANGLES) { 110 if (geo.Desc.Triangles.IndexFormat >= INDEX_FORMAT_END) { 111 IsSane = false; 112 } 113 else 114 { 115 if (geo.Desc.Triangles.IndexFormat != INDEX_FORMAT_NONE && geo.Desc.Triangles.IndexCount > 2) 116 { 117 IsSane = (geo.Desc.Triangles.VertexFormat < VERTEX_FORMAT_END) && 118 IsValid48bPtr((qword)geo.Desc.Triangles.pVertexBuffer) && 119 IsValid48bPtr((qword)geo.Desc.Triangles.pIndexBuffer); 120 } 121 else if (geo.Desc.Triangles.VertexCount > 2) 122 { 123 IsSane = 124 geo.Desc.Triangles.VertexFormat < VERTEX_FORMAT_END&& 125 IsValid48bPtr((qword)geo.Desc.Triangles.pVertexBuffer) != 0; 126 } 127 } 128 } 129 } 130 } 131 132 geo_insanity_error.offset_in_BVH = ID; 133 geo_insanity_error.when = phase; 134 if (!IsSane) { 135 commit_err(some_null, bvh, err, geo_insanity_error); 136 } 137 return; 138 } 139} 140 141__attribute__((reqd_work_group_size(16, 1, 1))) 142void kernel check_geos_vs_quads( 143 global BVHBase* bvh, 144 global GRL_RAYTRACING_GEOMETRY_DESC* geomDesc, 145 global uint* some_null, 146 global ERROR_INFO* err, 147 uint phase, 148 uint numGeos, 149 uint numThreads) 150{ 151 uint numQuads = BVHBase_GetNumQuads(bvh); 152 153 QuadLeaf* quads = BVHBase_GetQuadLeaves(bvh); 154 155 uint globalID = get_local_id(0) + get_group_id(0) * get_local_size(0); 156 uint qoffset = bvh->quadLeafStart; 157 158 if (err->type != error_t_no_error) return; 159 160 ERROR_INFO theErr = { error_t_no_error, 0 }; 161 162 for (uint ID = globalID; ID < numQuads; ID += numThreads * get_sub_group_size()) 163 { 164 ERROR_INFO quadErr = { error_t_quad_leaf_broken, qoffset + ID, phase }; 165 166 QuadLeaf quad = quads[ID]; 167 168 uint geoIdx = PrimLeaf_GetGeoIndex(&quad.leafDesc); 169 170 if (geoIdx > numGeos) { commit_err(some_null, bvh, err, quadErr); return; } 171 172 uint numPrimsInGeo = geomDesc[geoIdx].Desc.Triangles.IndexFormat != INDEX_FORMAT_NONE ? 173 geomDesc[geoIdx].Desc.Triangles.IndexCount / 3 : 174 geomDesc[geoIdx].Desc.Triangles.VertexCount / 3; 175 176 if(quad.primIndex0 >= numPrimsInGeo) { 177 commit_err(some_null, bvh, err, quadErr); 178 return; 179 } 180 181 if(!QuadLeaf_IsSingleTriangle(&quad) && 182 (quad.primIndex0 + QuadLeaf_GetPrimIndexDelta(&quad) >= numPrimsInGeo)) 183 { 184 commit_err(some_null, bvh, err, quadErr); 185 return; 186 } 187 } 188} 189 190__attribute__((reqd_work_group_size(16, 1, 1))) 191void kernel check_instances_linked_bvhs( 192 global uint* some_null, 193 global BVHBase* bvh, 194 global ERROR_INFO* err, 195 uint phase) 196{ 197 if (err->type != error_t_no_error) return; 198 199 uint instanceLeafStart = bvh->instanceLeafStart; 200 uint instanceLeafEnd = bvh->instanceLeafEnd; 201 uint numInstances = (instanceLeafEnd - instanceLeafStart) / 2; 202 203 uint globalID = get_local_id(0) + get_group_id(0) * get_local_size(0); 204 205 ERROR_INFO reterr = check_instances_linked_bvhs_helper(bvh, globalID, /*touchBlas*/true); 206 reterr.when = phase; 207 commit_err(some_null, bvh, err, reterr); 208} 209