From d9d0b4f03277027690a909a76b78d1622ac13498 Mon Sep 17 00:00:00 2001
From: Yong He <yonghe@outlook.com>
Date: Wed, 15 Jan 2025 15:49:29 -0800
Subject: Fix optix varying legalization. (#6089)

* Fix optix varying legalization.

* Add test.
---
 source/slang/slang-ir-legalize-varying-params.cpp |  2 ++
 tests/cuda/optix-hit-attributes.slang             | 15 +++++++++++++++
 2 files changed, 17 insertions(+)
 create mode 100644 tests/cuda/optix-hit-attributes.slang

diff --git a/source/slang/slang-ir-legalize-varying-params.cpp b/source/slang/slang-ir-legalize-varying-params.cpp
index 5a18b533a..a98290545 100644
--- a/source/slang/slang-ir-legalize-varying-params.cpp
+++ b/source/slang/slang-ir-legalize-varying-params.cpp
@@ -1053,6 +1053,8 @@ struct CUDAEntryPointVaryingParamLegalizeContext : EntryPointVaryingParamLegaliz
         IRType* typeToFetch,
         IRBuilder* builder)
     {
+        if (auto ptrValType = tryGetPointedToType(builder, typeToFetch))
+            typeToFetch = ptrValType;
         if (auto structType = as<IRStructType>(typeToFetch))
         {
             List<IRInst*> fieldVals;
diff --git a/tests/cuda/optix-hit-attributes.slang b/tests/cuda/optix-hit-attributes.slang
new file mode 100644
index 000000000..347401893
--- /dev/null
+++ b/tests/cuda/optix-hit-attributes.slang
@@ -0,0 +1,15 @@
+//TEST:SIMPLE(filecheck=CHECK): -target cuda
+//CHECK: __global__ void __closesthit__closestHitShaderA
+struct RayPayload
+{
+    float4 color;
+};
+
+[shader("closesthit")]
+void closestHitShaderA(inout RayPayload payload, in BuiltInTriangleIntersectionAttributes attr)
+{
+    uint primitiveIndex = PrimitiveIndex();
+    float4 color = float4(0, 0, 0, 1);
+    color[primitiveIndex] = 1;
+    payload.color = color;
+}
-- 
cgit v1.2.3