[BugFix][Arith] Fold float Min/Max commutatively on NaN

[sohampanda000]

Problem

T.max / T.min on two float compile-time constants produce a result that depends on argument order when one operand is NaN:

T.max(1.0, nan)  ->  1.0
T.max(nan, 1.0)  ->  nan

(same for T.min). The kernel compiles without error and the wrong value is baked into the emitted code as a constant.

The identical expression is order-independent once the operands are not compile-time constants: the runtime path lowers to fmaxf/fminf, which are NaN-quiet and return the non-NaN operand regardless of order. So the constant fold disagrees both with the runtime lowering of the same operation and with itself under operand swap.

Reported downstream as tile-ai/tilelang#2882.

Root cause

TryConstFold<tirx::Min> and TryConstFold<tirx::Max> in src/arith/const_fold.h use std::min / std::max on the float-constant branch:

if (fa && fb) return FloatImm(rtype, std::min(fa->value, fb->value));
...
if (fa && fb) return FloatImm(rtype, std::max(fa->value, fb->value));

std::max(a, b) is specified as a < b ? b : a. Every comparison against NaN is false, so it returns a — whichever operand happens to be first. std::fmax / std::fmin are the NaN-quiet counterparts and return the non-NaN operand in either order.

Fix

Use std::fmax / std::fmin on the float branch only. Two lines.

The integer branch is deliberately unchanged — integers have no NaN, and std::min/std::max are correct and cheaper there. <cmath> is already included by this header, so there is no new dependency.

Behaviour change

Two cases, both currently order-dependent — so no well-defined behaviour is being altered.

NaN. The motivating case:

before:  fold max(1.0, nan) = 1.0   fold max(nan, 1.0) = nan
after:   fold max(1.0, nan) = 1.0   fold max(nan, 1.0) = 1.0
runtime: fmaxf(1.0, nan)    = 1.0   fmaxf(nan, 1.0)    = 1.0

Signed zero. Raised in review, and it applies equally. +0.0 < -0.0 and -0.0 < +0.0 are both false, so std::max/std::min return whichever operand came first there too:

before:  fold max(+0.0, -0.0) = +0.0   fold max(-0.0, +0.0) = -0.0
         fold min(+0.0, -0.0) = +0.0   fold min(-0.0, +0.0) = -0.0
after:   fold max(+0.0, -0.0) = +0.0   fold max(-0.0, +0.0) = +0.0
         fold min(+0.0, -0.0) = -0.0   fold min(-0.0, +0.0) = -0.0

So the change makes both cases order-independent. Note the C standard leaves fmax(+0.0, -0.0) implementation-defined as to which zero is returned; what it guarantees, and what matters here, is that the result no longer depends on operand order.

[Copilot]

Pull request overview

This PR fixes incorrect constant folding of tirx::Min/tirx::Max for floating-point compile-time constants when NaN is involved, making the folded result order-independent and consistent with the runtime lowering (fminf/fmaxf).

Changes:

• Replace std::min/std::max with std::fmin/std::fmax for the float-constant folding path in TryConstFold<tirx::Min> and TryConstFold<tirx::Max>.
• Keep integer constant-folding behavior unchanged.

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[sohampanda000]

Thanks for the review — both points are fair.

Signed zero. Correct, and it applies to std::min/std::max for the same reason as NaN: +0.0 < -0.0 and -0.0 < +0.0 are both false, so the result is whichever operand came first. Confirmed against libc++:

std::max(+0,-0) = +0.0   std::max(-0,+0) = -0.0     <- order-dependent
std::min(+0,-0) = +0.0   std::min(-0,+0) = -0.0     <- order-dependent
std::fmax(+0,-0) = +0.0  std::fmax(-0,+0) = +0.0    <- order-independent
std::fmin(+0,-0) = -0.0  std::fmin(-0,+0) = -0.0    <- order-independent

So this case is also currently order-dependent and is also fixed. I have updated the PR description with it rather than claiming the change is NaN-only.

Regression test. Agreed it needs one. I would rather add a test I have actually executed than one written against a guessed API, so I will push it in a follow-up commit on this branch once I can run it — covering both operand orders for (1.0, nan) and (+0.0, -0.0), for both min and max.

Happy to hold the PR until that lands if you would prefer to review it as one piece.