From 154c00969c2e5cc18700220df7e16f855e5bea8f Mon Sep 17 00:00:00 2001
From: zhangyue <zhangyue@qiyuanlab.com>
Date: Tue, 24 Mar 2026 15:59:46 +0800
Subject: [PATCH] fix: mv `#include <musa_fp16.h>` ahead of `polyfills.cuh`

---
 src/common/cuda/kernel_commons.h | 1 +
 src/moore/add/kernel.h           | 8 --------
 src/moore/polyfills.cuh          | 8 ++------
 src/moore/swiglu/kernel.h        | 8 --------
 4 files changed, 3 insertions(+), 22 deletions(-)

diff --git a/src/common/cuda/kernel_commons.h b/src/common/cuda/kernel_commons.h
index e2ef107..be96b20 100644
--- a/src/common/cuda/kernel_commons.h
+++ b/src/common/cuda/kernel_commons.h
@@ -21,6 +21,7 @@ using cuda_bfloat162 = maca_bfloat162;
 #include <musa_bf16.h>
 #include <musa_fp16.h>
 #include <musa_runtime.h>
+#include "moore/polyfills.cuh"
 using cuda_bfloat16 = __mt_bfloat16;
 using cuda_bfloat162 = __mt_bfloat162;
 #endif
diff --git a/src/moore/add/kernel.h b/src/moore/add/kernel.h
index 21a51f6..ae08dbd 100644
--- a/src/moore/add/kernel.h
+++ b/src/moore/add/kernel.h
@@ -3,14 +3,6 @@
 
 #include <utility>
 
-// clang-format off
-#include <musa_runtime.h>
-// clang-format on
-
-// clang-format off
-#include "moore/polyfills.cuh"
-// clang-format on
-
 #include "cuda/add/kernel.h"
 
 namespace infini::ops {
diff --git a/src/moore/polyfills.cuh b/src/moore/polyfills.cuh
index b3c7e70..65dee5e 100644
--- a/src/moore/polyfills.cuh
+++ b/src/moore/polyfills.cuh
@@ -1,12 +1,6 @@
 #ifndef INFINI_OPS_MOORE_POLYFILLS_CUH_
 #define INFINI_OPS_MOORE_POLYFILLS_CUH_
 
-#include <type_traits>
-
-// clang-format off
-#include <musa_bf16.h>
-// clang-format on
-
 namespace infini::ops {
 
 template <typename T>
@@ -36,6 +30,8 @@ __device__ __forceinline__ T hrcp(const T& a) {
 
 }  // namespace infini::ops
 
+// Force unqualified `hrcp` calls to resolve to our polyfill, bypassing ADL
+// which would otherwise pick the MUSA SDK's declared-but-undefined hrcp(__half).
 #define hrcp infini::ops::hrcp
 
 #endif
diff --git a/src/moore/swiglu/kernel.h b/src/moore/swiglu/kernel.h
index a7759fb..b721dcd 100644
--- a/src/moore/swiglu/kernel.h
+++ b/src/moore/swiglu/kernel.h
@@ -3,14 +3,6 @@
 
 #include <utility>
 
-// clang-format off
-#include <musa_runtime.h>
-// clang-format on
-
-// clang-format off
-#include "moore/polyfills.cuh"
-// clang-format on
-
 #include "cuda/swiglu/kernel.h"
 
 namespace infini::ops {