refactor conditional usage; fix build on rocm6.1 where the reference didn't exist

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);