|
| 1 | +//===-- amdgpu floating point env manipulation functions --------*- C++ -*-===// |
| 2 | +// |
| 3 | +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. |
| 4 | +// See https://llvm.org/LICENSE.txt for license information. |
| 5 | +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception |
| 6 | +// |
| 7 | +//===----------------------------------------------------------------------===// |
| 8 | + |
| 9 | +#ifndef LLVM_LIBC_SRC___SUPPORT_FPUTIL_AMDGPU_FENVIMPL_H |
| 10 | +#define LLVM_LIBC_SRC___SUPPORT_FPUTIL_AMDGPU_FENVIMPL_H |
| 11 | + |
| 12 | +#include "src/__support/GPU/utils.h" |
| 13 | +#include "src/__support/macros/attributes.h" |
| 14 | +#include "src/__support/macros/properties/architectures.h" |
| 15 | + |
| 16 | +#if !defined(LIBC_TARGET_ARCH_IS_AMDGPU) |
| 17 | +#error "Invalid include" |
| 18 | +#endif |
| 19 | + |
| 20 | +#include <fenv.h> |
| 21 | +#include <stdint.h> |
| 22 | + |
| 23 | +namespace LIBC_NAMESPACE { |
| 24 | +namespace fputil { |
| 25 | + |
| 26 | +namespace internal { |
| 27 | +// Retuns the current status of the AMDGPU floating point environment. In |
| 28 | +// practice this is simply a 64-bit concatenation of the mode register and the |
| 29 | +// trap status register. |
| 30 | +// |
| 31 | +// The mode register controls the floating point behaviour of the device. It |
| 32 | +// can be read or written to by the kernel during runtime It is laid out as a |
| 33 | +// bit field with the following offsets and sizes listed for the relevant |
| 34 | +// entries. |
| 35 | +// |
| 36 | +// ┌─────┬─────────────┬─────┬─────────┬──────────┬─────────────┬────────────┐ |
| 37 | +// │ ... │ EXCP[20:12] │ ... │ IEEE[9] │ CLAMP[8] │ DENORM[7:4] │ ROUND[3:0] │ |
| 38 | +// └─────┴─────────────┴─────┴─────────┴──────────┴─────────────┴────────────┘ |
| 39 | +// |
| 40 | +// The rounding mode and denormal modes both control f64/f16 and f32 precision |
| 41 | +// operations separately with two bits. The accepted values for the rounding |
| 42 | +// mode are nearest, upward, downward, and toward given 0, 1, 2, and 3 |
| 43 | +// respectively. |
| 44 | +// |
| 45 | +// The CLAMP bit indicates that DirectX 10 handling of NaNs is enabled in the |
| 46 | +// vector ALU. When set this will clamp NaN values to zero and pass them |
| 47 | +// otherwise. A hardware bug causes this bit to prevent floating exceptions |
| 48 | +// from being recorded if this bit is set on all generations before GFX12. |
| 49 | +// |
| 50 | +// The IEEE bit controls whether or not floating point operations supporting |
| 51 | +// exception gathering are IEEE 754-2008 compliant. |
| 52 | +// |
| 53 | +// The EXCP field indicates which exceptions will cause the instruction to |
| 54 | +// take a trap if traps are enabled, see the status register. The bit layout |
| 55 | +// is identical to that in the trap status register. We are only concerned |
| 56 | +// with the first six bits and ignore the other three. |
| 57 | +// |
| 58 | +// The trap status register contains information about the status of the |
| 59 | +// exceptions. These bits are accumulated regarless of trap handling statuss |
| 60 | +// and are sticky until cleared. |
| 61 | +// |
| 62 | +// 5 4 3 2 1 0 |
| 63 | +// ┌─────────┬───────────┬──────────┬────────────────┬──────────┬─────────┐ |
| 64 | +// │ Inexact │ Underflow │ Overflow │ Divide by zero │ Denormal │ Invalid │ |
| 65 | +// └─────────┴───────────┴──────────┴────────────────┴──────────┴─────────┘ |
| 66 | +// |
| 67 | +// These exceptions indicate that at least one lane in the current wavefront |
| 68 | +// signalled an floating point exception. There is no way to increase the |
| 69 | +// granularity. |
| 70 | +// |
| 71 | +// The returned value has the following layout. |
| 72 | +// |
| 73 | +// ┌────────────────────┬─────────────────────┐ |
| 74 | +// │ Trap Status[38:32] │ Mode Register[31:0] │ |
| 75 | +// └────────────────────┴─────────────────────┘ |
| 76 | +LIBC_INLINE uint64_t get_fpenv() { return __builtin_amdgcn_get_fpenv(); } |
| 77 | + |
| 78 | +// Set the floating point environment using the same layout as above. |
| 79 | +LIBC_INLINE void set_fpenv(uint64_t env) { __builtin_amdgcn_set_fpenv(env); } |
| 80 | + |
| 81 | +// The six bits used to encode the standard floating point exceptions in the |
| 82 | +// trap status register. |
| 83 | +enum ExceptionFlags : uint32_t { |
| 84 | + EXCP_INVALID_F = 0x1, |
| 85 | + EXCP_DENORMAL_F = 0x2, |
| 86 | + EXCP_DIV_BY_ZERO_F = 0x4, |
| 87 | + EXCP_OVERFLOW_F = 0x8, |
| 88 | + EXCP_UNDERFLOW_F = 0x10, |
| 89 | + EXCP_INEXACT_F = 0x20, |
| 90 | +}; |
| 91 | + |
| 92 | +// The two bit encoded rounding modes used in the mode register. |
| 93 | +enum RoundingFlags : uint32_t { |
| 94 | + ROUND_TO_NEAREST = 0x0, |
| 95 | + ROUND_UPWARD = 0x1, |
| 96 | + ROUND_DOWNWARD = 0x2, |
| 97 | + ROUND_TOWARD_ZERO = 0x3, |
| 98 | +}; |
| 99 | + |
| 100 | +// Exception flags are individual bits in the corresponding hardware register. |
| 101 | +// This converts between the exported C standard values and the hardware values. |
| 102 | +LIBC_INLINE uint32_t get_status_value_for_except(uint32_t excepts) { |
| 103 | + return (excepts & FE_INVALID ? EXCP_INVALID_F : 0) | |
| 104 | +#ifdef __FE_DENORM |
| 105 | + (excepts & __FE_DENORM ? EXCP_DENORMAL_F : 0) | |
| 106 | +#endif // __FE_DENORM |
| 107 | + (excepts & FE_DIVBYZERO ? EXCP_DIV_BY_ZERO_F : 0) | |
| 108 | + (excepts & FE_OVERFLOW ? EXCP_OVERFLOW_F : 0) | |
| 109 | + (excepts & FE_UNDERFLOW ? EXCP_UNDERFLOW_F : 0) | |
| 110 | + (excepts & FE_INEXACT ? EXCP_INEXACT_F : 0); |
| 111 | +} |
| 112 | + |
| 113 | +LIBC_INLINE uint32_t get_except_value_for_status(uint32_t status) { |
| 114 | + return (status & EXCP_INVALID_F ? FE_INVALID : 0) | |
| 115 | +#ifdef __FE_DENORM |
| 116 | + (status & EXCP_DENORMAL_F ? __FE_DENORM : 0) | |
| 117 | +#endif // __FE_DENORM |
| 118 | + (status & EXCP_DIV_BY_ZERO_F ? FE_DIVBYZERO : 0) | |
| 119 | + (status & EXCP_OVERFLOW_F ? FE_OVERFLOW : 0) | |
| 120 | + (status & EXCP_UNDERFLOW_F ? FE_UNDERFLOW : 0) | |
| 121 | + (status & EXCP_INEXACT_F ? FE_INEXACT : 0); |
| 122 | +} |
| 123 | + |
| 124 | +// Access the four bits in the mode register's ROUND[3:0] field. The hardware |
| 125 | +// supports setting the f64/f16 and f32 precision rounding modes separately but |
| 126 | +// we will assume that these always match. |
| 127 | +LIBC_INLINE void set_rounding_mode(uint32_t flags) { |
| 128 | + uint64_t old = get_fpenv() & 0xfffffffffffffff0; |
| 129 | + set_fpenv(old | flags << 2 | flags); |
| 130 | +} |
| 131 | + |
| 132 | +// The control register can modify f32/f64 rounding modes individually. For our |
| 133 | +// purposes we assume that these always match as we do not expose this through |
| 134 | +// the C interface. |
| 135 | +LIBC_INLINE uint32_t get_rounding_mode() { return get_fpenv() & 0x3; } |
| 136 | + |
| 137 | +} // namespace internal |
| 138 | + |
| 139 | +// TODO: Not implemented yet. |
| 140 | +LIBC_INLINE int clear_except(int) { return 0; } |
| 141 | + |
| 142 | +// TODO: Not implemented yet. |
| 143 | +LIBC_INLINE int test_except(int) { return 0; } |
| 144 | + |
| 145 | +// TODO: Not implemented yet. |
| 146 | +LIBC_INLINE int get_except() { return 0; } |
| 147 | + |
| 148 | +// TODO: Not implemented yet. |
| 149 | +LIBC_INLINE int set_except(int) { return 0; } |
| 150 | + |
| 151 | +// TODO: Not implemented yet. |
| 152 | +LIBC_INLINE int enable_except(int) { return 0; } |
| 153 | + |
| 154 | +// TODO: Not implemented yet. |
| 155 | +LIBC_INLINE int disable_except(int) { return 0; } |
| 156 | + |
| 157 | +// TODO: Not implemented yet. |
| 158 | +LIBC_INLINE int raise_except(int) { return 0; } |
| 159 | + |
| 160 | +LIBC_INLINE int get_round() { |
| 161 | + switch (internal::get_rounding_mode()) { |
| 162 | + case internal::ROUND_TO_NEAREST: |
| 163 | + return FE_TONEAREST; |
| 164 | + case internal::ROUND_UPWARD: |
| 165 | + return FE_UPWARD; |
| 166 | + case internal::ROUND_DOWNWARD: |
| 167 | + return FE_DOWNWARD; |
| 168 | + case internal::ROUND_TOWARD_ZERO: |
| 169 | + return FE_TOWARDZERO; |
| 170 | + } |
| 171 | + __builtin_unreachable(); |
| 172 | +} |
| 173 | + |
| 174 | +LIBC_INLINE int set_round(int rounding_mode) { |
| 175 | + switch (rounding_mode) { |
| 176 | + case FE_TONEAREST: |
| 177 | + internal::set_rounding_mode(internal::ROUND_TO_NEAREST); |
| 178 | + break; |
| 179 | + case FE_UPWARD: |
| 180 | + internal::set_rounding_mode(internal::ROUND_UPWARD); |
| 181 | + break; |
| 182 | + case FE_DOWNWARD: |
| 183 | + internal::set_rounding_mode(internal::ROUND_DOWNWARD); |
| 184 | + break; |
| 185 | + case FE_TOWARDZERO: |
| 186 | + internal::set_rounding_mode(internal::ROUND_TOWARD_ZERO); |
| 187 | + break; |
| 188 | + default: |
| 189 | + return 1; |
| 190 | + } |
| 191 | + return 0; |
| 192 | +} |
| 193 | + |
| 194 | +LIBC_INLINE int get_env(fenv_t *env) { |
| 195 | + if (!env) |
| 196 | + return 1; |
| 197 | + |
| 198 | + env->__fpc = internal::get_fpenv(); |
| 199 | + return 0; |
| 200 | +} |
| 201 | + |
| 202 | +LIBC_INLINE int set_env(const fenv_t *env) { |
| 203 | + if (!env) |
| 204 | + return 1; |
| 205 | + |
| 206 | + internal::set_fpenv(env->__fpc); |
| 207 | + return 0; |
| 208 | +} |
| 209 | + |
| 210 | +} // namespace fputil |
| 211 | +} // namespace LIBC_NAMESPACE |
| 212 | + |
| 213 | +#endif // LLVM_LIBC_SRC___SUPPORT_FPUTIL_AMDGPU_FENVIMPL_H |
0 commit comments