[libc-commits] [libc] [libc] Implement basic 'fenv.h' utilities on the AMD GPU (PR #83500)

Matt Arsenault via libc-commits libc-commits at lists.llvm.org
Thu Feb 29 23:18:49 PST 2024


================
@@ -0,0 +1,302 @@
+//===-- amdgpu floating point env manipulation functions --------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIBC_SRC___SUPPORT_FPUTIL_AMDGPU_FENVIMPL_H
+#define LLVM_LIBC_SRC___SUPPORT_FPUTIL_AMDGPU_FENVIMPL_H
+
+#include "src/__support/GPU/utils.h"
+#include "src/__support/macros/attributes.h"
+#include "src/__support/macros/properties/architectures.h"
+
+#if !defined(LIBC_TARGET_ARCH_IS_AMDGPU)
+#error "Invalid include"
+#endif
+
+#include <fenv.h>
+#include <stdint.h>
+
+namespace LIBC_NAMESPACE {
+namespace fputil {
+
+namespace internal {
+
+// Gets the immediate argument to access the AMDGPU hardware register. The
+// register access is encoded in a 16-bit immediate value according to the
+// following layout.
+//
+// ┌──────────────┬──────────────┬───────────────┐
+// │  SIZE[15:11] │ OFFSET[10:6] │    ID[5:0]    │
+// └──────────────┴──────────────┴───────────────┘
+//
+// This will read the size number of bits starting at the offset bit from the
+// corresponding hardware register ID.
+LIBC_INLINE constexpr uint16_t get_register(uint32_t id, uint32_t offset,
+                                            uint32_t size) {
+  return static_cast<uint16_t>(size << 11 | offset << 6 | id);
+}
+
+// Integral identifiers for the relevant hardware registers.
+enum Register : uint16_t {
+  // The mode register controls the floating point behaviour of the device. It
+  // can be read or written to by the kernel during runtime It is laid out as a
+  // bit field with the following offsets and sizes listed for the relevant
+  // entries.
+  //
+  // ┌─────┬─────────────┬─────┬─────────┬──────────┬─────────────┬────────────┐
+  // │ ... │ EXCP[20:12] │ ... │ IEEE[9] │ CLAMP[8] │ DENORM[7:4] │ ROUND[3:0] │
+  // └─────┴─────────────┴─────┴─────────┴──────────┴─────────────┴────────────┘
+  //
+  // The rounding mode and denormal modes both control f64/f16 and f32 precision
+  // operations separately with two bits. The accepted values for the rounding
+  // mode are nearest, upward, downward, and toward given 0, 1, 2, and 3
+  // respectively.
+  //
+  // The CLAMP bit indicates that DirectX 10 handling of NaNs is enabled in the
+  // vector ALU. When set this will clamp NaN values to zero and pass them
+  // otherwise. A hardware bug causes this bit to prevent floating exceptions
+  // from being recorded if this bit is set on all generations before GFX12.
+  //
+  // The IEEE bit controls whether or not floating point operations supporting
+  // exception gathering are IEEE 754-2008 compliant.
+  //
+  // The EXCP field indicates which exceptions will cause the instruction to
+  // take a trap if traps are enabled, see the status register. The bit layout
+  // is identical to that in the trap status register. We are only concerned
+  // with the first six bits and ignore the other three.
+  HW_REG_MODE = 1,
+  HW_REG_MODE_ROUND = get_register(HW_REG_MODE, 0, 4),
+  HW_REG_MODE_CLAMP = get_register(HW_REG_MODE, 8, 1),
+  HW_REG_MODE_EXCP = get_register(HW_REG_MODE, 12, 6),
+
+  // The status register is a read-only register that contains information about
+  // how the kernel was launched. The sixth bit TRAP_EN[6] indicates whether or
+  // not traps are enabled for this kernel. If this bit is set along with the
+  // corresponding bit in the mode register then a trap will be taken.
+  HW_REG_STATUS = 2,
+  HW_REG_STATUS_TRAP_EN = get_register(HW_REG_STATUS, 6, 1),
+
+  // The trap status register contains information about the status of the
+  // exceptions. These bits are accumulated regarless of trap handling statuss
+  // and are sticky until cleared.
+  //
+  // 5         4           3          2                1          0
+  // ┌─────────┬───────────┬──────────┬────────────────┬──────────┬─────────┐
+  // │ Inexact │ Underflow │ Overflow │ Divide by zero │ Denormal │ Invalid │
+  // └─────────┴───────────┴──────────┴────────────────┴──────────┴─────────┘
+  //
+  // These exceptions indicate that at least one lane in the current wavefront
+  // signalled an floating point exception. There is no way to increase the
+  // granularity.
+  HW_REG_TRAPSTS = 3,
+  HW_REG_TRAPSTS_EXCP = get_register(HW_REG_TRAPSTS, 0, 6),
+};
+
+// The six bits used to encode the standard floating point exceptions in the
+// trap status register.
+enum ExceptionFlags : uint32_t {
+  EXCP_INVALID_F = 0x1,
+  EXCP_DENORMAL_F = 0x2,
+  EXCP_DIV_BY_ZERO_F = 0x4,
+  EXCP_OVERFLOW_F = 0x8,
+  EXCP_UNDERFLOW_F = 0x10,
+  EXCP_INEXACT_F = 0x20,
+};
+
+// The two bit encoded rounding modes used in the mode register.
+enum RoundingFlags : uint32_t {
+  ROUND_TO_NEAREST = 0x0,
+  ROUND_UPWARD = 0x1,
+  ROUND_DOWNWARD = 0x2,
+  ROUND_TOWARD_ZERO = 0x3,
+};
+
+// Exception flags are individual bits in the corresponding hardware register.
+// This converts between the exported C standard values and the hardware values.
+LIBC_INLINE uint32_t get_status_value_for_except(uint32_t excepts) {
+  return (excepts & FE_INVALID ? EXCP_INVALID_F : 0) |
+#ifdef __FE_DENORM
+         (excepts & __FE_DENORM ? EXCP_DENORMAL_F : 0) |
+#endif // __FE_DENORM
+         (excepts & FE_DIVBYZERO ? EXCP_DIV_BY_ZERO_F : 0) |
+         (excepts & FE_OVERFLOW ? EXCP_OVERFLOW_F : 0) |
+         (excepts & FE_UNDERFLOW ? EXCP_UNDERFLOW_F : 0) |
+         (excepts & FE_INEXACT ? EXCP_INEXACT_F : 0);
+}
+
+LIBC_INLINE uint32_t get_except_value_for_status(uint32_t status) {
+  return (status & EXCP_INVALID_F ? FE_INVALID : 0) |
+#ifdef __FE_DENORM
+         (status & EXCP_DENORMAL_F ? __FE_DENORM : 0) |
+#endif // __FE_DENORM
+         (status & EXCP_DIV_BY_ZERO_F ? FE_DIVBYZERO : 0) |
+         (status & EXCP_OVERFLOW_F ? FE_OVERFLOW : 0) |
+         (status & EXCP_UNDERFLOW_F ? FE_UNDERFLOW : 0) |
+         (status & EXCP_INEXACT_F ? FE_INEXACT : 0);
+}
+
+// FIXME: These require the 'noinline' attribute to pessimistically flush the
+//        state. Otherwise, reading from the register may return stale results.
+
+// Access the six bits in the trap status register for the floating point
+// exceptions.
+[[gnu::noinline]] LIBC_INLINE void set_trap_status(uint32_t status) {
+  uint32_t val = gpu::broadcast_value(gpu::get_lane_mask(), status);
+  __builtin_amdgcn_s_setreg(HW_REG_TRAPSTS_EXCP, val);
+}
+
+[[gnu::noinline]] LIBC_INLINE uint32_t get_trap_status() {
+  return __builtin_amdgcn_s_getreg(HW_REG_TRAPSTS_EXCP);
+}
+
+// Access the six bits in the mode register that control which exceptions will
+// result in a trap being taken. Uses the same flags as the status register.
+[[gnu::noinline]] LIBC_INLINE void set_enabled_trap(uint32_t flags) {
+  uint32_t val = gpu::broadcast_value(gpu::get_lane_mask(), flags);
+  __builtin_amdgcn_s_setreg(HW_REG_MODE_EXCP, val);
+}
+
+[[gnu::noinline]] LIBC_INLINE uint32_t get_enabled_trap() {
+  return __builtin_amdgcn_s_getreg(HW_REG_MODE_EXCP);
+}
+
+// Access the four bits in the mode register's ROUND[3:0] field. The hardware
+// supports setting the f64/f16 and f32 precision rounding modes separately but
+// we will assume that these always match.
+[[gnu::noinline]] LIBC_INLINE void set_rounding_mode(uint32_t flags) {
+  uint32_t val = gpu::broadcast_value(gpu::get_lane_mask(), flags);
+  __builtin_amdgcn_s_setreg(HW_REG_MODE_ROUND, val << 2 | val);
+}
+
+[[gnu::noinline]] LIBC_INLINE uint32_t get_rounding_mode() {
+  return __builtin_amdgcn_s_getreg(HW_REG_MODE_ROUND) & 0x3;
+}
+
+// NOTE: On architectures before GFX12 the DX10_CLAMP bit supresses all floating
+//       point exceptions. In order to get them to be presented we need to
+//       manually set if off.
+[[gnu::noinline]] LIBC_INLINE void set_clamp_low() {
+  __builtin_amdgcn_s_setreg(HW_REG_MODE_CLAMP, 0);
+}
+
+[[gnu::noinline]] LIBC_INLINE void set_clamp_high() {
----------------
arsenm wrote:

"high" should be "on"?

https://github.com/llvm/llvm-project/pull/83500


More information about the libc-commits mailing list