LngelKyo opened a new pull request, #20387:
URL: https://github.com/apache/tvm/pull/20387
### Bug: on `main`, the CUDA codegen prints bare `min()`/`max()` for
`T.min`/`T.max` on float32/float16/bfloat16, so a NaN operand is dropped
With #20054's test data, the CUDA target gives 4/8 bitwise matches: the NaN
lanes are discarded and the ±0 ties disagree with the C host.
### Fix
Same semantics as #20054: emit `(((a > b) || (a != a)) ? a : b)` for `max`
(and the `<` form for `min`), so a NaN on either side is preserved and ties
take `b`. Scalar form binds each operand once via SSA and references the
temporaries; the vector form expands the same expression per lane. Covers
float16/bfloat16/float32/float64. Integer and other non-float min/max keep the
existing base-`CodeGenC` path, including the per-lane vector expansion. Only
the CUDA codegen is touched; the host side is #20054's.
Generated source (scalar and per-lane vector), from the branch:
```
[float32/scalar] C_ptr[((int)threadIdx.x)] = (((v_ > v__1) || (v_ != v_)) ?
v_ : v__1);
[float32/vec4] __1.x = ((v_.x > v__1.x) || (v_.x != v_.x)) ? v_.x : v__1.x;
```
### Tests
Folded into `tests/python/codegen/test_target_codegen_cuda.py`: 4 dtypes
(float32/float64/float16/bfloat16) × 2 ops × {scalar, vec4}, a composite case
(`C[i] = max(A[i], B[i]) + 1.0` — catches the ternary being parsed as `(x +
cond) ? va : vb`; NaN lanes assert NaN, finite lanes bitwise), and an int32
vec4 regression against numpy. Both compile paths (nvcc, nvrtc) are covered by
the autouse fixture.
### Verification
RTX A6000, CUDA 13.0.88:
- positive run: **38/38 passed** (all four dtypes, scalar and vec4, both
compile paths);
- negative control (predicate forced to `return false`): **34 failed**, the
4 int cases stay green (`T7d`-style logs in the gist);
- restored: 38/38 again.
- `[email protected]` check/format and `clang-format 20.1.8 --dry-run --Werror`
clean.
Fixes #19579. Follows #20054.
cc @tlopex @yongwww @swjng — CI will likely need approval as before.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
