[BugFix] Reject round on a vectorized cast from a non-f32 source - #3350

Refs #3301 (F318).

Why

round and its rbits operand have a PTX lowering only for an f32 source, in the four codegen groups that consult cast_round themselves. Every other source dtype reaches a vectorized branch that emits the plain conversion helper and returns, before the rejection at the end of the cast lowering is reached. The requested rounding mode was dropped silently:

for i in T.vectorized(8):
    B[i] = T.cast(A[i], "float8_e4m3fn", round="rs", rbits=T.uint32(seed))

with A in float16 emitted __tl_cvt_half2_to_fp8x2, the round-to-nearest helper, and no rbits operand anywhere in the generated CUDA. The output was bit-identical across seed values 0x0, 0xFFFFFFFF and 0x12345678, and identical to the same cast with no rounding argument at all. The scalar form of the same cast was rejected:

Fatal: round 'rs' is not supported for cast from float16 to float8_e4m3fn
       (only supported f32 packed stochastic conversions are available)

So one request had two outcomes, decided only by whether the loop was vectorized.

What changed

A non-f32 source with a non-empty round is rejected before the vectorized branches, with the message the scalar path already produces. An f32 source is untouched and still takes the stochastic path through the groups that handle it.

Verification

• Two cases added to testing/python/language/test_tilelang_cast_rounding.py, the test module for this feature: the rejection of a vectorized rs cast from float16, and a control that the same vectorized cast with no rounding mode still lowers and still emits __tl_cvt_half2_to_fp8x2. The rejection fails on an unmodified tree with DID NOT RAISE Exception; the control passes in both states.

• test_tilelang_cast_rounding.py, test_tilelang_language_fp8.py, testing/python/quantize/ and test_tilelang_language_vectorized_cast.py: 134 passed, 1 skipped, 19 failed. Those 19 fail identically on an unmodified tree and are an environment limit, not a regression:

  tl_templates/cuda/cuda_fp4.h(313): static assertion failed with
    "Stochastic rounding f32-to-FP4 requires sm_100a or sm_103a"
  

  This machine is sm_120, so the tests that need sm_100a cannot build there. The count is 19 with and without the change.

• The kernel cache was cleared before each run.

• pre-commit run --files <changed files>: all 13 hooks pass.

Summary

Reject vectorized casts that request rounding when the source dtype is not float32. The check uses the existing scalar-path error message. float32 conversion paths remain unchanged.

Add tests that check rejection of vectorized float16 casts with round="rs" and verify that casts without a rounding mode still lower to the half-to-FP8 conversion helper.

Tests

The author reports 134 passed, 1 skipped, and 19 failed across the listed test suites. The author attributes the failures to the test machine's sm_120 architecture, which cannot build tests requiring sm_100a or sm_103a; the author reports the same failures on an unmodified tree. The author also reports that all 13 pre-commit hooks passed. These results were not independently verified.

C++ style / lint notes

The change touches C++ source covered by docs/developer_guide/cpp_style.md, but it does not change the documented style rules. The CI configuration includes the “C++ API Style Audit (warning only)” step. No audit result or current style findings were supplied.

Verified end to end, together with the other contributions from this batch: VERIFICATION.md