[Bug] CUDA: the NVRTC compile path (the default) always passes --use_fast_math; no opt-out like TVM_CUDA_NVCC_NO_FAST_MATH on the nvcc path

[KitKyoD]

Expected behavior

A way to compile CUDA kernels through the default NVRTC path without --use_fast_math, equivalent to what TVM_CUDA_NVCC_NO_FAST_MATH=1 does on the nvcc path. Whether that is the same variable, a new one, or a target attribute is up to you; I am not asking for the default to change. The use case is comparing GPU results bit-for-bit against a CPU reference, which the fast-math intrinsics make impossible.

Actual behavior

python/tvm/support/nvcc.py (main aafeaadc, same lines in the 0.27.0 wheel):

• nvcc path, line 219: --use_fast_math is added unless TVM_CUDA_NVCC_NO_FAST_MATH is set.
• NVRTC path, line 571: b"--use_fast_math" is added unconditionally; TVM_CUDA_NVRTC_EXTRA_OPTS (lines 588-590) is appended after it, so it can override -ftz/-prec-div/-prec-sqrt/-fmad but cannot remove the flag itself.
• line 947: the default compiler is NVRTC (TVM_CUDA_COMPILE_MODE defaults to nvrtc).

The part the extra options cannot undo is the source-level substitution --use_fast_math implies (expf -> __expf, logf -> __logf, x/y -> __fdividef, ...). Measured with the script below (five one-op kernels, 65536 float32 each, default cuda target, apache-tvm 0.27.0, NVRTC 13.0.88, sm_86):

kernel	PTX, default	PTX with TVM_CUDA_NVRTC_EXTRA_OPTS="--ftz=false --prec-div=true --prec-sqrt=true --fmad=false"	PTX with the flag removed
expf	ex2.approx.ftz.f32 mul.ftz.f32	still ex2.approx	fma.rn sequence (46 PTX lines)
logf	lg2.approx.ftz.f32 mul.ftz.f32	still lg2.approx	fma.rn sequence (74 lines)
sqrtf	sqrt.approx.ftz.f32	sqrt.rn.f32	sqrt.rn.f32
/	div.approx.ftz.f32	div.rn.f32	div.rn.f32
* (subnormal x 1)	mul.ftz.f32	mul.rn.f32	mul.f32

Values against NumPy (exp/log: float64 reference rounded to float32; sqrt/div/mul: NumPy float32), same hardware class on two different sm_86 GPUs, identical counts:

kernel	default: elements differing / max ULP	flag removed: differing / max ULP
expf, x in [-20, 20]	53792 / 15	19767 / 2
logf	32278 / 188	2320 / 1
sqrtf	11015 / 1	0 / 0
a / b	18911 / (see note)	0 / 0
a * 1.0f, a subnormal	65536 / (flushed to 0)	0 / 0

Note on the division row: for divisors in [2^126, 2^128) the default build returns 0 where NumPy gives 1.5, 1.0 and 5.877e-39 (sign-preserving zero, as the PTX ISA documents for div.approx.f32), and 3.0 where NumPy gives 3.0; the subnormal row is -ftz=true flushing inputs. With the flag removed all five rows are bit-exact except exp (<= 2 ULP) and log (<= 1 ULP), which is what NVIDIA documents for the non-fast-math expf/logf.

"Flag removed" in the script means filtering --use_fast_math out of the option list inside the script (--drop-flag); it is not a TVM feature, it only shows what an opt-out would give.

One thing to know when testing an opt-out: NVRTC's compilation cache (the driver's compute cache) keys the final cubin without the floating-point options, so on a machine that has already compiled the same source with --use_fast_math, a later compile without it gets the cached fast-math cubin back (the PTX changes, the values do not). The script therefore sets CUDA_CACHE_DISABLE=1 before touching the GPU.

Environment

apache-tvm 0.27.0 (PyPI wheel), apache-tvm-ffi 0.1.14.post1, cuda-bindings 13.0.3, NVRTC 13.0.88, CUDA runtime 13.0.96, Python 3.12.3, Linux x86_64, GPU sm_86. The nvcc.py lines quoted above were checked on main aafeaadc (2026-10-02); the measurements were taken with the 0.27.0 wheel. The nvcc path (TVM_CUDA_COMPILE_MODE=nvcc with TVM_CUDA_NVCC_NO_FAST_MATH=1, nvcc 13.0.88) gives the precise PTX, so the variable itself works there; I have not measured values on that path.

Steps to reproduce

python repro_nvrtc_fast_math.py                      # default: fast-math PTX + values vs NumPy
TVM_CUDA_NVCC_NO_FAST_MATH=1 python repro_nvrtc_fast_math.py        # no effect on the NVRTC path
TVM_CUDA_NVRTC_EXTRA_OPTS="--ftz=false --prec-div=true --prec-sqrt=true --fmad=false" python repro_nvrtc_fast_math.py
python repro_nvrtc_fast_math.py --drop-flag          # what an opt-out would give

Without a GPU the script prints only the PTX table; with one it also prints the value table.

repro_nvrtc_fast_math.py

"""What TVM's default CUDA build does to float32 math.

Five one-line kernels are built with the default CUDA target (fast math is not requested).
The script prints the float32 instructions in the PTX that tvm.support.nvcc produces for them
and, when a CUDA device is present, runs them and compares every element with NumPy.

    python repro_nvrtc_fast_math.py               # TVM as shipped
    python repro_nvrtc_fast_math.py --drop-flag   # --use_fast_math filtered out of the NVRTC options

Needs: apache-tvm, numpy, cuda-bindings, and a CUDA toolkit that tvm.support.nvcc can find.
"""

import os

# Both runs must compile from scratch, whatever was compiled on this machine before.
os.environ.setdefault("CUDA_CACHE_DISABLE", "1")

import argparse  # noqa: E402
import re  # noqa: E402

import numpy as np  # noqa: E402
import tvm  # noqa: E402
import tvm.support.nvcc as nvcc  # noqa: E402
from tvm.script import ir as I  # noqa: E402
from tvm.script import tirx as T  # noqa: E402

NB, NT = 256, 256
N = NB * NT


@I.ir_module
class Exp:
    @T.prim_func
    def main(A: T.Buffer((65536,), "float32"), C: T.Buffer((65536,), "float32")):
        T.func_attr({"tirx.noalias": True})
        for b in T.thread_binding(256, thread="blockIdx.x"):
            for t in T.thread_binding(256, thread="threadIdx.x"):
                C[b * 256 + t] = T.exp(A[b * 256 + t])


@I.ir_module
class Log:
    @T.prim_func
    def main(A: T.Buffer((65536,), "float32"), C: T.Buffer((65536,), "float32")):
        T.func_attr({"tirx.noalias": True})
        for b in T.thread_binding(256, thread="blockIdx.x"):
            for t in T.thread_binding(256, thread="threadIdx.x"):
                C[b * 256 + t] = T.log(A[b * 256 + t])


@I.ir_module
class Sqrt:
    @T.prim_func
    def main(A: T.Buffer((65536,), "float32"), C: T.Buffer((65536,), "float32")):
        T.func_attr({"tirx.noalias": True})
        for b in T.thread_binding(256, thread="blockIdx.x"):
            for t in T.thread_binding(256, thread="threadIdx.x"):
                C[b * 256 + t] = T.sqrt(A[b * 256 + t])


@I.ir_module
class Div:
    @T.prim_func
    def main(
        A: T.Buffer((65536,), "float32"),
        B: T.Buffer((65536,), "float32"),
        C: T.Buffer((65536,), "float32"),
    ):
        T.func_attr({"tirx.noalias": True})
        for b in T.thread_binding(256, thread="blockIdx.x"):
            for t in T.thread_binding(256, thread="threadIdx.x"):
                C[b * 256 + t] = A[b * 256 + t] / B[b * 256 + t]


@I.ir_module
class Mul:
    @T.prim_func
    def main(
        A: T.Buffer((65536,), "float32"),
        B: T.Buffer((65536,), "float32"),
        C: T.Buffer((65536,), "float32"),
    ):
        T.func_attr({"tirx.noalias": True})
        for b in T.thread_binding(256, thread="blockIdx.x"):
            for t in T.thread_binding(256, thread="threadIdx.x"):
                C[b * 256 + t] = A[b * 256 + t] * B[b * 256 + t]


def inputs():
    rng = np.random.default_rng(0)
    f32 = np.float32
    wide = (2.0 ** rng.uniform(-100, 100, N)).astype(f32)
    a = (rng.standard_normal(N) * 2.0 ** rng.uniform(-20, 20, N)).astype(f32)
    b = (rng.standard_normal(N) * 2.0 ** rng.uniform(-20, 20, N)).astype(f32)
    b[b == 0] = f32(1.0)
    # first lanes of the division: divisor in [2^126, 2^128)
    a[:4] = np.array([1.5 * 2.0**127, 3.0 * 2.0**126, 2.0**127, 1.0], dtype=f32)
    b[:4] = np.array([2.0**127, 2.0**126, 2.0**127, 2.0**127], dtype=f32)
    tiny = np.full(N, 1e-40, dtype=f32)  # subnormal
    tiny[1::2] = f32(3e-39)
    one = np.ones(N, dtype=f32)
    x = rng.uniform(-20, 20, N).astype(f32)
    with np.errstate(all="ignore"):
        return {
            "exp": (Exp, [x], np.exp(x.astype(np.float64)).astype(f32)),
            "log": (Log, [wide], np.log(wide.astype(np.float64)).astype(f32)),
            "sqrt": (Sqrt, [wide], np.sqrt(wide)),
            "div": (Div, [a, b], a / b),
            "mul (subnormal * 1)": (Mul, [tiny, one], tiny * one),
        }


def ordered(v):
    i = v.view(np.int32).astype(np.int64)
    return np.where(i < 0, -(i & 0x7FFFFFFF), i)


def compare(got, ref):
    finite = np.isfinite(got) & np.isfinite(ref)
    ulp = np.abs(ordered(got[finite]) - ordered(ref[finite]))
    other = int(np.sum(~finite & ~((got == ref) | (np.isnan(got) & np.isnan(ref)))))
    return int(np.sum(ulp > 0)), int(ulp.max()) if ulp.size else 0, other


def drop_fast_math():
    from cuda.bindings import nvrtc

    original = nvrtc.nvrtcCompileProgram

    def compile_without_flag(prog, num_options, options):
        kept = [o for o in options if o != b"--use_fast_math"]
        return original(prog, len(kept), kept)

    nvrtc.nvrtcCompileProgram = compile_without_flag


def main():
    ap = argparse.ArgumentParser()
    ap.add_argument("--drop-flag", action="store_true")
    ap.add_argument("--arch", default="sm_80", help="used when no CUDA device is present")
    args = ap.parse_args()
    if args.drop_flag:
        drop_fast_math()

    dev = tvm.cuda(0)
    arch = "sm_" + dev.compute_version.replace(".", "") if dev.exist else args.arch
    target = tvm.target.Target({"kind": "cuda", "arch": arch})
    print(
        f"tvm {tvm.__version__}  arch {arch}  device {dev.exist}  drop-flag {args.drop_flag}"
        f"  CUDA_CACHE_DISABLE={os.environ.get('CUDA_CACHE_DISABLE')}"
    )

    print(f"\n{'kernel':<22}{'generated source':<22}{'PTX lines':>10}  float32 instructions in the PTX")
    libs = {}
    for name, (mod, _, _) in inputs().items():
        lib = tvm.compile(mod, target=target)
        libs[name] = lib
        src = lib.mod.imports[0].inspect_source()
        stmt = [l for l in src[src.index("__global__") :].splitlines() if "C_ptr[" in l][0]
        rhs = re.sub(r"([AB])_ptr\[[^\]]*\]", r"\1[i]", stmt.split("=", 1)[1].strip(" ;"))
        with target:
            ptx = bytes(nvcc.compile_cuda(src, target_format="ptx", compiler="nvrtc")).decode()
        entry = ptx[ptx.index(".entry") :]
        ops = sorted(set(re.findall(r"\b((?:div|sqrt|ex2|lg2|mul|fma)\.[a-z0-9.]*f32)", entry)))
        print(f"{name:<22}{rhs:<22}{entry.count(chr(10)):>10}  {' '.join(ops)}")

    if not dev.exist:
        print("\nno CUDA device: PTX only")
        return

    print(f"\n{'kernel':<22}{'n':>7}{'differ':>9}{'max ulp':>12}{'inf/nan differ':>16}")
    for name, (_, arrays, ref) in inputs().items():
        ins = [tvm.runtime.tensor(v, dev) for v in arrays]
        out = tvm.runtime.empty((N,), "float32", dev)
        libs[name](*ins, out)
        dev.sync()
        got = out.numpy()
        differ, max_ulp, other = compare(got, ref)
        print(f"{name:<22}{N:>7}{differ:>9}{max_ulp:>12}{other:>16}")
        if name == "div":
            print(f"    divisor in [2^126, 2^128): got {got[:4]}  NumPy {ref[:4]}")
        if name.startswith("mul"):
            print(f"    subnormal inputs:          got {got[:2]}  NumPy {ref[:2]}")


if __name__ == "__main__":
    main()

Triage

• needs-triage
• backend:cuda

[KitKyoD]

The script above targets the 0.27.0 wheel. On current main, replace T.Buffer with T.Tensor (renamed in #20519); the PTX table is unchanged on main 043b57fbe9b8.