File: rocblas_bfloat16.h

package info (click to toggle)
rocblas 5.5.1%2Bdfsg-7
  • links: PTS, VCS
  • area: main
  • in suites: forky, sid, trixie
  • size: 565,372 kB
  • sloc: cpp: 198,491; python: 44,792; f90: 25,111; sh: 24,429; asm: 8,954; xml: 222; makefile: 147; ansic: 107; awk: 14
file content (296 lines) | stat: -rw-r--r-- 10,325 bytes parent folder | download
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
/* ************************************************************************
 * Copyright (C) 2016-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 cop-
 * ies 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 IM-
 * PLIED, 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 CONNE-
 * CTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
 *
 * ************************************************************************ */

/*!\file
 * \brief rocblas_bfloat16.h provides struct for rocblas_bfloat16 typedef
 */

#ifndef ROCBLAS_BFLOAT16_H
#define ROCBLAS_BFLOAT16_H

#if __cplusplus < 201103L || (!defined(__HCC__) && !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 rocblas_bfloat16

#include <stdint.h>
/*! \brief Struct to represent a 16 bit Brain floating-point number. */
typedef struct
{
    uint16_t data;
} rocblas_bfloat16;

#else // __cplusplus < 201103L || (!defined(__HCC__) && !defined(__HIPCC__))

#include <cmath>
#include <cstddef>
#include <cstdint>
#include <hip/hip_runtime.h>
#include <ostream>
#include <type_traits>

struct ROCBLAS_EXPORT rocblas_bfloat16
{
    uint16_t data;

    enum rocblas_truncate_t
    {
        rocblas_truncate
    };

    enum truncate_t
    {
        truncate __attribute__((deprecated(
            "From 3.0, The enum truncate_t and the value truncate will be "
            "removed and replaced by rocblas_truncate_t and rocblas_truncate, "
            "respectively. The new enum rocblas_truncate_t and the value rocblas_truncate could "
            "be used from this ROCm release for an easy transition.")))
    };

    __host__ __device__ rocblas_bfloat16() = default;

    // round upper 16 bits of IEEE float to convert to bfloat16
    explicit __host__ __device__ rocblas_bfloat16(float f)
        : data(float_to_bfloat16(f))
    {
    }

    explicit __host__ __device__ rocblas_bfloat16(float f, rocblas_truncate_t)
        : data(truncate_float_to_bfloat16(f))
    {
    }

    explicit __host__ __device__ rocblas_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;
    }

    explicit __host__ __device__ operator bool() const
    {
        return data & 0x7fff;
    }

private:
    static __host__ __device__ 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 uint16_t(u.int32 >> 16);
    }

    // Truncate instead of rounding, preserving SNaN
    static __host__ __device__ uint16_t truncate_float_to_bfloat16(float f)
    {
        union
        {
            float    fp32;
            uint32_t int32;
        } u = {f};
        return uint16_t(u.int32 >> 16) | (!(~u.int32 & 0x7f800000) && (u.int32 & 0xffff));
    }
};

typedef struct
{
    uint16_t data;
} rocblas_bfloat16_public;

static_assert(std::is_standard_layout<rocblas_bfloat16>{},
              "rocblas_bfloat16 is not a standard layout type, and thus is "
              "incompatible with C.");

static_assert(std::is_trivial<rocblas_bfloat16>{},
              "rocblas_bfloat16 is not a trivial type, and thus is "
              "incompatible with C.");

static_assert(sizeof(rocblas_bfloat16) == sizeof(rocblas_bfloat16_public)
                  && offsetof(rocblas_bfloat16, data) == offsetof(rocblas_bfloat16_public, data),
              "internal rocblas_bfloat16 does not match public rocblas_bfloat16");

inline std::ostream& operator<<(std::ostream& os, const rocblas_bfloat16& bf16)
{
    return os << float(bf16);
}
inline __host__ __device__ rocblas_bfloat16 operator+(rocblas_bfloat16 a)
{
    return a;
}
inline __host__ __device__ rocblas_bfloat16 operator-(rocblas_bfloat16 a)
{
    a.data ^= 0x8000;
    return a;
}
inline __host__ __device__ rocblas_bfloat16 operator+(rocblas_bfloat16 a, rocblas_bfloat16 b)
{
    return rocblas_bfloat16(float(a) + float(b));
}
inline __host__ __device__ rocblas_bfloat16 operator-(rocblas_bfloat16 a, rocblas_bfloat16 b)
{
    return rocblas_bfloat16(float(a) - float(b));
}
inline __host__ __device__ rocblas_bfloat16 operator*(rocblas_bfloat16 a, rocblas_bfloat16 b)
{
    return rocblas_bfloat16(float(a) * float(b));
}
inline __host__ __device__ rocblas_bfloat16 operator/(rocblas_bfloat16 a, rocblas_bfloat16 b)
{
    return rocblas_bfloat16(float(a) / float(b));
}
inline __host__ __device__ bool operator<(rocblas_bfloat16 a, rocblas_bfloat16 b)
{
    return float(a) < float(b);
}
inline __host__ __device__ bool operator==(rocblas_bfloat16 a, rocblas_bfloat16 b)
{
    return float(a) == float(b);
}
inline __host__ __device__ bool operator>(rocblas_bfloat16 a, rocblas_bfloat16 b)
{
    return b < a;
}
inline __host__ __device__ bool operator<=(rocblas_bfloat16 a, rocblas_bfloat16 b)
{
    return !(a > b);
}
inline __host__ __device__ bool operator!=(rocblas_bfloat16 a, rocblas_bfloat16 b)
{
    return !(a == b);
}
inline __host__ __device__ bool operator>=(rocblas_bfloat16 a, rocblas_bfloat16 b)
{
    return !(a < b);
}
inline __host__ __device__ rocblas_bfloat16& operator+=(rocblas_bfloat16& a, rocblas_bfloat16 b)
{
    return a = a + b;
}
inline __host__ __device__ rocblas_bfloat16& operator-=(rocblas_bfloat16& a, rocblas_bfloat16 b)
{
    return a = a - b;
}
inline __host__ __device__ rocblas_bfloat16& operator*=(rocblas_bfloat16& a, rocblas_bfloat16 b)
{
    return a = a * b;
}
inline __host__ __device__ rocblas_bfloat16& operator/=(rocblas_bfloat16& a, rocblas_bfloat16 b)
{
    return a = a / b;
}
inline __host__ __device__ rocblas_bfloat16& operator++(rocblas_bfloat16& a)
{
    return a += rocblas_bfloat16(1.0f);
}
inline __host__ __device__ rocblas_bfloat16& operator--(rocblas_bfloat16& a)
{
    return a -= rocblas_bfloat16(1.0f);
}
inline __host__ __device__ rocblas_bfloat16 operator++(rocblas_bfloat16& a, int)
{
    rocblas_bfloat16 orig = a;
    ++a;
    return orig;
}
inline __host__ __device__ rocblas_bfloat16 operator--(rocblas_bfloat16& a, int)
{
    rocblas_bfloat16 orig = a;
    --a;
    return orig;
}

namespace std
{
    constexpr __host__ __device__ bool isinf(rocblas_bfloat16 a)
    {
        return !(~a.data & 0x7f80) && !(a.data & 0x7f);
    }
    constexpr __host__ __device__ bool isnan(rocblas_bfloat16 a)
    {
        return !(~a.data & 0x7f80) && +(a.data & 0x7f);
    }
    constexpr __host__ __device__ bool iszero(rocblas_bfloat16 a)
    {
        return !(a.data & 0x7fff);
    }
    inline rocblas_bfloat16 sin(rocblas_bfloat16 a)
    {
        return rocblas_bfloat16(sinf(float(a)));
    }
    inline rocblas_bfloat16 cos(rocblas_bfloat16 a)
    {
        return rocblas_bfloat16(cosf(float(a)));
    }
    __device__ __host__ constexpr rocblas_bfloat16 real(const rocblas_bfloat16& a)
    {
        return a;
    }
}

#endif // __cplusplus < 201103L || (!defined(__HCC__) && !defined(__HIPCC__))

#endif // ROCBLAS_BFLOAT16_H