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

7 files changed, 257 insertions, 1 deletions

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;