diff --git a/Project.toml b/Project.toml index c2fe816..467d309 100644 --- a/Project.toml +++ b/Project.toml @@ -11,6 +11,7 @@ BFloat16s = "ab4f0b2a-ad5b-11e8-123f-65d77653426b" BitPacking = "b58c8408-13c4-4787-8733-7038ae624acf" Random = "9a3f8284-a2c9-5f02-9a11-845980a1fd5c" Republic = "27243419-9dde-4721-b67c-fd63626fea7f" +StaticArrays = "90137ffa-7385-5640-81b9-e52037218182" [weakdeps] CUDACore = "bd0ed864-bdfe-4181-a5ed-ce625a5fdea2" @@ -24,4 +25,5 @@ BitPacking = "0.2" CUDACore = "6.2.1" Random = "1" Republic = "2.1" +StaticArrays = "1.9" julia = "1.10" diff --git a/ext/CUDACoreExt/CUDACoreExt.jl b/ext/CUDACoreExt/CUDACoreExt.jl index f9e9ec9..0a3130c 100644 --- a/ext/CUDACoreExt/CUDACoreExt.jl +++ b/ext/CUDACoreExt/CUDACoreExt.jl @@ -1,5 +1,6 @@ module CUDACoreExt include("cudaDataType.jl") +include("conversion.jl") end diff --git a/ext/CUDACoreExt/conversion.jl b/ext/CUDACoreExt/conversion.jl new file mode 100644 index 0000000..e80b18e --- /dev/null +++ b/ext/CUDACoreExt/conversion.jl @@ -0,0 +1,227 @@ +# Device-side conversion overrides. +# +# Microfloats funnels every conversion through `cvt(T, x, mode, policy)` +# (scalar) and `cvt(NVector{T,N}, xs, mode, policy)` (vector), with the +# rounding mode and overflow policy as positional, dispatchable arguments. +# This file therefore only needs to override: +# +# 1. the error hooks — so the *generic* numeric kernels run unmodified on +# device (no duplicated conversion body); +# 2. narrow `cvt` signatures for exactly the (target, source, mode, policy) +# combinations that have native PTX conversion instructions, gated on +# compute capability. The gates fold at kernel compile time because +# `compute_capability()`/`target_feature_set()` are compile-time +# constants under GPUCompiler, so each kernel compiles to either the +# native instruction or the generic path with no runtime branch. +# +# PTX `cvt` into sub-byte float formats is only available as `.satfinite` +# (mandatory for fp8/fp6/fp4/ue8m0 destinations), so every native path +# implements the `SAT` overflow policy; `OVF` always takes the generic path. +# Instruction/operand-order conventions follow PTX ISA §9.7.9 ("cvt") as +# validated empirically in PTX.jl (H100/GB10): `cvt d, a, b` puts `a` in the +# UPPER lane and `b` in the LOWER lane of `d`. +# +# EXPERIMENTAL: exact parity between the native instructions and the generic +# path (rounding at the floatmax boundary, NaN payloads) has not been +# validated on hardware yet; run an on-device exhaustive parity sweep before +# relying on bit-exactness. + +using Microfloats +using Microfloats: Microfloat, cvt, cvt_generic, cvt_lanes, + OverflowPolicy, Overflowing, Saturating, SAT, OVF, + throw_negative_unsigned, throw_no_nan, + Float8_E4M3FN, Float8_E5M2, Float8_E8M0FNU, + Float6_E2M3FN, Float6_E3M2FN, Float4_E2M1FN +using BitPacking: NArray, NVector +using CUDACore: CUDACore, @device_override, compute_capability, target_feature_set + +# ───────────────────────── error hooks ────────────────────────── + +# With these four overrides the generic conversion kernels are device-safe +# as-is; everything below is optimization only. +# NB: the `where T` type variables are load-bearing — a bare `::Type` +# argument is left unspecialized by Julia, which turns these into dynamic +# calls in device code (InvalidIRError). +@device_override @noinline Microfloats.throw_negative_unsigned(::Type{T}, x) where T = + CUDACore.@gputhrow "DomainError" "negative input to unsigned microfloat" +@device_override @noinline Microfloats.throw_no_nan(::Type{T}, x) where T = + CUDACore.@gputhrow "DomainError" "microfloat format has no NaN" +@device_override @noinline Microfloats.throw_no_overflow_sentinel(::Type{T}, x) where T = + CUDACore.@gputhrow "DomainError" "microfloat format has no overflow sentinel; use overflow=SAT" +@device_override @noinline Microfloats.throw_unsupported_rounding(::Type{T}, mode) where T = + CUDACore.@gputhrow "ArgumentError" "unsupported rounding mode for microfloat conversion" + +# ───────────────────────── capability gates ────────────────────────── + +@inline function cc_ge(major::UInt32, minor::UInt32) + cc = compute_capability() + cc.major > major || (cc.major == major && cc.minor >= minor) +end + +@inline has_fp8_cvt() = cc_ge(UInt32(8), UInt32(9)) +@inline has_mxfp_cvt() = cc_ge(UInt32(10), UInt32(0)) && target_feature_set() === :arch + +# ───────────────────────── PTX cvt wrappers ────────────────────────── + +# Two Float32 lanes → one packed pair, low lane first. Inline asm rather +# than `llvm.nvvm.*` intrinsics so availability doesn't depend on the LLVM +# version Julia ships; a future PTX.jl-based extension can supersede these +# with intrinsic-backed lowering. +@generated function cvt_pair_bits(::Val{instr}, lo::Float32, hi::Float32) where instr + ir = """ + define i16 @entry(float %lo, float %hi) #0 { + %r = call i16 asm "$(String(instr)) \$0, \$1, \$2;", "=h,f,f"(float %hi, float %lo) + ret i16 %r + } + attributes #0 = { alwaysinline } + """ + :(Base.llvmcall(($ir, "entry"), UInt16, Tuple{Float32,Float32}, lo, hi)) +end + +# fp4 destinations are `.b8`, which has no NVPTX register-constraint letter; +# bridge through a 16-bit register (mirrors NVIDIA's shims and +# PTX.jl's hand-written e2m1x2 entries). +@generated function cvt_pair_bits_b8(::Val{instr}, lo::Float32, hi::Float32) where instr + ir = """ + define i16 @entry(float %lo, float %hi) #0 { + %r = call i16 asm "{ .reg .b8 t; $(String(instr)) t, \$1, \$2; mov.b16 \$0, {t, 0}; }", "=h,f,f"(float %hi, float %lo) + ret i16 %r + } + attributes #0 = { alwaysinline } + """ + :(Base.llvmcall(($ir, "entry"), UInt16, Tuple{Float32,Float32}, lo, hi)) +end + +# ───────────────────────── bit repacking ────────────────────────── + +# PTX returns byte-aligned lanes; NVector packs lanes densely (lane 1 at the +# LSB). 8-bit lanes coincide; 6-/4-bit lanes need compaction into the exact +# storage representation BitPacking's `pack` would choose. + +# 2×6-bit: byte-aligned b16 → dense 12 bits → NTuple{2,UInt8} storage. +@inline function fp6_pair_storage(bits::UInt16) + dense = (bits & 0x003f) | ((bits >> 2) & 0x0fc0) + (dense % UInt8, (dense >> 8) % UInt8) +end + +@inline pack2(::Type{T}, data::D) where {T,D} = NArray{T,1,Tuple{2},D}(data) +@inline pack4(::Type{T}, data::D) where {T,D} = NArray{T,1,Tuple{4},D}(data) + +# ───────────────────────── fp8: E4M3FN / E5M2 (sm_89+) ────────────────────────── + +for (T, instr) in ((Float8_E4M3FN, "cvt.rn.satfinite.e4m3x2.f32"), + (Float8_E5M2, "cvt.rn.satfinite.e5m2x2.f32")) + v = Val(Symbol(instr)) + @eval begin + @device_override @inline Microfloats.cvt(::Type{$T}, x::Float32, + mode::RoundingMode{:Nearest}, policy::Saturating) = + has_fp8_cvt() ? reinterpret($T, cvt_pair_bits($v, x, x) % UInt8) : + cvt_generic($T, x, mode, policy) + + @device_override @inline Microfloats.cvt(::Type{NVector{$T,2}}, xs::NTuple{2,Float32}, + mode::RoundingMode{:Nearest}, policy::Saturating) = + has_fp8_cvt() ? pack2($T, cvt_pair_bits($v, xs[1], xs[2])) : + cvt_lanes(NVector{$T,2}, xs, mode, policy) + + @device_override @inline Microfloats.cvt(::Type{NVector{$T,4}}, xs::NTuple{4,Float32}, + mode::RoundingMode{:Nearest}, policy::Saturating) = + has_fp8_cvt() ? pack4($T, UInt32(cvt_pair_bits($v, xs[1], xs[2])) | + (UInt32(cvt_pair_bits($v, xs[3], xs[4])) << 16)) : + cvt_lanes(NVector{$T,4}, xs, mode, policy) + end +end + +# ───────────────────────── fp6: E2M3FN / E3M2FN (sm_100a+) ────────────────────────── + +# FiniteOnly targets: the generic SAT path throws for NaN inputs, while +# hardware `.satfinite` silently maps NaN; guard first to keep semantics. +for (T, instr) in ((Float6_E2M3FN, "cvt.rn.satfinite.e2m3x2.f32"), + (Float6_E3M2FN, "cvt.rn.satfinite.e3m2x2.f32")) + v = Val(Symbol(instr)) + @eval begin + @device_override @inline function Microfloats.cvt(::Type{$T}, x::Float32, + mode::RoundingMode{:Nearest}, policy::Saturating) + has_mxfp_cvt() || return cvt_generic($T, x, mode, policy) + isnan(x) && throw_no_nan($T, x) + return reinterpret($T, (cvt_pair_bits($v, x, x) % UInt8) & 0x3f) + end + + @device_override @inline function Microfloats.cvt(::Type{NVector{$T,2}}, xs::NTuple{2,Float32}, + mode::RoundingMode{:Nearest}, policy::Saturating) + has_mxfp_cvt() || return cvt_lanes(NVector{$T,2}, xs, mode, policy) + (isnan(xs[1]) | isnan(xs[2])) && throw_no_nan($T, xs) + return pack2($T, fp6_pair_storage(cvt_pair_bits($v, xs[1], xs[2]))) + end + + @device_override @inline function Microfloats.cvt(::Type{NVector{$T,4}}, xs::NTuple{4,Float32}, + mode::RoundingMode{:Nearest}, policy::Saturating) + has_mxfp_cvt() || return cvt_lanes(NVector{$T,4}, xs, mode, policy) + (isnan(xs[1]) | isnan(xs[2]) | isnan(xs[3]) | isnan(xs[4])) && throw_no_nan($T, xs) + lo = fp6_pair_storage(cvt_pair_bits($v, xs[1], xs[2])) + hi = fp6_pair_storage(cvt_pair_bits($v, xs[3], xs[4])) + # dense 24-bit little-endian layout: lanes 1-2 in bits 0-11, 3-4 in 12-23 + return pack4($T, (lo[1], lo[2] | (hi[1] << 4), (hi[1] >> 4) | (hi[2] << 4))) + end + end +end + +# ───────────────────────── fp4: E2M1FN (sm_100a+) ────────────────────────── + +let T = Float4_E2M1FN, v = Val(Symbol("cvt.rn.satfinite.e2m1x2.f32")) + @eval begin + @device_override @inline function Microfloats.cvt(::Type{$T}, x::Float32, + mode::RoundingMode{:Nearest}, policy::Saturating) + has_mxfp_cvt() || return cvt_generic($T, x, mode, policy) + isnan(x) && throw_no_nan($T, x) + return reinterpret($T, (cvt_pair_bits_b8($v, x, x) % UInt8) & 0x0f) + end + + @device_override @inline function Microfloats.cvt(::Type{NVector{$T,2}}, xs::NTuple{2,Float32}, + mode::RoundingMode{:Nearest}, policy::Saturating) + has_mxfp_cvt() || return cvt_lanes(NVector{$T,2}, xs, mode, policy) + (isnan(xs[1]) | isnan(xs[2])) && throw_no_nan($T, xs) + return pack2($T, cvt_pair_bits_b8($v, xs[1], xs[2]) % UInt8) + end + + @device_override @inline function Microfloats.cvt(::Type{NVector{$T,4}}, xs::NTuple{4,Float32}, + mode::RoundingMode{:Nearest}, policy::Saturating) + has_mxfp_cvt() || return cvt_lanes(NVector{$T,4}, xs, mode, policy) + (isnan(xs[1]) | isnan(xs[2]) | isnan(xs[3]) | isnan(xs[4])) && throw_no_nan($T, xs) + return pack4($T, (cvt_pair_bits_b8($v, xs[1], xs[2]) & 0x00ff) | + (cvt_pair_bits_b8($v, xs[3], xs[4]) << 8)) + end + end +end + +# ───────────────────────── ue8m0: E8M0FNU (sm_100a+) ────────────────────────── + +# Hardware only converts to ue8m0 with .rz (and .rp); the type's default +# RoundNearest keeps the generic path. NaN maps to 0xff natively, matching +# `nan(Float8_E8M0FNU)` under SAT. The generic path throws for any negative +# input (including -0.0); guard to preserve that. +let T = Float8_E8M0FNU, v = Val(Symbol("cvt.rz.satfinite.ue8m0x2.f32")) + @eval begin + @device_override @inline function Microfloats.cvt(::Type{$T}, x::Float32, + mode::RoundingMode{:ToZero}, policy::Saturating) + has_mxfp_cvt() || return cvt_generic($T, x, mode, policy) + signbit(x) && throw_negative_unsigned($T, x) + return reinterpret($T, cvt_pair_bits($v, x, x) % UInt8) + end + + @device_override @inline function Microfloats.cvt(::Type{NVector{$T,2}}, xs::NTuple{2,Float32}, + mode::RoundingMode{:ToZero}, policy::Saturating) + has_mxfp_cvt() || return cvt_lanes(NVector{$T,2}, xs, mode, policy) + (signbit(xs[1]) | signbit(xs[2])) && throw_negative_unsigned($T, xs) + return pack2($T, cvt_pair_bits($v, xs[1], xs[2])) + end + + @device_override @inline function Microfloats.cvt(::Type{NVector{$T,4}}, xs::NTuple{4,Float32}, + mode::RoundingMode{:ToZero}, policy::Saturating) + has_mxfp_cvt() || return cvt_lanes(NVector{$T,4}, xs, mode, policy) + (signbit(xs[1]) | signbit(xs[2]) | signbit(xs[3]) | signbit(xs[4])) && + throw_negative_unsigned($T, xs) + return pack4($T, UInt32(cvt_pair_bits($v, xs[1], xs[2])) | + (UInt32(cvt_pair_bits($v, xs[3], xs[4])) << 16)) + end + end +end diff --git a/src/Microfloats.jl b/src/Microfloats.jl index 38baf12..2fc9dce 100644 --- a/src/Microfloats.jl +++ b/src/Microfloats.jl @@ -2,9 +2,8 @@ module Microfloats using Republic -import BFloat16s: BFloat16 - @republic import BitPacking: bitwidth +@republic import BFloat16s: BFloat16 include("utils.jl") @public sign_bits, exponent_bits, significand_bits @@ -20,6 +19,7 @@ export Microfloat include("conversion.jl") @public overflow_policy @public SAT, OVF +@public cvt, cvt_generic include("macro.jl") export @microfloat @@ -30,6 +30,19 @@ export Float8_E4M3FN, Float8_E8M0FNU export Float6_E2M3FN, Float6_E3M2FN export Float4_E2M1FN +include("vectorization.jl") +@public cvt_lanes + +include("specializations.jl") +@public Float16x2, Float16x4 +@public BFloat16x2, BFloat16x4 +@public Float8x2_E4M3FN, Float8x4_E4M3FN +@public Float8x2_E5M2, Float8x4_E5M2 +@public Float8x2_E8M0FNU, Float8x4_E8M0FNU +@public Float6x2_E2M3FN, Float6x4_E2M3FN +@public Float6x2_E3M2FN, Float6x4_E3M2FN +@public Float4x2_E2M1FN, Float4x4_E2M1FN + include("ops.jl") include("random.jl") diff --git a/src/conversion.jl b/src/conversion.jl index da8ae8b..2e71450 100644 --- a/src/conversion.jl +++ b/src/conversion.jl @@ -86,6 +86,24 @@ Microfloats.Saturating() overflow_policy(::Type{T}) where T<:Microfloat = error("$T must define `Microfloats.overflow_policy(::Type{$T})`") +# ───────────────────────── error hooks ────────────────────────── + +# Every error path in the conversion kernels routes through one of these +# `@noinline` hooks so device backends (e.g. CUDACoreExt) can override just +# the hooks — via `@device_override` — and run the *same* numeric kernels on +# device, instead of maintaining a duplicated device-safe copy of the whole +# conversion body. +@noinline throw_negative_unsigned(::Type{T}, x) where T = + throw(DomainError(x, "negative input to unsigned $T")) +@noinline throw_no_nan(::Type{T}, x) where T = + throw(DomainError(x, "$T has no NaN")) +@noinline throw_no_overflow_sentinel(::Type{T}, x) where T = + throw(DomainError(x, "$T has no overflow sentinel; use overflow=SAT")) +@noinline throw_unsupported_rounding(::Type{T}, mode) where T = + throw(ArgumentError("$T does not support rounding mode $mode")) + +# ───────────────────────── rounding shifts ────────────────────────── + function rshift_round_to_even(x::T, n::Int) where T<:Unsigned n <= 0 && return x >> n n > 8 * sizeof(T) && return zero(T) @@ -129,12 +147,12 @@ clamp_inf(x::T) where T<:Microfloat = signbit(x) ? -inf(T) : inf(T) function apply_overflow_policy(x::T, xf::Float32, mode::RoundingMode, ::Overflowing) where T<:Microfloat if isnan(xf) - return hasnan(T) ? nan(T) : throw(DomainError(xf, "$T has no NaN")) + return hasnan(T) ? nan(T) : throw_no_nan(T, xf) elseif isinf(xf) || is_outside_floatmax(xf, T) if mode_overflows_to_inf(mode, signbit(xf)) return hasinf(T) ? clamp_inf(x) : hasnan(T) ? nan(T) : - throw(DomainError(xf, "$T has no overflow sentinel; use overflow=SAT")) + throw_no_overflow_sentinel(T, xf) else return clamp_floatmax(x) end @@ -145,7 +163,7 @@ end function apply_overflow_policy(x::T, xf::Float32, ::RoundingMode, ::Saturating) where T<:Microfloat if isnan(xf) - return hasnan(T) ? nan(T) : throw(DomainError(xf, "$T has no NaN")) + return hasnan(T) ? nan(T) : throw_no_nan(T, xf) elseif isinf(xf) || is_outside_floatmax(xf, T) return clamp_floatmax(x) else @@ -158,7 +176,7 @@ function _round_to_microfloat(::Type{T}, x::Float32, rshift::F, mode::RoundingMode, policy::OverflowPolicy ) where {T<:Microfloat, F} if sign_bits(T) == 0 && signbit(x) - throw(DomainError(x, "negative input to unsigned $T")) + throw_negative_unsigned(T, x) end iszero(x) && return signbit(x) ? -zero(T) : zero(T) @@ -208,44 +226,83 @@ function _round_to_microfloat(::Type{T}, x::Float32, rshift::F, return apply_overflow_policy(reinterpret(T, t_raw), x, mode, policy) end -(::Type{T})(x::Float32, mode::RoundingMode{:Nearest}; - overflow::OverflowPolicy = overflow_policy(T)) where T<:Microfloat = - _round_to_microfloat(T, x, rshift_round_to_even, mode, overflow) -(::Type{T})(x::Float32, mode::RoundingMode{:NearestTiesAway}; - overflow::OverflowPolicy = overflow_policy(T)) where T<:Microfloat = - _round_to_microfloat(T, x, rshift_round_ties_away, mode, overflow) -(::Type{T})(x::Float32, mode::RoundingMode{:ToZero}; - overflow::OverflowPolicy = overflow_policy(T)) where T<:Microfloat = - _round_to_microfloat(T, x, rshift_truncate, mode, overflow) -(::Type{T})(x::Float32, mode::RoundingMode{:FromZero}; - overflow::OverflowPolicy = overflow_policy(T)) where T<:Microfloat = - _round_to_microfloat(T, x, rshift_round_up_magnitude, mode, overflow) +# ───────────────────────── conversion funnel ────────────────────────── + +""" + cvt(::Type{T}, x, mode::RoundingMode, policy::OverflowPolicy) -> T + +Central conversion funnel. Every scalar conversion into a +[`Microfloat`](@ref) — constructors, `convert`, broadcasts, and the packed +vector paths — reduces to a call of this function, with the rounding mode +and overflow policy as positional, dispatchable arguments. + +`cvt` is the extension surface for optimized conversions. To specialize, +add a method on any subset of `(T, typeof(x), mode, policy)`: + +- **Bit-twiddling / table specializations** add ordinary methods, e.g. + `Microfloats.cvt(::Type{Float8_E4M3}, x::Float4_E2M1FN, ::RoundingMode, + ::OverflowPolicy)`. See [`@cvt_table`](@ref) for a generated lookup-table + shortcut. +- **Device backends** (package extensions) use overlay method tables (e.g. + `CUDACore.@device_override`) on exactly the `(T, source, mode, policy)` + signatures the hardware supports natively; every other combination falls + through to the portable methods below. + +The always-correct reference path is [`cvt_generic`](@ref); specialized +methods that need a partial fallback should call it (not `cvt`, which on +overlay method tables would recurse into the override itself). +""" +@inline cvt(::Type{T}, x::Real, mode::RoundingMode, policy::OverflowPolicy) where T<:Microfloat = + cvt(T, Float32(x), mode, policy) +@inline cvt(::Type{T}, x::Float32, mode::RoundingMode, policy::OverflowPolicy) where T<:Microfloat = + cvt_generic(T, x, mode, policy) +@inline cvt(::Type{T}, x::Microfloat, mode::RoundingMode, policy::OverflowPolicy) where T<:Microfloat = + cvt_generic(T, Float32(x), mode, policy) + +""" + cvt_generic(::Type{T}, x::Float32, mode::RoundingMode, policy::OverflowPolicy) -> T + +The generic reference implementation behind [`cvt`](@ref): bit-level +rounding from `Float32` into any `Microfloat` layout, for every supported +rounding mode and overflow policy. Specialized `cvt` methods (and device +overrides) call this directly when their fast path does not apply. +""" +@inline cvt_generic(::Type{T}, x::Float32, mode::RoundingMode{:Nearest}, policy::OverflowPolicy) where T<:Microfloat = + _round_to_microfloat(T, x, rshift_round_to_even, mode, policy) +@inline cvt_generic(::Type{T}, x::Float32, mode::RoundingMode{:NearestTiesAway}, policy::OverflowPolicy) where T<:Microfloat = + _round_to_microfloat(T, x, rshift_round_ties_away, mode, policy) +@inline cvt_generic(::Type{T}, x::Float32, mode::RoundingMode{:ToZero}, policy::OverflowPolicy) where T<:Microfloat = + _round_to_microfloat(T, x, rshift_truncate, mode, policy) +@inline cvt_generic(::Type{T}, x::Float32, mode::RoundingMode{:FromZero}, policy::OverflowPolicy) where T<:Microfloat = + _round_to_microfloat(T, x, rshift_round_up_magnitude, mode, policy) # RoundUp/RoundDown are sign-dependent: "toward +∞" rounds the magnitude up # for positive inputs but truncates the magnitude for negative inputs (which # moves the value toward zero, i.e., closer to +∞). RoundDown is the mirror. -(::Type{T})(x::Float32, mode::RoundingMode{:Up}; - overflow::OverflowPolicy = overflow_policy(T)) where T<:Microfloat = - signbit(x) ? _round_to_microfloat(T, x, rshift_truncate, mode, overflow) : - _round_to_microfloat(T, x, rshift_round_up_magnitude, mode, overflow) -(::Type{T})(x::Float32, mode::RoundingMode{:Down}; - overflow::OverflowPolicy = overflow_policy(T)) where T<:Microfloat = - signbit(x) ? _round_to_microfloat(T, x, rshift_round_up_magnitude, mode, overflow) : - _round_to_microfloat(T, x, rshift_truncate, mode, overflow) - -# Errors on unsupported modes instead of recursing through the Real-level fallback below. -(::Type{T})(x::Float32, mode::RoundingMode; - overflow::OverflowPolicy = overflow_policy(T)) where T<:Microfloat = - throw(ArgumentError("$T does not support rounding mode $mode")) - +@inline cvt_generic(::Type{T}, x::Float32, mode::RoundingMode{:Up}, policy::OverflowPolicy) where T<:Microfloat = + signbit(x) ? _round_to_microfloat(T, x, rshift_truncate, mode, policy) : + _round_to_microfloat(T, x, rshift_round_up_magnitude, mode, policy) +@inline cvt_generic(::Type{T}, x::Float32, mode::RoundingMode{:Down}, policy::OverflowPolicy) where T<:Microfloat = + signbit(x) ? _round_to_microfloat(T, x, rshift_round_up_magnitude, mode, policy) : + _round_to_microfloat(T, x, rshift_truncate, mode, policy) + +cvt_generic(::Type{T}, x::Float32, mode::RoundingMode, ::OverflowPolicy) where T<:Microfloat = + throw_unsupported_rounding(T, mode) + +# ───────────────────────── constructors ────────────────────────── + +# Constructors are thin sugar over `cvt`: they only resolve defaults +# (RoundNearest, the type's registered overflow policy) and are never +# specialized or device-overridden themselves. +# # `Real` (not `Number`) avoids colliding with Base's # `(::Type{T})(::Real, ::RoundingMode) where T<:AbstractFloat`. (::Type{T})(x::Real; overflow::OverflowPolicy = overflow_policy(T)) where T<:Microfloat = - T(x, RoundNearest; overflow=overflow) + cvt(T, x, RoundNearest, overflow) (::Type{T})(x::Real, mode::RoundingMode; overflow::OverflowPolicy = overflow_policy(T)) where T<:Microfloat = - T(Float32(x), mode; overflow=overflow) + cvt(T, x, mode, overflow) # Returns the BFloat16 encoding as raw UInt16 bits. Internal plumbing stays # in bits because on Julia >= 1.12 a BFloat16 *value* crossing a function-call @@ -303,8 +360,7 @@ function _to_bfloat16_bits(x::T) where {T<:Microfloat} return bf16_sign_bit | UInt16(sub_q & 0x7f) end else - bf16_raw_out = bf16_sign_bit | (UInt16(bf16_exponent_field & 0xff) << 7) | UInt16((bf16_significand_total - 0x80) & 0x7f) - return bf16_raw_out + return bf16_sign_bit | (UInt16(bf16_exponent_field & 0xff) << 7) | UInt16((bf16_significand_total - 0x80) & 0x7f) end end @@ -363,6 +419,76 @@ function decimal_string end BFloat16(x::T) where T<:Microfloat = to_bfloat16(x) (::Type{T})(x::Microfloat) where T<:Number = T(reinterpret(Float32, UInt32(to_bfloat16_bits(x)) << 16)) -# Microfloat → Microfloat: route through Float32 (matches the Real-input path -# and avoids the BFloat16 intermediate's narrower exponent dynamic range). -(::Type{T})(x::Microfloat) where T<:Microfloat = T(Float32(x)) +# Microfloat → Microfloat: disambiguates the two methods above; the default +# `cvt` route goes through Float32 (matching the Real-input path and avoiding +# the BFloat16 intermediate's narrower exponent dynamic range) unless a +# specialized `cvt` method — e.g. one registered by `@cvt_table` — applies. +(::Type{T})(x::Microfloat; + overflow::OverflowPolicy = overflow_policy(T)) where T<:Microfloat = + cvt(T, x, RoundNearest, overflow) + +# ───────────────────────── @cvt_table ────────────────────────── + +# Generator backend for `@cvt_table`: builds the complete raw-bits lookup +# table for one (Dst, Src, mode, policy) combination by running every +# possible source bit pattern through `cvt_generic`, so the table is correct +# by construction. If any entry throws (e.g. negative source values into an +# unsigned target), the whole combination falls back to the runtime generic +# path so error behavior is preserved exactly. +function table_cvt_expr(::Type{T}, ::Type{S}, ::Type{M}, ::Type{P} + ) where {T<:Microfloat, S<:Microfloat, M<:RoundingMode, P<:OverflowPolicy} + mode, policy = M.instance, P.instance + n = 1 << bitwidth(S) + vals = UInt8[] + for raw in UInt8(0):UInt8(n - 1) + y = try + cvt_generic(T, Float32(reinterpret(S, raw)), mode, policy) + catch + return :($cvt_generic($T, Float32(x), mode, policy)) + end + push!(vals, reinterpret(UInt8, y)) + end + table = Tuple(vals) + mask = UInt8(n - 1) + return :(reinterpret($T, $table[Int(reinterpret(UInt8, x) & $mask) + 1])) +end + +""" + @cvt_table Src => Dst + +Register an optimized lookup-table method on the conversion funnel +[`cvt`](@ref) for converting microfloat `Src` values to microfloat `Dst`. + +Expands to a `@generated` method of `Microfloats.cvt` whose lookup table is +computed lazily — once per `(mode, policy)` combination actually used — by +running every `Src` bit pattern through [`cvt_generic`](@ref), so results +are identical to the generic path but cost a single `2^bitwidth(Src)`-entry +table lookup. Combinations where the generic path throws (e.g. signed +source into unsigned target) keep the runtime path and its errors. + +Invoke *after* both types are defined. Microfloats registers tables for all +pairs of built-in types; user-defined `@microfloat` types can opt in: + +```julia +@microfloat MyFloat6 exponent=3 significand=2 +Microfloats.@cvt_table MyFloat6 => Float8_E4M3 +Microfloats.@cvt_table Float8_E4M3 => MyFloat6 +``` + +To hand-optimize a pair instead (e.g. branch-free bit-twiddling), define +the `Microfloats.cvt` method for it directly rather than invoking +`@cvt_table` for that pair. +""" +macro cvt_table(pair) + (pair isa Expr && pair.head === :call && pair.args[1] === :(=>)) || + throw(ArgumentError("@cvt_table expects `Src => Dst`, got `$pair`")) + S, T = pair.args[2], pair.args[3] + ex = quote + Base.@generated function $(@__MODULE__).cvt(::Type{$T}, x::$S, + mode::$RoundingMode, policy::$OverflowPolicy) + $table_cvt_expr($T, $S, mode, policy) + end + nothing + end + return esc(ex) +end diff --git a/src/specializations.jl b/src/specializations.jl new file mode 100644 index 0000000..58632ff --- /dev/null +++ b/src/specializations.jl @@ -0,0 +1,52 @@ +using BitPacking: NArray + +# Hand-optimized bit-twiddling `cvt` specializations. +# +# Every Float4_E2M1FN value is exactly representable in the E4M3 layouts +# (M: 1 ≤ 3; exponent range ⊂ target range; subnormal 0.5 becomes normal), +# so the conversion is independent of rounding mode and overflow policy and +# reduces to a few branch-free ALU ops on the raw bits. Unlike the +# `@cvt_table` lookup (a gather in vectorized loops), these autovectorize. +# +# Pairs implemented here are excluded from the table-registration loop in +# variants.jl (`TWIDDLED_PAIRS`) so the definitions don't collide. + +# 3-bit magnitude m = e₁e₀f: 0 → 0; 1 (subnormal, 0.5) → 0x30 (2⁻¹); +# normals are linear: (e+6) << 3 | f << 2 == (m << 2) + 0x30. +@inline function _e2m1_to_e4m3_mag(m::UInt8) + t = (m << 2) + 0x30 + t = ifelse(m == 0x01, 0x30, t) + return ifelse(m == 0x00, 0x00, t) +end + +@inline _e2m1_to_e4m3_byte(n::UInt8) = + _e2m1_to_e4m3_mag(n & 0x07) | ((n & 0x08) << 4) + +for T in (:Float8_E4M3, :Float8_E4M3FN) + @eval begin + @inline cvt(::Type{$T}, x::Float4_E2M1FN, + ::RoundingMode, ::OverflowPolicy) = + reinterpret($T, _e2m1_to_e4m3_byte(reinterpret(UInt8, x))) + + # Packed → packed: twiddle directly on the storage bits, so loops + # over packed buffers never materialize individual lanes. + @inline function cvt(::Type{NVector{$T,2}}, xs::NVector{Float4_E2M1FN,2}, + ::RoundingMode, ::OverflowPolicy) + raw = reinterpret(UInt8, xs) # lane 1 in the low nibble + lo = _e2m1_to_e4m3_byte(raw & 0x0f) + hi = _e2m1_to_e4m3_byte(raw >> 4) + return NArray{$T,1,Tuple{2},UInt16}(UInt16(lo) | (UInt16(hi) << 8)) + end + + @inline function cvt(::Type{NVector{$T,4}}, xs::NVector{Float4_E2M1FN,4}, + ::RoundingMode, ::OverflowPolicy) + raw = reinterpret(UInt16, xs) + b1 = _e2m1_to_e4m3_byte(raw % UInt8 & 0x0f) + b2 = _e2m1_to_e4m3_byte((raw >> 4) % UInt8 & 0x0f) + b3 = _e2m1_to_e4m3_byte((raw >> 8) % UInt8 & 0x0f) + b4 = _e2m1_to_e4m3_byte((raw >> 12) % UInt8) + return NArray{$T,1,Tuple{4},UInt32}( + UInt32(b1) | (UInt32(b2) << 8) | (UInt32(b3) << 16) | (UInt32(b4) << 24)) + end + end +end diff --git a/src/variants.jl b/src/variants.jl index b26eb3d..9037fa2 100644 --- a/src/variants.jl +++ b/src/variants.jl @@ -18,12 +18,28 @@ @microfloat Float6_E3M2FN exponent=3 significand=2 nonfinite=FiniteOnly @microfloat Float4_E2M1FN exponent=2 significand=1 nonfinite=FiniteOnly -for T in ( +const BUILTIN_TYPES = ( :Float8_E5M2, :Float8_E4M3, :Float8_E3M4, :Float8_E4M3FN, :Float8_E8M0FNU, :Float6_E2M3FN, :Float6_E3M2FN, :Float4_E2M1FN, ) + +# Register lookup-table `cvt` methods for every ordered pair of built-in +# types (including identity, which normalizes NaN encodings like the generic +# path does). Tables are built lazily per (mode, policy) on first use. +# Pairs with hand-written bit-twiddling methods (specializations.jl) are +# excluded so the definitions don't collide. +const TWIDDLED_PAIRS = ( + (:Float4_E2M1FN, :Float8_E4M3), + (:Float4_E2M1FN, :Float8_E4M3FN), +) +for S in BUILTIN_TYPES, T in BUILTIN_TYPES + (S, T) in TWIDDLED_PAIRS && continue + @eval @cvt_table $S => $T +end + +for T in BUILTIN_TYPES @eval @doc """ $($T) diff --git a/src/vectorization.jl b/src/vectorization.jl new file mode 100644 index 0000000..9db4058 --- /dev/null +++ b/src/vectorization.jl @@ -0,0 +1,83 @@ +using BitPacking: NVector +using StaticArrays: SVector, StaticArray + +# 16-bit + +const Float16x2 = SVector{2,Float16} +const Float16x4 = SVector{4,Float16} + +const BFloat16x2 = SVector{2,BFloat16} +const BFloat16x4 = SVector{4,BFloat16} + +# 8-bit + +const Float8x2_E4M3FN = NVector{Float8_E4M3FN,2} +const Float8x4_E4M3FN = NVector{Float8_E4M3FN,4} + +const Float8x2_E5M2 = NVector{Float8_E5M2,2} +const Float8x4_E5M2 = NVector{Float8_E5M2,4} + +const Float8x2_E8M0FNU = NVector{Float8_E8M0FNU,2} +const Float8x4_E8M0FNU = NVector{Float8_E8M0FNU,4} + +# 6-bit + +const Float6x2_E2M3FN = NVector{Float6_E2M3FN,2} +const Float6x2_E3M2FN = NVector{Float6_E3M2FN,2} + +const Float6x4_E2M3FN = NVector{Float6_E2M3FN,4} +const Float6x4_E3M2FN = NVector{Float6_E3M2FN,4} + +# 4-bit + +const Float4x2_E2M1FN = NVector{Float4_E2M1FN,2} +const Float4x4_E2M1FN = NVector{Float4_E2M1FN,4} + +# ───────────────────────── vector conversion funnel ────────────────────────── + +""" + cvt_lanes(::Type{NVector{T,N}}, xs::NTuple{N,Any}, mode, policy) -> NVector{T,N} + +Reference lanewise implementation of the vector conversion funnel: converts +each lane through the scalar [`cvt`](@ref) funnel (so per-lane +specializations and device overrides still apply), then packs. Vectorized +`cvt` specializations call this when their fast path does not cover the +requested combination. +""" +@inline cvt_lanes(::Type{NVector{T,N}}, xs::NTuple{N,Any}, + mode::RoundingMode, policy::OverflowPolicy) where {T<:Microfloat,N} = + NVector{T,N}(ntuple(i -> cvt(T, xs[i], mode, policy), Val(N))) + +""" + cvt(::Type{NVector{T,N}}, xs::NTuple{N,Any}, mode, policy) -> NVector{T,N} + +Vector form of the conversion funnel: convert `N` source lanes into a packed +`BitPacking.NVector` in one call. The default is lanewise +([`cvt_lanes`](@ref)); specialize on `(T, N, lane type, mode, policy)` for +multi-lane hardware conversions (e.g. PTX `cvt` x2 instructions in device +overlays) or SIMD bit-twiddling over packed sources. + +Source containers (`SVector`, `NVector`, any `StaticArray` vector) normalize +to `NTuple` first; packed→packed specializations may intercept the +`NVector`-source signature before it is unpacked. +""" +@inline cvt(::Type{NVector{T,N}}, xs::NTuple{N,Any}, + mode::RoundingMode, policy::OverflowPolicy) where {T<:Microfloat,N} = + cvt_lanes(NVector{T,N}, xs, mode, policy) +@inline cvt(::Type{NVector{T,N}}, xs::StaticArray{Tuple{N},<:Any,1}, + mode::RoundingMode, policy::OverflowPolicy) where {T<:Microfloat,N} = + cvt(NVector{T,N}, Tuple(xs), mode, policy) + +# Entry points: like the scalar constructors, these only resolve defaults and +# hand off to the funnel. `StaticArray{Tuple{N},<:Real,1}` covers both +# `SVector` and packed `NVector` sources. +(::Type{NVector{T,N}})(xs::StaticArray{Tuple{N},<:Real,1}, mode::RoundingMode; + overflow::OverflowPolicy = overflow_policy(T)) where {T<:Microfloat,N} = + cvt(NVector{T,N}, xs, mode, overflow) +(::Type{NVector{T,N}})(xs::StaticArray{Tuple{N},<:Real,1}; + overflow::OverflowPolicy = overflow_policy(T)) where {T<:Microfloat,N} = + cvt(NVector{T,N}, xs, RoundNearest, overflow) +# Same-eltype repacking involves no rounding; this also disambiguates against +# BitPacking's exact-eltype StaticArray constructor. +(::Type{NVector{T,N}})(xs::StaticArray{Tuple{N},T,1}) where {T<:Microfloat,N} = + NVector{T,N}(Tuple(xs)) diff --git a/test/cuda_extension.jl b/test/cuda_extension.jl new file mode 100644 index 0000000..4c27dcd --- /dev/null +++ b/test/cuda_extension.jl @@ -0,0 +1,144 @@ +using Test +using Microfloats +using Microfloats: + Float8x2_E4M3FN, Float8x4_E4M3FN, + Float8x2_E5M2, Float8x4_E5M2, + Float8x2_E8M0FNU, Float8x4_E8M0FNU, + Float6x2_E2M3FN, Float6x4_E2M3FN, + Float6x2_E3M2FN, Float6x4_E3M2FN, + Float4x2_E2M1FN, Float4x4_E2M1FN + +using CUDACore +using CUDACore: CuArray + +samebits(xs, ys) = reinterpret.(UInt8, xs) == reinterpret.(UInt8, ys) +sametuples(xs, ys) = Tuple.(xs) == Tuple.(ys) + +gpu_broadcast(::Type{T}, xs) where T = Array(T.(CuArray(xs))) + +struct GPUConvertWithMode{T,M,O} end +@inline (::GPUConvertWithMode{T,M,O})(x) where {T,M,O} = + T(x, M(); overflow=O()) + +gpu_broadcast(::Type{T}, xs, mode::M; + overflow = Microfloats.overflow_policy(T)) where {T,M} = + Array(GPUConvertWithMode{T,M,typeof(overflow)}().(CuArray(xs))) + +# Vector constructor with a saturating policy; on sm_89+ the fp8 targets +# lower to native cvt.rn.satfinite instructions. +struct GPUVecSAT{V} end +@inline (::GPUVecSAT{V})(xs) where V = V(xs; overflow=Microfloats.SAT) + +# NaN payloads of native conversions are hardware-defined; compare NaN lanes +# NaN-aware, everything else bit-exact. +naneq(a, b) = (isnan(a) && isnan(b)) || a === b +naneq_tuples(xs, ys) = all(map((x, y) -> all(naneq.(Tuple(x), Tuple(y))), xs, ys)) + +const SCALAR_TARGETS = ( + Float8_E5M2, Float8_E4M3, Float8_E3M4, Float8_E4M3FN, Float8_E8M0FNU, + Float6_E2M3FN, Float6_E3M2FN, Float4_E2M1FN, +) + +lane(::Type{T}, x) where T = T(x) +lane(::Type{Microfloats.BFloat16}, x) = Microfloats.BFloat16(x) + +svector2(::Type{T}, a, b) where T = + Microfloats.SVector{2,T}(lane(T, a), lane(T, b)) +svector4(::Type{T}, a, b, c, d) where T = + Microfloats.SVector{4,T}(lane(T, a), lane(T, b), lane(T, c), lane(T, d)) + +@testset "CUDACore extension" begin + if CUDACore.functional() + values = Float32[0, 0.5, 1, 1.5, 2, 3] + f16s = Float16.(values) + bf16s = Microfloats.BFloat16.(values) + f32s = values + f64s = Float64.(values) + e4m3s = Float8_E4M3.(values) + + @testset "scalar broadcast" begin + for xs in (f16s, bf16s, f32s, f64s, e4m3s), T in SCALAR_TARGETS + @test samebits(gpu_broadcast(T, xs), T.(xs)) + end + + for T in SCALAR_TARGETS + @test samebits(gpu_broadcast(T, f32s, RoundToZero), + T.(f32s, Ref(RoundToZero))) + end + end + + f16x2 = [svector2(Float16, 0f0, 1f0), svector2(Float16, 1.5f0, 2f0)] + bf16x2 = [svector2(Microfloats.BFloat16, 0f0, 1f0), svector2(Microfloats.BFloat16, 1.5f0, 2f0)] + f32x2 = [svector2(Float32, 0f0, 1f0), svector2(Float32, 1.5f0, 2f0)] + f64x2 = [svector2(Float64, 0f0, 1f0), svector2(Float64, 1.5f0, 2f0)] + f4x2 = Float4x2_E2M1FN.(f32x2) + + @testset "x2 broadcast" begin + @test sametuples(gpu_broadcast(Float8x2_E4M3FN, f16x2), Float8x2_E4M3FN.(f16x2)) + @test sametuples(gpu_broadcast(Float8x2_E5M2, bf16x2), Float8x2_E5M2.(bf16x2)) + @test sametuples(gpu_broadcast(Float8x2_E8M0FNU, f64x2), Float8x2_E8M0FNU.(f64x2)) + @test sametuples(gpu_broadcast(Float6x2_E2M3FN, f32x2), Float6x2_E2M3FN.(f32x2)) + @test sametuples(gpu_broadcast(Float6x2_E3M2FN, f4x2), Float6x2_E3M2FN.(f4x2)) + @test sametuples(gpu_broadcast(Float4x2_E2M1FN, f32x2), Float4x2_E2M1FN.(f32x2)) + end + + f16x4 = [ + svector4(Float16, 0f0, 1f0, 1.5f0, 2f0), + svector4(Float16, 2f0, 3f0, 4f0, 6f0), + ] + bf16x4 = [ + svector4(Microfloats.BFloat16, 0f0, 1f0, 1.5f0, 2f0), + svector4(Microfloats.BFloat16, 2f0, 3f0, 4f0, 6f0), + ] + f32x4 = [ + svector4(Float32, 0f0, 1f0, 1.5f0, 2f0), + svector4(Float32, 2f0, 3f0, 4f0, 6f0), + ] + f64x4 = [ + svector4(Float64, 0f0, 1f0, 1.5f0, 2f0), + svector4(Float64, 2f0, 3f0, 4f0, 6f0), + ] + f4x4 = Float4x4_E2M1FN.(f32x4) + + @testset "x4 broadcast" begin + @test sametuples(gpu_broadcast(Float8x4_E4M3FN, f16x4), Float8x4_E4M3FN.(f16x4)) + @test sametuples(gpu_broadcast(Float8x4_E5M2, bf16x4), Float8x4_E5M2.(bf16x4)) + @test sametuples(gpu_broadcast(Float8x4_E8M0FNU, f64x4), Float8x4_E8M0FNU.(f64x4)) + @test sametuples(gpu_broadcast(Float6x4_E2M3FN, f32x4), Float6x4_E2M3FN.(f32x4)) + @test sametuples(gpu_broadcast(Float6x4_E3M2FN, f4x4), Float6x4_E3M2FN.(f4x4)) + @test sametuples(gpu_broadcast(Float4x4_E2M1FN, f32x4), Float4x4_E2M1FN.(f32x4)) + end + + # Saturating policy: exercises the native `.satfinite` overrides on + # capable hardware (fp8 on sm_89+, fp6/fp4/ue8m0 on sm_100a) and the + # generic fallback elsewhere — device must match host either way. + sat_scalar = Float32[0, -0.0, 0.5, 1.5, -2, 447, 448, 449, -1e9, + 57344, 6e4, 1e9, Inf, -Inf, NaN] + @testset "saturating scalar broadcast" begin + for T in (Float8_E4M3FN, Float8_E5M2) + got = gpu_broadcast(T, sat_scalar, RoundNearest; overflow=Microfloats.SAT) + want = T.(sat_scalar, Ref(RoundNearest); overflow=Microfloats.SAT) + @test all(naneq.(got, want)) + end + e8m0_vals = Float32[0.5, 1, 3, 2f0^-127, 1e30, Inf, NaN] + got = gpu_broadcast(Float8_E8M0FNU, e8m0_vals, RoundToZero; overflow=Microfloats.SAT) + want = Float8_E8M0FNU.(e8m0_vals, Ref(RoundToZero); overflow=Microfloats.SAT) + @test all(naneq.(got, want)) + end + + sat2 = [svector2(Float32, 448f0, 1f9), svector2(Float32, -1f9, 0.5f0), + svector2(Float32, NaN32, 2f0), svector2(Float32, 6f4, -6f4)] + sat4 = [svector4(Float32, 448f0, 1f9, -1f9, 0.5f0), + svector4(Float32, NaN32, 2f0, 6f4, -6f4)] + @testset "saturating x2/x4 broadcast" begin + for V in (Float8x2_E4M3FN, Float8x2_E5M2) + @test naneq_tuples(Array(GPUVecSAT{V}().(CuArray(sat2))), GPUVecSAT{V}().(sat2)) + end + for V in (Float8x4_E4M3FN, Float8x4_E5M2) + @test naneq_tuples(Array(GPUVecSAT{V}().(CuArray(sat4))), GPUVecSAT{V}().(sat4)) + end + end + else + @test_skip "CUDACore.functional() == false" + end +end diff --git a/test/cvt.jl b/test/cvt.jl new file mode 100644 index 0000000..2994cec --- /dev/null +++ b/test/cvt.jl @@ -0,0 +1,126 @@ +using Microfloats: cvt, cvt_generic + +# Result-or-throw comparator: conversions must agree on values *and* on +# whether they throw. +_cvt_outcome(f) = try f() catch e; (e isa DomainError || e isa ArgumentError) ? :threw : rethrow() end + +@testset "conversion funnel" begin + @testset "cvt ≡ constructors" begin + for T in TYPES, x in (0.0, -0.0, 0.4, 1.0, 1.5, -2.75, 100.0, 1e6, -1e6, Inf, -Inf, NaN), + mode in (RoundNearest, RoundToZero, RoundUp, RoundDown, RoundFromZero, RoundNearestTiesAway), + pol in (Microfloats.SAT, Microfloats.OVF) + + a = _cvt_outcome(() -> T(x, mode; overflow=pol)) + b = _cvt_outcome(() -> cvt(T, Float32(x), mode, pol)) + @test a === b + end + end + + @testset "table specializations ≡ generic path" begin + # Built-in pairs have @cvt_table-generated methods; they must agree + # with the runtime generic path for every bit pattern, mode, policy. + bad = [] + for S in TYPES_BUILTIN, T in TYPES_BUILTIN, + mode in (RoundNearest, RoundToZero, RoundUp, RoundDown, RoundFromZero, RoundNearestTiesAway), + pol in (Microfloats.SAT, Microfloats.OVF) + + for raw in 0x00:UInt8(2^bitwidth(S) - 1) + x = reinterpret(S, raw) + a = _cvt_outcome(() -> cvt(T, x, mode, pol)) + b = _cvt_outcome(() -> cvt_generic(T, Float32(x), mode, pol)) + a === b || push!(bad, (S, T, raw, mode, pol, a, b)) + end + end + @test isempty(bad) + end + + @testset "@cvt_table on user-defined types" begin + # Types defined long after Microfloats loaded — exercises the + # world-age design of the generated table methods. + Microfloats.@cvt_table UFloat5_E2M3 => Float8_E4M3 + Microfloats.@cvt_table Float8_E4M3 => UFloat5_E2M3 + for raw in 0x00:UInt8(2^bitwidth(UFloat5_E2M3) - 1) + x = reinterpret(UFloat5_E2M3, raw) + @test cvt(Float8_E4M3, x, RoundNearest, Microfloats.SAT) === + cvt_generic(Float8_E4M3, Float32(x), RoundNearest, Microfloats.SAT) + end + # signed → unsigned: table generation hits throwing entries and must + # preserve the runtime error path + for raw in 0x00:0xff + x = reinterpret(Float8_E4M3, UInt8(raw)) + a = _cvt_outcome(() -> cvt(UFloat5_E2M3, x, RoundNearest, Microfloats.SAT)) + b = _cvt_outcome(() -> cvt_generic(UFloat5_E2M3, Float32(x), RoundNearest, Microfloats.SAT)) + @test a === b + end + end + + @testset "bit-twiddling specializations ≡ generic path" begin + # scalar twiddle: exact widening, so every mode/policy must agree + for T in (Float8_E4M3, Float8_E4M3FN), + mode in (RoundNearest, RoundToZero, RoundUp, RoundDown), + pol in (Microfloats.SAT, Microfloats.OVF) + + for raw in 0x00:0x0f + x = reinterpret(Float4_E2M1FN, raw) + @test cvt(T, x, mode, pol) === cvt_generic(T, Float32(x), mode, pol) + end + end + # packed → packed twiddle vs lanewise reference, all storage patterns + NV, NA = Microfloats.NVector, Microfloats.NArray + for T in (Float8_E4M3, Float8_E4M3FN) + bad = 0 + for raw in 0x00:0xff + xs = NA{Float4_E2M1FN,1,Tuple{2},UInt8}(raw) + a = cvt(NV{T,2}, xs, RoundNearest, Microfloats.SAT) + b = Microfloats.cvt_lanes(NV{T,2}, Tuple(xs), RoundNearest, Microfloats.SAT) + bad += a !== b + end + for raw in 0x0000:0xffff + xs = NA{Float4_E2M1FN,1,Tuple{4},UInt16}(UInt16(raw)) + a = cvt(NV{T,4}, xs, RoundNearest, Microfloats.SAT) + b = Microfloats.cvt_lanes(NV{T,4}, Tuple(xs), RoundNearest, Microfloats.SAT) + bad += a !== b + end + @test bad == 0 + end + # entry constructor reaches the packed path + f4 = Microfloats.Float4x2_E2M1FN((Float4_E2M1FN(0.5), Float4_E2M1FN(-6))) + @test Tuple(Microfloats.Float8x2_E4M3FN(f4)) == (Float8_E4M3FN(0.5), Float8_E4M3FN(-6)) + end + + @testset "error hooks" begin + @test_throws DomainError Float8_E8M0FNU(-1.0) + @test_throws DomainError Float8_E8M0FNU(-0.0) + @test_throws DomainError Float4_E2M1FN(NaN) + @test_throws DomainError Float4_E2M1FN(Inf; overflow=Microfloats.OVF) + @test_throws ArgumentError Float8_E4M3(1.0, RoundNearestTiesUp) + end + + @testset "vector funnel" begin + SV, NV = Microfloats.SVector, Microfloats.NVector + + # entry constructors resolve defaults and route through cvt + v = Microfloats.Float8x2_E4M3FN(SV{2,Float32}(1.0f0, 0.99f0)) + @test Tuple(v) == (Float8_E4M3FN(1), Float8_E4M3FN(1)) + + # rounding mode and overflow policy per call + v = Microfloats.Float8x2_E4M3FN(SV{2,Float32}(0.99f0, 1f9), RoundToZero; overflow=Microfloats.SAT) + @test Tuple(v) == (Float8_E4M3FN(0.9375), floatmax(Float8_E4M3FN)) + + # sources: Float64/Float16 SVectors, packed NVectors, same-type repack + @test Tuple(Microfloats.Float4x2_E2M1FN(SV{2,Float64}(1.0, 2.0))) == + (Float4_E2M1FN(1), Float4_E2M1FN(2)) + f4 = Microfloats.Float4x2_E2M1FN(SV{2,Float16}(1, 2)) + @test Tuple(Microfloats.Float6x2_E2M3FN(f4)) == (Float6_E2M3FN(1), Float6_E2M3FN(2)) + @test Microfloats.Float4x2_E2M1FN(f4) === f4 + + # lanewise default agrees with scalar funnel, including policy + xs = (10000f0, -1.5f0, 0.1f0, NaN32) + a = Microfloats.cvt_lanes(NV{Float8_E5M2,4}, xs, RoundNearest, Microfloats.OVF) + @test Tuple(a) === map(x -> Float8_E5M2(x; overflow=Microfloats.OVF), xs) + + # x4 entry constructor + v4 = Microfloats.Float6x4_E2M3FN(SV{4,Float32}(0f0, 1f0, 1.5f0, 2f0)) + @test Tuple(v4) == (Float6_E2M3FN(0), Float6_E2M3FN(1), Float6_E2M3FN(1.5), Float6_E2M3FN(2)) + end +end diff --git a/test/runtests.jl b/test/runtests.jl index fe95e95..0703231 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -45,6 +45,14 @@ const UNSIGNED_TYPES = ( const TYPES = (SIGNED_TYPES..., UNSIGNED_TYPES...) +# Shipped types, which have @cvt_table lookup methods registered +const TYPES_BUILTIN = ( + Float8_E5M2, Float8_E4M3, Float8_E3M4, + Float8_E4M3FN, Float8_E8M0FNU, + Float6_E2M3FN, Float6_E3M2FN, + Float4_E2M1FN, +) + # OCP Microscaling Formats v1.0 aliases const MX_E5M2 = Float8_E5M2 const MX_E4M3 = Float8_E4M3FN @@ -53,15 +61,20 @@ const MX_E2M3 = Float6_E2M3FN const MX_E2M1 = Float4_E2M1FN const MX_E8M0 = Float8_E8M0FNU -@testset "Microfloats" begin +import CUDACore + +@testset "Microfloats.jl" begin include("basic.jl") + include("cvt.jl") include("overflow.jl") include("floatmin.jl") include("rounding_modes.jl") + include("vectorization.jl") include("mx_compliance.jl") include("mx_properties.jl") include("dlfp8_parity.jl") - + # extensions include("extensions/CUDACore.jl") + include("cuda_extension.jl") end diff --git a/test/vectorization.jl b/test/vectorization.jl new file mode 100644 index 0000000..e430cc2 --- /dev/null +++ b/test/vectorization.jl @@ -0,0 +1,37 @@ +@testset "Vector aliases" begin + @test Microfloats.Float16x2 === Microfloats.SVector{2,Float16} + @test Microfloats.BFloat16x4 === Microfloats.SVector{4,BFloat16} + + @test Microfloats.Float8x2_E4M3FN === Microfloats.NVector{Float8_E4M3FN,2} + @test Microfloats.Float8x4_E5M2 === Microfloats.NVector{Float8_E5M2,4} + @test Microfloats.Float8x2_E8M0FNU === Microfloats.NVector{Float8_E8M0FNU,2} + + @test Microfloats.Float6x2_E2M3FN === Microfloats.NVector{Float6_E2M3FN,2} + @test Microfloats.Float6x4_E3M2FN === Microfloats.NVector{Float6_E3M2FN,4} + @test Microfloats.Float4x4_E2M1FN === Microfloats.NVector{Float4_E2M1FN,4} + + f4 = Microfloats.Float4x2_E2M1FN((Float4_E2M1FN(1), Float4_E2M1FN(2))) + @test bitwidth(f4) == 8 + @test reinterpret(UInt8, f4) == + (reinterpret(UInt8, Float4_E2M1FN(1)) | + (reinterpret(UInt8, Float4_E2M1FN(2)) << 4)) + + f6 = Microfloats.Float6x4_E2M3FN(( + Float6_E2M3FN(1), Float6_E2M3FN(2), + Float6_E2M3FN(3), Float6_E2M3FN(4), + )) + @test bitwidth(f6) == 24 + @test Tuple(f6) == ( + Float6_E2M3FN(1), Float6_E2M3FN(2), + Float6_E2M3FN(3), Float6_E2M3FN(4), + ) + + f4_from_svec = Microfloats.Float4x2_E2M1FN(Microfloats.SVector{2,Float16}(1, 2)) + @test Tuple(f4_from_svec) == (Float4_E2M1FN(1), Float4_E2M1FN(2)) + + f8_from_svec = Microfloats.Float8x2_E5M2(Microfloats.SVector{2,BFloat16}(BFloat16(1), BFloat16(2))) + @test Tuple(f8_from_svec) == (Float8_E5M2(1), Float8_E5M2(2)) + + f6_from_nvec = Microfloats.Float6x2_E2M3FN(f4) + @test Tuple(f6_from_nvec) == (Float6_E2M3FN(1), Float6_E2M3FN(2)) +end