From f40258cb827d2272625a9960cbc59a74d52d71c0 Mon Sep 17 00:00:00 2001
From: Illia Silin <98187287+illsilin@users.noreply.github.com>
Date: Thu, 11 Dec 2025 08:09:29 -0800
Subject: [PATCH] Fix compilation errors with latest clang22 version. (#3396)

* remove target attributes from deduction guides

* switch CK_TILE_HOST_DEVICE_EXTERN based on clang version

[ROCm/composable_kernel commit: b2925ee207e6179be727ff6e12748d69d0b5df73]
---
 include/ck_tile/core/config.hpp               |  4 ++++
 include/ck_tile/core/numeric/math.hpp         | 21 +++++++------------
 .../core/utility/unary_element_function.hpp   |  3 +--
 3 files changed, 12 insertions(+), 16 deletions(-)

diff --git a/include/ck_tile/core/config.hpp b/include/ck_tile/core/config.hpp
index 678a2fbfff..0e7d1def75 100644
--- a/include/ck_tile/core/config.hpp
+++ b/include/ck_tile/core/config.hpp
@@ -39,8 +39,12 @@
 #define CK_TILE_DEVICE inline __device__
 #define CK_TILE_HOST_DEVICE inline __host__ __device__
 #define CK_TILE_DEVICE_EXTERN __device__
+#if __clang_major__ < 22
 #define CK_TILE_HOST_DEVICE_EXTERN __host__ __device__
 #else
+#define CK_TILE_HOST_DEVICE_EXTERN
+#endif
+#else
 #define CK_TILE_HOST inline
 #define CK_TILE_DEVICE inline
 #define CK_TILE_HOST_DEVICE inline
diff --git a/include/ck_tile/core/numeric/math.hpp b/include/ck_tile/core/numeric/math.hpp
index 57f3953514..8a0e3b3408 100644
--- a/include/ck_tile/core/numeric/math.hpp
+++ b/include/ck_tile/core/numeric/math.hpp
@@ -41,9 +41,8 @@ struct scales
     Scale lhs_;
 };
 
-/// FIXME: create macro to replace '__host__ __device__' and nothing more
 template <typename Scale>
-__host__ __device__ scales(Scale) -> scales<Scale>;
+CK_TILE_HOST_DEVICE_EXTERN scales(Scale) -> scales<Scale>;
 
 template <typename Left = void, typename Right = Left>
 struct plus
@@ -66,8 +65,7 @@ struct plus<void, void>
     }
 };
 
-/// FIXME: create macro to replace '__host__ __device__' and nothing more
-__host__ __device__ plus() -> plus<void, void>;
+CK_TILE_HOST_DEVICE_EXTERN plus() -> plus<void, void>;
 
 template <typename Left = void, typename Right = Left>
 struct minus
@@ -90,8 +88,7 @@ struct minus<void, void>
     }
 };
 
-/// FIXME: create macro to replace '__host__ __device__' and nothing more
-__host__ __device__ minus() -> minus<void, void>;
+CK_TILE_HOST_DEVICE_EXTERN minus() -> minus<void, void>;
 
 template <typename Left = void, typename Right = Left>
 struct multiplies
@@ -114,8 +111,7 @@ struct multiplies<void, void>
     }
 };
 
-/// FIXME: create macro to replace '__host__ __device__' and nothing more
-__host__ __device__ multiplies() -> multiplies<void, void>;
+CK_TILE_HOST_DEVICE_EXTERN multiplies() -> multiplies<void, void>;
 
 template <typename T>
 struct maximize
@@ -345,8 +341,7 @@ struct equal<void, void>
     }
 };
 
-/// FIXME: create macro to replace '__host__ __device__' and nothing more
-__host__ __device__ equal() -> equal<void, void>;
+CK_TILE_HOST_DEVICE_EXTERN equal() -> equal<void, void>;
 
 template <>
 struct equal<float, float>
@@ -387,8 +382,7 @@ struct less<void, void>
     }
 };
 
-/// FIXME: create macro to replace '__host__ __device__' and nothing more
-__host__ __device__ less() -> less<void, void>;
+CK_TILE_HOST_DEVICE_EXTERN less() -> less<void, void>;
 
 template <typename Left = void, typename Right = Left>
 struct less_equal
@@ -411,8 +405,7 @@ struct less_equal<void, void>
     }
 };
 
-/// FIXME: create macro to replace '__host__ __device__' and nothing more
-__host__ __device__ less_equal() -> less_equal<void, void>;
+CK_TILE_HOST_DEVICE_EXTERN less_equal() -> less_equal<void, void>;
 
 template <>
 struct less_equal<float, float>
diff --git a/include/ck_tile/core/utility/unary_element_function.hpp b/include/ck_tile/core/utility/unary_element_function.hpp
index b195275bdc..595b8522da 100644
--- a/include/ck_tile/core/utility/unary_element_function.hpp
+++ b/include/ck_tile/core/utility/unary_element_function.hpp
@@ -47,9 +47,8 @@ struct composes<F>
     F f_;
 };
 
-/// FIXME: create macro to replace '__host__ __device__' and nothing more
 template <typename... Ts>
-__host__ __device__ composes(Ts&&...) -> composes<remove_cvref_t<Ts>...>;
+CK_TILE_HOST_DEVICE_EXTERN composes(Ts&&...) -> composes<remove_cvref_t<Ts>...>;
 
 template <typename SaturateType>
 struct saturates