From d4fecfb0beb1a31c9e109c2c41b4131df0b63c7e Mon Sep 17 00:00:00 2001
From: Arm Patinyasakdikul <arm.patinyasakdikul@amd.com>
Date: Tue, 19 Aug 2025 20:24:31 -0500
Subject: [PATCH] Remove noinline attribute from reduceCopyPacks and (#1864)

reduceCopyPacksWithBias.

[ROCm/rccl commit: fb882e80f6e06d512885ebf7b8231c1be50e0eb2]
---
 projects/rccl/src/device/common_kernel.h | 8 --------
 1 file changed, 8 deletions(-)

diff --git a/projects/rccl/src/device/common_kernel.h b/projects/rccl/src/device/common_kernel.h
index 98113702fc..31dd2b05d3 100644
--- a/projects/rccl/src/device/common_kernel.h
+++ b/projects/rccl/src/device/common_kernel.h
@@ -32,11 +32,7 @@ template<typename RedFn, typename T, int Unroll, int BytePerPack,
          int MultimemSrcs, int MinSrcs, int MaxSrcs,
          int MultimemDsts, int MinDsts, int MaxDsts, int PreOpSrcs,
          typename IntBytes, typename SrcPtrFn, typename DstPtrFn>
-#if defined(__gfx942__) || defined(__gfx950__)
 __device__ __forceinline__ void reduceCopyPacks(
-#else
-__device__ __attribute__((noinline)) void reduceCopyPacks(
-#endif
     int nThreads, int &thread,
     uint64_t redArg, uint64_t *preOpArgs, bool postOp,
     int nSrcs, SrcPtrFn const &srcPtrFn, int nDsts, DstPtrFn const &dstPtrFn,
@@ -428,11 +424,7 @@ template<typename RedFn, typename T, int Unroll, int BytePerPack,
          int MultimemSrcs, int MinSrcs, int MaxSrcs,
          int MultimemDsts, int MinDsts, int MaxDsts, int PreOpSrcs,
          typename IntBytes, typename SrcPtrFn, typename DstPtrFn, typename AccPtrFn>
-#if defined(__gfx942__) || defined(__gfx950__)
 __device__ __forceinline__ void reduceCopyPacksWithBias(
-#else
-__device__ __attribute__((noinline)) void reduceCopyPacksWithBias(
-#endif
     int nThreads, int &thread,
     uint64_t redArg, uint64_t *preOpArgs, bool postOp,
     int nSrcs, SrcPtrFn const &srcPtrFn, int nDsts, DstPtrFn const &dstPtrFn,