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