Skip to content

Commit d446e9c

Browse files
committed
sycl: cache committed candidate type in hit writeback path
rtcIntersectRTHW queried intel_get_hit_candidate(committed_hit) multiple times in the final writeback section. Cache the committed candidate once and reuse it for: - triangle primID/barycentric extraction - triangle normal fetch This removes a redundant hardware query in a hot post-traversal path without changing behavior.
1 parent 99097b6 commit d446e9c

1 file changed

Lines changed: 3 additions & 2 deletions

File tree

kernels/sycl/rthwif_embree.cpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -736,7 +736,8 @@ SYCL_EXTERNAL __attribute__((always_inline)) void rtcIntersectRTHW(sycl::global_
736736
unsigned int geomID = intel_get_hit_geometry_id(query, intel_hit_type_committed_hit);
737737

738738
unsigned int primID = ray.primID;
739-
if (intel_get_hit_candidate(query, intel_hit_type_committed_hit) == intel_candidate_type_triangle) {
739+
const intel_candidate_type_t committed_candidate = intel_get_hit_candidate(query, intel_hit_type_committed_hit);
740+
if (committed_candidate == intel_candidate_type_triangle) {
740741
primID = intel_get_hit_triangle_primitive_id(query, intel_hit_type_committed_hit);
741742
uv = intel_get_hit_barycentrics (query, intel_hit_type_committed_hit);
742743
}
@@ -768,7 +769,7 @@ SYCL_EXTERNAL __attribute__((always_inline)) void rtcIntersectRTHW(sycl::global_
768769
#endif
769770

770771
/* calculate geometry normal for hardware accelerated triangles */
771-
if (intel_get_hit_candidate(query, intel_hit_type_committed_hit) == intel_candidate_type_triangle)
772+
if (committed_candidate == intel_candidate_type_triangle)
772773
ray.Ng = intel_get_hit_triangle_normal(query, intel_hit_type_committed_hit);
773774

774775
rayhit_i->hit.Ng_x = ray.Ng.x;

0 commit comments

Comments
 (0)