NVIDIA / NVIDIA/cutlass

numeric_limits< float_e2m3_t / float_e3m2_t >::lowest() returns a small negative value instead of the format minimum

Open Beginner friendly
#3,493 0 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

CUTLASS C++
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

Description

numeric_limits<cutlass::float_e2m3_t>::lowest() and numeric_limits<cutlass::float_e3m2_t>::lowest() return values near -1.8 / -1.75 instead of the most negative finite value of each format. The same wrong constant is used in the std:: specializations, the cutlass::platform:: specializations, and the float_e2m3_unpack8bits_t / float_e3m2_unpack8bits_t variants, six call sites in total.

The type's own metadata disagrees with the hardcoded constant: BitRepresentation::MIN_VALUE for both formats is 0x3F, and max() (which uses MAX_VALUE = 0x1F) is correct.

Evidence

For E2M3 (bias 1) the most negative finite pattern is 0x3F (s=1, exp=11, man=111, value -(1+7/8)*2^2 = -7.5). The code returns bitcast(0x2f), which decodes to s=1, exp=01, man=111 = -1.875.

For E3M2 (bias 3) the most negative finite pattern is 0x3F (s=1, exp=111, man=11, value -(1+3/4)*2^4 = -28). The returned 0x2f decodes to s=1, exp=011, man=11 = -1.75.

Repro (g++ -std=c++17 -O2 against include/ only):

#include "cutlass/float_subbyte.h"
#include <cstdio>
#include <limits>

int main() {
    printf("e2m3: max=%g lowest=%g\n",
        float(std::numeric_limits<cutlass::float_e2m3_t>::max()),
        float(std::numeric_limits<cutlass::float_e2m3_t>::lowest()));
    printf("e3m2: max=%g lowest=%g\n",
        float(std::numeric_limits<cutlass::float_e3m2_t>::max()),
        float(std::numeric_limits<cutlass::float_e3m2_t>::lowest()));
}

Actual output:

e2m3: max=7.5 lowest=-1.875
e3m2: max=28 lowest=-1.75

Expected: -7.5 and -28. Note lowest() for float_e2m1_t (0xf) is correct, which shows the intent was sign bit | MAX_VALUE.

Suggested fix

Replace bitcast(0x2f) with bitcast(0x3f) in all six places (include/cutlass/float_subbyte.h lines 628, 640, 715, 727, 738, and the e3m2_unpack8bits equivalent), or better, derive it as type::bitcast(type::Base::BitRepresentation::MIN_VALUE) so it cannot drift from the encoding again. epsilon() on these types checks out and needs no change.

No in-tree caller currently reads lowest() on these types, so the defect is latent but silently wrong for any external user relying on numeric_limits.

Contributor guide

No contributing guide indexed for this repository

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start in include/cutlass/float_subbyte.h at the numeric_limits specializations and the float_e2m3_unpack8bits_t/float_e3m2_unpack8bits_t variants; inspect all six lowest() call sites and compare them with BitRepresentation::MIN_VALUE. Verify the reported g++ reproduction, and consider the issue done when both formats and variants return their most negative finite values without changing epsilon().

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
backend
Issue type
Bug
Difficulty
2/5
Estimated time
1-3 hours
Activity status
Active
Clarity
Clearly specified
Newbie friendliness
88/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.