summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--docs/command-line-slangc-reference.md3
-rw-r--r--docs/user-guide/a3-02-reference-capability-atoms.md6
-rw-r--r--prelude/slang-cuda-prelude.h49
-rw-r--r--source/slang/hlsl.meta.slang208
-rw-r--r--source/slang/slang-capabilities.capdef12
-rw-r--r--tests/cuda/lss-test.slang34
-rw-r--r--tests/hlsl-intrinsic/ray-tracing/rt-lss-intrinsics-chit.slang45
-rw-r--r--tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-ray.slang17
8 files changed, 374 insertions, 0 deletions
diff --git a/docs/command-line-slangc-reference.md b/docs/command-line-slangc-reference.md
index 4ecaac546..b976d3ed3 100644
--- a/docs/command-line-slangc-reference.md
+++ b/docs/command-line-slangc-reference.md
@@ -1128,6 +1128,7 @@ A capability describes an optional feature that a target may or may not support.
* `SPV_NV_ray_tracing_motion_blur` : enables the SPV_NV_ray_tracing_motion_blur extension
* `SPV_NV_shader_invocation_reorder` : enables the SPV_NV_shader_invocation_reorder extension
* `SPV_NV_cluster_acceleration_structure` : enables the SPV_NV_cluster_acceleration_structure extension
+* `SPV_NV_linear_swept_spheres` : enables the SPV_NV_linear_swept_spheres extension
* `SPV_NV_shader_image_footprint` : enables the SPV_NV_shader_image_footprint extension
* `SPV_KHR_compute_shader_derivatives` : enables the SPV_KHR_compute_shader_derivatives extension
* `SPV_GOOGLE_user_type` : enables the SPV_GOOGLE_user_type extension
@@ -1168,6 +1169,7 @@ A capability describes an optional feature that a target may or may not support.
* `spvRayQueryPositionFetchKHR`
* `spvShaderInvocationReorderNV`
* `spvRayTracingClusterAccelerationStructureNV`
+* `spvRayTracingLinearSweptSpheresGeometryNV`
* `spvShaderClockKHR`
* `spvShaderNonUniformEXT`
* `spvShaderNonUniform`
@@ -1496,6 +1498,7 @@ A capability describes an optional feature that a target may or may not support.
* `raytracing_anyhit`
* `raytracing_intersection`
* `raytracing_anyhit_closesthit`
+* `raytracing_lss`
* `raytracing_anyhit_closesthit_intersection`
* `raytracing_raygen_closesthit_miss`
* `raytracing_anyhit_closesthit_intersection_miss`
diff --git a/docs/user-guide/a3-02-reference-capability-atoms.md b/docs/user-guide/a3-02-reference-capability-atoms.md
index 0766fdf82..78affcb7d 100644
--- a/docs/user-guide/a3-02-reference-capability-atoms.md
+++ b/docs/user-guide/a3-02-reference-capability-atoms.md
@@ -432,6 +432,9 @@ Extensions
`SPV_NV_cluster_acceleration_structure`
> Represents the SPIR-V extension for cluster acceleration structure.
+`SPV_NV_linear_swept_spheres`
+> Represents the SPIR-V extension for linear swept spheres.
+
`SPV_NV_shader_image_footprint`
> Represents the SPIR-V extension for shader image footprint.
@@ -555,6 +558,9 @@ Extensions
`spvRayTracingClusterAccelerationStructureNV`
> Represents the SPIR-V capability for cluster acceleration structure.
+`spvRayTracingLinearSweptSpheresGeometryNV`
+> Represents the SPIR-V capability for linear swept spheres.
+
`spvShaderClockKHR`
> Represents the SPIR-V capability for shader clock.
diff --git a/prelude/slang-cuda-prelude.h b/prelude/slang-cuda-prelude.h
index d2c9fce9d..d1160cdd3 100644
--- a/prelude/slang-cuda-prelude.h
+++ b/prelude/slang-cuda-prelude.h
@@ -31,6 +31,7 @@
#ifdef SLANG_CUDA_ENABLE_OPTIX
#include <optix.h>
+#include <optix_device.h>
#endif
// Define slang offsetof implementation
@@ -3221,6 +3222,54 @@ __forceinline__ __device__ void* optixTrace(
r1);
}
+__forceinline__ __device__ float4 optixGetSpherePositionAndRadius()
+{
+ float4 data[1];
+ optixGetSphereData(data);
+ return data;
+}
+
+__forceinline__ __device__ float4 optixHitObjectGetSpherePositionAndRadius()
+{
+ float4 data[1];
+ optixHitObjectGetSphereData(data);
+ return data;
+}
+
+__forceinline__ __device__ Matrix<float, 2, 4> optixGetSpherePositionAndRadius()
+{
+ float4 data[2];
+ optixGetLinearCurveVertexData(data);
+ return Matrix<float, 2, 4>(data[0], data[1]);
+}
+
+__forceinline__ __device__ float2x4 optixHitObjectGetSpherePositionAndRadius()
+{
+ float4 data[2];
+ optixHitObjectGetLinearCurveVertexData(data);
+ return Matrix<float, 2, 4>(data[0], data[1]);
+}
+
+__forceinline__ __device__ bool optixIsSphereHit()
+{
+ return optixGetPrimitiveType() == OPTIX_PRIMITIVE_TYPE_SPHERE;
+}
+
+__forceinline__ __device__ bool optixHitObjectIsSphereHit()
+{
+ return optixGetPrimitiveType(optixHitObjectGetHitKind()) == OPTIX_PRIMITIVE_TYPE_SPHERE;
+}
+
+__forceinline__ __device__ bool optixIsLSSHit()
+{
+ return optixGetPrimitiveType() == OPTIX_PRIMITIVE_TYPE_ROUND_LINEAR;
+}
+
+__forceinline__ __device__ bool optixHitObjectIsLSSHit()
+{
+ return optixGetPrimitiveType(optixHitObjectGetHitKind()) == OPTIX_PRIMITIVE_TYPE_ROUND_LINEAR;
+}
+
template<typename T>
__forceinline__ __device__ void* optixTraverse(
OptixTraversableHandle AccelerationStructure,
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
index fd7c7cfc7..9b15392ec 100644
--- a/source/slang/hlsl.meta.slang
+++ b/source/slang/hlsl.meta.slang
@@ -17948,6 +17948,114 @@ int GetClusterID()
};
}
}
+
+/// @category raytracing
+[__requiresNVAPI]
+[NonUniformReturn]
+[require(cuda_hlsl_spirv, raytracing_lss)]
+float4 GetSpherePositionAndRadius()
+{
+ __target_switch
+ {
+ case hlsl: __intrinsic_asm "NvRtSphereObjectPositionAndRadius";
+ case cuda:
+ {
+ __intrinsic_asm "optixGetSpherePositionAndRadius";
+ }
+ case spirv:
+ return spirv_asm
+ {
+ OpExtension "SPV_NV_linear_swept_spheres";
+ OpCapability RayTracingLinearSweptSpheresGeometryNV;
+ OpCapability RayTracingSpheresGeometryNV;
+ %pos:$$float3 = OpLoad builtin(HitSpherePositionNV:float3);
+ %rad:$$float = OpLoad builtin(HitSphereRadiusNV:float);
+ result:$$float4 = OpCompositeConstruct %pos %rad;
+ };
+ }
+}
+
+/// @category raytracing
+[__requiresNVAPI]
+[NonUniformReturn]
+[require(cuda_hlsl_spirv, raytracing_lss)]
+float2x4 GetLssPositionsAndRadii()
+{
+ __target_switch
+ {
+ case hlsl: __intrinsic_asm "NvRtLssObjectPositionsAndRadii";
+ case cuda:
+ {
+ __intrinsic_asm "optixObjectPositionsAndRadii";
+ }
+ case spirv:
+ return spirv_asm
+ {
+ OpExtension "SPV_NV_linear_swept_spheres";
+ OpCapability RayTracingLinearSweptSpheresGeometryNV;
+ OpCapability RayTracingSpheresGeometryNV;
+ %positions:$$float3[2] = OpLoad builtin(HitLSSPositionsNV:float3[2]);
+ %radii:$$float[2] = OpLoad builtin(HitLSSRadiiNV:float[2]);
+ %r0:$$float = OpCompositeExtract %radii 0;
+ %r1:$$float = OpCompositeExtract %radii 1;
+ %p0:$$float3 = OpCompositeExtract %positions 0;
+ %p1:$$float3 = OpCompositeExtract %positions 1;
+ %a:$$float4 = OpCompositeConstruct %p0 %r0;
+ %b:$$float4 = OpCompositeConstruct %p1 %r1;
+ result:$$float2x4 = OpCompositeConstruct %a %b;
+ };
+
+ }
+}
+
+/// @category raytracing
+[__requiresNVAPI]
+[NonUniformReturn]
+[require(cuda_hlsl_spirv, raytracing_lss)]
+bool IsSphereHit()
+{
+ __target_switch
+ {
+ case hlsl: __intrinsic_asm "NvRtIsSphereHit";
+ case cuda:
+ {
+ __intrinsic_asm "optixIsSphereHit";
+ }
+ case spirv:
+ return spirv_asm
+ {
+ OpExtension "SPV_NV_linear_swept_spheres";
+ OpCapability RayTracingLinearSweptSpheresGeometryNV;
+ OpCapability RayTracingSpheresGeometryNV;
+ result:$$bool = OpLoad builtin(HitIsSphereNV:bool);
+ };
+ }
+}
+
+/// @category raytracing
+[__requiresNVAPI]
+[NonUniformReturn]
+[require(cuda_hlsl_spirv, raytracing_lss)]
+bool IsLssHit()
+{
+ __target_switch
+ {
+ case hlsl: __intrinsic_asm "NvRtIsLssHit";
+ case cuda:
+ {
+ __intrinsic_asm "optixIsLSSHit";
+ }
+ case spirv:
+ return spirv_asm
+ {
+ OpExtension "SPV_NV_linear_swept_spheres";
+ OpCapability RayTracingLinearSweptSpheresGeometryNV;
+ OpCapability RayTracingSpheresGeometryNV;
+ result:$$bool = OpLoad builtin(HitIsLSSNV:bool);
+ };
+ }
+}
+
// Note: The provisional DXR spec included these unadorned
// `ObjectToWorld()` and `WorldToObject()` functions, so
// we will forward them to the new names as a convience
@@ -20529,6 +20637,106 @@ struct HitObject
}
[__requiresNVAPI]
+ [NonUniformReturn]
+ [require(cuda_hlsl_spirv, raytracing_lss)]
+ float4 GetSpherePositionAndRadius()
+ {
+ __target_switch
+ {
+ case hlsl: __intrinsic_asm "NvRtSphereObjectPositionAndRadius";
+ case cuda:
+ {
+ __intrinsic_asm "optixHitObjectGetSpherePositionAndRadius";
+ }
+ case spirv:
+ return spirv_asm
+ {
+ OpExtension "SPV_NV_linear_swept_spheres";
+ OpCapability RayTracingLinearSweptSpheresGeometryNV;
+ OpCapability RayTracingSpheresGeometryNV;
+ %position:$$float3 = OpHitObjectGetSpherePositionNV &this;
+ %radius:$$float = OpHitObjectGetSphereRadiusNV &this;
+ result:$$float4 = OpCompositeConstruct %position %radius;
+ };
+ }
+ }
+
+ [__requiresNVAPI]
+ [NonUniformReturn]
+ [require(cuda_hlsl_spirv, raytracing_lss)]
+ float2x4 GetLssPositionsAndRadii()
+ {
+ __target_switch
+ {
+ case hlsl: __intrinsic_asm "NvRtLssObjectPositionsAndRadii";
+ case cuda:
+ {
+ __intrinsic_asm "optixHitObjectGetSpherePositionAndRadius";
+ }
+ case spirv:
+ return spirv_asm
+ {
+ OpExtension "SPV_NV_linear_swept_spheres";
+ OpCapability RayTracingLinearSweptSpheresGeometryNV;
+ OpCapability RayTracingSpheresGeometryNV;
+ %positions:$$float3[2] = OpHitObjectGetLSSPositionsNV &this;
+ %radii:$$float[2] = OpHitObjectGetLSSRadiiNV &this;
+ %r0:$$float = OpCompositeExtract %radii 0;
+ %r1:$$float = OpCompositeExtract %radii 1;
+ %p0:$$float3 = OpCompositeExtract %positions 0;
+ %p1:$$float3 = OpCompositeExtract %positions 1;
+ %a:$$float4 = OpCompositeConstruct %p0 %r0;
+ %b:$$float4 = OpCompositeConstruct %p1 %r1;
+ result:$$float2x4 = OpCompositeConstruct %a %b;
+ };
+ }
+ }
+
+ [__requiresNVAPI]
+ [NonUniformReturn]
+ [require(cuda_hlsl_spirv, raytracing_lss)]
+ bool IsSphereHit()
+ {
+ __target_switch
+ {
+ case hlsl: __intrinsic_asm "NvRtIsSphereHit";
+ case cuda:
+ {
+ __intrinsic_asm "optixHitObjectIsSphereHit";
+ }
+ case spirv:
+ return spirv_asm
+ {
+ OpExtension "SPV_NV_linear_swept_spheres";
+ OpCapability RayTracingLinearSweptSpheresGeometryNV;
+ result:$$bool = OpHitObjectIsSphereHitNV &this;
+ };
+ }
+ }
+
+ [__requiresNVAPI]
+ [NonUniformReturn]
+ [require(cuda_hlsl_spirv, raytracing_lss)]
+ bool IsLssHit()
+ {
+ __target_switch
+ {
+ case hlsl: __intrinsic_asm "NvRtIsLssHit";
+ case cuda:
+ {
+ __intrinsic_asm "optixHitObjectIsLSSHit";
+ }
+ case spirv:
+ return spirv_asm
+ {
+ OpExtension "SPV_NV_linear_swept_spheres";
+ OpCapability RayTracingLinearSweptSpheresGeometryNV;
+ result:$$bool = OpHitObjectIsLSSHitNV &this;
+ };
+ }
+ }
+
+ [__requiresNVAPI]
__glsl_extension(GL_EXT_ray_tracing)
[ForceInline]
[require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)]
diff --git a/source/slang/slang-capabilities.capdef b/source/slang/slang-capabilities.capdef
index 48617c54d..28fa211e4 100644
--- a/source/slang/slang-capabilities.capdef
+++ b/source/slang/slang-capabilities.capdef
@@ -547,6 +547,10 @@ def SPV_NV_shader_invocation_reorder : _spirv_1_5 + SPV_KHR_ray_tracing;
/// [EXT]
def SPV_NV_cluster_acceleration_structure : _spirv_1_0;
+/// Represents the SPIR-V extension for linear swept spheres.
+/// [EXT]
+def SPV_NV_linear_swept_spheres : _spirv_1_0;
+
/// Represents the SPIR-V extension for shader image footprint.
/// [EXT]
def SPV_NV_shader_image_footprint : _spirv_1_0;
@@ -714,6 +718,10 @@ def spvShaderInvocationReorderNV : SPV_NV_shader_invocation_reorder;
/// [EXT]
def spvRayTracingClusterAccelerationStructureNV : SPV_NV_cluster_acceleration_structure;
+/// Represents the SPIR-V capability for cluster acceleration structure.
+/// [EXT]
+def spvRayTracingLinearSweptSpheresGeometryNV : SPV_NV_linear_swept_spheres;
+
/// Represents the SPIR-V capability for shader clock.
/// [EXT]
def spvShaderClockKHR : SPV_KHR_shader_clock;
@@ -2246,6 +2254,10 @@ alias raytracing_intersection = intersection + raytracing;
/// [Compound]
alias raytracing_anyhit_closesthit = anyhit_closesthit + raytracing;
+/// Collection of capabilities for linear swept spheres.
+/// [Compound]
+alias raytracing_lss = raytracing_anyhit_closesthit | spvRayTracingLinearSweptSpheresGeometryNV;
+
/// Collection of capabilities for raytracing with the shader stages of anyhit, closesthit, and intersection.
/// [Compound]
alias raytracing_anyhit_closesthit_intersection = anyhit_closesthit_intersection + raytracing;
diff --git a/tests/cuda/lss-test.slang b/tests/cuda/lss-test.slang
new file mode 100644
index 000000000..4b0512cb1
--- /dev/null
+++ b/tests/cuda/lss-test.slang
@@ -0,0 +1,34 @@
+//TEST:SIMPLE(filecheck=CHECK): -target cuda
+//CHECK_: __global__ void __closesthit__closestHitShaderLss
+//CHECK: optixGetSpherePositionAndRadius
+//CHECK: optixObjectPositionsAndRadii
+//CHECK: optixIsSphereHit
+//CHECK: optixIsLSSHit
+
+struct RayPayload
+{
+ float4 color;
+ float2x4 lssData;
+ bool isSphere;
+ bool isLss;
+};
+
+[shader("closesthit")]
+void closestHitShaderLss(inout RayPayload payload, in BuiltInTriangleIntersectionAttributes attr)
+{
+ // Test TraceRay API functions
+ payload.color = GetSpherePositionAndRadius();
+ payload.lssData = GetLssPositionsAndRadii();
+ payload.isSphere = IsSphereHit();
+ payload.isLss = IsLssHit();
+
+// TODO: This will be enabled once issue #6647 is completed.
+#if 0
+ // Test HitObject API functions
+ HitObject hitObj;
+ float4 sphereData = hitObj.GetSphereObjectPositionAndRadius();
+ float2x4 lssData = hitObj.GetLssObjectPositionsAndRadii();
+ bool isSphereHit = hitObj.IsSphereHit();
+ bool isLssHit = hitObj.IsLssHit();
+#endif
+} \ No newline at end of file
diff --git a/tests/hlsl-intrinsic/ray-tracing/rt-lss-intrinsics-chit.slang b/tests/hlsl-intrinsic/ray-tracing/rt-lss-intrinsics-chit.slang
new file mode 100644
index 000000000..b4d181c7d
--- /dev/null
+++ b/tests/hlsl-intrinsic/ray-tracing/rt-lss-intrinsics-chit.slang
@@ -0,0 +1,45 @@
+//TEST:SIMPLE(filecheck=HLSL): -target hlsl
+//TEST:SIMPLE(filecheck=SPIRV): -target spirv-asm
+//HLSL: NvRtSphereObjectPositionAndRadius
+//HLSL: NvRtLssObjectPositionsAndRadii
+//HLSL: NvRtIsSphereHit
+//HLSL: NvRtIsLssHit
+
+//SPIRV: HitSpherePositionNV
+//SPIRV: HitSphereRadiusNV
+//SPIRV: HitLSSPositionsNV
+//SPIRV: HitIsSphereNV
+//SPIRV: HitIsLSSNV
+
+// Hit object variants
+//SPIRV: OpHitObjectGetSpherePositionNV
+//SPIRV: OpHitObjectGetSphereRadiusNV
+//SPIRV: OpHitObjectGetLSSPositionsNV
+//SPIRV: OpHitObjectGetLSSRadiiNV
+//SPIRV: OpHitObjectIsSphereHitNV
+//SPIRV: OpHitObjectIsLSSHitNV
+
+struct RayPayload
+{
+ float4 color;
+ float2x4 lssData;
+ bool isSphere;
+ bool isLss;
+};
+
+[shader("closesthit")]
+void closestHitShaderLss(inout RayPayload payload, in BuiltInTriangleIntersectionAttributes attr)
+{
+ // Test TraceRay API functions
+ payload.color = GetSpherePositionAndRadius();
+ payload.lssData = GetLssPositionsAndRadii();
+ payload.isSphere = IsSphereHit();
+ payload.isLss = IsLssHit();
+
+ // Test HitObject API functions
+ HitObject hitObj;
+ float4 sphereData = hitObj.GetSpherePositionAndRadius();
+ float2x4 lssData = hitObj.GetLssPositionsAndRadii();
+ bool isSphereHit = hitObj.IsSphereHit();
+ bool isLssHit = hitObj.IsLssHit();
+} \ No newline at end of file
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 71c113934..be38915b1 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
@@ -32,6 +32,10 @@ uint calcValue(HitObject hit)
uint geometryIndex = hit.GetGeometryIndex();
uint primitiveIndex = hit.GetPrimitiveIndex();
int clusterID = hit.GetClusterID();
+ float4 posRadius = hit.GetSpherePositionAndRadius();
+ float2x4 positionsRadii = hit.GetLssPositionsAndRadii();
+ uint isSphereHit = uint(hit.IsSphereHit());
+ uint isLssHit = uint(hit.IsLssHit());
SomeValues objSomeValues = hit.GetAttributes<SomeValues>();
@@ -41,6 +45,19 @@ uint calcValue(HitObject hit)
r += primitiveIndex;
r += objSomeValues.a;
r += clusterID;
+
+ r += int(posRadius.x);
+ r += int(posRadius.y);
+ r += int(posRadius.z);
+ r += int(posRadius.w);
+
+ r += int(positionsRadii[0].x);
+ r += int(positionsRadii[0].y);
+ r += int(positionsRadii[0].z);
+ r += int(positionsRadii[0].w);
+
+ r += isSphereHit;
+ r += isLssHit;
}
return r;
generated by cgit v1.2.3 (git 2.39.1) at 2026-06-01 05:43:54 +0000