Prevent copy-elision pass for CudaDeviceExport (#8941)

The copy-elision optimization pass was incorrectly transforming struct
parameters in functions marked with [CudaDeviceExport] from
pass-by-value to pass-by-pointer, breaking the externally visible
function signatures that clients depend on.

This fix adds CudaDeviceExportDecoration to the exclusion list in
shouldProcessFunction(), ensuring that exported CUDA device functions
preserve their original signatures.

Fixes: #8874

Author: jkwak-work

source/slang/slang-ir-transform-params-to-constref.cpp

@@ -222,6 +222,11 @@ struct TransformParamsToConstRefContext
             if (as<IREntryPointDecoration>(decoration) || as<IRCudaKernelDecoration>(decoration) ||
                 as<IRAutoPyBindCudaDecoration>(decoration))
                 return false;
+
+            // Skip functions with CudaDeviceExport decoration.
+            // These functions have externally visible signatures that should not be changed.
+            if (func->findDecorationImpl(kIROp_CudaDeviceExportDecoration))
+                return false;
         }
 
         // Skip functions with `kIROp_GenericAsm` since

tests/autodiff/cuda-device-export-struct-param.slang

@@ -0,0 +1,27 @@
+//TEST:SIMPLE(filecheck=CUDA): -target cuda -line-directive-mode none
+
+// Verify that struct parameters in [CudaDeviceExport] functions are passed by value,
+// not transformed to pointers by the copy-elision pass.
+// This is a regression test for issue #8874.
+
+struct CommonParameters
+{
+    int id;
+    float value;
+}
+
+struct Parameters
+{
+    float3 position;
+    float scale;
+}
+
+// CUDA: __device__ Parameters_[[#]] processParameters(uint idx_[[#]], CommonParameters_[[#]] commonParams_[[#]])
+[CudaDeviceExport]
+Parameters processParameters(uint idx, CommonParameters commonParams)
+{
+    Parameters result;
+    result.position = float3(commonParams.value, commonParams.id, 0.0f);
+    result.scale = commonParams.value + float(idx);
+    return result;
+}