107 lines
4.6 KiB
C
107 lines
4.6 KiB
C
/* SPDX-License-Identifier: AGPL-3.0-or-later */
|
|
/**
|
|
* @file ft_bfp.h
|
|
* @brief Host-side block-floating-point tile packer: fp32 tiles to and from the
|
|
* Tenstorrent BFP8/BFP4 byte layout the unpacker reads from L1.
|
|
*
|
|
* FORMATS (tt-isa-documentation, FloatBitPatterns.md "BFP"): every group of 16
|
|
* datums shares one 8-bit exponent. A datum is a sign bit followed by an
|
|
* unsigned magnitude with no implicit leading one: 7 magnitude bits for BFP8,
|
|
* 3 for BFP4. The "b" formats store an FP32-style biased exponent (bias 127);
|
|
* the "a" formats store an FP16-style 5-bit exponent (bias 15) zero-extended
|
|
* to a byte.
|
|
*
|
|
* TILE LAYOUT: callers hand in 32x32 tiles stored row-major within each tile.
|
|
* The packer walks the tile as four 16x16 faces (top-left, top-right,
|
|
* bottom-left, bottom-right) and, inside each face, row by row; each face row
|
|
* of 16 values is one block. A packed tile is the 64 shared-exponent bytes
|
|
* (one per block, Packers/FormatConversion.md: "a separate output stream ...
|
|
* one byte per 16 datums") followed by the datums in the same block order: one
|
|
* byte each for BFP8, two per byte (low nibble first) for BFP4. Unpacking
|
|
* restores the row-major tile.
|
|
*
|
|
* ROUNDING: fp32 to datum takes the 24-bit significand (implicit one
|
|
* included), aligns it to the block's shared exponent with a truncating right
|
|
* shift, then rounds the aligned value to the datum width, nearest with ties
|
|
* to even, saturating at the largest magnitude instead of carrying. Zeros and
|
|
* denormals pack to datum 0, and a datum whose magnitude rounds to 0 never
|
|
* carries a sign. For the "a" formats a value below the 5-bit exponent range
|
|
* is first raised to +-2^-15. The shared exponent is the largest exponent
|
|
* field in the block (clamped to 0..31 for "a"); infinity and NaN take part
|
|
* like any other value.
|
|
*
|
|
* DECODING follows FloatBitPatterns.md's BFP-to-float routine, except that a
|
|
* sign-only datum (magnitude 0) decodes to -0 rather than -2^128/-2^16, and an
|
|
* exponent that underflows wraps rather than clamps; both match the reference.
|
|
*
|
|
* @status Complete. Bit-exact against reference/ft_bfp_reference_vectors.h.
|
|
* @issues None
|
|
* @todo None
|
|
*/
|
|
|
|
#ifndef FT_BFP_H
|
|
#define FT_BFP_H
|
|
|
|
#include <stddef.h>
|
|
#include <stdint.h>
|
|
|
|
/** The block-float formats this packer produces and consumes. */
|
|
typedef enum {
|
|
FT_BFP_BFP8_B = 0, /**< 8-bit datums, 8-bit FP32-biased shared exponent. */
|
|
FT_BFP_BFP4_B = 1, /**< 4-bit datums, 8-bit FP32-biased shared exponent. */
|
|
FT_BFP_BFP8_A = 2, /**< 8-bit datums, 5-bit FP16-biased shared exponent. */
|
|
FT_BFP_BFP4_A = 3 /**< 4-bit datums, 5-bit FP16-biased shared exponent. */
|
|
} ft_bfp_format;
|
|
|
|
/** Return codes of the tile-level calls. */
|
|
enum {
|
|
FT_BFP_OK = 0,
|
|
FT_BFP_EINVAL = -3
|
|
};
|
|
|
|
/** Datums sharing one exponent byte. */
|
|
#define FT_BFP_BLOCK_VALUES 16
|
|
|
|
/** Exponent bytes at the head of every packed 32x32 tile (1024 / 16). */
|
|
#define FT_BFP_TILE_EXPONENT_BYTES 64
|
|
|
|
/** Bytes of one packed 32x32 tile; 0 for an unknown format. */
|
|
size_t ft_bfp_tile_bytes(ft_bfp_format format);
|
|
|
|
/** Magnitude bits per datum (the sign bit excluded); 0 for an unknown format. */
|
|
uint32_t ft_bfp_mantissa_bits(ft_bfp_format format);
|
|
|
|
/**
|
|
* Pack `tile_count` row-major 32x32 fp32 tiles into `out`, which must hold
|
|
* tile_count * ft_bfp_tile_bytes(format) bytes. Zero tiles is a no-op.
|
|
* Returns FT_BFP_OK, or FT_BFP_EINVAL for an unknown format or a NULL buffer.
|
|
*/
|
|
int ft_bfp_pack_tiles(ft_bfp_format format, const float *tiles, size_t tile_count, void *out);
|
|
|
|
/** Inverse of ft_bfp_pack_tiles: decode packed tiles back to row-major fp32. */
|
|
int ft_bfp_unpack_tiles(ft_bfp_format format, const void *packed, size_t tile_count, float *tiles_out);
|
|
|
|
/**
|
|
* Encode one fp32 bit pattern as a datum under `shared_exp`, which is in the
|
|
* format's own exponent encoding (8-bit for "b", 5-bit for "a"). Counts ties.
|
|
*/
|
|
uint8_t ft_bfp_convert_float_bits(ft_bfp_format format, uint32_t input_bits, uint32_t shared_exp);
|
|
|
|
/** Decode one datum under `shared_exp` to fp32 bits, as the hardware reads it. */
|
|
uint32_t ft_bfp_convert_datum(ft_bfp_format format, uint8_t datum, uint32_t shared_exp);
|
|
|
|
/** The shared exponent byte for 16 fp32 bit patterns, in the format's encoding. */
|
|
uint8_t ft_bfp_block_shared_exponent(ft_bfp_format format, const uint32_t *block_bits);
|
|
|
|
/**
|
|
* Diagnostic counters since the last reset: rounding ties met by
|
|
* ft_bfp_convert_float_bits (and therefore by packing), and how many of them
|
|
* rounded the magnitude up. Either pointer may be NULL. Not thread-safe.
|
|
*/
|
|
void ft_bfp_tie_counters(uint64_t *ties_seen, uint64_t *ties_rounded_up);
|
|
|
|
/** Zero both tie counters. */
|
|
void ft_bfp_tie_counters_reset(void);
|
|
|
|
#endif
|