From f570e109c4039e15526af38e17f350c115327489 Mon Sep 17 00:00:00 2001
From: Mukund Keshava <mkeshava@nvidia.com>
Date: Tue, 27 May 2025 17:46:04 +0530
Subject: Add LSS intrinsics (#7200)

* WiP: LSS intrinsics: initial commit

* format code

* Fix CI failures

* Address review comment

---------

Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
---
 prelude/slang-cuda-prelude.h | 49 ++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 49 insertions(+)

(limited to 'prelude')

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,
-- 
cgit v1.2.3