From e595743b5aa4f6bd88800cfbcfd6eead3cc3d01b Mon Sep 17 00:00:00 2001
From: "Harsh Aggarwal (NVIDIA)" <haaggarwal@nvidia.com>
Date: Fri, 8 Aug 2025 03:13:25 +0530
Subject: Fix intrinsic LoadLocalRootTableConstant for optix (#7949)

Due to an older version of spec referred there was an inconsitency v1.29
2/20/2025 - [HitObject LoadLocalRootArgumentsConstant]

Latest spec

https://microsoft.github.io/DirectX-Specs/d3d/Raytracing.html#hitobject-loadlocalroottableconstant

Refer:
OptiX backend support for Shader Execution Reordering (SER) features as
outlined in issue #6647. -
---
 prelude/slang-cuda-prelude.h | 35 +++++++++++++++++------------------
 1 file changed, 17 insertions(+), 18 deletions(-)

(limited to 'prelude')

diff --git a/prelude/slang-cuda-prelude.h b/prelude/slang-cuda-prelude.h
index a66fa15cb..c366ae856 100644
--- a/prelude/slang-cuda-prelude.h
+++ b/prelude/slang-cuda-prelude.h
@@ -3644,22 +3644,23 @@ __forceinline__ __device__ void* optixTraverse(
         r1);
 }
 
-static __forceinline__ __device__ bool optixHitObjectIsHit(OptixTraversableHandle* hitObj)
+static __forceinline__ __device__ bool slangOptixHitObjectIsHit(OptixTraversableHandle* hitObj)
 {
     return optixHitObjectIsHit();
 }
 
-static __forceinline__ __device__ bool optixHitObjectIsMiss(OptixTraversableHandle* hitObj)
+static __forceinline__ __device__ bool slangOptixHitObjectIsMiss(OptixTraversableHandle* hitObj)
 {
     return optixHitObjectIsMiss();
 }
 
-static __forceinline__ __device__ bool optixHitObjectIsNop(OptixTraversableHandle* hitObj)
+static __forceinline__ __device__ bool slangOptixHitObjectIsNop(OptixTraversableHandle* hitObj)
 {
     return optixHitObjectIsNop();
 }
 
-static __forceinline__ __device__ uint optixHitObjectGetClusterId(OptixTraversableHandle* hitObj)
+static __forceinline__ __device__ uint
+slangOptixHitObjectGetClusterId(OptixTraversableHandle* hitObj)
 {
     return optixHitObjectGetClusterId();
 }
@@ -3809,7 +3810,7 @@ static __forceinline__ __device__ void optixMakeHitObject(
         0 /*numTransforms */);
 }
 
-static __forceinline__ __device__ void optixMakeNopHitObject(OptixTraversableHandle* Obj)
+static __forceinline__ __device__ void slangOptixMakeNopHitObject(OptixTraversableHandle* Obj)
 {
     optixMakeNopHitObject();
 }
@@ -3834,22 +3835,25 @@ static __forceinline__ __device__ RayDesc optixHitObjectGetRayDesc(OptixTraversa
     return ray;
 }
 
-static __forceinline__ __device__ uint optixHitObjectGetInstanceIndex(OptixTraversableHandle* Obj)
+static __forceinline__ __device__ uint
+slangOptixHitObjectGetInstanceIndex(OptixTraversableHandle* Obj)
 {
     return optixHitObjectGetInstanceIndex();
 }
 
-static __forceinline__ __device__ uint optixHitObjectGetInstanceId(OptixTraversableHandle* Obj)
+static __forceinline__ __device__ uint slangOptixHitObjectGetInstanceId(OptixTraversableHandle* Obj)
 {
     return optixHitObjectGetInstanceId();
 }
 
-static __forceinline__ __device__ uint optixHitObjectGetSbtGASIndex(OptixTraversableHandle* Obj)
+static __forceinline__ __device__ uint
+slangOptixHitObjectGetSbtGASIndex(OptixTraversableHandle* Obj)
 {
     return optixHitObjectGetSbtGASIndex();
 }
 
-static __forceinline__ __device__ uint optixHitObjectGetPrimitiveIndex(OptixTraversableHandle* Obj)
+static __forceinline__ __device__ uint
+slangOptixHitObjectGetPrimitiveIndex(OptixTraversableHandle* Obj)
 {
     return optixHitObjectGetPrimitiveIndex();
 }
@@ -3888,22 +3892,17 @@ static __forceinline__ __device__ T optixHitObjectGetAttribute(OptixTraversableH
     return result;
 }
 
-static __forceinline__ __device__ uint optixHitObjectGetSbtRecordIndex(OptixTraversableHandle* Obj)
+static __forceinline__ __device__ uint
+slangOptixHitObjectGetSbtRecordIndex(OptixTraversableHandle* Obj)
 {
     return optixHitObjectGetSbtRecordIndex();
 }
 
 static __forceinline__ __device__ uint
-optixHitObjectSetSbtRecordIndex(OptixTraversableHandle* Obj, uint sbtRecordIndex)
+slangOptixHitObjectSetSbtRecordIndex(OptixTraversableHandle* Obj, uint sbtRecordIndex)
 {
     optixHitObjectSetSbtRecordIndex(sbtRecordIndex); // returns void
-    return 0;
-}
-static __forceinline__ __device__ uint
-optixHitObjectGetSbtDataPointer(OptixTraversableHandle* Obj, uint sbtRecordIndex)
-{
-    optixHitObjectGetSbtDataPointer(); // returns void
-    return 0;
+    return sbtRecordIndex;
 }
 #endif
 static const int kSlangTorchTensorMaxDim = 5;
-- 
cgit v1.2.3