From 243d4dacd19aef839f950637b8ed3b4a895086c6 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Tue, 15 Nov 2022 14:42:30 +0000 Subject: [PATCH] SWDEV-366306 - SWDEV-366314 - Support hip vector types and hip_bfloat16 with hiprtc Also, moves the hip_bfloat16 header implemenation from hip to hipamd Change-Id: I9980126fa1ee8160ff9f5f8241f64699a9702c95 [ROCm/clr commit: 5f4aa2e9947a6ba05f706a48eb10ad04d8f721b4] --- .../include/hip/amd_detail/amd_hip_bfloat16.h | 290 ++++++++++++++++++ .../include/hip/amd_detail/host_defines.h | 87 ++++++ 2 files changed, 377 insertions(+) create mode 100644 projects/clr/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h new file mode 100644 index 0000000000..59f8f11b09 --- /dev/null +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h @@ -0,0 +1,290 @@ +/** + * MIT License + * + * Copyright (c) 2019 - 2022 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ + +/*!\file + * \brief hip_bfloat16.h provides struct for hip_bfloat16 typedef + */ + +#ifndef _HIP_INCLUDE_HIP_AMD_DETAIL_HIP_BFLOAT16_H_ +#define _HIP_INCLUDE_HIP_AMD_DETAIL_HIP_BFLOAT16_H_ + +#if defined(__HIPCC_RTC__) + #define __HOST_DEVICE__ __device__ +#else + #define __HOST_DEVICE__ __host__ __device__ +#endif + +#if __cplusplus < 201103L || !defined(__HIPCC__) + +// If this is a C compiler, C++ compiler below C++11, or a host-only compiler, we only +// include a minimal definition of hip_bfloat16 + +#include +/*! \brief Struct to represent a 16 bit brain floating point number. */ +typedef struct +{ + uint16_t data; +} hip_bfloat16; + +#else // __cplusplus < 201103L || !defined(__HIPCC__) + +#include "host_defines.h" + +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wshadow" +struct hip_bfloat16 +{ + __hip_uint16_t data; + + enum truncate_t + { + truncate + }; + + __HOST_DEVICE__ hip_bfloat16() = default; + + // round upper 16 bits of IEEE float to convert to bfloat16 + explicit __HOST_DEVICE__ hip_bfloat16(float f) + : data(float_to_bfloat16(f)) + { + } + + explicit __HOST_DEVICE__ hip_bfloat16(float f, truncate_t) + : data(truncate_float_to_bfloat16(f)) + { + } + + // zero extend lower 16 bits of bfloat16 to convert to IEEE float + __HOST_DEVICE__ operator float() const + { + union + { + uint32_t int32; + float fp32; + } u = {uint32_t(data) << 16}; + return u.fp32; + } + + __HOST_DEVICE__ hip_bfloat16 &operator=(const float& f) + { + data = float_to_bfloat16(f); + return *this; + } + + static __HOST_DEVICE__ hip_bfloat16 round_to_bfloat16(float f) + { + hip_bfloat16 output; + output.data = float_to_bfloat16(f); + return output; + } + + static __HOST_DEVICE__ hip_bfloat16 round_to_bfloat16(float f, truncate_t) + { + hip_bfloat16 output; + output.data = truncate_float_to_bfloat16(f); + return output; + } + +private: + static __HOST_DEVICE__ __hip_uint16_t float_to_bfloat16(float f) + { + union + { + float fp32; + uint32_t int32; + } u = {f}; + if(~u.int32 & 0x7f800000) + { + // When the exponent bits are not all 1s, then the value is zero, normal, + // or subnormal. We round the bfloat16 mantissa up by adding 0x7FFF, plus + // 1 if the least significant bit of the bfloat16 mantissa is 1 (odd). + // This causes the bfloat16's mantissa to be incremented by 1 if the 16 + // least significant bits of the float mantissa are greater than 0x8000, + // or if they are equal to 0x8000 and the least significant bit of the + // bfloat16 mantissa is 1 (odd). This causes it to be rounded to even when + // the lower 16 bits are exactly 0x8000. If the bfloat16 mantissa already + // has the value 0x7f, then incrementing it causes it to become 0x00 and + // the exponent is incremented by one, which is the next higher FP value + // to the unrounded bfloat16 value. When the bfloat16 value is subnormal + // with an exponent of 0x00 and a mantissa of 0x7F, it may be rounded up + // to a normal value with an exponent of 0x01 and a mantissa of 0x00. + // When the bfloat16 value has an exponent of 0xFE and a mantissa of 0x7F, + // incrementing it causes it to become an exponent of 0xFF and a mantissa + // of 0x00, which is Inf, the next higher value to the unrounded value. + u.int32 += 0x7fff + ((u.int32 >> 16) & 1); // Round to nearest, round to even + } + else if(u.int32 & 0xffff) + { + // When all of the exponent bits are 1, the value is Inf or NaN. + // Inf is indicated by a zero mantissa. NaN is indicated by any nonzero + // mantissa bit. Quiet NaN is indicated by the most significant mantissa + // bit being 1. Signaling NaN is indicated by the most significant + // mantissa bit being 0 but some other bit(s) being 1. If any of the + // lower 16 bits of the mantissa are 1, we set the least significant bit + // of the bfloat16 mantissa, in order to preserve signaling NaN in case + // the bloat16's mantissa bits are all 0. + u.int32 |= 0x10000; // Preserve signaling NaN + } + return __hip_uint16_t(u.int32 >> 16); + } + + // Truncate instead of rounding, preserving SNaN + static __HOST_DEVICE__ __hip_uint16_t truncate_float_to_bfloat16(float f) + { + union + { + float fp32; + uint32_t int32; + } u = {f}; + return __hip_uint16_t(u.int32 >> 16) | (!(~u.int32 & 0x7f800000) && (u.int32 & 0xffff)); + } +}; +#pragma clang diagnostic pop + +typedef struct +{ + __hip_uint16_t data; +} hip_bfloat16_public; + +static_assert(__hip_internal::is_standard_layout{}, + "hip_bfloat16 is not a standard layout type, and thus is " + "incompatible with C."); + +static_assert(__hip_internal::is_trivial{}, + "hip_bfloat16 is not a trivial type, and thus is " + "incompatible with C."); +#if !defined(__HIPCC_RTC__) +static_assert(sizeof(hip_bfloat16) == sizeof(hip_bfloat16_public) + && offsetof(hip_bfloat16, data) == offsetof(hip_bfloat16_public, data), + "internal hip_bfloat16 does not match public hip_bfloat16"); +#endif +inline __hip_internal::ostream& operator<<(__hip_internal::ostream& os, const hip_bfloat16& bf16) +{ + return os << bf16; +} +inline __HOST_DEVICE__ hip_bfloat16 operator+(hip_bfloat16 a) +{ + return a; +} +inline __HOST_DEVICE__ hip_bfloat16 operator-(hip_bfloat16 a) +{ + a.data ^= 0x8000; + return a; +} +inline __HOST_DEVICE__ hip_bfloat16 operator+(hip_bfloat16 a, hip_bfloat16 b) +{ + return hip_bfloat16(float(a) + float(b)); +} +inline __HOST_DEVICE__ hip_bfloat16 operator-(hip_bfloat16 a, hip_bfloat16 b) +{ + return hip_bfloat16(float(a) - float(b)); +} +inline __HOST_DEVICE__ hip_bfloat16 operator*(hip_bfloat16 a, hip_bfloat16 b) +{ + return hip_bfloat16(float(a) * float(b)); +} +inline __HOST_DEVICE__ hip_bfloat16 operator/(hip_bfloat16 a, hip_bfloat16 b) +{ + return hip_bfloat16(float(a) / float(b)); +} +inline __HOST_DEVICE__ bool operator<(hip_bfloat16 a, hip_bfloat16 b) +{ + return float(a) < float(b); +} +inline __HOST_DEVICE__ bool operator==(hip_bfloat16 a, hip_bfloat16 b) +{ + return float(a) == float(b); +} +inline __HOST_DEVICE__ bool operator>(hip_bfloat16 a, hip_bfloat16 b) +{ + return b < a; +} +inline __HOST_DEVICE__ bool operator<=(hip_bfloat16 a, hip_bfloat16 b) +{ + return !(a > b); +} +inline __HOST_DEVICE__ bool operator!=(hip_bfloat16 a, hip_bfloat16 b) +{ + return !(a == b); +} +inline __HOST_DEVICE__ bool operator>=(hip_bfloat16 a, hip_bfloat16 b) +{ + return !(a < b); +} +inline __HOST_DEVICE__ hip_bfloat16& operator+=(hip_bfloat16& a, hip_bfloat16 b) +{ + return a = a + b; +} +inline __HOST_DEVICE__ hip_bfloat16& operator-=(hip_bfloat16& a, hip_bfloat16 b) +{ + return a = a - b; +} +inline __HOST_DEVICE__ hip_bfloat16& operator*=(hip_bfloat16& a, hip_bfloat16 b) +{ + return a = a * b; +} +inline __HOST_DEVICE__ hip_bfloat16& operator/=(hip_bfloat16& a, hip_bfloat16 b) +{ + return a = a / b; +} +inline __HOST_DEVICE__ hip_bfloat16& operator++(hip_bfloat16& a) +{ + return a += hip_bfloat16(1.0f); +} +inline __HOST_DEVICE__ hip_bfloat16& operator--(hip_bfloat16& a) +{ + return a -= hip_bfloat16(1.0f); +} +inline __HOST_DEVICE__ hip_bfloat16 operator++(hip_bfloat16& a, int) +{ + hip_bfloat16 orig = a; + ++a; + return orig; +} +inline __HOST_DEVICE__ hip_bfloat16 operator--(hip_bfloat16& a, int) +{ + hip_bfloat16 orig = a; + --a; + return orig; +} + +namespace std +{ + constexpr __HOST_DEVICE__ bool isinf(hip_bfloat16 a) + { + return !(~a.data & 0x7f80) && !(a.data & 0x7f); + } + constexpr __HOST_DEVICE__ bool isnan(hip_bfloat16 a) + { + return !(~a.data & 0x7f80) && +(a.data & 0x7f); + } + constexpr __HOST_DEVICE__ bool iszero(hip_bfloat16 a) + { + return !(a.data & 0x7fff); + } +} + +#endif // __cplusplus < 201103L || !defined(__HIPCC__) + +#endif // _HIP_BFLOAT16_H_ diff --git a/projects/clr/hipamd/include/hip/amd_detail/host_defines.h b/projects/clr/hipamd/include/hip/amd_detail/host_defines.h index 5a89433c04..e66dc7a4cc 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/host_defines.h +++ b/projects/clr/hipamd/include/hip/amd_detail/host_defines.h @@ -39,6 +39,93 @@ THE SOFTWARE. #if defined(__clang__) && defined(__HIP__) +namespace __hip_internal { +typedef unsigned short uint16_t; +template struct integral_constant { + static constexpr const _Tp value = __v; + typedef _Tp value_type; + typedef integral_constant type; + constexpr operator value_type() const { return value; } + constexpr value_type operator()() const { return value; } +}; +template constexpr const _Tp integral_constant<_Tp, __v>::value; + +typedef integral_constant true_type; +typedef integral_constant false_type; + +template using bool_constant = integral_constant; +typedef bool_constant true_type; +typedef bool_constant false_type; + +template struct enable_if {}; +template struct enable_if { typedef __T type; }; + +template struct true_or_false_type : public false_type {}; +template<> struct true_or_false_type : public true_type {}; + +template struct is_integral : public false_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; + +template struct is_arithmetic : public false_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; + +template struct is_floating_point : public false_type {}; +template<> struct is_floating_point : public true_type {}; +template<> struct is_floating_point : public true_type {}; +template<> struct is_floating_point : public true_type {}; + +template struct is_same : public false_type {}; +template struct is_same<__T, __T> : public true_type {}; + +template::value> + struct is_signed : public false_type {}; +template + struct is_signed<_Tp, true> : public true_or_false_type<_Tp(-1) < _Tp(0)> {}; + +template struct char_traits; +template> class basic_istream; +template> class basic_ostream; +typedef basic_istream istream; +typedef basic_ostream ostream; + +template + struct is_standard_layout + : public integral_constant + { }; + +template + struct is_trivial + : public integral_constant + { }; +} +typedef __hip_internal::uint16_t __hip_uint16_t; + #if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #define __host__ __attribute__((host)) #define __device__ __attribute__((device))