diff --git a/src/include/rccl_float8.h b/src/include/rccl_float8.h index 19304c74b0..5db894394a 100644 --- a/src/include/rccl_float8.h +++ b/src/include/rccl_float8.h @@ -24,8 +24,9 @@ #define ROCBLAS_FLOAT8_H #include +#include -#if __cplusplus < 201103L || (!defined(__HCC__) && !defined(__HIPCC__)) +#if __cplusplus < 201103L || (!defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__)) /*! \brief Struct to represent a 8 bit floating-point number. */ typedef struct @@ -38,7 +39,55 @@ typedef struct uint8_t data; } rccl_bfloat8; -#else // __cplusplus < 201103L || (!defined(__HCC__) && !defined(__HIPCC__)) +// __cplusplus < 201103L || (!defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__)) +#elif HIP_VERSION >= 60200000 + +#include + +#if __HIP_DEVICE_COMPILE__ && (defined(__gfx950__) || defined(__gfx1200__) || defined(__gfx1201__) || (defined(__gfx1100__) || defined(__gfx1101__)))//HIP_FP8_TYPE_OCP is enabled. +typedef __hip_fp8_e4m3 rccl_float8; +typedef __hip_fp8_e5m2 rccl_bfloat8; +#elif __HIP_DEVICE_COMPILE__ && (defined(__gfx942__)) +typedef __hip_fp8_e4m3_fnuz rccl_float8; +typedef __hip_fp8_e5m2_fnuz rccl_bfloat8; +#else +typedef __hip_fp8_e4m3 rccl_float8; +typedef __hip_fp8_e5m2 rccl_bfloat8; +#endif + +inline std::ostream& operator<<(std::ostream& os, const rccl_float8& f8) +{ + return os << float(f8); +} + +inline std::ostream& operator<<(std::ostream& os, const rccl_bfloat8& bf8) +{ + return os << float(bf8); +} + +inline __host__ __device__ float operator*(rccl_float8 a, rccl_float8 b) +{ + return float(a) * float(b); +} + +inline __host__ __device__ float operator*(rccl_bfloat8 a, rccl_bfloat8 b) +{ + return float(a) * float(b); +} + +inline __host__ __device__ float operator*(rccl_float8 a, float b) +{ + return float(a) * float(b); +} + +inline __host__ __device__ float operator*(rccl_bfloat8 a, float b) +{ + return float(a) * float(b); +} + +// For older versions of ROCm that do not include hip_fp8.h, +// we provide a local version of the header file as a fallback. +#else #define HIP_HOST_DEVICE __host__ __device__ #define HIP_HOST __host__ @@ -1016,6 +1065,6 @@ inline __host__ __device__ T explicit_downcast(Ta a, uint32_t rng) // ================================================================================================= -#endif // __cplusplus < 201103L || (!defined(__HCC__) && !defined(__HIPCC__)) +#endif #endif // ROCBLAS_FLOAT8_H diff --git a/test/common/StandaloneUtils.hpp b/test/common/StandaloneUtils.hpp index 8bb3347c09..119ab6e1e9 100644 --- a/test/common/StandaloneUtils.hpp +++ b/test/common/StandaloneUtils.hpp @@ -26,7 +26,7 @@ } while(0) // should be 112, temp fix to make CI pass -#define MAX_STACK_SIZE 480 +#define MAX_STACK_SIZE 496 #ifdef ENABLE_LL128 #define MAX_STACK_SIZE_gfx90a 360