Fix intrinsic LoadLocalRootTableConstant for optix (#7949)

szihs · web-flow · commit e595743b5aa4 · 2025-08-07T21:43:25.000Z
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. -
diff --git 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;
diff --git a/source/core/slang-platform.cpp b/source/core/slang-platform.cpp
@@ -174,12 +174,27 @@ SLANG_COMPILE_TIME_ASSERT(E_OUTOFMEMORY == SLANG_E_OUT_OF_MEMORY);
 }
 
 #else // _WIN32
-
 /* static */ SlangResult PlatformUtil::getInstancePath([[maybe_unused]] StringBuilder& out)
 {
-    // On non Windows it's typically hard to get the instance path, so we'll say not implemented.
-    // The meaning is also somewhat more ambiguous - is it the exe or the shared library path?
+#if defined(__linux__) || defined(__CYGWIN__)
+    char path[PATH_MAX];
+    ssize_t len = readlink("/proc/self/exe", path, sizeof(path) - 1);
+    if (len == -1)
+    {
+        return SLANG_FAIL;
+    }
+
+    path[len] = '\0';
+    String pathString(path);
+
+    // We don't want the instance name, just the path to it
+    out.clear();
+    out.append(Path::getParentDirectory(pathString));
+
+    return out.getLength() > 0 ? SLANG_OK : SLANG_FAIL;
+#else
     return SLANG_E_NOT_IMPLEMENTED;
+#endif
 }
 
 /* static */ SlangResult PlatformUtil::appendResult(
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
@@ -20472,7 +20472,7 @@ struct HitObject
             __intrinsic_asm "($0 = NvMakeNop())";
         case glsl:
             __glslMakeNop(__return_val);
-        case cuda: __intrinsic_asm "optixMakeNopHitObject";
+        case cuda: __intrinsic_asm "slangOptixMakeNopHitObject";
         case spirv:
             spirv_asm 
             {
@@ -20561,7 +20561,7 @@ struct HitObject
         {
         case hlsl: __intrinsic_asm ".IsMiss";
         case glsl: __intrinsic_asm "hitObjectIsMissNV($0)";
-        case cuda: __intrinsic_asm "optixHitObjectIsMiss";
+        case cuda: __intrinsic_asm "slangOptixHitObjectIsMiss";
         case spirv:
             return spirv_asm 
             {
@@ -20582,7 +20582,7 @@ struct HitObject
         {
         case hlsl: __intrinsic_asm ".IsHit";
         case glsl: __intrinsic_asm "hitObjectIsHitNV($0)";
-        case cuda: __intrinsic_asm "optixHitObjectIsHit";
+        case cuda: __intrinsic_asm "slangOptixHitObjectIsHit";
         case spirv: 
             return spirv_asm 
             {
@@ -20603,7 +20603,7 @@ struct HitObject
         {
         case hlsl: __intrinsic_asm ".IsNop";
         case glsl: __intrinsic_asm "hitObjectIsEmptyNV($0)";
-        case cuda: __intrinsic_asm "optixHitObjectIsNop";
+        case cuda: __intrinsic_asm "slangOptixHitObjectIsNop";
         case spirv: 
             return spirv_asm 
             {
@@ -20655,7 +20655,7 @@ struct HitObject
         {
         case hlsl: __intrinsic_asm ".GetShaderTableIndex";
         case glsl: __intrinsic_asm "hitObjectGetShaderBindingTableRecordIndexNV($0)";
-        case cuda: __intrinsic_asm "optixHitObjectGetSbtRecordIndex";
+        case cuda: __intrinsic_asm "slangOptixHitObjectGetSbtRecordIndex";
         case spirv: 
             return spirv_asm 
             {
@@ -20675,7 +20675,7 @@ struct HitObject
         __target_switch
         {
         case hlsl: __intrinsic_asm ".SetShaderTableIndex";
-        case cuda: __intrinsic_asm "optixHitObjectSetSbtRecordIndex";
+        case cuda: __intrinsic_asm "slangOptixHitObjectSetSbtRecordIndex";
         }
     }
         /// Returns the instance index of a hit. Valid if the hit object represents a hit.
@@ -20689,7 +20689,7 @@ struct HitObject
         {
         case hlsl: __intrinsic_asm ".GetInstanceIndex";
         case glsl: __intrinsic_asm "hitObjectGetInstanceIdNV($0)";
-        case cuda: __intrinsic_asm "optixHitObjectGetInstanceIndex";
+        case cuda: __intrinsic_asm "slangOptixHitObjectGetInstanceIndex";
         case spirv: 
             return spirv_asm 
             {
@@ -20711,7 +20711,7 @@ struct HitObject
         {
         case hlsl: __intrinsic_asm ".GetInstanceID";
         case glsl: __intrinsic_asm "hitObjectGetInstanceCustomIndexNV($0)";
-        case cuda: __intrinsic_asm "optixHitObjectGetInstanceId";
+        case cuda: __intrinsic_asm "slangOptixHitObjectGetInstanceId";
         case spirv:
             return spirv_asm 
             {
@@ -20733,7 +20733,7 @@ struct HitObject
         {
         case hlsl: __intrinsic_asm ".GetGeometryIndex";
         case glsl: __intrinsic_asm "hitObjectGetGeometryIndexNV($0)";
-        case cuda: __intrinsic_asm "optixHitObjectGetSbtGASIndex";
+        case cuda: __intrinsic_asm "slangOptixHitObjectGetSbtGASIndex";
         case spirv: 
             return spirv_asm 
             {
@@ -20755,7 +20755,7 @@ struct HitObject
         {
         case hlsl: __intrinsic_asm ".GetPrimitiveIndex";
         case glsl: __intrinsic_asm "hitObjectGetPrimitiveIndexNV($0)";
-        case cuda: __intrinsic_asm "optixHitObjectGetPrimitiveIndex";
+        case cuda: __intrinsic_asm "slangOptixHitObjectGetPrimitiveIndex";
         case spirv: 
             return spirv_asm 
             {
@@ -20799,7 +20799,7 @@ struct HitObject
         {
         case hlsl: __intrinsic_asm ".GetClusterID";
         case glsl: __intrinsic_asm "hitObjectGetClusterIdNV($0)";
-        case cuda: __intrinsic_asm "optixHitObjectGetClusterId";
+        case cuda: __intrinsic_asm "slangOptixHitObjectGetClusterId";
         case spirv:
             return spirv_asm
             {
diff --git a/tests/cuda/optix-cluster.slang b/tests/cuda/optix-cluster.slang
@@ -1,4 +1,5 @@
 //TEST:SIMPLE(filecheck=CHECK): -target cuda
+//TEST:SIMPLE(filecheck=CHECK-PTX): -target ptx -Xnvrtc -I"./external/optix-dev/include/" -entry closestHitShaderA
 //CHECK: __global__ void __closesthit__closestHitShaderA
 //CHECK: optixGetClusterId
 struct RayPayload
@@ -9,6 +10,7 @@ struct RayPayload
 [shader("closesthit")]
 void closestHitShaderA(inout RayPayload payload, in BuiltInTriangleIntersectionAttributes attr)
 {
+    //CHECK-PTX:_optix_get_cluster_id 
     int clusterId = GetClusterID();
     float4 color = float4(0, 0, 0, 1);
     if (clusterId >= 0)
diff --git a/tests/cuda/optix-coopvec.slang b/tests/cuda/optix-coopvec.slang
@@ -1,5 +1,7 @@
 //TEST:SIMPLE(filecheck=CHECK): -target cuda -capability optix_coopvec
+//TEST:SIMPLE(filecheck=CHECK-PTX): -target ptx -Xnvrtc -I"./external/optix-dev/include/"
 
+// CHECK-PTX: add.f32
 // CHECK: optixCoopVecLoad
 // CHECK: OptixCoopVec
 // CHECK: optixCoopVecTanh
diff --git a/tests/cuda/optix-hit-attributes.slang b/tests/cuda/optix-hit-attributes.slang
@@ -1,5 +1,6 @@
 //TEST:SIMPLE(filecheck=CHECK): -target cuda
 //CHECK: __global__ void __closesthit__closestHitShaderA
+//TEST:SIMPLE(filecheck=CHECK-PTX): -target ptx -Xnvrtc -I"./external/optix-dev/include/"
 struct RayPayload
 {
     float4 color;
@@ -8,6 +9,7 @@ struct RayPayload
 [shader("closesthit")]
 void closestHitShaderA(inout RayPayload payload, in BuiltInTriangleIntersectionAttributes attr)
 {
+    //CHECK-PTX: _optix_read_primitive_idx
     uint primitiveIndex = PrimitiveIndex();
     float4 color = float4(0, 0, 0, 1);
     color[primitiveIndex] = 1;
diff --git a/tests/cuda/optix-ignore-hit.slang b/tests/cuda/optix-ignore-hit.slang
@@ -1,8 +1,10 @@
 // optix-ignore-hit.slang
 //TEST:SIMPLE(filecheck=CHECK): -target cuda -entry anyHitShader
+//TEST:SIMPLE(filecheck=CHECK-PTX): -target ptx -Xnvrtc -I"./external/optix-dev/include/"
 //CHECK: HitBuffer_insert_0(((HitBuffer_0 *)getOptiXRayPayloadPtr()), hit_0.t_0);
 //CHECK: optixIgnoreIntersection
 
+//CHECK-PTX: _optix_get_ray_tmax
 
 struct HitBuffer
 {
@@ -28,6 +30,7 @@ void anyHitShader(inout HitBuffer rayHitBuffer)
     // Modify the inout parameter
     rayHitBuffer.insert(hit.t);
     
+    // CHECK-PTX: _optix_ignore_intersection
     // Early exit - should not lose the modification to inout rayHitBuffer
     if (hit.t < rayHitBuffer.last)
         IgnoreHit();
diff --git a/tests/cuda/optix-ser.slang b/tests/cuda/optix-ser.slang

Original file line number	Diff line number	Diff line change
`@@ -3644,22 +3644,23 @@ __forceinline__ __device__ void* optixTraverse(`
`3644`	`3644`	`r1);`
`3645`	`3645`	`}`
`3646`	`3646`
`3647`		`-static __forceinline__ __device__ bool optixHitObjectIsHit(OptixTraversableHandle* hitObj)`
	`3647`	`+static __forceinline__ __device__ bool slangOptixHitObjectIsHit(OptixTraversableHandle* hitObj)`
`3648`	`3648`	`{`
`3649`	`3649`	`return optixHitObjectIsHit();`
`3650`	`3650`	`}`
`3651`	`3651`
`3652`		`-static __forceinline__ __device__ bool optixHitObjectIsMiss(OptixTraversableHandle* hitObj)`
	`3652`	`+static __forceinline__ __device__ bool slangOptixHitObjectIsMiss(OptixTraversableHandle* hitObj)`
`3653`	`3653`	`{`
`3654`	`3654`	`return optixHitObjectIsMiss();`
`3655`	`3655`	`}`
`3656`	`3656`
`3657`		`-static __forceinline__ __device__ bool optixHitObjectIsNop(OptixTraversableHandle* hitObj)`
	`3657`	`+static __forceinline__ __device__ bool slangOptixHitObjectIsNop(OptixTraversableHandle* hitObj)`
`3658`	`3658`	`{`
`3659`	`3659`	`return optixHitObjectIsNop();`
`3660`	`3660`	`}`
`3661`	`3661`
`3662`		`-static __forceinline__ __device__ uint optixHitObjectGetClusterId(OptixTraversableHandle* hitObj)`
	`3662`	`+static __forceinline__ __device__ uint`
	`3663`	`+slangOptixHitObjectGetClusterId(OptixTraversableHandle* hitObj)`
`3663`	`3664`	`{`
`3664`	`3665`	`return optixHitObjectGetClusterId();`
`3665`	`3666`	`}`
`@@ -3809,7 +3810,7 @@ static __forceinline__ __device__ void optixMakeHitObject(`
`3809`	`3810`	`0 /numTransforms /);`
`3810`	`3811`	`}`
`3811`	`3812`
`3812`		`-static __forceinline__ __device__ void optixMakeNopHitObject(OptixTraversableHandle* Obj)`
	`3813`	`+static __forceinline__ __device__ void slangOptixMakeNopHitObject(OptixTraversableHandle* Obj)`
`3813`	`3814`	`{`
`3814`	`3815`	`optixMakeNopHitObject();`
`3815`	`3816`	`}`
`@@ -3834,22 +3835,25 @@ static __forceinline__ __device__ RayDesc optixHitObjectGetRayDesc(OptixTraversa`
`3834`	`3835`	`return ray;`
`3835`	`3836`	`}`
`3836`	`3837`
`3837`		`-static __forceinline__ __device__ uint optixHitObjectGetInstanceIndex(OptixTraversableHandle* Obj)`
	`3838`	`+static __forceinline__ __device__ uint`
	`3839`	`+slangOptixHitObjectGetInstanceIndex(OptixTraversableHandle* Obj)`
`3838`	`3840`	`{`
`3839`	`3841`	`return optixHitObjectGetInstanceIndex();`
`3840`	`3842`	`}`
`3841`	`3843`
`3842`		`-static __forceinline__ __device__ uint optixHitObjectGetInstanceId(OptixTraversableHandle* Obj)`
	`3844`	`+static __forceinline__ __device__ uint slangOptixHitObjectGetInstanceId(OptixTraversableHandle* Obj)`
`3843`	`3845`	`{`
`3844`	`3846`	`return optixHitObjectGetInstanceId();`
`3845`	`3847`	`}`
`3846`	`3848`
`3847`		`-static __forceinline__ __device__ uint optixHitObjectGetSbtGASIndex(OptixTraversableHandle* Obj)`
	`3849`	`+static __forceinline__ __device__ uint`
	`3850`	`+slangOptixHitObjectGetSbtGASIndex(OptixTraversableHandle* Obj)`
`3848`	`3851`	`{`
`3849`	`3852`	`return optixHitObjectGetSbtGASIndex();`
`3850`	`3853`	`}`
`3851`	`3854`
`3852`		`-static __forceinline__ __device__ uint optixHitObjectGetPrimitiveIndex(OptixTraversableHandle* Obj)`
	`3855`	`+static __forceinline__ __device__ uint`
	`3856`	`+slangOptixHitObjectGetPrimitiveIndex(OptixTraversableHandle* Obj)`
`3853`	`3857`	`{`
`3854`	`3858`	`return optixHitObjectGetPrimitiveIndex();`
`3855`	`3859`	`}`
`@@ -3888,22 +3892,17 @@ static __forceinline__ __device__ T optixHitObjectGetAttribute(OptixTraversableH`
`3888`	`3892`	`return result;`
`3889`	`3893`	`}`
`3890`	`3894`
`3891`		`-static __forceinline__ __device__ uint optixHitObjectGetSbtRecordIndex(OptixTraversableHandle* Obj)`
	`3895`	`+static __forceinline__ __device__ uint`
	`3896`	`+slangOptixHitObjectGetSbtRecordIndex(OptixTraversableHandle* Obj)`
`3892`	`3897`	`{`
`3893`	`3898`	`return optixHitObjectGetSbtRecordIndex();`
`3894`	`3899`	`}`
`3895`	`3900`
`3896`	`3901`	`static __forceinline__ __device__ uint`
`3897`		`-optixHitObjectSetSbtRecordIndex(OptixTraversableHandle* Obj, uint sbtRecordIndex)`
	`3902`	`+slangOptixHitObjectSetSbtRecordIndex(OptixTraversableHandle* Obj, uint sbtRecordIndex)`
`3898`	`3903`	`{`
`3899`	`3904`	`optixHitObjectSetSbtRecordIndex(sbtRecordIndex); // returns void`
`3900`		`- return 0;`
`3901`		`-}`
`3902`		`-static __forceinline__ __device__ uint`
`3903`		`-optixHitObjectGetSbtDataPointer(OptixTraversableHandle* Obj, uint sbtRecordIndex)`
`3904`		`-{`
`3905`		`- optixHitObjectGetSbtDataPointer(); // returns void`
`3906`		`- return 0;`
	`3905`	`+ return sbtRecordIndex;`
`3907`	`3906`	`}`
`3908`	`3907`	`#endif`
`3909`	`3908`	`static const int kSlangTorchTensorMaxDim = 5;`
Original file line number	Diff line number	Diff line change
`@@ -20472,7 +20472,7 @@ struct HitObject`
`20472`	`20472`	`__intrinsic_asm "($0 = NvMakeNop())";`
`20473`	`20473`	`case glsl:`
`20474`	`20474`	`__glslMakeNop(__return_val);`
`20475`		`- case cuda: __intrinsic_asm "optixMakeNopHitObject";`
	`20475`	`+ case cuda: __intrinsic_asm "slangOptixMakeNopHitObject";`
`20476`	`20476`	`case spirv:`
`20477`	`20477`	`spirv_asm`
`20478`	`20478`	`{`
`@@ -20561,7 +20561,7 @@ struct HitObject`
`20561`	`20561`	`{`
`20562`	`20562`	`case hlsl: __intrinsic_asm ".IsMiss";`
`20563`	`20563`	`case glsl: __intrinsic_asm "hitObjectIsMissNV($0)";`
`20564`		`- case cuda: __intrinsic_asm "optixHitObjectIsMiss";`
	`20564`	`+ case cuda: __intrinsic_asm "slangOptixHitObjectIsMiss";`
`20565`	`20565`	`case spirv:`
`20566`	`20566`	`return spirv_asm`
`20567`	`20567`	`{`
`@@ -20582,7 +20582,7 @@ struct HitObject`
`20582`	`20582`	`{`
`20583`	`20583`	`case hlsl: __intrinsic_asm ".IsHit";`
`20584`	`20584`	`case glsl: __intrinsic_asm "hitObjectIsHitNV($0)";`
`20585`		`- case cuda: __intrinsic_asm "optixHitObjectIsHit";`
	`20585`	`+ case cuda: __intrinsic_asm "slangOptixHitObjectIsHit";`
`20586`	`20586`	`case spirv:`
`20587`	`20587`	`return spirv_asm`
`20588`	`20588`	`{`
`@@ -20603,7 +20603,7 @@ struct HitObject`
`20603`	`20603`	`{`
`20604`	`20604`	`case hlsl: __intrinsic_asm ".IsNop";`
`20605`	`20605`	`case glsl: __intrinsic_asm "hitObjectIsEmptyNV($0)";`
`20606`		`- case cuda: __intrinsic_asm "optixHitObjectIsNop";`
	`20606`	`+ case cuda: __intrinsic_asm "slangOptixHitObjectIsNop";`
`20607`	`20607`	`case spirv:`
`20608`	`20608`	`return spirv_asm`
`20609`	`20609`	`{`
`@@ -20655,7 +20655,7 @@ struct HitObject`
`20655`	`20655`	`{`
`20656`	`20656`	`case hlsl: __intrinsic_asm ".GetShaderTableIndex";`
`20657`	`20657`	`case glsl: __intrinsic_asm "hitObjectGetShaderBindingTableRecordIndexNV($0)";`
`20658`		`- case cuda: __intrinsic_asm "optixHitObjectGetSbtRecordIndex";`
	`20658`	`+ case cuda: __intrinsic_asm "slangOptixHitObjectGetSbtRecordIndex";`
`20659`	`20659`	`case spirv:`
`20660`	`20660`	`return spirv_asm`
`20661`	`20661`	`{`
`@@ -20675,7 +20675,7 @@ struct HitObject`
`20675`	`20675`	`__target_switch`
`20676`	`20676`	`{`
`20677`	`20677`	`case hlsl: __intrinsic_asm ".SetShaderTableIndex";`
`20678`		`- case cuda: __intrinsic_asm "optixHitObjectSetSbtRecordIndex";`
	`20678`	`+ case cuda: __intrinsic_asm "slangOptixHitObjectSetSbtRecordIndex";`
`20679`	`20679`	`}`
`20680`	`20680`	`}`
`20681`	`20681`	`/// Returns the instance index of a hit. Valid if the hit object represents a hit.`
`@@ -20689,7 +20689,7 @@ struct HitObject`
`20689`	`20689`	`{`
`20690`	`20690`	`case hlsl: __intrinsic_asm ".GetInstanceIndex";`
`20691`	`20691`	`case glsl: __intrinsic_asm "hitObjectGetInstanceIdNV($0)";`
`20692`		`- case cuda: __intrinsic_asm "optixHitObjectGetInstanceIndex";`
	`20692`	`+ case cuda: __intrinsic_asm "slangOptixHitObjectGetInstanceIndex";`
`20693`	`20693`	`case spirv:`
`20694`	`20694`	`return spirv_asm`
`20695`	`20695`	`{`
`@@ -20711,7 +20711,7 @@ struct HitObject`
`20711`	`20711`	`{`
`20712`	`20712`	`case hlsl: __intrinsic_asm ".GetInstanceID";`
`20713`	`20713`	`case glsl: __intrinsic_asm "hitObjectGetInstanceCustomIndexNV($0)";`
`20714`		`- case cuda: __intrinsic_asm "optixHitObjectGetInstanceId";`
	`20714`	`+ case cuda: __intrinsic_asm "slangOptixHitObjectGetInstanceId";`
`20715`	`20715`	`case spirv:`
`20716`	`20716`	`return spirv_asm`
`20717`	`20717`	`{`
`@@ -20733,7 +20733,7 @@ struct HitObject`
`20733`	`20733`	`{`
`20734`	`20734`	`case hlsl: __intrinsic_asm ".GetGeometryIndex";`
`20735`	`20735`	`case glsl: __intrinsic_asm "hitObjectGetGeometryIndexNV($0)";`
`20736`		`- case cuda: __intrinsic_asm "optixHitObjectGetSbtGASIndex";`
	`20736`	`+ case cuda: __intrinsic_asm "slangOptixHitObjectGetSbtGASIndex";`
`20737`	`20737`	`case spirv:`
`20738`	`20738`	`return spirv_asm`
`20739`	`20739`	`{`
`@@ -20755,7 +20755,7 @@ struct HitObject`
`20755`	`20755`	`{`
`20756`	`20756`	`case hlsl: __intrinsic_asm ".GetPrimitiveIndex";`
`20757`	`20757`	`case glsl: __intrinsic_asm "hitObjectGetPrimitiveIndexNV($0)";`
`20758`		`- case cuda: __intrinsic_asm "optixHitObjectGetPrimitiveIndex";`
	`20758`	`+ case cuda: __intrinsic_asm "slangOptixHitObjectGetPrimitiveIndex";`
`20759`	`20759`	`case spirv:`
`20760`	`20760`	`return spirv_asm`
`20761`	`20761`	`{`
`@@ -20799,7 +20799,7 @@ struct HitObject`
`20799`	`20799`	`{`
`20800`	`20800`	`case hlsl: __intrinsic_asm ".GetClusterID";`
`20801`	`20801`	`case glsl: __intrinsic_asm "hitObjectGetClusterIdNV($0)";`
`20802`		`- case cuda: __intrinsic_asm "optixHitObjectGetClusterId";`
	`20802`	`+ case cuda: __intrinsic_asm "slangOptixHitObjectGetClusterId";`
`20803`	`20803`	`case spirv:`
`20804`	`20804`	`return spirv_asm`
`20805`	`20805`	`{`
Original file line number	Diff line number	Diff line change
`@@ -1,5 +1,6 @@`
`1`	`1`	`//TEST:SIMPLE(filecheck=CHECK): -target cuda`
`2`	`2`	`//CHECK: __global__ void __closesthit__closestHitShaderA`
	`3`	`+//TEST:SIMPLE(filecheck=CHECK-PTX): -target ptx -Xnvrtc -I"./external/optix-dev/include/"`
`3`	`4`	`struct RayPayload`
`4`	`5`	`{`
`5`	`6`	`float4 color;`
`@@ -8,6 +9,7 @@ struct RayPayload`
`8`	`9`	`[shader("closesthit")]`
`9`	`10`	`void closestHitShaderA(inout RayPayload payload, in BuiltInTriangleIntersectionAttributes attr)`
`10`	`11`	`{`
	`12`	`+ //CHECK-PTX: _optix_read_primitive_idx`
`11`	`13`	`uint primitiveIndex = PrimitiveIndex();`
`12`	`14`	`float4 color = float4(0, 0, 0, 1);`
`13`	`15`	`color[primitiveIndex] = 1;`