ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-14 10:09:41 +00:00
Code Issues Packages Projects Releases Wiki Activity

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

Browse Source 
[ROCm/composable_kernel commit: 6ef8d3c295]
This commit is contained in:
Max Podkorytov
2024-12-12 19:47:57 +00:00
parent b0cd5fea51
commit 477e028e58
1 changed files with 10 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

14
include/ck/utility/amd_ck_fp8.hpp
View File

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