From c3b8de4ec8b4cb4efc926f313eb71a7f5b29852a Mon Sep 17 00:00:00 2001
From: Nilesh M Negi <Nilesh.Negi@amd.com>
Date: Fri, 15 Aug 2025 15:15:02 -0500
Subject: [PATCH] [DEVICE] Use noinline for LLGenericOp only on gfx950 (#1849)

---
 src/device/prims_ll.h | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/src/device/prims_ll.h b/src/device/prims_ll.h
index 2709df2dfc..0995730fd3 100644
--- a/src/device/prims_ll.h
+++ b/src/device/prims_ll.h
@@ -432,7 +432,11 @@ private:
   }
 
   template <int RECV, int SEND, int SrcBuf, int DstBuf>
+#if defined(__gfx950__)
   __device__ __attribute__((noinline)) void LLGenericOp(intptr_t srcIx, intptr_t dstIx, int nelem, bool postOp) {
+#else
+  __device__ void LLGenericOp(intptr_t srcIx, intptr_t dstIx, int nelem, bool postOp) {
+#endif
     constexpr int SRC = SrcBuf != -1 ? 1 : 0;
     constexpr int DST = DstBuf != -1 ? 1 : 0;
     T *srcElts = SrcBuf == -1 ? nullptr : userBufs[SrcBuf] + srcIx;