From 477e028e58e016788df9837915e01cfa1cc3cacc Mon Sep 17 00:00:00 2001
From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com>
Date: Thu, 12 Dec 2024 19:47:57 +0000
Subject: [PATCH] refactor conditional usage; fix build on rocm6.1 where the
 reference didn't exist

[ROCm/composable_kernel commit: 6ef8d3c295686b872d7e7a86621b68f765d98572]
---
 include/ck/utility/amd_ck_fp8.hpp | 14 ++++++++++----
 1 file changed, 10 insertions(+), 4 deletions(-)

diff --git a/include/ck/utility/amd_ck_fp8.hpp b/include/ck/utility/amd_ck_fp8.hpp
index 7b21ad6464..1bdb1d078e 100644
--- a/include/ck/utility/amd_ck_fp8.hpp
+++ b/include/ck/utility/amd_ck_fp8.hpp
@@ -18,6 +18,12 @@
 #define CK_USE_OCP_FP8 0
 #endif
 
+namespace {
+// https://en.cppreference.com/w/cpp/types/conditional
+template <bool B, class T, class F> struct conditional { using type = T; };
+template <class T, class F> struct conditional<false, T, F> { using type = F; };
+}
+
 namespace ck {
 
 using f8_fnuz_t  = _BitInt(8);
@@ -191,10 +197,10 @@ __host__ __device__ static inline T cast_from_f8(fp8_storage_t x)
         }
     }
 
-    typename __hip_internal::conditional<
+    typename conditional<
         sizeof(T) == 2,
         unsigned short int,
-        typename __hip_internal::conditional<sizeof(T) == 4, unsigned int, unsigned long long>::
+        typename conditional<sizeof(T) == 4, unsigned int, unsigned long long>::
             type>::type retval;
 
     if constexpr(we == 5 && is_half && !is_fnuz)
@@ -538,10 +544,10 @@ __host__ __device__ static inline fp8_storage_t cast_to_f8(T _x, unsigned int rn
 
     constexpr int mfmt = (sizeof(T) == 8) ? 52 : ((sizeof(T) == 4) ? 23 : 10);
 
-    using T_bitwise = typename __hip_internal::conditional<
+    using T_bitwise = typename conditional<
         sizeof(T) == 2,
         unsigned short int,
-        typename __hip_internal::conditional<sizeof(T) == 4, unsigned int, unsigned long long>::
+        typename conditional<sizeof(T) == 4, unsigned int, unsigned long long>::
             type>::type;
     T_bitwise x_bitwise = bit_cast<T_bitwise>(_x);