From 64a43b17a76d7a8269f55cbe67ef708cd464774c Mon Sep 17 00:00:00 2001 From: "Harsh Aggarwal (NVIDIA)" Date: Mon, 12 May 2025 13:05:34 +0530 Subject: 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 --- tests/cuda/optix-cluster.slang | 17 +++++ .../raytracing/glsl-rayAnyhit-cluster.slang | 34 +++++++++ .../raytracing/ray-query-cluster-intrinsics.slang | 80 ++++++++++++++++++++++ .../ray-tracing/ray-query-cluster-intrinsics.slang | 79 +++++++++++++++++++++ .../ray-tracing/ray-query-intrinsics.slang | 2 +- .../ray-tracing/rt-cluster-intrinsics-chit.slang | 41 +++++++++++ .../hit-object-trace-ray.slang | 5 ++ 7 files changed, 257 insertions(+), 1 deletion(-) create mode 100644 tests/cuda/optix-cluster.slang create mode 100644 tests/glsl-intrinsic/raytracing/glsl-rayAnyhit-cluster.slang create mode 100644 tests/glsl-intrinsic/raytracing/ray-query-cluster-intrinsics.slang create mode 100644 tests/hlsl-intrinsic/ray-tracing/ray-query-cluster-intrinsics.slang create mode 100644 tests/hlsl-intrinsic/ray-tracing/rt-cluster-intrinsics-chit.slang (limited to 'tests') 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(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 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 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(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(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 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 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(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(); @@ -37,6 +41,7 @@ uint calcValue(HitObject hit) r += geometryIndex; r += primitiveIndex; r += objSomeValues.a; + r += clusterID; } return r; -- cgit v1.2.3