diff --git a/.github/copilot-instructions.md b/.github/copilot-instructions.md index 47cb0be1..57fdb035 100644 --- a/.github/copilot-instructions.md +++ b/.github/copilot-instructions.md @@ -16,7 +16,6 @@ The library has CPU and CUDA backends. HIP support is architecturally possible b --- ## Repository Layout - ``` FusedKernelLibrary/ ├── .clang-format # LLVM-based style, 4-space indent, 120-char column limit @@ -44,7 +43,6 @@ FusedKernelLibrary/ ├── utests/ # Unit tests (header .h files, auto-discovered) └── benchmarks/ # Benchmarks (disabled by default, ENABLE_BENCHMARK=ON) ``` - --- ## Build System @@ -56,15 +54,13 @@ FusedKernelLibrary/ - **MSVC**: Visual Studio 2022 or Visual Studio 2026 (MSVC_VERSION >= 1930) required; ### Configure and Build (typical) -```bash # Linux (Ninja) cmake -G "Ninja" -B build -DCMAKE_BUILD_TYPE=Release -S . -cmake --build build --config Release +cmake --build build --config Release --parallel 32 # Windows (Ninja, inside VS Developer Shell) cmake -G "Ninja" -B build -DCMAKE_BUILD_TYPE=Release -S . -cmake --build build --config Release -``` +cmake --build build --config Release --parallel 32 ### Key CMake Options | Option | Default | Description | @@ -91,12 +87,8 @@ cmake --build build --config Release --- ## Running Tests - -```bash cd build ctest --build-config Release --output-junit test_results.xml -``` - Tests are registered with CTest automatically. Individual targets follow the naming pattern `_cpp` (CPU) and `_cu` (CUDA). --- @@ -116,7 +108,9 @@ Tests in `tests/` and `utests/` are **not** written with a traditional test fram 6. Use `// ONLY_CPU` in a test header to suppress the `_cu` target. ### Test File Structure + Every test header must define a `launch()` function returning `int`: + ```cpp #include #include @@ -163,13 +157,15 @@ In CPU-only mode (no NVCC, no CLANG_HOST_DEVICE), these macros degrade to standa The `FK_STATIC_STRUCT(StructName, StructAlias)` macro marks a struct as non-constructible and non-copyable (deletes default/copy/move constructors and assignment operators). ### Type Aliases + The library defines CUDA-compatible type aliases (also available in CPU mode): + ```cpp -using uchar = unsigned char; -using schar = signed char; -using uint = unsigned int; -using ushort = unsigned short; -using ulong = unsigned long; +using uchar = unsigned char; +using schar = signed char; +using uint = unsigned int; +using ushort = unsigned short; +using ulong = unsigned long; using longlong = long long; using ulonglong = unsigned long long; ``` @@ -190,7 +186,9 @@ Run `clang-format` using the `.clang-format` file at the repo root: ## Core API Patterns ### Executing Fused Operations + The primary entry point is `fk::executeOperations(stream, op1, op2, ...)`: + ```cpp #include using namespace fk; @@ -263,3 +261,5 @@ See existing operations like `Mul`, `Add`, `SaturateCast` in `include/fused_kern 1. **Windows Ninja + NVCC path**: After CMake configure on Windows with Ninja, `/CMakeFiles/rules.ninja` may contain an incorrect path to `nvcc.exe`. The CI workflow patches this with PowerShell `Set-Content`. If you hit this locally, check that `CUDACXX` env var is set before invoking CMake and verify the generated `rules.ninja`. +2. **CUDA Compilation Availability**: CUDA compilation is supported and should not be declared blocked based on an unsuccessful ad hoc nvcc invocation; use the repository's documented VS Developer Shell/CMake build configuration and correct architecture/toolchain settings before concluding CUDA is unavailable. + diff --git a/.github/workflows/cmake-linux-amd64.yml b/.github/workflows/cmake-linux-amd64.yml index d1565122..ff90f5af 100644 --- a/.github/workflows/cmake-linux-amd64.yml +++ b/.github/workflows/cmake-linux-amd64.yml @@ -8,7 +8,7 @@ on: jobs: build: - runs-on: [self-hosted, linux, x64] + runs-on: [self-hosted, linux, x64, cuda] strategy: # Set fail-fast to false to ensure that feedback is delivered for all matrix combinations. diff --git a/.github/workflows/cmake-linux-arm64.yml b/.github/workflows/cmake-linux-arm64.yml index 15a15981..6f52ec19 100644 --- a/.github/workflows/cmake-linux-arm64.yml +++ b/.github/workflows/cmake-linux-arm64.yml @@ -8,7 +8,7 @@ on: jobs: build: - runs-on: [self-hosted, linux, arm64] + runs-on: [self-hosted, linux, arm64, cuda] strategy: # Set fail-fast to false to ensure that feedback is delivered for all matrix combinations. diff --git a/.github/workflows/cmake-windows-amd64.yml b/.github/workflows/cmake-windows-amd64.yml index 78da62f2..5b62ccd0 100644 --- a/.github/workflows/cmake-windows-amd64.yml +++ b/.github/workflows/cmake-windows-amd64.yml @@ -8,7 +8,7 @@ on: jobs: build: - runs-on: [self-hosted, windows, x64] + runs-on: [self-hosted, windows, x64, cuda] strategy: # Set fail-fast to false to ensure that feedback is delivered for all matrix combinations. Consider changing this to true when your workflow is stable. diff --git a/.github/workflows/copilot-setup-steps.yml b/.github/workflows/copilot-setup-steps.yml index f1ebfd88..0b8e1084 100644 --- a/.github/workflows/copilot-setup-steps.yml +++ b/.github/workflows/copilot-setup-steps.yml @@ -6,7 +6,7 @@ on: jobs: copilot-setup-steps: # Run Copilot setup on the self-hosted Linux x64 runner - runs-on: [self-hosted, linux, x64] + runs-on: [self-hosted, linux, x64, cuda] permissions: contents: read diff --git a/include/fused_kernel/algorithms/basic_ops/math.h b/include/fused_kernel/algorithms/basic_ops/math.h index 5268133b..10250236 100644 --- a/include/fused_kernel/algorithms/basic_ops/math.h +++ b/include/fused_kernel/algorithms/basic_ops/math.h @@ -28,7 +28,8 @@ namespace fk { struct AbsFunc { using InstanceType = UnaryType; template FK_HOST_DEVICE_FUSE auto exec(const ST& s) { - if constexpr (std::is_signed_v) return s < ST(0) ? static_cast(-s) : s; + if constexpr (isReducedFloat) return ReducedFloatTraits::abs(s); + else if constexpr (std::is_signed_v) return s < ST(0) ? static_cast(-s) : s; else return s; } }; diff --git a/include/fused_kernel/algorithms/basic_ops/vector_ops.h b/include/fused_kernel/algorithms/basic_ops/vector_ops.h index f4c97584..59c1456c 100644 --- a/include/fused_kernel/algorithms/basic_ops/vector_ops.h +++ b/include/fused_kernel/algorithms/basic_ops/vector_ops.h @@ -32,7 +32,7 @@ namespace fk { static_assert(std::is_same_v, VBase>, "Base types should be the same"); const auto result = cxp::discard>::f(input); - if constexpr (std::is_fundamental_v) { + if constexpr (validScalar) { return result.x; } else { return result; diff --git a/include/fused_kernel/algorithms/image_processing/color_conversion.h b/include/fused_kernel/algorithms/image_processing/color_conversion.h index 31166b44..4d049eb3 100644 --- a/include/fused_kernel/algorithms/image_processing/color_conversion.h +++ b/include/fused_kernel/algorithms/image_processing/color_conversion.h @@ -70,6 +70,9 @@ namespace fk { #endif } else if constexpr (std::is_floating_point_v) { return compute_luminance(input); + } else { + static_assert(isReducedFloat, "RGB2Gray: unsupported output type"); + return static_cast(compute_luminance(input)); } } private: diff --git a/include/fused_kernel/core/constexpr_libs/constexpr_cmath.h b/include/fused_kernel/core/constexpr_libs/constexpr_cmath.h index 4fb7653d..f4c1a117 100644 --- a/include/fused_kernel/core/constexpr_libs/constexpr_cmath.h +++ b/include/fused_kernel/core/constexpr_libs/constexpr_cmath.h @@ -31,7 +31,21 @@ namespace cxp { constexpr T maxValue = std::numeric_limits::max(); template - constexpr T smallestPositiveValue = std::is_floating_point_v ? std::numeric_limits::min() : static_cast(1); + constexpr T smallestPositiveValue = fk::validFloatingPoint ? std::numeric_limits::min() : static_cast(1); + + namespace detail { + // Reduced floats promote to float for generic comparisons/arithmetic: their mixed + // implicit conversions are ambiguous (fp16/bf16) or non existent (fp8/fp4). + template + FK_HOST_DEVICE_CNST auto promoteReduced(const T& value) { + if constexpr (fk::isReducedFloat) { + return static_cast(value); + } else { + return value; + } + } + } // namespace detail + using detail::promoteReduced; #define CXP_F_FUNC \ template \ @@ -44,7 +58,13 @@ namespace cxp { using InstanceType = fk::UnaryType; template FK_HOST_DEVICE_FUSE bool exec(const ST& s) { - return s != s; + if constexpr (fk::isReducedFloat) { + // fp8/fp4 have no comparison operators, and e4m3/e2m1 NaN is not + // detectable via s != s anyway: classify by bit pattern. + return fk::ReducedFloatTraits::isNaN(s); + } else { + return s != s; + } } }; CXP_F_FUNC @@ -55,7 +75,11 @@ namespace cxp { using InstanceType = fk::UnaryType; template FK_HOST_DEVICE_FUSE bool exec(const ST& s) { - return s == s && s != ST(0) && s + s == s; + if constexpr (fk::isReducedFloat) { + return fk::ReducedFloatTraits::isInf(s); + } else { + return s == s && s != ST(0) && s + s == s; + } } }; CXP_F_FUNC @@ -79,23 +103,29 @@ namespace cxp { using InstanceType = fk::BinaryType; template FK_HOST_DEVICE_FUSE bool exec(const ST1& s1, const ST2& s2) { - static_assert(!std::is_same_v && std::is_fundamental_v, - "First parameter must be a fundamental type other than bool"); - static_assert(!std::is_same_v && std::is_fundamental_v, - "Second parameter must be a fundamental type other than bool"); - constexpr bool isAnyFloatingPoint = std::is_floating_point_v || std::is_floating_point_v; - constexpr bool areBothSigned = std::is_signed_v == std::is_signed_v; - if constexpr (isAnyFloatingPoint || areBothSigned) { - // Safe comparison cases - return s1 == s2; - } else if constexpr (std::is_signed_v) { - // T is signed, U is unsigned, both are integers - if (s1 < 0) return false; // Negative cannot equal any unsigned. - return static_cast>(s1) == s2; + if constexpr (fk::isReducedFloat || fk::isReducedFloat) { + // Promote reduced floats explicitly: mixed implicit conversions between + // float and half-like types are ambiguous, and fp8/fp4 have no operators. + return exec(promoteReduced(s1), promoteReduced(s2)); } else { - // T is unsigned, U is signed, both are integers - if (s2 < 0) return false; // Negative cannot equal any unsigned. - return s1 == static_cast>(s2); + static_assert(!std::is_same_v && std::is_fundamental_v, + "First parameter must be a fundamental type other than bool"); + static_assert(!std::is_same_v && std::is_fundamental_v, + "Second parameter must be a fundamental type other than bool"); + constexpr bool isAnyFloatingPoint = std::is_floating_point_v || std::is_floating_point_v; + constexpr bool areBothSigned = std::is_signed_v == std::is_signed_v; + if constexpr (isAnyFloatingPoint || areBothSigned) { + // Safe comparison cases + return s1 == s2; + } else if constexpr (std::is_signed_v) { + // T is signed, U is unsigned, both are integers + if (s1 < 0) return false; // Negative cannot equal any unsigned. + return static_cast>(s1) == s2; + } else { + // T is unsigned, U is signed, both are integers + if (s2 < 0) return false; // Negative cannot equal any unsigned. + return s1 == static_cast>(s2); + } } } }; @@ -120,23 +150,27 @@ namespace cxp { using InstanceType = fk::BinaryType; template FK_HOST_DEVICE_FUSE bool exec(const ST1& s1, const ST2& s2) { - static_assert(!std::is_same_v && std::is_fundamental_v, - "First parameter must be a fundamental type other than bool"); - static_assert(!std::is_same_v && std::is_fundamental_v, - "Second parameter must be a fundamental type other than bool"); - constexpr bool isAnyFloatingPoint = std::is_floating_point_v || std::is_floating_point_v; - constexpr bool areBothSigned = std::is_signed_v == std::is_signed_v; - if constexpr (isAnyFloatingPoint || areBothSigned) { - // Safe comparison cases - return s1 < s2; - } else if constexpr (std::is_signed_v) { - // T is signed, U is unsigned, both are integers - if (s1 < 0) return true; // Signed negative is always less than unsigned. - return static_cast>(s1) < s2; + if constexpr (fk::isReducedFloat || fk::isReducedFloat) { + return exec(promoteReduced(s1), promoteReduced(s2)); } else { - // T is unsigned, U is signed, both are integers - if (s2 < 0) return false; // Unsigned is never less than a signed negative. - return s1 < static_cast>(s2); + static_assert(!std::is_same_v && std::is_fundamental_v, + "First parameter must be a fundamental type other than bool"); + static_assert(!std::is_same_v && std::is_fundamental_v, + "Second parameter must be a fundamental type other than bool"); + constexpr bool isAnyFloatingPoint = std::is_floating_point_v || std::is_floating_point_v; + constexpr bool areBothSigned = std::is_signed_v == std::is_signed_v; + if constexpr (isAnyFloatingPoint || areBothSigned) { + // Safe comparison cases + return s1 < s2; + } else if constexpr (std::is_signed_v) { + // T is signed, U is unsigned, both are integers + if (s1 < 0) return true; // Signed negative is always less than unsigned. + return static_cast>(s1) < s2; + } else { + // T is unsigned, U is signed, both are integers + if (s2 < 0) return false; // Unsigned is never less than a signed negative. + return s1 < static_cast>(s2); + } } } }; @@ -296,14 +330,19 @@ namespace cxp { using InstanceType = fk::BinaryType; template FK_HOST_DEVICE_FUSE auto exec(const ST& s1, const ST& s2) - -> std::enable_if_t, ST> { - return s1 >= s2 ? s1 : s2; + -> std::enable_if_t, ST> { + if constexpr (fk::isReducedFloat) { + // fp8/fp4 have no comparison operators: compare in float. + return static_cast(s1) >= static_cast(s2) ? s1 : s2; + } else { + return s1 >= s2 ? s1 : s2; + } } }; CXP_F_FUNC template FK_HOST_DEVICE_FUSE ST f(const ST& s) { - return s; + return s; } }; @@ -311,15 +350,19 @@ namespace cxp { struct BaseFunc { using InstanceType = fk::BinaryType; template - FK_HOST_DEVICE_FUSE auto exec(const ST& s1, const ST& s2) - -> std::enable_if_t, ST> { - return s1 <= s2 ? s1 : s2; + FK_HOST_DEVICE_FUSE auto exec(const ST& s1, const ST& s2) + -> std::enable_if_t, ST> { + if constexpr (fk::isReducedFloat) { + return static_cast(s1) <= static_cast(s2) ? s1 : s2; + } else { + return s1 <= s2 ? s1 : s2; + } } }; CXP_F_FUNC template FK_HOST_DEVICE_FUSE ST f(const ST& value) { - return value; + return value; } }; @@ -328,8 +371,11 @@ namespace cxp { using InstanceType = fk::UnaryType; template FK_HOST_DEVICE_FUSE auto exec(const ST& s) { - static_assert(std::is_fundamental_v, "abs does not support non fundamental types"); - if constexpr (std::is_signed_v) { + static_assert(fk::validScalar, "abs does not support non fundamental types"); + if constexpr (fk::isReducedFloat) { + // Sign-magnitude formats: clearing the sign bit is exact (and NaN safe). + return fk::ReducedFloatTraits::abs(s); + } else if constexpr (std::is_signed_v) { // For signed integrals, when x is std::numerical_limits::lowest(), // the result is undefined behavior in C++. So, for the sake of performance, // we will not do any special treatment for those cases. diff --git a/include/fused_kernel/core/constexpr_libs/constexpr_saturate.h b/include/fused_kernel/core/constexpr_libs/constexpr_saturate.h index 70b8f27c..d9f62bb0 100644 --- a/include/fused_kernel/core/constexpr_libs/constexpr_saturate.h +++ b/include/fused_kernel/core/constexpr_libs/constexpr_saturate.h @@ -27,6 +27,12 @@ namespace cxp { using InstanceType = fk::UnaryType; template FK_HOST_DEVICE_FUSE auto exec(const ST& s) { + if constexpr (fk::isReducedFloat) { + // Promote reduced float sources once: every reduced value is exactly + // representable in float, and the float path already handles rounding + // and clamping towards any output type. + return exec(static_cast(s)); + } else { constexpr auto maxValOutput = maxValue>; constexpr auto minValueOutput = minValue>; if (cxp::cmp_greater::BaseFunc::exec(s, maxValOutput)) { @@ -44,6 +50,7 @@ namespace cxp { return static_cast>(s); } } + } } }; public: diff --git a/include/fused_kernel/core/constexpr_libs/constexpr_vector_exec.h b/include/fused_kernel/core/constexpr_libs/constexpr_vector_exec.h index 382a6151..f9011274 100644 --- a/include/fused_kernel/core/constexpr_libs/constexpr_vector_exec.h +++ b/include/fused_kernel/core/constexpr_libs/constexpr_vector_exec.h @@ -26,7 +26,7 @@ namespace cxp { struct Exec>> { template FK_HOST_DEVICE_FUSE auto exec(const T& val) { - if constexpr (std::is_fundamental_v) { + if constexpr (fk::validScalar) { return Op::exec(val); } else { static_assert(fk::validCUDAVec, "Type not supported in Unary operation execution."); diff --git a/include/fused_kernel/core/data/array.h b/include/fused_kernel/core/data/array.h index c8ad04ec..23cc0359 100644 --- a/include/fused_kernel/core/data/array.h +++ b/include/fused_kernel/core/data/array.h @@ -69,7 +69,7 @@ namespace fk { template union ArrayVector { - static_assert(std::is_fundamental_v, "ArrayVector can only be used with fundamental types"); + static_assert(fk::validScalar, "ArrayVector can only be used with fundamental or reduced float types"); enum { size = 1 }; T at[1]; struct { @@ -94,7 +94,7 @@ namespace fk { template union ArrayVector { - static_assert(std::is_fundamental_v, "ArrayVector can only be used with fundamental types"); + static_assert(fk::validScalar, "ArrayVector can only be used with fundamental or reduced float types"); enum { size = 2 }; T at[2]; struct { @@ -128,7 +128,7 @@ namespace fk { template union ArrayVector { - static_assert(std::is_fundamental_v, "ArrayVector can only be used with fundamental types"); + static_assert(fk::validScalar, "ArrayVector can only be used with fundamental or reduced float types"); enum { size = 3 }; T at[3]; struct { @@ -163,7 +163,7 @@ namespace fk { template union ArrayVector { - static_assert(std::is_fundamental_v, "ArrayVector can only be used with fundamental types"); + static_assert(fk::validScalar, "ArrayVector can only be used with fundamental or reduced float types"); enum { size = 4 }; T at[4]; struct { diff --git a/include/fused_kernel/core/data/ptr_nd.h b/include/fused_kernel/core/data/ptr_nd.h index 0e70f253..3504ac22 100644 --- a/include/fused_kernel/core/data/ptr_nd.h +++ b/include/fused_kernel/core/data/ptr_nd.h @@ -539,7 +539,7 @@ namespace fk { return *At::cr_point(p, ptr_pinned); } else { throw std::runtime_error("Cannot access data in Device memory from host code"); - return make_set(0); + return T{}; } } diff --git a/include/fused_kernel/core/data/reduced_float_types.h b/include/fused_kernel/core/data/reduced_float_types.h new file mode 100644 index 00000000..81dce339 --- /dev/null +++ b/include/fused_kernel/core/data/reduced_float_types.h @@ -0,0 +1,545 @@ +/* Copyright 2026 Oscar Amoros Huguet, Johnny Nunez + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. */ + +#ifndef FK_REDUCED_FLOAT_TYPES +#define FK_REDUCED_FLOAT_TYPES + +#include +#include + +#if defined(__NVCC__) && !defined(NVRTC_COMPILER) +// Native types are used only as runtime fast paths and for interop; the fk types below are +// the element types of the library on every backend. cuda_fp8.h/cuda_fp4.h are deliberately +// NOT included: fp8/fp4 conversions are pure software on all paths (see PR notes), which keeps +// the API independent of the installed toolkit version. +#include +#include +#endif + +#if !defined(NVRTC_COMPILER) +#include +#endif + +namespace fk { + + namespace detail_rf { + // Description of a reduced floating point format plus its conversion policy. + // HAS_INF: the format encodes infinities (fp16, bf16, e5m2). + // HAS_NAN: the format encodes NaN (all but e2m1). + // SAT_FINITE: out of range values (including infinities) saturate to the maximum finite + // value instead of producing Inf, matching NVIDIA's __NV_SATFINITE fp8/fp4 + // converters. NaN handling under SAT_FINITE is canonical (sign dropped). + template + struct MiniFloatFormat { + static constexpr int EXP_BITS = EXP_BITS_; + static constexpr int MAN_BITS = MAN_BITS_; + static constexpr int BIAS = (1 << (EXP_BITS - 1)) - 1; + static constexpr bool HAS_INF = HAS_INF_; + static constexpr bool HAS_NAN = HAS_NAN_; + static constexpr bool SAT_FINITE = SAT_FINITE_; + static constexpr unsigned int SIGN_MASK = 1u << (EXP_BITS + MAN_BITS); + static constexpr unsigned int EXP_MASK = ((1u << EXP_BITS) - 1u) << MAN_BITS; + static constexpr unsigned int MAN_MASK = (1u << MAN_BITS) - 1u; + // Maximum finite value (sign bit not included): + // - Formats with Inf: largest exponent field is reserved, so exp = max-1, mantissa all ones. + // - e4m3 (no Inf, NaN = exp and mantissa all ones): everything but the NaN code is finite. + // - e2m1 (no Inf, no NaN): every code is finite. + static constexpr unsigned int MAX_FINITE = + HAS_INF ? ((((EXP_MASK >> MAN_BITS) - 1u) << MAN_BITS) | MAN_MASK) + : (HAS_NAN ? (EXP_MASK | (MAN_MASK - 1u)) : (EXP_MASK | MAN_MASK)); + // Canonical NaN produced when converting a NaN into this format. NVIDIA's fp8 + // converters produce 0x7F for both e4m3 and e5m2; fp16/bf16 use a quiet NaN. + // For e2m1 (no NaN) NVIDIA converts NaN to +MAX_FINITE (+6.0). + static constexpr unsigned int NAN_CODE = + HAS_NAN ? (SAT_FINITE ? (EXP_MASK | MAN_MASK) : (EXP_MASK | (1u << (MAN_BITS - 1)))) + : MAX_FINITE; + static constexpr unsigned int INF_CODE = EXP_MASK; + static constexpr unsigned int MIN_NORMAL = 1u << MAN_BITS; + }; + + using FmtF16 = MiniFloatFormat<5, 10, true, true, false>; + using FmtBF16 = MiniFloatFormat<8, 7, true, true, false>; + using FmtE4M3 = MiniFloatFormat<4, 3, false, true, true>; + using FmtE5M2 = MiniFloatFormat<5, 2, true, true, true>; + using FmtE2M1 = MiniFloatFormat<2, 1, false, false, true>; + + // Round a 64 bit significand to the lowest "shift" bits using round to nearest, ties to even. + FK_HOST_DEVICE_CNST unsigned long long roundShiftRNE(const unsigned long long value, const int shift) { + if (shift <= 0) { + return value << (-shift); + } else if (shift >= 64) { + return 0ull; + } + const unsigned long long truncated = value >> shift; + const unsigned long long guard = (value >> (shift - 1)) & 1ull; + const unsigned long long stickyMask = (1ull << (shift - 1)) - 1ull; + const bool sticky = (value & stickyMask) != 0ull; + const bool roundUp = guard && (sticky || ((truncated & 1ull) != 0ull)); + return truncated + (roundUp ? 1ull : 0ull); + } + + // Encode an IEEE 754 binary32/binary64 value into the target reduced format, with a + // single rounding step (round to nearest even) and the format's overflow/NaN policy. + template + FK_HOST_DEVICE_CNST unsigned int encode(const FloatType value) { + static_assert(std::is_same_v || std::is_same_v, + "encode only accepts float or double sources"); + constexpr bool IS_F32 = std::is_same_v; + using UIntT = std::conditional_t; + constexpr int SRC_MAN_BITS = IS_F32 ? 23 : 52; + constexpr int SRC_EXP_BITS = IS_F32 ? 8 : 11; + constexpr int SRC_BIAS = IS_F32 ? 127 : 1023; + constexpr UIntT SRC_MAN_MASK = (UIntT(1) << SRC_MAN_BITS) - 1; + constexpr unsigned int SRC_EXP_MAX = (1u << SRC_EXP_BITS) - 1u; + + const UIntT srcBits = __builtin_bit_cast(UIntT, value); + const unsigned int sign = static_cast(srcBits >> (SRC_MAN_BITS + SRC_EXP_BITS)) & 1u; + const unsigned int srcExp = static_cast(srcBits >> SRC_MAN_BITS) & SRC_EXP_MAX; + const unsigned long long srcMan = static_cast(srcBits & SRC_MAN_MASK); + const unsigned int signBits = sign << (Fmt::EXP_BITS + Fmt::MAN_BITS); + + if (srcExp == SRC_EXP_MAX) { + if (srcMan != 0ull) { // NaN + if constexpr (!Fmt::HAS_NAN) { + return Fmt::NAN_CODE; // e2m1: NaN converts to +MAX_FINITE, sign dropped + } else if constexpr (Fmt::SAT_FINITE) { + return Fmt::NAN_CODE; // fp8: canonical NaN, sign dropped + } else { + // fp16/bf16: quiet NaN preserving the truncated payload + const unsigned int payload = + static_cast(srcMan >> (SRC_MAN_BITS - Fmt::MAN_BITS)) & Fmt::MAN_MASK; + return signBits | Fmt::EXP_MASK | (1u << (Fmt::MAN_BITS - 1)) | payload; + } + } else { // Inf + if constexpr (Fmt::HAS_INF && !Fmt::SAT_FINITE) { + return signBits | Fmt::INF_CODE; + } else { + return signBits | Fmt::MAX_FINITE; + } + } + } + + // Finite input: build (unbiased exponent, 1.xxx significand) normalizing subnormal sources + unsigned long long significand = srcMan; + int exponent = static_cast(srcExp) - SRC_BIAS; + if (srcExp == 0u) { + if (significand == 0ull) { + return signBits; // +/- 0 + } + exponent = 1 - SRC_BIAS; + while ((significand & (UIntT(1) << SRC_MAN_BITS)) == 0ull) { + significand <<= 1; + --exponent; + } + } else { + significand |= (UIntT(1) << SRC_MAN_BITS); + } + + const int targetBiasedExp = exponent + Fmt::BIAS; + unsigned long long candidate = 0ull; + if (targetBiasedExp >= 1) { + // Normal range: round the significand to MAN_BITS fractional bits + const unsigned long long rounded = roundShiftRNE(significand, SRC_MAN_BITS - Fmt::MAN_BITS); + // rounded holds the implicit bit at position MAN_BITS (or MAN_BITS+1 after a carry); + // adding (exp-1)<(targetBiasedExp - 1) << Fmt::MAN_BITS); + } else { + // Subnormal range for the target: shift out extra bits before rounding. + // A round up out of the subnormal range naturally produces the minimum normal. + const int extraShift = 1 - targetBiasedExp; + candidate = roundShiftRNE(significand, (SRC_MAN_BITS - Fmt::MAN_BITS) + extraShift); + } + if (candidate > static_cast(Fmt::MAX_FINITE)) { + if constexpr (Fmt::HAS_INF && !Fmt::SAT_FINITE) { + return signBits | Fmt::INF_CODE; + } else { + return signBits | Fmt::MAX_FINITE; + } + } + return signBits | static_cast(candidate); + } + + // Decode a reduced format value into IEEE 754 binary32. Every finite value of every + // supported format is exactly representable in binary32. + template + FK_HOST_DEVICE_CNST float decode(const unsigned int bits) { + const unsigned int sign = (bits & Fmt::SIGN_MASK) ? 1u : 0u; + const unsigned int expField = (bits & Fmt::EXP_MASK) >> Fmt::MAN_BITS; + const unsigned int mantissa = bits & Fmt::MAN_MASK; + const unsigned int maxExpField = (1u << Fmt::EXP_BITS) - 1u; + const unsigned int signBit32 = sign << 31; + + if constexpr (Fmt::HAS_INF) { + if (expField == maxExpField) { + if (mantissa == 0u) { + return __builtin_bit_cast(float, signBit32 | 0x7F800000u); + } + // Quiet NaN preserving the payload in the top mantissa bits + return __builtin_bit_cast(float, signBit32 | 0x7FC00000u | (mantissa << (23 - Fmt::MAN_BITS))); + } + } else if constexpr (Fmt::HAS_NAN) { + if ((bits & (Fmt::EXP_MASK | Fmt::MAN_MASK)) == (Fmt::EXP_MASK | Fmt::MAN_MASK)) { + return __builtin_bit_cast(float, signBit32 | 0x7FC00000u); + } + } + + if (expField == 0u) { + if (mantissa == 0u) { + return __builtin_bit_cast(float, signBit32); // +/- 0 + } + // Subnormal: normalize. The result is usually a binary32 normal value, except + // for the smallest bf16 subnormals (down to 2^-133), which land below the + // binary32 normal range (2^-126) and must be emitted as binary32 subnormals. + int exponent = 1 - Fmt::BIAS; + unsigned int man = mantissa; + while ((man & Fmt::MIN_NORMAL) == 0u) { + man <<= 1; + --exponent; + } + man &= Fmt::MAN_MASK; + const int exp32 = exponent + 127; + if (exp32 >= 1) { + return __builtin_bit_cast(float, signBit32 | (static_cast(exp32) << 23) + | (man << (23 - Fmt::MAN_BITS))); + } + // Binary32 subnormal: shift the full significand (implicit bit included) into + // place. The shift never discards set bits: the smallest representable input + // (bf16 2^-133) still has 16 trailing zero bits available. + const unsigned int sig32 = (1u << 23) | (man << (23 - Fmt::MAN_BITS)); + return __builtin_bit_cast(float, signBit32 | (sig32 >> (1 - exp32))); + } + const int exponent = static_cast(expField) - Fmt::BIAS; + return __builtin_bit_cast(float, signBit32 | (static_cast(exponent + 127) << 23) + | (mantissa << (23 - Fmt::MAN_BITS))); + } + } // namespace detail_rf + + // Reduced precision floating point types, usable as element types in any FKL pipeline on + // every backend. They are the same fk-owned types under nvcc, g++, clang and MSVC (identical + // layout, traits and semantics), fully constexpr through a software conversion core, with + // native CUDA fast paths taken automatically at runtime under nvcc (fp16/bf16). + // Layouts are bit-compatible with the CUDA types (__half, __nv_bfloat16, __nv_fp8_e4m3, + // __nv_fp8_e5m2, __nv_fp4_e2m1), so device buffers can be reinterpreted freely. +// VALUE_MASK canonicalizes the stored bits on construction from raw bits: for fp4 the upper +// nibble is unspecified storage (NVIDIA converters write 0 and ignore it on read), and keeping +// it would make numerically equal values compare bit-different through toBits(). +#define FK_RF_COMMON(TypeName, StorageT, Fmt, VALUE_MASK) \ + private: \ + StorageT bits_; \ + struct FromBitsTag {}; \ + FK_HOST_DEVICE_CNST TypeName(const StorageT bits, const FromBitsTag&) \ + : bits_(static_cast(bits & VALUE_MASK)) {} \ + public: \ + using Format = Fmt; \ + using StorageType = StorageT; \ + TypeName() = default; \ + FK_HOST_DEVICE_STATIC constexpr TypeName fromBits(const StorageT bits) { \ + return TypeName(bits, FromBitsTag{}); \ + } \ + FK_HOST_DEVICE_CNST StorageT bits() const { return bits_; } + + struct Fp16 { + FK_RF_COMMON(Fp16, unsigned short, detail_rf::FmtF16, 0xFFFFu) + FK_HOST_DEVICE_CNST explicit Fp16(const float value) { +#if defined(__NVCC__) && !defined(NVRTC_COMPILER) + if (!__builtin_is_constant_evaluated()) { + bits_ = __builtin_bit_cast(unsigned short, __half(value)); + return; + } +#endif + bits_ = static_cast(detail_rf::encode(value)); + } + FK_HOST_DEVICE_CNST explicit Fp16(const double value) + : bits_(static_cast(detail_rf::encode(value))) {} + template >> + FK_HOST_DEVICE_CNST explicit Fp16(const I value) : Fp16(static_cast(value)) {} + FK_HOST_DEVICE_CNST operator float() const { +#if defined(__NVCC__) && !defined(NVRTC_COMPILER) + if (!__builtin_is_constant_evaluated()) { + return static_cast(__builtin_bit_cast(__half, bits_)); + } +#endif + return detail_rf::decode(bits_); + } + }; + + struct Bf16 { + FK_RF_COMMON(Bf16, unsigned short, detail_rf::FmtBF16, 0xFFFFu) + FK_HOST_DEVICE_CNST explicit Bf16(const float value) { +#if defined(__NVCC__) && !defined(NVRTC_COMPILER) + if (!__builtin_is_constant_evaluated()) { + bits_ = __builtin_bit_cast(unsigned short, __nv_bfloat16(value)); + return; + } +#endif + bits_ = static_cast(detail_rf::encode(value)); + } + FK_HOST_DEVICE_CNST explicit Bf16(const double value) + : bits_(static_cast(detail_rf::encode(value))) {} + template >> + FK_HOST_DEVICE_CNST explicit Bf16(const I value) : Bf16(static_cast(value)) {} + FK_HOST_DEVICE_CNST operator float() const { +#if defined(__NVCC__) && !defined(NVRTC_COMPILER) + if (!__builtin_is_constant_evaluated()) { + return static_cast(__builtin_bit_cast(__nv_bfloat16, bits_)); + } +#endif + return detail_rf::decode(bits_); + } + }; + + struct Fp8E4m3 { + FK_RF_COMMON(Fp8E4m3, unsigned char, detail_rf::FmtE4M3, 0xFFu) + FK_HOST_DEVICE_CNST explicit Fp8E4m3(const float value) + : bits_(static_cast(detail_rf::encode(value))) {} + FK_HOST_DEVICE_CNST explicit Fp8E4m3(const double value) + : bits_(static_cast(detail_rf::encode(value))) {} + FK_HOST_DEVICE_CNST explicit Fp8E4m3(const Fp16 value) : Fp8E4m3(static_cast(value)) {} + FK_HOST_DEVICE_CNST explicit Fp8E4m3(const Bf16 value) : Fp8E4m3(static_cast(value)) {} + template >> + FK_HOST_DEVICE_CNST explicit Fp8E4m3(const I value) : Fp8E4m3(static_cast(value)) {} + FK_HOST_DEVICE_CNST explicit operator float() const { return detail_rf::decode(bits_); } + }; + + struct Fp8E5m2 { + FK_RF_COMMON(Fp8E5m2, unsigned char, detail_rf::FmtE5M2, 0xFFu) + FK_HOST_DEVICE_CNST explicit Fp8E5m2(const float value) + : bits_(static_cast(detail_rf::encode(value))) {} + FK_HOST_DEVICE_CNST explicit Fp8E5m2(const double value) + : bits_(static_cast(detail_rf::encode(value))) {} + FK_HOST_DEVICE_CNST explicit Fp8E5m2(const Fp16 value) : Fp8E5m2(static_cast(value)) {} + FK_HOST_DEVICE_CNST explicit Fp8E5m2(const Bf16 value) : Fp8E5m2(static_cast(value)) {} + template >> + FK_HOST_DEVICE_CNST explicit Fp8E5m2(const I value) : Fp8E5m2(static_cast(value)) {} + FK_HOST_DEVICE_CNST explicit operator float() const { return detail_rf::decode(bits_); } + }; + + // One element per byte, value in the low nibble, matching CUDA's scalar __nv_fp4_e2m1. + // The truly packed two-per-byte format (__nv_fp4x2_e2m1) cannot be an element type in the + // RawPtr model (one addressable element per T); use a fused dequantizing Read for that. + struct Fp4E2m1 { + FK_RF_COMMON(Fp4E2m1, unsigned char, detail_rf::FmtE2M1, 0x0Fu) + FK_HOST_DEVICE_CNST explicit Fp4E2m1(const float value) + : bits_(static_cast(detail_rf::encode(value))) {} + FK_HOST_DEVICE_CNST explicit Fp4E2m1(const double value) + : bits_(static_cast(detail_rf::encode(value))) {} + FK_HOST_DEVICE_CNST explicit Fp4E2m1(const Fp16 value) : Fp4E2m1(static_cast(value)) {} + FK_HOST_DEVICE_CNST explicit Fp4E2m1(const Bf16 value) : Fp4E2m1(static_cast(value)) {} + template >> + FK_HOST_DEVICE_CNST explicit Fp4E2m1(const I value) : Fp4E2m1(static_cast(value)) {} + // bits_ is canonical (upper nibble masked on construction), so no masking is needed here. + FK_HOST_DEVICE_CNST explicit operator float() const { return detail_rf::decode(bits_); } + }; + +#undef FK_RF_COMMON + + using fp16 = Fp16; + using bf16 = Bf16; + using fp8_e4m3 = Fp8E4m3; + using fp8_e5m2 = Fp8E5m2; + using fp4_e2m1 = Fp4E2m1; + + static_assert(sizeof(fp16) == 2 && alignof(fp16) == 2, "fp16 layout must match __half"); + static_assert(sizeof(bf16) == 2 && alignof(bf16) == 2, "bf16 layout must match __nv_bfloat16"); + static_assert(sizeof(fp8_e4m3) == 1 && alignof(fp8_e4m3) == 1, "fp8_e4m3 layout must match __nv_fp8_e4m3"); + static_assert(sizeof(fp8_e5m2) == 1 && alignof(fp8_e5m2) == 1, "fp8_e5m2 layout must match __nv_fp8_e5m2"); + static_assert(sizeof(fp4_e2m1) == 1 && alignof(fp4_e2m1) == 1, "fp4_e2m1 layout must match __nv_fp4_e2m1"); + static_assert(std::is_trivially_copyable_v && std::is_trivially_copyable_v && + std::is_trivially_copyable_v && std::is_trivially_copyable_v && + std::is_trivially_copyable_v, "Reduced float types must be trivially copyable"); + + // Type classification. validScalar/validFloatingPoint are the library-wide replacements for + // std::is_fundamental_v / std::is_floating_point_v wherever reduced floats must be admitted. + using ReducedFloatTypes = TypeList; + template + constexpr bool isReducedFloat = one_of_v; + // Types with a full arithmetic and comparison operator surface (fp8/fp4 are conversion-only, + // exactly like the CUDA types, which define no operators for them). + template + constexpr bool isArithmeticReducedFloat = one_of_v>; + template + constexpr bool validScalar = std::is_fundamental_v || isReducedFloat; + template + constexpr bool validFloatingPoint = std::is_floating_point_v || isReducedFloat; + + template + struct ReducedFloatTraits { + using Format = typename T::Format; + using StorageType = typename T::StorageType; + static constexpr StorageType maxBits = static_cast(Format::MAX_FINITE); + static constexpr StorageType lowestBits = static_cast(Format::SIGN_MASK | Format::MAX_FINITE); + static constexpr StorageType minNormalBits = static_cast(Format::MIN_NORMAL); + static constexpr StorageType minSubnormalBits = static_cast(1u); + static constexpr StorageType quietNaNBits = static_cast(Format::NAN_CODE); + static constexpr StorageType infBits = static_cast(Format::INF_CODE); + static constexpr bool hasInf = Format::HAS_INF; + static constexpr bool hasNaN = Format::HAS_NAN; + FK_HOST_DEVICE_FUSE bool isNaN(const T& value) { + if constexpr (Format::HAS_INF) { + return (value.bits() & Format::EXP_MASK) == Format::EXP_MASK && + (value.bits() & Format::MAN_MASK) != 0u; + } else if constexpr (Format::HAS_NAN) { + return (value.bits() & (Format::EXP_MASK | Format::MAN_MASK)) == + (Format::EXP_MASK | Format::MAN_MASK); + } else { + return false; + } + } + FK_HOST_DEVICE_FUSE bool isInf(const T& value) { + if constexpr (Format::HAS_INF) { + return (value.bits() & (Format::EXP_MASK | Format::MAN_MASK)) == Format::EXP_MASK; + } else { + return false; + } + } + FK_HOST_DEVICE_FUSE T abs(const T& value) { + return T::fromBits(static_cast(value.bits() & ~Format::SIGN_MASK)); + } + }; + + template >> + FK_HOST_DEVICE_CNST auto toBits(const T& value) { + return value.bits(); + } + +#if defined(__NVCC__) && !defined(NVRTC_COMPILER) + // Zero cost interop with the CUDA native types (bit-identical layouts). + // NOTE: mixed expressions between fk::fp16 and __half (or fk::bf16 and __nv_bfloat16) are + // ambiguous by design - both types convert implicitly to several built in arithmetic types. + // Convert one operand explicitly with toNative()/fromNative() (both are free, bit casts). + FK_HOST_DEVICE_CNST __half toNative(const fp16& value) { + return __builtin_bit_cast(__half, value.bits()); + } + FK_HOST_DEVICE_CNST __nv_bfloat16 toNative(const bf16& value) { + return __builtin_bit_cast(__nv_bfloat16, value.bits()); + } + FK_HOST_DEVICE_CNST fp16 fromNative(const __half& value) { + return fp16::fromBits(__builtin_bit_cast(unsigned short, value)); + } + FK_HOST_DEVICE_CNST bf16 fromNative(const __nv_bfloat16& value) { + return bf16::fromBits(__builtin_bit_cast(unsigned short, value)); + } +#endif + +} // namespace fk + +// Arithmetic and comparisons for fp16/bf16: promote to float, compute, demote (RN). +// binary32 has more than 2p+2 mantissa bits for both formats, so the results are the +// correctly rounded reduced precision results (identical to CUDA's host/device operators). +// Declared at global scope like the vector operators in vector_utils.h: declaring them inside +// namespace fk would hide every global operator from unqualified lookup within fk. +#define FK_RF_ARITHMETIC(TypeName) \ + FK_HOST_DEVICE_CNST TypeName operator+(const TypeName& a, const TypeName& b) { \ + return TypeName(static_cast(a) + static_cast(b)); \ + } \ + FK_HOST_DEVICE_CNST TypeName operator-(const TypeName& a, const TypeName& b) { \ + return TypeName(static_cast(a) - static_cast(b)); \ + } \ + FK_HOST_DEVICE_CNST TypeName operator*(const TypeName& a, const TypeName& b) { \ + return TypeName(static_cast(a) * static_cast(b)); \ + } \ + FK_HOST_DEVICE_CNST TypeName operator/(const TypeName& a, const TypeName& b) { \ + return TypeName(static_cast(a) / static_cast(b)); \ + } \ + FK_HOST_DEVICE_CNST TypeName operator-(const TypeName& a) { \ + return TypeName::fromBits(static_cast( \ + a.bits() ^ TypeName::Format::SIGN_MASK)); \ + } \ + FK_HOST_DEVICE_CNST TypeName& operator+=(TypeName& a, const TypeName& b) { a = a + b; return a; } \ + FK_HOST_DEVICE_CNST TypeName& operator-=(TypeName& a, const TypeName& b) { a = a - b; return a; } \ + FK_HOST_DEVICE_CNST TypeName& operator*=(TypeName& a, const TypeName& b) { a = a * b; return a; } \ + FK_HOST_DEVICE_CNST TypeName& operator/=(TypeName& a, const TypeName& b) { a = a / b; return a; } \ + FK_HOST_DEVICE_CNST bool operator==(const TypeName& a, const TypeName& b) { \ + return static_cast(a) == static_cast(b); \ + } \ + FK_HOST_DEVICE_CNST bool operator!=(const TypeName& a, const TypeName& b) { \ + return static_cast(a) != static_cast(b); \ + } \ + FK_HOST_DEVICE_CNST bool operator<(const TypeName& a, const TypeName& b) { \ + return static_cast(a) < static_cast(b); \ + } \ + FK_HOST_DEVICE_CNST bool operator>(const TypeName& a, const TypeName& b) { \ + return static_cast(a) > static_cast(b); \ + } \ + FK_HOST_DEVICE_CNST bool operator<=(const TypeName& a, const TypeName& b) { \ + return static_cast(a) <= static_cast(b); \ + } \ + FK_HOST_DEVICE_CNST bool operator>=(const TypeName& a, const TypeName& b) { \ + return static_cast(a) >= static_cast(b); \ + } + +FK_RF_ARITHMETIC(fk::Fp16) +FK_RF_ARITHMETIC(fk::Bf16) +#undef FK_RF_ARITHMETIC + +#if !defined(NVRTC_COMPILER) +// std::numeric_limits for the reduced float types: makes fk::maxValue/minValue (vlimits.h) and +// cxp::minValue/maxValue work unmodified, and is the canonical source for the limit values. +#define FK_RF_NUMERIC_LIMITS(TypeName, DIGITS_, DIGITS10_, MAX_DIGITS10_, MAX_EXP_, MAX_EXP10_, \ + MIN_EXP_, MIN_EXP10_) \ + template <> \ + class std::numeric_limits { \ + public: \ + using Traits = fk::ReducedFloatTraits; \ + static constexpr bool is_specialized = true; \ + static constexpr bool is_signed = true; \ + static constexpr bool is_integer = false; \ + static constexpr bool is_exact = false; \ + static constexpr bool has_infinity = Traits::hasInf; \ + static constexpr bool has_quiet_NaN = Traits::hasNaN; \ + static constexpr bool has_signaling_NaN = false; \ + static constexpr std::float_denorm_style has_denorm = std::denorm_present; \ + static constexpr bool has_denorm_loss = false; \ + static constexpr std::float_round_style round_style = std::round_to_nearest; \ + static constexpr bool is_iec559 = false; \ + static constexpr bool is_bounded = true; \ + static constexpr bool is_modulo = false; \ + static constexpr bool traps = false; \ + static constexpr bool tinyness_before = false; \ + static constexpr int digits = DIGITS_; \ + static constexpr int digits10 = DIGITS10_; \ + static constexpr int max_digits10 = MAX_DIGITS10_; \ + static constexpr int max_exponent = MAX_EXP_; \ + static constexpr int max_exponent10 = MAX_EXP10_; \ + static constexpr int min_exponent = MIN_EXP_; \ + static constexpr int min_exponent10 = MIN_EXP10_; \ + static constexpr int radix = 2; \ + static constexpr TypeName max() noexcept { return TypeName::fromBits(Traits::maxBits); } \ + static constexpr TypeName lowest() noexcept { return TypeName::fromBits(Traits::lowestBits); } \ + static constexpr TypeName min() noexcept { return TypeName::fromBits(Traits::minNormalBits); } \ + static constexpr TypeName denorm_min() noexcept { \ + return TypeName::fromBits(Traits::minSubnormalBits); \ + } \ + static constexpr TypeName epsilon() noexcept { \ + return TypeName(1.0f / static_cast(1 << (DIGITS_ - 1))); \ + } \ + static constexpr TypeName round_error() noexcept { return TypeName(0.5f); } \ + static constexpr TypeName infinity() noexcept { return TypeName::fromBits(Traits::infBits); } \ + static constexpr TypeName quiet_NaN() noexcept { \ + return TypeName::fromBits(Traits::quietNaNBits); \ + } \ + static constexpr TypeName signaling_NaN() noexcept { \ + return TypeName::fromBits(Traits::quietNaNBits); \ + } \ + }; + +FK_RF_NUMERIC_LIMITS(fk::Fp16, 11, 3, 5, 16, 4, -13, -4) +FK_RF_NUMERIC_LIMITS(fk::Bf16, 8, 2, 4, 128, 38, -125, -37) +FK_RF_NUMERIC_LIMITS(fk::Fp8E4m3, 4, 0, 3, 9, 2, -5, -1) +FK_RF_NUMERIC_LIMITS(fk::Fp8E5m2, 3, 0, 2, 16, 4, -13, -4) +FK_RF_NUMERIC_LIMITS(fk::Fp4E2m1, 2, 0, 2, 3, 0, 1, 0) +#undef FK_RF_NUMERIC_LIMITS +#endif // !NVRTC_COMPILER + +#endif // FK_REDUCED_FLOAT_TYPES diff --git a/include/fused_kernel/core/data/vector_types.h b/include/fused_kernel/core/data/vector_types.h index 3ebd2258..30dab2b5 100644 --- a/include/fused_kernel/core/data/vector_types.h +++ b/include/fused_kernel/core/data/vector_types.h @@ -18,6 +18,7 @@ #include #include +#include namespace fk { @@ -257,6 +258,119 @@ namespace fk { struct alignas(16) Double4 { double x, y, z, w; }; + + // Reduced precision vector types. Like the Bool vectors, they are fk-owned aggregates on + // every backend: CUDA provides no 3/4 channel variants, and its 2 channel types (__half2, + // __nv_bfloat162) are non-aggregate classes that would diverge per backend. Layouts match + // the CUDA counterparts where one exists (__half2/__nv_bfloat162: 4 bytes, alignment 4; + // __nv_fp8x2_*: 2/2; __nv_fp8x4_*: 4/4). They live in the fk namespace only, because CUDA + // headers define global half/half2/nv_bfloat162 typedefs that must not be shadowed. + struct Fp16_1 { + fp16 x; + }; + + struct alignas(4) Fp16_2 { + fp16 x, y; + }; + + struct Fp16_3 { + fp16 x, y, z; + }; + + struct alignas(8) Fp16_4 { + fp16 x, y, z, w; + }; + + struct Bf16_1 { + bf16 x; + }; + + struct alignas(4) Bf16_2 { + bf16 x, y; + }; + + struct Bf16_3 { + bf16 x, y, z; + }; + + struct alignas(8) Bf16_4 { + bf16 x, y, z, w; + }; + + struct Fp8E4m3_1 { + fp8_e4m3 x; + }; + + struct alignas(2) Fp8E4m3_2 { + fp8_e4m3 x, y; + }; + + struct Fp8E4m3_3 { + fp8_e4m3 x, y, z; + }; + + struct alignas(4) Fp8E4m3_4 { + fp8_e4m3 x, y, z, w; + }; + + struct Fp8E5m2_1 { + fp8_e5m2 x; + }; + + struct alignas(2) Fp8E5m2_2 { + fp8_e5m2 x, y; + }; + + struct Fp8E5m2_3 { + fp8_e5m2 x, y, z; + }; + + struct alignas(4) Fp8E5m2_4 { + fp8_e5m2 x, y, z, w; + }; + + struct Fp4E2m1_1 { + fp4_e2m1 x; + }; + + struct alignas(2) Fp4E2m1_2 { + fp4_e2m1 x, y; + }; + + struct Fp4E2m1_3 { + fp4_e2m1 x, y, z; + }; + + struct alignas(4) Fp4E2m1_4 { + fp4_e2m1 x, y, z, w; + }; + + using fp16_1 = Fp16_1; + using fp16_2 = Fp16_2; + using fp16_3 = Fp16_3; + using fp16_4 = Fp16_4; + using bf16_1 = Bf16_1; + using bf16_2 = Bf16_2; + using bf16_3 = Bf16_3; + using bf16_4 = Bf16_4; + using fp8_e4m3_1 = Fp8E4m3_1; + using fp8_e4m3_2 = Fp8E4m3_2; + using fp8_e4m3_3 = Fp8E4m3_3; + using fp8_e4m3_4 = Fp8E4m3_4; + using fp8_e5m2_1 = Fp8E5m2_1; + using fp8_e5m2_2 = Fp8E5m2_2; + using fp8_e5m2_3 = Fp8E5m2_3; + using fp8_e5m2_4 = Fp8E5m2_4; + using fp4_e2m1_1 = Fp4E2m1_1; + using fp4_e2m1_2 = Fp4E2m1_2; + using fp4_e2m1_3 = Fp4E2m1_3; + using fp4_e2m1_4 = Fp4E2m1_4; + + static_assert(sizeof(fp16_2) == 4 && alignof(fp16_2) == 4, "fp16_2 layout must match __half2"); + static_assert(sizeof(bf16_2) == 4 && alignof(bf16_2) == 4, "bf16_2 layout must match __nv_bfloat162"); + static_assert(sizeof(fp16_4) == 8 && sizeof(fp16_3) == 6, "fp16_3/fp16_4 unexpected layout"); + static_assert(sizeof(fp8_e4m3_2) == 2 && alignof(fp8_e4m3_2) == 2, "fp8x2 layout must match __nv_fp8x2_e4m3"); + static_assert(sizeof(fp8_e4m3_4) == 4 && alignof(fp8_e4m3_4) == 4, "fp8x4 layout must match __nv_fp8x4_e4m3"); } // namespace fk #if defined(__NVCC__) diff --git a/include/fused_kernel/core/execution_model/thread_fusion.h b/include/fused_kernel/core/execution_model/thread_fusion.h index 8c4d139b..a9da71e3 100644 --- a/include/fused_kernel/core/execution_model/thread_fusion.h +++ b/include/fused_kernel/core/execution_model/thread_fusion.h @@ -47,17 +47,69 @@ namespace fk { Times bigger can be: 1, 2, 4 */ - using TFSourceTypes = TypeListCat_t; - using TFBiggerTypes = TypeList; + // The reduced float entries are appended as a self contained block AFTER the standard + // types, so the positional correspondence of the existing 65 pairs cannot shift. + // Mappings follow the size precedents above: 2 byte scalars fuse x2 (short -> short2), + // 1 byte scalars fuse x4 (char -> char4), everything of 4+ bytes stays as is. + using RFTFSourceTypes = TypeList; + using RFTFBiggerTypes = TypeList; + + using TFSourceTypes = TypeListCat_t; + using TFBiggerTypes = TypeListCat_t< + TypeList, + RFTFBiggerTypes /*scalars*/, RFTFBiggerTypes /*x1*/, RFTFBiggerTypes /*x2*/, + RFVThree /*x3: no fusion*/, RFVFour /*x4: no fusion*/>; template using FilteredType_t = std::conditional_t, typename VectorTraits::base, T>; template using TFBiggerType_t = EquivalentType_t, TFSourceTypes, TFBiggerTypes>; + // The source/bigger mapping is positional: a misaligned insertion corrupts vectorized loads + // silently. These asserts pin every reduced float mapping and the legacy row boundaries. + static_assert(TFSourceTypes::size == TFBiggerTypes::size, + "TFSourceTypes and TFBiggerTypes must have the same number of entries"); + static_assert(std::is_same_v, bool4> && std::is_same_v, double> && + std::is_same_v, bool4> && std::is_same_v, double> && + std::is_same_v, bool4> && std::is_same_v, double2> && + std::is_same_v, bool3> && std::is_same_v, double3> && + std::is_same_v, bool4> && std::is_same_v, double4>, + "Legacy thread fusion row boundaries shifted"); + static_assert(std::is_same_v, fp16_2> && + std::is_same_v, bf16_2> && + std::is_same_v, fp8_e4m3_4> && + std::is_same_v, fp8_e5m2_4> && + std::is_same_v, fp4_e2m1_4>, + "Reduced float scalar thread fusion mappings shifted"); + static_assert(std::is_same_v, fp16_2> && + std::is_same_v, bf16_2> && + std::is_same_v, fp8_e4m3_4> && + std::is_same_v, fp8_e5m2_4> && + std::is_same_v, fp4_e2m1_4>, + "Reduced float x1 thread fusion mappings shifted"); + static_assert(std::is_same_v, fp16_2> && + std::is_same_v, bf16_2> && + std::is_same_v, fp8_e4m3_4> && + std::is_same_v, fp8_e5m2_4> && + std::is_same_v, fp4_e2m1_4>, + "Reduced float x2 thread fusion mappings shifted"); + static_assert(std::is_same_v, fp16_3> && + std::is_same_v, bf16_3> && + std::is_same_v, fp8_e4m3_3> && + std::is_same_v, fp8_e5m2_3> && + std::is_same_v, fp4_e2m1_3>, + "Reduced float x3 thread fusion mappings shifted"); + static_assert(std::is_same_v, fp16_4> && + std::is_same_v, bf16_4> && + std::is_same_v, fp8_e4m3_4> && + std::is_same_v, fp8_e5m2_4> && + std::is_same_v, fp4_e2m1_4>, + "Reduced float x4 thread fusion mappings shifted"); + constexpr std::integer_sequence validChannelsSequence; template diff --git a/include/fused_kernel/core/utils/vector_utils.h b/include/fused_kernel/core/utils/vector_utils.h index 91e43f17..4b59180f 100644 --- a/include/fused_kernel/core/utils/vector_utils.h +++ b/include/fused_kernel/core/utils/vector_utils.h @@ -69,6 +69,53 @@ namespace fk { template <> struct VectorType { using type = char4; using type_v = type; }; + // Reduced precision types: hand-written like char/schar because their vector aliases do not + // follow the BaseType##N token-pasting pattern for fp8/fp4. + template <> + struct VectorType { using type = fp16; using type_v = fp16_1; }; + template <> + struct VectorType { using type = fp16_2; using type_v = type; }; + template <> + struct VectorType { using type = fp16_3; using type_v = type; }; + template <> + struct VectorType { using type = fp16_4; using type_v = type; }; + + template <> + struct VectorType { using type = bf16; using type_v = bf16_1; }; + template <> + struct VectorType { using type = bf16_2; using type_v = type; }; + template <> + struct VectorType { using type = bf16_3; using type_v = type; }; + template <> + struct VectorType { using type = bf16_4; using type_v = type; }; + + template <> + struct VectorType { using type = fp8_e4m3; using type_v = fp8_e4m3_1; }; + template <> + struct VectorType { using type = fp8_e4m3_2; using type_v = type; }; + template <> + struct VectorType { using type = fp8_e4m3_3; using type_v = type; }; + template <> + struct VectorType { using type = fp8_e4m3_4; using type_v = type; }; + + template <> + struct VectorType { using type = fp8_e5m2; using type_v = fp8_e5m2_1; }; + template <> + struct VectorType { using type = fp8_e5m2_2; using type_v = type; }; + template <> + struct VectorType { using type = fp8_e5m2_3; using type_v = type; }; + template <> + struct VectorType { using type = fp8_e5m2_4; using type_v = type; }; + + template <> + struct VectorType { using type = fp4_e2m1; using type_v = fp4_e2m1_1; }; + template <> + struct VectorType { using type = fp4_e2m1_2; using type_v = type; }; + template <> + struct VectorType { using type = fp4_e2m1_3; using type_v = type; }; + template <> + struct VectorType { using type = fp4_e2m1_4; using type_v = type; }; + template using VectorType_t = typename VectorType::type; @@ -85,6 +132,11 @@ namespace fk { template using longlong_ = VectorType_t; template using float_ = VectorType_t; template using double_ = VectorType_t; + template using fp16_ = VectorType_t; + template using bf16_ = VectorType_t; + template using fp8_e4m3_ = VectorType_t; + template using fp8_e5m2_ = VectorType_t; + template using fp4_e2m1_ = VectorType_t; template using VectorTypeList = TypeList, uchar_, char_, ushort_, short_, uint_, int_, @@ -99,7 +151,13 @@ namespace fk { using VTwo = VectorTypeList<2>; using VThree = VectorTypeList<3>; using VFour = VectorTypeList<4>; - using VAll = TypeListCat_t; + // Reduced precision vector lists are kept separate so the standard lists (and everything + // positionally derived from them) stay untouched. + using RFVOne = TypeList; + using RFVTwo = TypeList; + using RFVThree = TypeList; + using RFVFour = TypeList; + using VAll = TypeListCat_t; template constexpr bool validCUDAVec = one_of::value; @@ -109,14 +167,15 @@ namespace fk { template FK_HOST_DEVICE_CNST int Channels() { - if constexpr (one_of_v || !validCUDAVec) { + if constexpr (one_of_v || one_of_v || !validCUDAVec) { return 1; - } else if constexpr (one_of_v) { + } else if constexpr (one_of_v || one_of_v) { return 2; - } else if constexpr (one_of_v) { + } else if constexpr (one_of_v || one_of_v) { return 3; } else { - static_assert(one_of_v, "Type T must be a valid CUDA vector type (1, 2, 3, or 4 channels)"); + static_assert(one_of_v || one_of_v, + "Type T must be a valid CUDA vector type (1, 2, 3, or 4 channels)"); return 4; } } @@ -167,6 +226,26 @@ namespace fk { #undef VECTOR_TRAITS +#define VECTOR_TRAITS_RF(BaseType, V1, V2, V3, V4) \ + template <> \ + struct VectorTraits { using base = BaseType; enum { bytes = sizeof(base) }; }; \ + template <> \ + struct VectorTraits { using base = BaseType; enum { bytes = sizeof(base) }; }; \ + template <> \ + struct VectorTraits { using base = BaseType; enum { bytes = sizeof(base) * 2 }; }; \ + template <> \ + struct VectorTraits { using base = BaseType; enum { bytes = sizeof(base) * 3 }; }; \ + template <> \ + struct VectorTraits { using base = BaseType; enum { bytes = sizeof(base) * 4 }; }; + + VECTOR_TRAITS_RF(fp16, fp16_1, fp16_2, fp16_3, fp16_4) + VECTOR_TRAITS_RF(bf16, bf16_1, bf16_2, bf16_3, bf16_4) + VECTOR_TRAITS_RF(fp8_e4m3, fp8_e4m3_1, fp8_e4m3_2, fp8_e4m3_3, fp8_e4m3_4) + VECTOR_TRAITS_RF(fp8_e5m2, fp8_e5m2_1, fp8_e5m2_2, fp8_e5m2_3, fp8_e5m2_4) + VECTOR_TRAITS_RF(fp4_e2m1, fp4_e2m1_1, fp4_e2m1_2, fp4_e2m1_3, fp4_e2m1_4) + +#undef VECTOR_TRAITS_RF + template using VBase = typename VectorTraits::base; @@ -188,7 +267,7 @@ namespace fk { struct vector_at { template FK_HOST_DEVICE_FUSE auto f(const int& idx, const VT& v) - -> std::enable_if_t, VT> { + -> std::enable_if_t, VT> { return v; } template @@ -324,19 +403,19 @@ namespace fk { struct AreSV : public std::false_type {}; template - struct AreSV && fk::validCUDAVec, void>> : public std::true_type {}; + struct AreSV && fk::validCUDAVec, void>> : public std::true_type {}; template struct AreVS : public std::false_type {}; template - struct AreVS && std::is_fundamental_v, void>> : public std::true_type {}; + struct AreVS && fk::validScalar, void>> : public std::true_type {}; template struct AreSS : public std::false_type {}; template - struct AreSS && std::is_fundamental_v, void>> : public std::true_type {}; + struct AreSS && fk::validScalar, void>> : public std::true_type {}; // Utils to check if the type or combination of types can be used with a particular operator template @@ -369,12 +448,34 @@ namespace fk { std::enable_if_t<(AreVVEqCN::value || AreVS::value) && BothIntegrals::value, void>> : public std::true_type {}; + // Closed form availability of the ELEMENT level compound operator (x op= y). Expressed as + // a trait instead of probing the expression: an expression probe would re-enter the vector + // operator templates and create a self dependent constraint. Fundamentals compound among + // themselves and with the arithmetic reduced floats (through their implicit float + // conversion); fp16/bf16 compound only with themselves; fp8/fp4 have no operators at all. + template + constexpr bool hasCompoundBaseOp = + (std::is_fundamental_v && std::is_fundamental_v) || + (std::is_fundamental_v && isArithmeticReducedFloat) || + (isArithmeticReducedFloat && std::is_same_v); + template struct CanCompound : public std::false_type {}; + // Vector op= vector: both bases are registered, so VBase is safe to evaluate on both sides. template struct CanCompound::value || AreVS::value, void>> : public std::true_type {}; + std::enable_if_t::value && + hasCompoundBaseOp, VBase>, void>> : public std::true_type {}; + + // Vector op= scalar: VBase must NOT be evaluated for plain fundamentals - types like + // long double or wchar_t are fundamental (and compound with any fundamental base through + // the built in operators) but have no VectorTraits specialization. + template + struct CanCompound::value && + ((std::is_fundamental_v && std::is_fundamental_v>) || + (isReducedFloat && hasCompoundBaseOp, I2>)), void>> : public std::true_type {}; template struct CanCompoundLogical : public std::false_type {}; @@ -441,11 +542,14 @@ inline constexpr typename std::enable_if_t, std::ostream&> o // ####################### VECTOR OPERATORS ########################## // Implemented in a way that the return types follow the c++ standard, for each vector component // The user is responsible for knowing the type conversion hazards, inherent to the C++ language. +// The Can* gates live in a default template argument (not in the return type): they must be +// checked BEFORE the decltype over the base types is substituted, otherwise resolving the +// element level operator for class type scalars (fp16/bf16) re-enters this same template and +// recurses infinitely. #define VEC_UNARY_UNIVERSAL(op) \ -template \ +template ::value>> \ FK_HOST_DEVICE_CNST auto operator op(const T& a) -> \ - std::enable_if_t::value, \ - fk::VectorType_t>()), fk::cn>> { \ + fk::VectorType_t>()), fk::cn> { \ using O = fk::VectorType_t>()), fk::cn>; \ if constexpr (fk::cn == 1) { \ return fk::make_(op a.x); \ @@ -514,11 +618,10 @@ VEC_COMPOUND_LOGICAL(|=) // We don't need to check for I2 being a vector type, because the enable_if condition ensures it is a cuda vector if the two previous conditions are false #define VEC_BINARY(op) \ -template \ +template ::value>> \ FK_HOST_DEVICE_CNST auto operator op(const I1& a, const I2& b) \ - -> std::enable_if_t::value, \ - typename fk::VectorType>() op std::declval>()), \ - (fk::cn > fk::cn ? fk::cn : fk::cn)>::type_v> { \ + -> typename fk::VectorType>() op std::declval>()), \ + (fk::cn > fk::cn ? fk::cn : fk::cn)>::type_v { \ using O = typename fk::VectorType>() op std::declval>()), \ (fk::cn > fk::cn ? fk::cn : fk::cn)>::type_v; \ if constexpr (fk::validCUDAVec && fk::validCUDAVec) { \ @@ -571,11 +674,10 @@ VEC_BINARY(||) #undef VEC_BINARY #define VEC_BINARY_BITWISE(op) \ -template \ +template ::value>> \ FK_HOST_DEVICE_CNST auto operator op(const I1& a, const I2& b) \ - -> std::enable_if_t::value, \ - typename fk::VectorType>() op std::declval>()), \ - (fk::cn > fk::cn ? fk::cn : fk::cn)>::type_v> { \ + -> typename fk::VectorType>() op std::declval>()), \ + (fk::cn > fk::cn ? fk::cn : fk::cn)>::type_v { \ using O = typename fk::VectorType>() op std::declval>()), \ (fk::cn > fk::cn ? fk::cn : fk::cn)>::type_v; \ if constexpr (fk::validCUDAVec && fk::validCUDAVec) { \ diff --git a/include/fused_kernel/core/utils/vlimits.h b/include/fused_kernel/core/utils/vlimits.h index 2d569863..d3a44d08 100644 --- a/include/fused_kernel/core/utils/vlimits.h +++ b/include/fused_kernel/core/utils/vlimits.h @@ -20,18 +20,35 @@ #include namespace fk { + namespace vlimits_detail { + // std::numeric_limits' primary template silently returns value initialized (zero) + // limits for unspecialized types: turn that into a loud compile error. + template + FK_HOST_DEVICE_CNST T checkedMax() { + static_assert(std::numeric_limits::is_specialized, + "fk::maxValue: std::numeric_limits is not specialized for this type"); + return std::numeric_limits::max(); + } + template + FK_HOST_DEVICE_CNST T checkedLowest() { + static_assert(std::numeric_limits::is_specialized, + "fk::minValue: std::numeric_limits is not specialized for this type"); + return std::numeric_limits::lowest(); + } + } // namespace vlimits_detail + // Limits template constexpr T maxValue{}; template - constexpr T maxValue && !std::is_aggregate_v>> = std::numeric_limits::max(); + constexpr T maxValue && !std::is_aggregate_v>> = vlimits_detail::checkedMax(); template constexpr T maxValue >> = make_set(maxValue>); template constexpr T minValue{}; template - constexpr T minValue && !std::is_aggregate_v>> = std::numeric_limits::lowest(); + constexpr T minValue && !std::is_aggregate_v>> = vlimits_detail::checkedLowest(); template constexpr T minValue >> = make_set(minValue>); diff --git a/tests/operation_test_utils.h b/tests/operation_test_utils.h index 4c62a2ec..16d958bc 100644 --- a/tests/operation_test_utils.h +++ b/tests/operation_test_utils.h @@ -88,7 +88,12 @@ return correct ? 0 : -1; template constexpr inline bool equalValues(const T & val1, const T & val2) { - if constexpr (std::is_floating_point_v) { + if constexpr (fk::isReducedFloat) { + // Bit exact comparison: a float tolerance would quantize to 0 on fp8/fp4 grids, + // and two NaNs (any payload) must compare equal. + return fk::toBits(val1) == fk::toBits(val2) || + (fk::ReducedFloatTraits::isNaN(val1) && fk::ReducedFloatTraits::isNaN(val2)); + } else if constexpr (std::is_floating_point_v) { return std::abs(val1 - val2) < static_cast(0.0001); } else { return val1 == val2; @@ -394,6 +399,41 @@ struct TestCaseBuilder +struct TestCaseBuilder::value && + !fk::validCUDAVec && + !fk::validCUDAVec && + (fk::isReducedFloat || + fk::isReducedFloat), void>> { + template + static inline void addTest(std::map>& testCases, + const std::array& inputElems, + const std::array& expectedElems) { + const std::string testName = fk::typeToString(); + testCases[testName] = [testName, inputElems, expectedElems]() { + const auto outputPtr = test_case_builder::detail::launchUnary(testName, inputElems); + bool result{ true }; + for (size_t i = 0; i < N; ++i) { + const auto generated = outputPtr.at(fk::Point{ static_cast(i), 0, 0 }); + static_assert(std::is_same_v, std::decay_t>, "Output and Expected types are not the same"); + const auto resultV = equalValues(generated, expectedElems[i]); + if (!resultV) { + std::cout << "\033[31m" << "FAIL!!" << "\033[0m" << std::endl; + std::cout << "\033[31m Mismatch at test element index " << i << ": Expected value " + << static_cast(expectedElems[i]) << ", got " + << static_cast(generated) << "\033[0m" << std::endl; + } + result &= resultV; + } + if (result) + std::cout << "\033[32m" << "Success!!" << "\033[0m" << std::endl; + return result; + }; + } +}; + template struct TestCaseBuilder::value && (fk::validCUDAVec || diff --git a/utests/core/data/utest_reduced_float_native_parity.h b/utests/core/data/utest_reduced_float_native_parity.h new file mode 100644 index 00000000..d87d4308 --- /dev/null +++ b/utests/core/data/utest_reduced_float_native_parity.h @@ -0,0 +1,213 @@ +/* Copyright 2026 Oscar Amoros Huguet, Johnny Nunez + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. */ + +// Verifies bit exact parity between the fk software conversions and the NVIDIA native +// converters, for every representable code and a large sample of float inputs. +#define __ONLY_CU__ + +#include + +#include +#include + +#include +#include +#if __has_include() +#include +#define UTEST_HAS_FP8 1 +#endif +#if __has_include() +#include +#define UTEST_HAS_FP4 1 +#endif + +#include +#include + +using namespace fk; + +// The constexpr software path and the runtime native path must agree inside the same TU. +static_assert(fp16(1.5f).bits() == 0x3E00); +static_assert(bf16(1.5f).bits() == 0x3FC0); +static_assert(sizeof(fp16) == sizeof(__half) && alignof(fp16) == alignof(__half)); +static_assert(sizeof(bf16) == sizeof(__nv_bfloat16) && alignof(bf16) == alignof(__nv_bfloat16)); +static_assert(sizeof(fp16_2) == sizeof(__half2) && alignof(fp16_2) == alignof(__half2)); +static_assert(sizeof(bf16_2) == sizeof(__nv_bfloat162) && alignof(bf16_2) == alignof(__nv_bfloat162)); +#if defined(UTEST_HAS_FP8) +static_assert(sizeof(fp8_e4m3) == sizeof(__nv_fp8_e4m3) && alignof(fp8_e4m3) == alignof(__nv_fp8_e4m3)); +static_assert(sizeof(fp8_e5m2) == sizeof(__nv_fp8_e5m2) && alignof(fp8_e5m2) == alignof(__nv_fp8_e5m2)); +static_assert(sizeof(fp8_e4m3_2) == sizeof(__nv_fp8x2_e4m3) && alignof(fp8_e4m3_2) == alignof(__nv_fp8x2_e4m3)); +static_assert(sizeof(fp8_e4m3_4) == sizeof(__nv_fp8x4_e4m3) && alignof(fp8_e4m3_4) == alignof(__nv_fp8x4_e4m3)); +#endif +#if defined(UTEST_HAS_FP4) +static_assert(sizeof(fp4_e2m1) == sizeof(__nv_fp4_e2m1) && alignof(fp4_e2m1) == alignof(__nv_fp4_e2m1)); +#endif + +namespace { + int failures = 0; + + void reportFailure(const char* what, const unsigned int detail) { + if (failures < 50) { + std::cout << what << " (0x" << std::hex << detail << std::dec << ")" << std::endl; + } + ++failures; + } + + void checkEncodeFp16(const float f) { + const unsigned int sw = detail_rf::encode(f); + const unsigned short nat = __builtin_bit_cast(unsigned short, __half(f)); + if (std::isnan(f)) { + const bool swNan = ((sw & 0x7C00u) == 0x7C00u) && ((sw & 0x3FFu) != 0u); + const bool natNan = ((nat & 0x7C00u) == 0x7C00u) && ((nat & 0x3FFu) != 0u); + if (!swNan || !natNan) reportFailure("fp16 NaN encode parity", sw); + } else if (sw != nat) { + reportFailure("fp16 encode parity", __builtin_bit_cast(unsigned int, f)); + } + } + + void checkEncodeBf16(const float f) { + const unsigned int sw = detail_rf::encode(f); + const unsigned short nat = __builtin_bit_cast(unsigned short, __nv_bfloat16(f)); + if (std::isnan(f)) { + const bool swNan = ((sw & 0x7F80u) == 0x7F80u) && ((sw & 0x7Fu) != 0u); + const bool natNan = ((nat & 0x7F80u) == 0x7F80u) && ((nat & 0x7Fu) != 0u); + if (!swNan || !natNan) reportFailure("bf16 NaN encode parity", sw); + } else if (sw != nat) { + reportFailure("bf16 encode parity", __builtin_bit_cast(unsigned int, f)); + } + } +} // namespace + +int launch() { + // Exhaustive decode parity for the 2 byte formats. detail_rf::decode is called directly: + // fp16/bf16 operator float takes the native fast path at runtime under nvcc, and the point + // here is to compare the SOFTWARE decoder (the one CPU builds and constexpr use) against + // the native converter. + for (unsigned int code = 0; code < 0x10000u; ++code) { + const float nativeF16 = __half2float(__builtin_bit_cast(__half, static_cast(code))); + const float fkF16 = detail_rf::decode(code); + if (std::isnan(nativeF16) ? !std::isnan(fkF16) + : (nativeF16 != fkF16 || std::signbit(nativeF16) != std::signbit(fkF16))) { + reportFailure("fp16 decode parity", code); + } + const float nativeBf16 = __bfloat162float(__builtin_bit_cast(__nv_bfloat16, static_cast(code))); + const float fkBf16 = detail_rf::decode(code); + if (std::isnan(nativeBf16) ? !std::isnan(fkBf16) + : (nativeBf16 != fkBf16 || std::signbit(nativeBf16) != std::signbit(fkBf16))) { + reportFailure("bf16 decode parity", code); + } + } + + // Encode parity: directed corpus plus 2M sampled bit patterns + const float corpus[] = { 0.0f, -0.0f, 1.0f, -1.0f, 0.5f, 1.5f, 2.5f, 65504.0f, 65505.0f, 65520.0f, + 65536.0f, 1e20f, -1e20f, 5.96e-8f, 2.98e-8f, 1e-40f, -1e-45f, 3.4e38f, + INFINITY, -INFINITY, NAN, 448.0f, 449.0f, 464.0f, 465.0f, 57344.0f, + 6.0f, 5.0f, 7.0f, 0.25f, 0.75f, 3.3895314e38f, 1.7014118e38f }; + for (const float f : corpus) { + checkEncodeFp16(f); + checkEncodeBf16(f); + } + unsigned int lcg = 12345u; + for (int i = 0; i < 2000000; ++i) { + lcg = lcg * 1664525u + 1013904223u; + const float f = __builtin_bit_cast(float, lcg); + checkEncodeFp16(f); + checkEncodeBf16(f); + } + + // Native interop roundtrip + { + const fp16 h(3.25f); + const __half nh = toNative(h); + if (toBits(fromNative(nh)) != h.bits() || __half2float(nh) != 3.25f) { + reportFailure("fp16 toNative/fromNative roundtrip", h.bits()); + } + const bf16 b(-7.5f); + if (toBits(fromNative(toNative(b))) != b.bits()) { + reportFailure("bf16 toNative/fromNative roundtrip", b.bits()); + } + } + +#if defined(UTEST_HAS_FP8) + for (unsigned int code = 0; code < 256u; ++code) { + __nv_fp8_e4m3 n43; n43.__x = static_cast<__nv_fp8_storage_t>(code); + const float nat43 = static_cast(n43); + const float fk43 = static_cast(fp8_e4m3::fromBits(static_cast(code))); + if (std::isnan(nat43) ? !std::isnan(fk43) + : (nat43 != fk43 || std::signbit(nat43) != std::signbit(fk43))) { + reportFailure("fp8_e4m3 decode parity", code); + } + __nv_fp8_e5m2 n52; n52.__x = static_cast<__nv_fp8_storage_t>(code); + const float nat52 = static_cast(n52); + const float fk52 = static_cast(fp8_e5m2::fromBits(static_cast(code))); + if (std::isnan(nat52) ? !std::isnan(fk52) + : (nat52 != fk52 || std::signbit(nat52) != std::signbit(fk52))) { + reportFailure("fp8_e5m2 decode parity", code); + } + } + lcg = 999u; + for (int i = 0; i < 500000; ++i) { + lcg = lcg * 1664525u + 1013904223u; + const float f = __builtin_bit_cast(float, lcg); + if (detail_rf::encode(f) != + static_cast(__nv_cvt_float_to_fp8(f, __NV_SATFINITE, __NV_E4M3))) { + reportFailure("fp8_e4m3 encode parity", lcg); + } + if (detail_rf::encode(f) != + static_cast(__nv_cvt_float_to_fp8(f, __NV_SATFINITE, __NV_E5M2))) { + reportFailure("fp8_e5m2 encode parity", lcg); + } + // Also sample the fp8-relevant exponent range densely + const float g = std::ldexp(static_cast(static_cast(lcg % 4096u)) / 256.0f - 8.0f, + static_cast(lcg % 40u) - 20); + if (detail_rf::encode(g) != + static_cast(__nv_cvt_float_to_fp8(g, __NV_SATFINITE, __NV_E4M3))) { + reportFailure("fp8_e4m3 encode parity (ranged)", __builtin_bit_cast(unsigned int, g)); + } + } +#endif + +#if defined(UTEST_HAS_FP4) + for (unsigned int code = 0; code < 16u; ++code) { + __nv_fp4_e2m1 n4; n4.__x = static_cast<__nv_fp4_storage_t>(code); + const float nat = static_cast(n4); + const float fkv = static_cast(fp4_e2m1::fromBits(static_cast(code))); + if (nat != fkv || std::signbit(nat) != std::signbit(fkv)) { + reportFailure("fp4_e2m1 decode parity", code); + } + } + for (int i = -20000; i <= 20000; ++i) { + const float f = static_cast(i) / 1000.0f; + const __nv_fp4_e2m1 n4(f); + if (detail_rf::encode(f) != static_cast(n4.__x & 0x0Fu)) { + reportFailure("fp4_e2m1 encode parity", static_cast(i)); + } + } + { + const __nv_fp4_e2m1 nNan(NAN); + if (detail_rf::encode(NAN) != static_cast(nNan.__x & 0x0Fu)) { + reportFailure("fp4_e2m1 NaN encode parity", 0u); + } + const __nv_fp4_e2m1 nInf(INFINITY); + if (detail_rf::encode(INFINITY) != static_cast(nInf.__x & 0x0Fu)) { + reportFailure("fp4_e2m1 Inf encode parity", 0u); + } + } +#endif + + if (failures == 0) { + std::cout << "utest_reduced_float_native_parity: all checks passed" << std::endl; + } + return failures == 0 ? 0 : -1; +} diff --git a/utests/core/data/utest_reduced_float_types.h b/utests/core/data/utest_reduced_float_types.h new file mode 100644 index 00000000..a96df3ac --- /dev/null +++ b/utests/core/data/utest_reduced_float_types.h @@ -0,0 +1,225 @@ +/* Copyright 2026 Oscar Amoros Huguet, Johnny Nunez + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. */ + +#include + +#include +#include +#include +#include +#include + +#include +#include + +// The CPU backend must never see a CUDA header: cuda_fp16.h compiles silently under plain +// g++ when the toolkit is installed, so this tripwire is the only reliable guard. +#if !defined(__NVCC__) && defined(__CUDA_FP16_TYPES_EXIST__) +#error "cuda_fp16.h leaked into a CPU-only translation unit" +#endif + +using namespace fk; + +// ---- Layout: bit compatible with the CUDA types ---- +static_assert(sizeof(fp16) == 2 && alignof(fp16) == 2); +static_assert(sizeof(bf16) == 2 && alignof(bf16) == 2); +static_assert(sizeof(fp8_e4m3) == 1 && alignof(fp8_e4m3) == 1); +static_assert(sizeof(fp8_e5m2) == 1 && alignof(fp8_e5m2) == 1); +static_assert(sizeof(fp4_e2m1) == 1 && alignof(fp4_e2m1) == 1); +static_assert(sizeof(fp16_2) == 4 && alignof(fp16_2) == 4); // __half2 is alignas(4) +static_assert(sizeof(bf16_2) == 4 && alignof(bf16_2) == 4); // __nv_bfloat162 is alignas(4) +static_assert(sizeof(fp16_3) == 6 && sizeof(fp16_4) == 8 && alignof(fp16_4) == 8); +static_assert(sizeof(fp8_e4m3_2) == 2 && alignof(fp8_e4m3_2) == 2); // __nv_fp8x2_e4m3 +static_assert(sizeof(fp8_e4m3_4) == 4 && alignof(fp8_e4m3_4) == 4); // __nv_fp8x4_e4m3 + +// ---- Trait registration ---- +static_assert(validScalar && validScalar && validScalar && + validScalar && validScalar); +static_assert(validFloatingPoint && !validFloatingPoint); +static_assert(isArithmeticReducedFloat && isArithmeticReducedFloat && + !isArithmeticReducedFloat); +static_assert(validCUDAVec && validCUDAVec && validCUDAVec); +static_assert(cn == 1 && cn == 2 && cn == 3 && cn == 4); +static_assert(std::is_same_v, fp16> && std::is_same_v, fp8_e5m2>); +static_assert(std::is_same_v, fp16_2> && + std::is_same_v, fp4_e2m1_4>); + +// ---- Constexpr conversions (the software path is always taken at constant evaluation) ---- +static_assert(fp16(1.0f).bits() == 0x3C00 && fp16(-2.0f).bits() == 0xC000); +static_assert(fp16(65504.0f).bits() == 0x7BFF && fp16(1e20f).bits() == 0x7C00); +static_assert(static_cast(fp16::fromBits(0x0001)) == 5.960464477539063e-08f); // min subnormal +static_assert(bf16(1.0f).bits() == 0x3F80); +static_assert(fp8_e4m3(448.0f).bits() == 0x7E && fp8_e4m3(1000.0f).bits() == 0x7E); // SATFINITE +static_assert(fp8_e5m2(1e6f).bits() == 0x7B); // SATFINITE 57344 +static_assert(fp4_e2m1(1.5f).bits() == 0x03 && fp4_e2m1(-0.5f).bits() == 0x09); +static_assert(fp4_e2m1(1000.0f).bits() == 0x07); // saturates to 6.0 +static_assert(static_cast(fp4_e2m1::fromBits(0xF2)) == 1.0f); // upper nibble ignored +// Round to nearest even ties +static_assert(fp4_e2m1(5.0f).bits() == 0x06); // tie between 4 (even mantissa) and 6 -> 4 +static_assert(fp4_e2m1(0.25f).bits() == 0x00); // tie between 0 and 0.5 -> 0 +static_assert(fp16(65505.0f).bits() == 0x7BFF); // rounds down to max, not Inf +// Single rounding from double: 2049.0000001 is above the 2048/2050 tie, so it must round to +// 2050 (0x6801). Converting through float first would round to 2049.0f and then to 2048 (tie +// to even) - a double rounding error this assert would catch. +static_assert(fp16(2049.0000001).bits() == 0x6801); +// NaN handling +static_assert(fp8_e4m3(std::numeric_limits::quiet_NaN()).bits() == 0x7F); +static_assert(fp4_e2m1(std::numeric_limits::quiet_NaN()).bits() == 0x07); // e2m1 has no NaN + +// ---- Constexpr arithmetic and comparisons (fp16/bf16) ---- +static_assert((fp16(2.0f) * fp16(3.0f)).bits() == fp16(6.0f).bits()); +static_assert((bf16(2.0f) + bf16(3.0f)).bits() == bf16(5.0f).bits()); +static_assert((-fp16(1.5f)).bits() == fp16(-1.5f).bits()); +static_assert(fp16(1.0f) < fp16(2.0f) && bf16(-1.0f) <= bf16(1.0f) && fp16(3.0f) == fp16(3.0f)); + +// ---- Limits: fk::maxValue/minValue and cxp limits work through std::numeric_limits ---- +static_assert(toBits(maxValue) == 0x7BFF && toBits(minValue) == 0xFBFF); +static_assert(static_cast(maxValue) == 448.0f); +static_assert(static_cast(maxValue) == 57344.0f); +static_assert(static_cast(maxValue) == 6.0f); +static_assert(static_cast(maxValue) == 3.3895313892515355e+38f); +static_assert(toBits(maxValue.x) == 0x7BFF && toBits(maxValue.y) == 0x7BFF); +static_assert(static_cast(cxp::maxValue) == 65504.0f); + +// ---- cxp classification and saturate_cast ---- +static_assert(cxp::isnan::f(fp16::fromBits(0x7E00)) && !cxp::isnan::f(fp16(1.0f))); +static_assert(cxp::isinf::f(fp16::fromBits(0x7C00)) && !cxp::isinf::f(fp16(1.0f))); +static_assert(cxp::isnan::f(fp8_e4m3::fromBits(0x7F)) && !cxp::isnan::f(fp8_e4m3::fromBits(0x7E))); +static_assert(!cxp::isinf::f(fp4_e2m1::fromBits(0x07))); +static_assert(toBits(cxp::abs::f(fp16(-3.5f))) == fp16(3.5f).bits()); +static_assert(toBits(cxp::abs::f(fp8_e5m2(-2.0f))) == fp8_e5m2(2.0f).bits()); +static_assert(toBits(cxp::saturate_cast::f(1e20f)) == 0x7BFF); // clamps, no Inf +static_assert(cxp::saturate_cast::f(fp16(3.7f)) == 4); // rounds via float +static_assert(toBits(cxp::saturate_cast::f(fp16(1000.0f))) == 0x7E); +static_assert(cxp::saturate_cast::f(fp16(300.0f)) == 255); +static_assert(toBits(cxp::max::f(fp8_e4m3(2.0f), fp8_e4m3(-3.0f))) == fp8_e4m3(2.0f).bits()); + +int launch() { + int failures = 0; + + // Exhaustive roundtrip on the software path: decode every code, encode it back. + for (unsigned int code = 0; code < 0x10000u; ++code) { + const auto h = fp16::fromBits(static_cast(code)); + if (!ReducedFloatTraits::isNaN(h)) { + const float f = static_cast(h); + if (toBits(fp16(f)) != code) { + std::cout << "fp16 roundtrip failed for code " << code << std::endl; + ++failures; + } + } + const auto b = bf16::fromBits(static_cast(code)); + if (!ReducedFloatTraits::isNaN(b)) { + const float f = static_cast(b); + if (toBits(bf16(f)) != code) { + std::cout << "bf16 roundtrip failed for code " << code << std::endl; + ++failures; + } + } + } + for (unsigned int code = 0; code < 256u; ++code) { + const auto e43 = fp8_e4m3::fromBits(static_cast(code)); + if (!ReducedFloatTraits::isNaN(e43)) { + if (toBits(fp8_e4m3(static_cast(e43))) != code) { + std::cout << "fp8_e4m3 roundtrip failed for code " << code << std::endl; + ++failures; + } + } + const auto e52 = fp8_e5m2::fromBits(static_cast(code)); + if (!ReducedFloatTraits::isNaN(e52) && !ReducedFloatTraits::isInf(e52)) { + if (toBits(fp8_e5m2(static_cast(e52))) != code) { + std::cout << "fp8_e5m2 roundtrip failed for code " << code << std::endl; + ++failures; + } + } + } + for (unsigned int code = 0; code < 16u; ++code) { + const auto e21 = fp4_e2m1::fromBits(static_cast(code)); + // -0.0 encodes back to +0.0? No: sign is preserved for zero, so identity must hold. + if (toBits(fp4_e2m1(static_cast(e21))) != code) { + std::cout << "fp4_e2m1 roundtrip failed for code " << code << std::endl; + ++failures; + } + } + + // Arithmetic double oracle: the float promote-compute-demote path must produce the + // correctly rounded reduced result (binary32 has > 2p+2 mantissa bits for both formats). + unsigned int lcg = 42u; + for (int i = 0; i < 200000; ++i) { + lcg = lcg * 1664525u + 1013904223u; + const auto a16 = fp16::fromBits(static_cast(lcg & 0xFFFFu)); + const auto b16 = fp16::fromBits(static_cast((lcg >> 16) & 0xFFFFu)); + if (!ReducedFloatTraits::isNaN(a16) && !ReducedFloatTraits::isNaN(b16)) { + const auto sum = a16 + b16; + const auto oracle = fp16(static_cast(static_cast(a16)) + + static_cast(static_cast(b16))); + if (toBits(sum) != toBits(oracle) && + !(ReducedFloatTraits::isNaN(sum) && ReducedFloatTraits::isNaN(oracle))) { + std::cout << "fp16 add oracle mismatch: " << toBits(a16) << " + " << toBits(b16) << std::endl; + ++failures; + } + const auto prod = a16 * b16; + const auto prodOracle = fp16(static_cast(static_cast(a16)) * + static_cast(static_cast(b16))); + if (toBits(prod) != toBits(prodOracle) && + !(ReducedFloatTraits::isNaN(prod) && ReducedFloatTraits::isNaN(prodOracle))) { + std::cout << "fp16 mul oracle mismatch: " << toBits(a16) << " * " << toBits(b16) << std::endl; + ++failures; + } + } + const auto a8 = bf16::fromBits(static_cast(lcg & 0xFFFFu)); + const auto b8 = bf16::fromBits(static_cast((lcg >> 16) & 0xFFFFu)); + if (!ReducedFloatTraits::isNaN(a8) && !ReducedFloatTraits::isNaN(b8)) { + const auto sum = a8 + b8; + const auto oracle = bf16(static_cast(static_cast(a8)) + + static_cast(static_cast(b8))); + if (toBits(sum) != toBits(oracle) && + !(ReducedFloatTraits::isNaN(sum) && ReducedFloatTraits::isNaN(oracle))) { + std::cout << "bf16 add oracle mismatch: " << toBits(a8) << " + " << toBits(b8) << std::endl; + ++failures; + } + } + } + + // Vector operators over reduced vectors + { + const fp16_2 a{ fp16(1.0f), fp16(2.0f) }; + const fp16_2 b{ fp16(0.5f), fp16(0.25f) }; + const auto c = a * b; + static_assert(std::is_same_v, fp16_2>); + if (toBits(c.x) != fp16(0.5f).bits() || toBits(c.y) != fp16(0.5f).bits()) { + std::cout << "fp16_2 vector multiply failed" << std::endl; + ++failures; + } + const auto d = a * fp16(2.0f); // vector x scalar + if (toBits(d.x) != fp16(2.0f).bits() || toBits(d.y) != fp16(4.0f).bits()) { + std::cout << "fp16_2 vector x scalar multiply failed" << std::endl; + ++failures; + } + const auto eq = (a == a); + if (!Bool::vAnd(eq)) { + std::cout << "fp16_2 vector comparison failed" << std::endl; + ++failures; + } + const auto ms = make_set(bf16(1.5f)); + if (toBits(ms.x) != bf16(1.5f).bits() || toBits(ms.w) != bf16(1.5f).bits()) { + std::cout << "make_set failed" << std::endl; + ++failures; + } + } + + if (failures == 0) { + std::cout << "utest_reduced_float_types: all checks passed" << std::endl; + } + return failures == 0 ? 0 : -1; +} diff --git a/utests/core/execution_model/utest_reduced_float_pipelines.h b/utests/core/execution_model/utest_reduced_float_pipelines.h new file mode 100644 index 00000000..90e502d2 --- /dev/null +++ b/utests/core/execution_model/utest_reduced_float_pipelines.h @@ -0,0 +1,211 @@ +/* Copyright 2026 Oscar Amoros Huguet, Johnny Nunez + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. */ + +// Reduced float element types in full executeOperations pipelines, on both backends: +// the same source builds as _cpp (ParArch::CPU) and _cu (ParArch::GPU_NVIDIA). + +#include +#include + +#include +#include +#include +#include +#include + +#include + +using namespace fk; + +namespace { + constexpr int WIDTH = 67; // odd width: exercises the non thread-divisible path too + constexpr int HEIGHT = 23; + + // Read T -> Cast to float -> x2 -> SaturateCast back to T. The expected value is computed + // element by element with the same scalar operations, so the comparison is bit exact. + template + bool scalarPipeline() { + Stream stream; + Ptr2D input(WIDTH, HEIGHT); + Ptr2D output(WIDTH, HEIGHT); + for (int y = 0; y < HEIGHT; ++y) { + for (int x = 0; x < WIDTH; ++x) { + // Values exactly representable in every format under test + input.at(Point(x, y)) = T(static_cast((x % 3) - 1)); + } + } + input.upload(stream); + executeOperations>(stream, + PerThreadRead::build(input.ptr()), + Cast::build(), + Mul::build(2.0f), + SaturateCast::build(), + PerThreadWrite::build(output.ptr())); + output.download(stream); + stream.sync(); + for (int y = 0; y < HEIGHT; ++y) { + for (int x = 0; x < WIDTH; ++x) { + const T expected = cxp::saturate_cast::f(static_cast(input.at(Point(x, y))) * 2.0f); + if (!equalValues(output.at(Point(x, y)), expected)) { + std::cout << "scalar pipeline mismatch at (" << x << "," << y << ") for " + << typeToString() << std::endl; + return false; + } + } + } + return true; + } + + // Same pipeline expressed as loose operations and as a .then() fused operation: the two + // must produce bit identical results (quantization applies equally to both). + bool fusedVsLoose() { + Stream stream; + Ptr2D input(WIDTH, HEIGHT); + Ptr2D outLoose(WIDTH, HEIGHT); + Ptr2D outFused(WIDTH, HEIGHT); + for (int y = 0; y < HEIGHT; ++y) { + for (int x = 0; x < WIDTH; ++x) { + input.at(Point(x, y)) = fp16(0.25f * static_cast(x - 30)); + } + } + input.upload(stream); + executeOperations>(stream, + PerThreadRead::build(input.ptr()), + Cast::build(), + Mul::build(3.0f), + SaturateCast::build(), + PerThreadWrite::build(outLoose.ptr())); + const auto fusedRead = PerThreadRead::build(input.ptr()) + .then(Cast::build()) + .then(Mul::build(3.0f)) + .then(SaturateCast::build()); + executeOperations>(stream, fusedRead, + PerThreadWrite::build(outFused.ptr())); + outLoose.download(stream); + outFused.download(stream); + stream.sync(); + for (int y = 0; y < HEIGHT; ++y) { + for (int x = 0; x < WIDTH; ++x) { + if (toBits(outLoose.at(Point(x, y))) != toBits(outFused.at(Point(x, y)))) { + std::cout << "fused vs loose mismatch at (" << x << "," << y << ")" << std::endl; + return false; + } + } + } + return true; + } + + // Vector element type end to end + bool vectorPipeline() { + Stream stream; + Ptr2D input(WIDTH, HEIGHT); + Ptr2D output(WIDTH, HEIGHT); + for (int y = 0; y < HEIGHT; ++y) { + for (int x = 0; x < WIDTH; ++x) { + input.at(Point(x, y)) = fp16_2{ fp16(static_cast(x % 5)), fp16(-0.5f * (y % 4)) }; + } + } + input.upload(stream); + executeOperations>(stream, + PerThreadRead::build(input.ptr()), + Cast::build(), + Mul::build(float2{ 2.0f, 4.0f }), + SaturateCast::build(), + PerThreadWrite::build(output.ptr())); + output.download(stream); + stream.sync(); + for (int y = 0; y < HEIGHT; ++y) { + for (int x = 0; x < WIDTH; ++x) { + const auto in = input.at(Point(x, y)); + const auto out = output.at(Point(x, y)); + const fp16 expectedX = cxp::saturate_cast::f(static_cast(in.x) * 2.0f); + const fp16 expectedY = cxp::saturate_cast::f(static_cast(in.y) * 4.0f); + if (!equalValues(out.x, expectedX) || !equalValues(out.y, expectedY)) { + std::cout << "vector pipeline mismatch at (" << x << "," << y << ")" << std::endl; + return false; + } + } + } + return true; + } + + // Thread fusion enabled: fp16 maps to fp16_2 (two elements per thread), including the + // non divisible tail path (WIDTH is odd). + bool threadFusionPipeline() { + Stream stream; + Ptr2D input(WIDTH, HEIGHT); + Ptr2D output(WIDTH, HEIGHT); + for (int y = 0; y < HEIGHT; ++y) { + for (int x = 0; x < WIDTH; ++x) { + input.at(Point(x, y)) = fp16(0.5f * static_cast(x)); + } + } + input.upload(stream); + executeOperations>( + stream, + PerThreadRead::build(input.ptr()), + Cast::build(), + Mul::build(2.0f), + SaturateCast::build(), + PerThreadWrite::build(output.ptr())); + output.download(stream); + stream.sync(); + for (int y = 0; y < HEIGHT; ++y) { + for (int x = 0; x < WIDTH; ++x) { + const fp16 expected = cxp::saturate_cast::f(static_cast(input.at(Point(x, y))) * 2.0f); + if (!equalValues(output.at(Point(x, y)), expected)) { + std::cout << "thread fusion pipeline mismatch at (" << x << "," << y << ")" << std::endl; + return false; + } + } + } + return true; + } +} // namespace + +int launch() { + bool correct = true; + + correct &= scalarPipeline(); + correct &= scalarPipeline(); + correct &= scalarPipeline(); + correct &= scalarPipeline(); + correct &= scalarPipeline(); + correct &= fusedVsLoose(); + correct &= vectorPipeline(); + correct &= threadFusionPipeline(); + + // Exactly representable values: expectations hold bit exactly on every backend. + TestCaseBuilder>::addTest(testCases, + std::array{ fp16(1.0f), fp16(-2.5f), fp16(0.0f), fp16(65504.0f) }, + std::array{ 1.0f, -2.5f, 0.0f, 65504.0f }); + TestCaseBuilder>::addTest(testCases, + std::array{ 1.0f, 1e20f, -1e20f, 0.5f }, + std::array{ fp16(1.0f), fp16(65504.0f), fp16(-65504.0f), fp16(0.5f) }); + TestCaseBuilder>::addTest(testCases, + std::array{ 1.0f, 1000.0f, -1000.0f, 0.25f }, + std::array{ fp8_e4m3(1.0f), fp8_e4m3(448.0f), fp8_e4m3(-448.0f), fp8_e4m3(0.25f) }); + TestCaseBuilder>::addTest(testCases, + std::array{ fp16(1.0f), fp16(300.0f), fp16(-5.0f), fp16(254.5f) }, + std::array{ 1, 255, 0, 254 }); + for (const auto& [testName, testFunc] : testCases) { + correct &= testFunc(); + } + testCases.clear(); + + if (correct) { + std::cout << "utest_reduced_float_pipelines: all checks passed" << std::endl; + } + return correct ? 0 : -1; +}