Fix min_float32 for parsing float literals

[cchan]

The minimum normal 2**-126, but the minimum denormal is 2**-23 times smaller than that.

[jlebar]

On second thought, I'm not sure about this PR.

Subnormal values lose precision. I think if we want to accept them as float32's, we should also check that they are exactly representable as float32?

[jlebar]

exactly representable as float32?

I grant this is also a little weird because if your input is e.g. 1/3, that will be represented as a python f64, and it's not exactly-representable as a f32. But we'll still cast it to a Triton f32. So why have a special case for subnormals? Except I feel like it's one thing to lose half the mantissa bits of 1/3, while it's another thing to lose potentially all but one of the mantissa bits of a f64 that is represented as an f32 subnormal.

[ThomasRaoux]

I agree with Justin, in addition to that one thing to consider is that this breaks backward compatibility and may break existing kernels in a non trivial way.

[cchan]

Hmmm, those are fair points. There should be a way to produce custom constexpr types somehow, though.

e.g. what about something like this?

x : tl.constexpr[tl.float32] = FLOAT32_MIN_POSITIVE
y : tl.constexpr[tl.float64] = 1/3

[ThomasRaoux]

Hmmm, those are fair points. There should be a way to produce custom constexpr types somehow, though.

e.g. what about something like this?

x : tl.constexpr[tl.float32] = FLOAT32_MIN_POSITIVE
y : tl.constexpr[tl.float64] = 1/3

can you just assign the constexpr and do tl.cast to get to the type you want?

[cchan]

I did what y'all suggested, and it seems to have been smart enough to do the right thing, turning it into an immediate.

        mul.f32         %f9, %f8, 0f07800000;
        mul.f32         %f10, %f7, 0f07800000;
        mul.f32         %f11, %f6, 0f07800000;
        mul.f32         %f12, %f5, 0f07800000;
        mul.f32         %f13, %f4, 0f07800000;
        mul.f32         %f14, %f3, 0f07800000;
        mul.f32         %f15, %f2, 0f07800000;
        mul.f32         %f16, %f1, 0f07800000;

Thanks both :)