cluster acceleration structure optix 6431 (#7028)

* Add cluster geometry intrinsics for ray tracing

- Added GetClusterID() method to HitObject class
- Added CandidateClusterID() and CommittedClusterID() methods to
RayQuery class
- Added SPV_NV_cluster_acceleration_structure extension support
- Added GL_NV_cluster_acceleration_structure extension support
- Added test files for RayQuery and HitObject cluster methods

Fixes #6431

* OpRayQueryGetIntersectionClusterIdNV - unrecognized spirv
Disabling spirv backend for SPV_NV_cluster_acceleration_structure

hlsl.meta.slang(18674): error 29100: unrecognized spirv opcode:
OpRayQueryGetIntersectionClusterIdNV
                result:$$int = OpRayQueryGetIntersectionClusterIdNV
&this $iCandidateOrCommitted;
                ^~~~~~
hlsl.meta.slang(18670): error 30019: expected an expression of type
'int', got 'void'
            return spirv_asm
                   ^~~~~~~~~
ninja: build stopped: subcommand failed.

* 6431 - Fix spirv opcode

* Remove tests

* Add relevant tests

* Review - Simplify tests

12 files changed, 409 insertions, 2 deletions

diff --git a/docs/user-guide/a3-02-reference-capability-atoms.md b/docs/user-guide/a3-02-reference-capability-atoms.md
index b6aa7aba2..f0345679c 100644
--- a/docs/user-guide/a3-02-reference-capability-atoms.md
+++ b/docs/user-guide/a3-02-reference-capability-atoms.md
@@ -423,6 +423,9 @@ Extensions
 > Represents the SPIR-V extension for shader invocation reorder.
 > Requires SPV_KHR_ray_tracing.
 
+`SPV_NV_cluster_acceleration_structure`
+> Represents the SPIR-V extension for cluster acceleration structure.
+
 `SPV_NV_shader_image_footprint`
 > Represents the SPIR-V extension for shader image footprint.
 
@@ -534,6 +537,9 @@ Extensions
 `spvShaderInvocationReorderNV`
 > Represents the SPIR-V capability for shader invocation reorder.
 
+`spvRayTracingClusterAccelerationStructureNV`
+> Represents the SPIR-V capability for cluster acceleration structure.
+
 `spvShaderClockKHR`
 > Represents the SPIR-V capability for shader clock.
 
@@ -741,6 +747,9 @@ Extensions
 `GL_NV_shader_texture_footprint`
 > Represents the GL_NV_shader_texture_footprint extension.
 
+`GL_NV_cluster_acceleration_structure`
+> Represents the GL_NV_cluster_acceleration_structure extension.
+
 Compound Capabilities
 ----------------------
 *Capabilities to specify capabilities created by other capabilities (`raytracing`, `meshshading`...)*
diff --git a/source/slang/glsl.meta.slang b/source/slang/glsl.meta.slang
index 00ad88add..05cef836c 100644
--- a/source/slang/glsl.meta.slang
+++ b/source/slang/glsl.meta.slang
@@ -4893,6 +4893,35 @@ public property uint gl_HitKindEXT
     }
 }
 
+public property int gl_ClusterIDNV
+{
+    [require(glsl_spirv, raytracing_anyhit_closesthit)]
+    get
+    {
+        setupExtForRayTracingBuiltIn();
+        __target_switch
+        {
+        case glsl:
+        {
+            __requireTargetExtension("GL_NV_cluster_acceleration_structure");
+            __intrinsic_asm "(gl_ClusterIDNV)";
+        }
+        case spirv:
+        {
+            return spirv_asm
+            {
+                OpCapability RayTracingClusterAccelerationStructureNV;
+                OpExtension "SPV_NV_cluster_acceleration_structure";
+                result:$$int = OpLoad builtin(ClusterIDNV:int);
+            };
+        }
+        }
+    }
+}
+// Constant to indicate that a cluster acceleration structure was not hit.
+// Corresponds to VK_GEOMETRY_INSTANCE_CLUSTER_ID_NONE_NV in Vulkan
+public static const int gl_ClusterIDNoneNV = -1;
+
 public property mat4x3 gl_ObjectToWorldEXT 
 {
     [require(glsl_spirv, raytracing_anyhit_closesthit_intersection)]
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
index 99eba0a42..34423d4f3 100644
--- a/source/slang/hlsl.meta.slang
+++ b/source/slang/hlsl.meta.slang
@@ -17889,6 +17889,26 @@ float RayCurrentTime()
     }
 }
 
+/// @category raytracing
+[__requiresNVAPI]
+__glsl_extension(GL_NV_cluster_acceleration_structure)
+__glsl_extension(GL_EXT_ray_tracing)
+[NonUniformReturn]
+[require(cuda_glsl_hlsl_spirv, raytracing_anyhit_closesthit)]
+int GetClusterID()
+{
+    __target_switch
+    {
+    case hlsl:  __intrinsic_asm "NvRtGetClusterID";
+    case glsl:  __intrinsic_asm "(gl_ClusterIDNV)";
+    case cuda: __intrinsic_asm "optixGetClusterId";
+    case spirv:
+        return spirv_asm
+        {
+            result:$$int = OpLoad builtin(ClusterIDNV:int);
+        };
+    }
+}
 // Note: The provisional DXR spec included these unadorned
 // `ObjectToWorld()` and `WorldToObject()` functions, so
 // we will forward them to the new names as a convience
@@ -18683,6 +18703,31 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE>
         }
     }
 
+    /// Gets the cluster ID of the candidate hit.
+    /// @return Cluster ID for the candidate hit
+    /// @remarks HLSL/GLSL/SPIRV
+    [__requiresNVAPI]
+    __glsl_extension(GL_NV_cluster_acceleration_structure)
+    __glsl_extension(GL_EXT_ray_query)
+    [__NoSideEffect]
+    [require(glsl_hlsl_spirv, rayquery)]
+    int CandidateClusterID()
+    {
+        __target_switch
+        {
+        case hlsl: __intrinsic_asm "NvRtGetCandidateClusterID";
+        case glsl: __intrinsic_asm "rayQueryGetIntersectionClusterIdNV($0, false)";
+        case spirv:
+            uint iCandidateOrCommitted = 0;
+            return spirv_asm
+            {
+                OpExtension "SPV_NV_cluster_acceleration_structure";
+                OpCapability RayTracingClusterAccelerationStructureNV;
+                result:$$int = OpRayQueryGetClusterIdNV &this $iCandidateOrCommitted;
+            };
+        }
+    }
+
     /// Gets the custom index of the instance containing the committed hit.
     /// @return User-provided instance identifier
     /// @remarks GLSL/SPIRV only
@@ -18698,7 +18743,32 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE>
             uint iCandidateOrCommitted = 1;
             return spirv_asm
             {
-                result:$$int = OpRayQueryGetIntersectionInstanceCustomIndexKHR &this $iCandidateOrCommitted;
+                result:$$int = OpRayQueryGetClusterIdNV &this $iCandidateOrCommitted;
+            };
+        }
+    }
+
+    /// Gets the cluster ID of the committed hit.
+    /// @return Cluster ID for the committed hit
+    /// @remarks HLSL/GLSL/SPIRV
+    [__requiresNVAPI]
+    __glsl_extension(GL_NV_cluster_acceleration_structure)
+    __glsl_extension(GL_EXT_ray_query)
+    [__NoSideEffect]
+    [require(glsl_hlsl_spirv, rayquery)]
+    int CommittedClusterID()
+    {
+        __target_switch
+        {
+        case hlsl: __intrinsic_asm "NvRtGetCommittedClusterID";
+        case glsl: __intrinsic_asm "rayQueryGetIntersectionClusterIdNV($0, true)";
+        case spirv:
+            uint iCandidateOrCommitted = 1; // 1 for committed
+            return spirv_asm
+            {
+                OpExtension "SPV_NV_cluster_acceleration_structure";
+                OpCapability RayTracingClusterAccelerationStructureNV;
+                result:$$int = OpRayQueryGetClusterIdNV &this $iCandidateOrCommitted;
             };
         }
     }
@@ -20354,6 +20424,29 @@ struct HitObject
         }
     }
 
+    /// Returns the cluster ID of the current hit. Valid if the hit object represents a hit.
+    [__requiresNVAPI]
+    __glsl_extension(GL_NV_cluster_acceleration_structure)
+    __glsl_extension(GL_EXT_ray_tracing)
+    [ForceInline]
+    [require(cuda_glsl_hlsl_spirv, ser_raygen_closesthit_miss)]
+    int GetClusterID()
+    {
+        __target_switch
+        {
+        case hlsl: __intrinsic_asm ".GetClusterID";
+        case glsl: __intrinsic_asm "hitObjectGetClusterIdNV($0)";
+        case cuda: __intrinsic_asm "optixHitObjectGetClusterId";
+        case spirv:
+            return spirv_asm
+            {
+                OpExtension "SPV_NV_cluster_acceleration_structure";
+                OpCapability RayTracingClusterAccelerationStructureNV;
+                result:$$int = OpHitObjectGetClusterIdNV &this;
+            };
+        }
+    }
+
     [__requiresNVAPI]
     __glsl_extension(GL_EXT_ray_tracing)
     [ForceInline]
diff --git a/source/slang/slang-capabilities.capdef b/source/slang/slang-capabilities.capdef
index b50519a72..fafcb214f 100644
--- a/source/slang/slang-capabilities.capdef
+++ b/source/slang/slang-capabilities.capdef
@@ -539,6 +539,10 @@ def SPV_NV_ray_tracing_motion_blur : _spirv_1_0;
 /// [EXT]
 def SPV_NV_shader_invocation_reorder : _spirv_1_5 + SPV_KHR_ray_tracing;
 
+/// Represents the SPIR-V extension for cluster acceleration structure.
+/// [EXT]
+def SPV_NV_cluster_acceleration_structure : _spirv_1_0;
+
 /// Represents the SPIR-V extension for shader image footprint.
 /// [EXT]
 def SPV_NV_shader_image_footprint : _spirv_1_0;
@@ -689,6 +693,10 @@ def spvRayQueryPositionFetchKHR : SPV_KHR_ray_tracing_position_fetch + spvRayQue
 /// [EXT]
 def spvShaderInvocationReorderNV : SPV_NV_shader_invocation_reorder;
 
+/// Represents the SPIR-V capability for cluster acceleration structure.
+/// [EXT]
+def spvRayTracingClusterAccelerationStructureNV : SPV_NV_cluster_acceleration_structure;
+
 /// Represents the SPIR-V capability for shader clock.
 /// [EXT]
 def spvShaderClockKHR : SPV_KHR_shader_clock;
@@ -809,6 +817,7 @@ def _GL_NV_shader_atomic_fp16_vector : _GL_NV_gpu_shader5;
 def _GL_NV_shader_invocation_reorder : _GLSL_460;
 def _GL_NV_shader_subgroup_partitioned : _GLSL_140;
 def _GL_NV_shader_texture_footprint : _GLSL_450;
+def _GL_NV_cluster_acceleration_structure : _GLSL_460;
 
 // GLSL extension and SPV extension associations.
 
@@ -1044,6 +1053,10 @@ alias GL_NV_shader_subgroup_partitioned = _GL_NV_shader_subgroup_partitioned | s
 /// [EXT]
 alias GL_NV_shader_texture_footprint = _GL_NV_shader_texture_footprint | spvImageFootprintNV;
 
+/// Represents the GL_NV_cluster_acceleration_structure extension.
+/// [EXT]
+alias GL_NV_cluster_acceleration_structure = _GL_NV_cluster_acceleration_structure | spvRayTracingClusterAccelerationStructureNV;
+
 // Define feature names not reliant on shader stages
 
 /// NVAPI capability for HLSL
diff --git a/source/slang/slang-emit-spirv.cpp b/source/slang/slang-emit-spirv.cpp
index bd9b23b2d..096e7d8bc 100644
--- a/source/slang/slang-emit-spirv.cpp
+++ b/source/slang/slang-emit-spirv.cpp
@@ -1433,6 +1433,13 @@ struct SPIRVEmitContext : public SourceEmitterBase, public SPIRVEmitSharedContex
         return ensureExtensionDeclaration(name);
     }
 
+    // Ensure cluster acceleration structure extensions and capabilities are declared
+    void requireRayTracingClusterAccelerationStructure()
+    {
+        requireSPIRVCapability(SpvCapabilityRayTracingClusterAccelerationStructureNV);
+        ensureExtensionDeclaration(UnownedStringSlice("SPV_NV_cluster_acceleration_structure"));
+    }
+
     bool hasExtensionDeclaration(const UnownedStringSlice& name)
     {
         return m_extensionInsts.containsKey(name);
diff --git a/tests/cuda/optix-cluster.slang b/tests/cuda/optix-cluster.slang
new file mode 100644
index 000000000..151b9c7dd
--- /dev/null
+++ b/tests/cuda/optix-cluster.slang
@@ -0,0 +1,17 @@
+//TEST:SIMPLE(filecheck=CHECK): -target cuda
+//CHECK: __global__ void __closesthit__closestHitShaderA
+//CHECK: optixGetClusterId
+struct RayPayload
+{
+    float4 color;
+};
+
+[shader("closesthit")]
+void closestHitShaderA(inout RayPayload payload, in BuiltInTriangleIntersectionAttributes attr)
+{
+    int clusterId = GetClusterID();
+    float4 color = float4(0, 0, 0, 1);
+    if (clusterId >= 0)
+        color[clusterId] = 1;
+    payload.color = color;
+}
diff --git a/tests/glsl-intrinsic/raytracing/glsl-rayAnyhit-cluster.slang b/tests/glsl-intrinsic/raytracing/glsl-rayAnyhit-cluster.slang
new file mode 100644
index 000000000..4f5d1c05b
--- /dev/null
+++ b/tests/glsl-intrinsic/raytracing/glsl-rayAnyhit-cluster.slang
@@ -0,0 +1,34 @@
+//TEST:SIMPLE(filecheck=CHECK_GLSL): -allow-glsl -stage anyhit -entry main -target glsl -capability GL_EXT_ray_tracing -capability GL_NV_cluster_acceleration_structure
+
+//TEST:SIMPLE(filecheck=CHECK_SPV): -allow-glsl -emit-spirv-directly -stage anyhit -entry main -target spirv-assembly -capability SPV_KHR_ray_tracing -capability SPV_NV_cluster_acceleration_structure
+
+// CHECK_GLSL: #extension GL_EXT_ray_tracing : require
+// CHECK_GLSL: #extension GL_NV_cluster_acceleration_structure : require
+__glsl_extension(GL_EXT_ray_tracing)
+__glsl_extension(GL_NV_cluster_acceleration_structure)
+
+layout(binding = 0) uniform accelerationStructureEXT as;
+
+//TEST_INPUT:ubuffer(data=[0], stride=4):out,name=outBuffer
+//TEST_INPUT:end
+
+// This test demonstrates the usage of cluster acceleration structure intrinsics
+// in ray tracing for the NVIDIA extension.
+
+void main()
+{
+    // CHECK_GLSL-DAG: gl_ClusterIDNV
+    // CHECK_SPV-DAG: ClusterIDNV
+
+    // Get the cluster ID of the current hit
+    int clusterId = gl_ClusterIDNV;
+
+    // Check if the hit is part of a cluster
+    bool isClusterHit = (clusterId != gl_ClusterIDNoneNV);
+
+    // Skip hits for certain cluster IDs (example: skip cluster ID 5)
+    if (isClusterHit && clusterId == 5)
+    {
+        ignoreIntersectionEXT;
+    }
+}
diff --git a/tests/glsl-intrinsic/raytracing/ray-query-cluster-intrinsics.slang b/tests/glsl-intrinsic/raytracing/ray-query-cluster-intrinsics.slang
new file mode 100644
index 000000000..b2c7d25b0
--- /dev/null
+++ b/tests/glsl-intrinsic/raytracing/ray-query-cluster-intrinsics.slang
@@ -0,0 +1,80 @@
+//TEST:SIMPLE(filecheck=CHECK_GLSL): -entry computeMain -target glsl -profile cs_6_5 
+//TEST:SIMPLE(filecheck=CHECK_SPV): -entry computeMain -target spirv-assembly -profile cs_6_5 
+
+//CHECK_GLSL:rayQueryGetIntersectionClusterIdNV 
+//CHECK_SPV:OpRayQueryGetClusterIdNV  
+
+uniform RaytracingAccelerationStructure accelStruct;
+
+float GetRayT<let RAY_QUERY_FLAGS: uint>(uint rayInlineFlags)
+{
+    RayDesc ray;
+    ray.Origin = float3(0.1f, 0.1f, 0.0f);
+    ray.Direction = float3(0.0f, 0.0f, 1.0f);
+    ray.TMin = 0.01f;
+    ray.TMax = 1e4f;
+
+    RayQuery<RAY_QUERY_FLAGS> rq;
+    rq.TraceRayInline(accelStruct, rayInlineFlags, 0xff, ray);
+    bool proceed = rq.Proceed();
+
+    if( proceed )
+    {
+        switch( rq.CandidateType() )
+        {
+            case CANDIDATE_NON_OPAQUE_TRIANGLE:
+                int candidateClusterId = rq.CandidateClusterID();
+                if (candidateClusterId >= 0)
+                {
+                    rq.CommitNonOpaqueTriangleHit();
+                }
+                rq.Abort();
+                return rq.CommittedRayT();
+
+            case CANDIDATE_PROCEDURAL_PRIMITIVE:
+                rq.CommitProceduralPrimitiveHit(0.5f);
+                rq.Abort();
+                return rq.CommittedRayT();
+
+            default:
+                rq.Abort();
+                return 0.0f;
+        }
+    }
+    else 
+    {
+        if( rq.CommittedStatus() == COMMITTED_TRIANGLE_HIT )
+        {  
+            rq.Abort();
+            return rq.CommittedRayT();
+        }
+    }
+    
+    int committedClusterId = rq.CommittedClusterID();
+    ray.TMin = 0.01f + committedClusterId;
+    return 0.0f;
+}
+
+//TEST_INPUT: ubuffer(data=[0], stride=4):out,name outputBuffer
+RWStructuredBuffer<float> outputBuffer;
+
+[shader("compute")]
+[numthreads(1, 1, 1)]
+void computeMain(int3 dispatchThreadID: SV_DispatchThreadID)
+{
+    int idx = dispatchThreadID.x;
+    float val = 0.0f;
+
+    const uint rayFlags[] = {
+        RAY_FLAG_NONE,
+    };
+    const uint numRayFlags = sizeof(rayFlags) / sizeof(uint);
+
+    // RAY_FLAG_FORCE_NON_OPAQUE, RAY_FLAG_CULL_OPAQUE, RAY_FLAG_SKIP_TRIANGLES returns 0 instead of 0.5
+    for( uint i = 0; i < numRayFlags; ++i )
+    {
+        val += GetRayT<RAY_FLAG_NONE>(rayFlags[i]);
+    }
+
+    outputBuffer[idx] = val;
+}
diff --git a/tests/hlsl-intrinsic/ray-tracing/ray-query-cluster-intrinsics.slang b/tests/hlsl-intrinsic/ray-tracing/ray-query-cluster-intrinsics.slang
new file mode 100644
index 000000000..49acfe878
--- /dev/null
+++ b/tests/hlsl-intrinsic/ray-tracing/ray-query-cluster-intrinsics.slang
@@ -0,0 +1,79 @@
+//TEST:SIMPLE(filecheck=CHECK_HLSL): -entry computeMain -target hlsl -profile cs_6_5 
+
+//CHECK_HLSL: NvRtGetCandidateClusterID
+//CHECK_HLSL: NvRtGetCommittedClusterID
+
+uniform RaytracingAccelerationStructure accelStruct;
+
+float GetRayT<let RAY_QUERY_FLAGS: uint>(uint rayInlineFlags)
+{
+    RayDesc ray;
+    ray.Origin = float3(0.1f, 0.1f, 0.0f);
+    ray.Direction = float3(0.0f, 0.0f, 1.0f);
+    ray.TMin = 0.01f;
+    ray.TMax = 1e4f;
+
+    RayQuery<RAY_QUERY_FLAGS> rq;
+    rq.TraceRayInline(accelStruct, rayInlineFlags, 0xff, ray);
+    bool proceed = rq.Proceed();
+
+    if( proceed )
+    {
+        switch( rq.CandidateType() )
+        {
+            case CANDIDATE_NON_OPAQUE_TRIANGLE:
+                int candidateClusterId = rq.CandidateClusterID();
+                if (candidateClusterId >= 0)
+                {
+                    rq.CommitNonOpaqueTriangleHit();
+                }
+                rq.Abort();
+                return rq.CommittedRayT();
+
+            case CANDIDATE_PROCEDURAL_PRIMITIVE:
+                rq.CommitProceduralPrimitiveHit(0.5f);
+                rq.Abort();
+                return rq.CommittedRayT();
+
+            default:
+                rq.Abort();
+                return 0.0f;
+        }
+    }
+    else 
+    {
+        if( rq.CommittedStatus() == COMMITTED_TRIANGLE_HIT )
+        {  
+            rq.Abort();
+            return rq.CommittedRayT();
+        }
+    }
+    
+    int committedClusterId = rq.CommittedClusterID();
+    ray.TMin = 0.01f + committedClusterId;
+    return 0.0f;
+}
+
+//TEST_INPUT: ubuffer(data=[0], stride=4):out,name outputBuffer
+RWStructuredBuffer<float> outputBuffer;
+
+[shader("compute")]
+[numthreads(1, 1, 1)]
+void computeMain(int3 dispatchThreadID: SV_DispatchThreadID)
+{
+    int idx = dispatchThreadID.x;
+    float val = 0.0f;
+
+    const uint rayFlags[] = {
+        RAY_FLAG_NONE,
+    };
+    const uint numRayFlags = sizeof(rayFlags) / sizeof(uint);
+
+    // RAY_FLAG_FORCE_NON_OPAQUE, RAY_FLAG_CULL_OPAQUE, RAY_FLAG_SKIP_TRIANGLES returns 0 instead of 0.5
+    for( uint i = 0; i < numRayFlags; ++i )
+    {
+        val += GetRayT<RAY_FLAG_NONE>(rayFlags[i]);
+    }
+
+    outputBuffer[idx] = val;
+}
diff --git a/tests/hlsl-intrinsic/ray-tracing/ray-query-intrinsics.slang b/tests/hlsl-intrinsic/ray-tracing/ray-query-intrinsics.slang
index c3b251053..6744b87e7 100644
--- a/tests/hlsl-intrinsic/ray-tracing/ray-query-intrinsics.slang
+++ b/tests/hlsl-intrinsic/ray-tracing/ray-query-intrinsics.slang
@@ -249,4 +249,4 @@ void computeMain(int3 dispatchThreadID: SV_DispatchThreadID)
 }
 
 // DX12: 88.5
-// VK: 88.5
\ No newline at end of file
+// VK: 88.5
diff --git a/tests/hlsl-intrinsic/ray-tracing/rt-cluster-intrinsics-chit.slang b/tests/hlsl-intrinsic/ray-tracing/rt-cluster-intrinsics-chit.slang
new file mode 100644
index 000000000..37e6aaf40
--- /dev/null
+++ b/tests/hlsl-intrinsic/ray-tracing/rt-cluster-intrinsics-chit.slang
@@ -0,0 +1,41 @@
+//TEST:SIMPLE(filecheck=DXIL):-target dxil-assembly -entry main -stage closesthit -profile sm_6_5 -DNV_SHADER_EXTN_SLOT=u0
+
+// DXIL: main
+
+// DXIL call void @dx.op.rawBufferStore.i32(i32 140, %dx.types.Handle %80, i32 %79, i32 0, i32 93, i32 undef, i32 undef, i32 undef, i8 1, i32 4)  ; RawBufferStore(uav,index,elementOffset,value0,value1,value2,value3,mask,alignment)
+
+[[vk::binding(0)]]
+uniform RaytracingAccelerationStructure accelStruct : register(t0);
+
+struct RayPayload
+{
+    float RayHitT;
+};
+
+struct CallableParams
+{
+    float value;
+};
+
+#define T_MIN 0.01f
+#define T_MAX 1e4f
+
+[shader("closesthit")]
+void main( inout RayPayload payload, in BuiltInTriangleIntersectionAttributes attribs )
+{
+    float2 dir = (DispatchRaysIndex().xy / DispatchRaysDimensions().xy) * 2.0f - 1.0f;
+    float aspectRatio = DispatchRaysDimensions().x / DispatchRaysDimensions().y;
+
+    RayDesc rayDesc;
+    rayDesc.Origin = float3(0.0f, 0.0f, 0.0f);
+    rayDesc.Direction = normalize(float3(dir.x * aspectRatio, -dir.y, 1)); ;
+    rayDesc.TMin = T_MIN;
+    rayDesc.TMax = T_MAX;
+
+    //RayPayload payload;
+    payload.RayHitT = T_MAX;
+
+    val += GetClusterID();
+
+    payload.RayHitT = val;
+}
diff --git a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-ray.slang b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-ray.slang
index 0f60e3cf5..877e41977 100644
--- a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-ray.slang
+++ b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-ray.slang
@@ -4,6 +4,9 @@
 //TEST:SIMPLE(filecheck=SPIRV): -target spirv -entry rayGenerationMain -stage raygeneration -profile sm_6_5 -line-directive-mode none 
 //TEST:SIMPLE(filecheck=SPIRV): -target spirv -entry rayGenerationMain -stage raygeneration -emit-spirv-directly
 
+// Note: HitObject::TraceRay is not supported in raygen stage for cuda target
+//DISABLE_TEST:SIMPLE: -target cuda -entry rayGenerationMain -stage raygeneration
+
 //DISABLE_TEST(compute):COMPARE_COMPUTE:-d3d12 -output-using-type -use-dxil -profile sm_6_6 -render-feature ray-query
 //DISABLE_TEST(compute):COMPARE_COMPUTE:-vk -output-using-type -render-feature ray-query
 
@@ -29,6 +32,7 @@ uint calcValue(HitObject hit)
         uint instanceID = hit.GetInstanceID();
         uint geometryIndex = hit.GetGeometryIndex();
         uint primitiveIndex = hit.GetPrimitiveIndex();
+        int clusterID = hit.GetClusterID();
 
         SomeValues objSomeValues = hit.GetAttributes<SomeValues>();
         
@@ -37,6 +41,7 @@ uint calcValue(HitObject hit)
         r += geometryIndex;
         r += primitiveIndex;
         r += objSomeValues.a;
+        r += clusterID;
     }
     
     return r;