From b3b4aa8d57e5635c7b11fc29f93d1639efa3f7b5 Mon Sep 17 00:00:00 2001 From: Xiao Li Date: Tue, 24 Jun 2025 21:46:15 -0700 Subject: [PATCH] Fix amd_ck_fp8.hpp macro definitions (#2325) * Fix amd_ck_fp8.hpp macro definitions 1. Define CK_USE_FNUZ_FP8 and CK_USE_OCP_FP8 definitions only if they were not defined before. 2. Prefix __assert_fnuz_support and __assert_ocp_support with namespace fp8_impl to avoid redefined error when building with rocm 6.4+ (rocm/6.4.0/include/hip/amd_detail/amd_hip_fp8.h) Co-authored-by: Andriy Roshchenko [ROCm/composable_kernel commit: bac51b6ec0d8e3e5f333a42af999cc097306a394] --- include/ck/utility/amd_ck_fp8.hpp | 16 ++++++---------- 1 file changed, 6 insertions(+), 10 deletions(-) diff --git a/include/ck/utility/amd_ck_fp8.hpp b/include/ck/utility/amd_ck_fp8.hpp index cdc2a4fbda..b7af32d3dc 100644 --- a/include/ck/utility/amd_ck_fp8.hpp +++ b/include/ck/utility/amd_ck_fp8.hpp @@ -10,15 +10,11 @@ #include "ck/utility/functional.hpp" #include "ck/utility/type.hpp" -#ifdef CK_USE_FNUZ_FP8 -#define CK_USE_FNUZ_FP8 1 -#else +#ifndef CK_USE_FNUZ_FP8 #define CK_USE_FNUZ_FP8 0 #endif -#ifdef CK_USE_OCP_FP8 -#define CK_USE_OCP_FP8 1 -#else +#ifndef CK_USE_OCP_FP8 #define CK_USE_OCP_FP8 0 #endif @@ -432,7 +428,7 @@ __host__ __device__ inline constexpr bool fp8_is_inf(bf8_ocp_t a) namespace fp8_impl { // Assertions to check for supported conversion types -#define __assert_ocp_support(interp) \ +#define __fp8_impl_assert_ocp_support(interp) \ { \ if(interp != ck_fp8_interpretation_t::CK_E4M3_OCP && \ interp != ck_fp8_interpretation_t::CK_E5M2_OCP) \ @@ -440,7 +436,7 @@ namespace fp8_impl { __hip_assert(false && "type is unsupported by current target device"); \ } \ } -#define __assert_fnuz_support(interp) \ +#define __fp8_impl_assert_fnuz_support(interp) \ { \ if(interp != ck_fp8_interpretation_t::CK_E4M3_FNUZ && \ interp != ck_fp8_interpretation_t::CK_E5M2_FNUZ) \ @@ -454,10 +450,10 @@ __is_interpret_supported([[maybe_unused]] ck_fp8_interpretation_t interp) { #if defined(__HIP_DEVICE_COMPILE__) && __HIP_DEVICE_COMPILE__ #if CK_USE_OCP_FP8 - __assert_ocp_support(interp); + __fp8_impl_assert_ocp_support(interp); #endif #if CK_USE_FNUZ_FP8 - __assert_fnuz_support(interp); + __fp8_impl_assert_fnuz_support(interp); #endif #endif }