This is an automated email from the ASF dual-hosted git repository.
tlopex pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm-tirx.git
The following commit(s) were added to refs/heads/main by this push:
new 05961665 [FIX][CUDA] Preserve NaNs in floating-point min and max (#23)
05961665 is described below
commit 05961665c7b97feb473be07b8ff49629d76af4ed
Author: LngelKyo <[email protected]>
AuthorDate: Sat Oct 24 05:32:14 2026 +0900
[FIX][CUDA] Preserve NaNs in floating-point min and max (#23)
This is apache/tvm#20387 re-targeted at this repository after the TIRx
sources moved here in apache/tvm#20803. The diff is the same apart from
the test file path and the module paths in the tests (`tvm.compile` ->
`tvm_tirx.tirx.compile`; `tvm.testing`, `tvm.runtime`, `tvm.cuda` -> the
`tvm_tirx.` forms the rest of the file uses).
### 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 apache/tvm#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 the host-side proposal apache/tvm#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 bound names; the vector
form expands per lane.
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;
```
bfloat16 constants: `PrintConst` printed a bfloat16 NaN or Inf constant
as `__float2bfloat16_rn(nanf/*nan*/)` / `(inff/*inf*/)`, which does not
compile (the float32 branch prints `CUDART_NAN_F` / `CUDART_INF_F`, and
the float16 branch delegates to it). The bfloat16 branch now delegates
to the float32 printer in the same way; finite bfloat16 constants print
as before. The bfloat16 rows of the test need it; it is a separate
commit.
### Tests
Folded into `tests/python/backend/cuda/codegen/test_basic.py` as one
parametrized test: for each dtype in {float32, float64, float16,
bfloat16} × {scalar, vec4}, a kernel computes the 15-row matrix of the
host-side test in #31 — min and max of (a, b), (0, a), (a, 0), (nan, a),
(a, nan); `max(a, b) + 1` and `max(0, a) + 1` (which catch the ternary
being parsed as `(x + cond) ? a : b`); the chained `C = 0; C = max(C,
3); C = max(C, 2)` row; `min(a, +inf)` and `max(a, −inf)` — and an int32
kernel computes min and max and must keep the base path. 10 cases, all
requesting the `backend_config` fixture, so they run on both the nvcc
and NVRTC paths (20 in total). For float32 the generated source is also
checked: the general rows emit the `||` form, the constant NaN rows
`CUDART_NAN_F`, a scalar constant rhs the reversed non-strict compare,
and no bare `min()`/`max()`.
The scalar path's SSA bindings must not stay alive across statements:
`SSAGetID` caches by printed text within the live scope and the
base-codegen `BufferStore` dispatch does not invalidate it, so `C[v] =
max(C[v], A[v]); C[v] = max(C[v], B[v])` would reuse the first read of
`C[v]`. Row 12 of the matrix (`C = 0; C = max(C, 3); C = max(C, 2)`)
pins this: a stale cached read gives 2.
Constant-operand fast paths (from the review of apache/tvm#20054): a
constant NaN lhs is the result (both operands are bound before the
classification, so the rhs is still evaluated); a constant NaN rhs needs
only the `(a != a)` test; a constant non-NaN lhs keeps the ordered
compare; a constant non-NaN rhs reverses it to the non-strict form (`<`
-> `>=`, `>` -> `<=`) so ties still take `b`. The three ternary forms
are parenthesized so they survive nesting in a compound expression (row
11 of the matrix); the constant-NaN-lhs form prints the constant alone.
The tests are written as plain TIRx (`@T.function` with `T.constexpr`
selections) for the transpiler parser.
### Verification
Before the move, in apache/tvm:
- RTX A6000 (sm_86) and L40 (sm_89), CUDA 13.0.88: `-k min_max` **74
passed** (nvcc + nvrtc); the whole `test_target_codegen_cuda.py` of the
time **440 passed, 6 skipped**; `test_gpu_codegen_allreduce.py` **328
passed, 330 skipped**; negative control with the predicate forced to
`return false`: **68 failed / 6 passed** (the 6 are the 4 int-vector and
2 chained cases; the chained ones are green by construction because the
control bypasses the new path); restored: **74/74**.
- apache/tvm Jenkins GPU job on the 2026-10-14 merge with main:
`test_min_max` **74 passed** (37 nvcc + 37 NVRTC), 0 failed.
- `[email protected]` check/format and `clang-format 20.1.8 --dry-run
--Werror` clean.
NaN constants and NaN results are only guaranteed to be NaN, not a
particular payload (on sm_89 nvcc normalizes some forms to
`0x7FFFFFFF`); the tests assert `isnan` on the NaN lanes and bitwise
equality on the finite lanes.
After the move: the same patch applied onto this repository's `main`
(the C++ hunks apply unchanged at shifted offsets; the tests change only
the module paths above). I have not re-run the GPU tests against this
repository's build, so the numbers above are from apache/tvm; the
Jenkins GPU pipeline here runs this file on both compiler paths.
Partially addresses apache/tvm#19579 (CUDA side only). The host side (C
host and LLVM codegen, plus the constant folding) is #31 in this
repository; this PR does not touch those paths. The constant
classification here is local to the CUDA codegen (scalar `FloatImm`); it
can move to `src/codegen/min_max_utils.h` once #31 lands.
---
src/backend/cuda/codegen/codegen_cuda.cc | 123 +++++++++++++++++++++-
src/backend/cuda/codegen/codegen_cuda.h | 6 ++
tests/python/backend/cuda/codegen/test_basic.py | 133 ++++++++++++++++++++++++
3 files changed, 259 insertions(+), 3 deletions(-)
diff --git a/src/backend/cuda/codegen/codegen_cuda.cc
b/src/backend/cuda/codegen/codegen_cuda.cc
index 068779bd..32a91899 100644
--- a/src/backend/cuda/codegen/codegen_cuda.cc
+++ b/src/backend/cuda/codegen/codegen_cuda.cc
@@ -73,6 +73,25 @@ TVM_FFI_INLINE bool IsPackedFloat(const PrimType& ty) {
return IsFloat8(ty) || IsFloat6(ty) || IsFloat4(ty);
}
+// Constant-operand classification for min/max, the same classification as
+// the host side's min_max_utils.h (#31), scalar FloatImm only here. A
+// constant NaN lhs is the result (the rhs is still bound and evaluated); a
+// constant NaN rhs keeps the lhs on any comparison outcome; a constant
+// non-NaN operand lets the ternary skip its NaN clause (an ordered compare
+// against a constant is never NaN). A vector constant (Broadcast of a
+// FloatImm) takes the general per-lane path.
+TVM_FFI_INLINE const FloatImmObj* AsFloatImm(const PrimExpr& e) { return
e.as<FloatImmObj>(); }
+
+// NaN-preserving min/max is emitted for float, half and bfloat16: the C
+// ternary below keeps the NaN operand instead of discarding it (apache/tvm
+// PR #20054). Integer min/max stays as it is.
+TVM_FFI_INLINE bool IsFloatMinMaxNanPreserving(const PrimType& ty) {
+ bool is_fp = ty.MatchesCode(DLDataTypeCode::kDLFloat) &&
+ (ty.bits() == 16 || ty.bits() == 32 || ty.bits() == 64);
+ bool is_bf16 = ty.MatchesCode(DLDataTypeCode::kDLBfloat) && ty.bits() == 16;
+ return is_fp || is_bf16;
+}
+
} // namespace
std::string GetFP8Type(const PrimType& type_ty) {
@@ -736,6 +755,104 @@ void CodeGenCUDA::PrintVecConstructor(const PrimType& t,
std::ostream& os) {
PrintType(t, os);
}
+// min/max for CUDA, keeping a NaN operand (apache/tvm PR #20054), for
+// float16/bfloat16/float32/float64. For a NaN-free pair this is min()/max();
+// the lhs-NaN clause fires on the first operand and the rhs-NaN case is
+// covered because the ordering compare is false when either operand is NaN.
+template <typename T>
+void CodeGenCUDA::PrintMinMaxNanPreservingImpl(const T* op, const char* opstr,
+ std::ostream& os) { //
NOLINT(*)
+ PrimType op_ty = op->ty.template as_or_throw<PrimType>();
+ const char* cmp = (opstr[0] == 'm' && opstr[1] == 'i') ? "<" : ">";
+ if (!IsFloatMinMaxNanPreserving(op_ty)) {
+ // Integer / non-float min/max keeps the base-codegen path, including the
+ // vectorized per-lane expansion for lanes > 1.
+ this->CodeGenC::Dispatch_(op, os);
+ return;
+ }
+ if (op_ty.lanes() > 1) {
+ this->PrintVecBinaryOpNanPreserving(std::string(opstr), op_ty, op->a,
op->b, cmp, os);
+ return;
+ }
+ // Bind both operands once (SSAGetID, in the expression's own scope: the
+ // same binding the host side and PrintVecBinaryOp use), then reference the
+ // temporaries: the expressions must not be re-evaluated per clause. Both
+ // are bound before the constant classification below, so a constant NaN
+ // lhs still evaluates the rhs and only the select is dropped. A
+ // float32/float64 constant comes back from SSAGetID as itself (PrintConst
+ // registers its text with MarkConst); a float16/bfloat16 constant is
+ // wrapped in a conversion call that is not registered and gets a temporary
+ // like any other operand. The bindings live in their own scope so a later
+ // statement that prints the same text (e.g.
+ // `red_buf[0] = max(red_buf[0], shuffle_down(...))` repeated by warp
+ // reduction) does not hit the cache and read the pre-write value.
+ int ssa_scope = BeginScope();
+ std::string va = SSAGetID(PrintExpr(op->a), op->a.ty());
+ std::string vb = SSAGetID(PrintExpr(op->b), op->b.ty());
+ const FloatImmObj* ca = AsFloatImm(op->a);
+ const FloatImmObj* cb = AsFloatImm(op->b);
+ if (ca != nullptr && std::isnan(ca->value)) {
+ // a constant NaN lhs is the result; the rhs stays bound above.
+ os << va;
+ } else if (ca != nullptr) {
+ // lhs is a constant non-NaN: ordered compare, b's NaN makes the compare
+ // false and b is kept.
+ os << "((" << va << ' ' << cmp << ' ' << vb << ") ? " << va << " : " << vb
<< ')';
+ } else if (cb != nullptr && std::isnan(cb->value)) {
+ // rhs is a constant NaN: keep a on any comparison outcome.
+ os << "((" << va << " != " << va << ") ? " << va << " : " << vb << ')';
+ } else if (cb != nullptr) {
+ // rhs is a constant non-NaN: reverse to a non-strict compare so a's NaN
+ // makes the compare false and a is kept; ties take b.
+ const char* rcmp = (cmp[0] == '<') ? ">=" : "<=";
+ os << "((" << va << ' ' << rcmp << ' ' << vb << ") ? " << vb << " : " <<
va << ')';
+ } else {
+ os << "(((" << va << ' ' << cmp << ' ' << vb << ") || (" << va << " != "
<< va << ")) ? " << va
+ << " : " << vb << ')';
+ }
+ EndScope(ssa_scope);
+}
+
+void CodeGenCUDA::Dispatch_(const prim::MinObj* op, std::ostream& os) { //
NOLINT(*)
+ this->PrintMinMaxNanPreservingImpl(op, "min", os);
+}
+
+void CodeGenCUDA::Dispatch_(const prim::MaxObj* op, std::ostream& os) { //
NOLINT(*)
+ this->PrintMinMaxNanPreservingImpl(op, "max", os);
+}
+
+void CodeGenCUDA::PrintVecBinaryOpNanPreserving(const std::string& op, const
PrimType& t,
+ PrimExpr lhs, PrimExpr rhs,
const char* cmp,
+ std::ostream& os) { //
NOLINT(*)
+ std::string sret = name_supply_->FreshName("_");
+ this->PrintIndent();
+ this->PrintType(t, stream);
+ stream << ' ' << sret << ";\n";
+ int ssa_scope = BeginScope();
+ {
+ std::string vlhs = SSAGetID(PrintExpr(lhs), lhs.ty());
+ std::string vrhs = SSAGetID(PrintExpr(rhs), rhs.ty());
+ for (int i = 0, lanes = t.lanes(); i < lanes; ++i) {
+ std::ostringstream value_temp;
+ value_temp << "((";
+ PrintVecElemLoad(vlhs, lhs.ty(), i, value_temp);
+ value_temp << ' ' << cmp << ' ';
+ PrintVecElemLoad(vrhs, rhs.ty(), i, value_temp);
+ value_temp << ") || (";
+ PrintVecElemLoad(vlhs, lhs.ty(), i, value_temp);
+ value_temp << " != ";
+ PrintVecElemLoad(vlhs, lhs.ty(), i, value_temp);
+ value_temp << ")) ? ";
+ PrintVecElemLoad(vlhs, lhs.ty(), i, value_temp);
+ value_temp << " : ";
+ PrintVecElemLoad(vrhs, rhs.ty(), i, value_temp);
+ PrintVecElemStore(sret, t, i, value_temp.str());
+ }
+ }
+ EndScope(ssa_scope);
+ os << sret;
+}
+
void CodeGenCUDA::PrintVecBinaryOp(const std::string& op, const PrimType& t,
PrimExpr lhs,
PrimExpr rhs, std::ostream& os) { //
NOLINT(*)
// Declare the result.
@@ -1503,9 +1620,9 @@ inline void PrintConst(const FloatImmObj* op,
std::ostream& os, CodeGenCUDA* p)
PrimType op_ty = op->ty.as_or_throw<PrimType>();
// Type code is kBFloat
if (op_ty.MatchesElementType(DLDataTypeCode::kDLBfloat, 16)) {
- os << "__float2bfloat16_rn";
- os << '(' << std::hexfloat << op->value << 'f';
- os << "/*" << std::scientific << op->value << "*/";
+ os << "__float2bfloat16_rn" << '(';
+ FloatImm const_f32 = FloatImm(PrimType::Float(32), op->value);
+ PrintConst(const_f32.get(), os, p);
os << ')';
return;
}
diff --git a/src/backend/cuda/codegen/codegen_cuda.h
b/src/backend/cuda/codegen/codegen_cuda.h
index 133a6f70..f8292bab 100644
--- a/src/backend/cuda/codegen/codegen_cuda.h
+++ b/src/backend/cuda/codegen/codegen_cuda.h
@@ -83,6 +83,12 @@ class CodeGenCUDA final : public CodeGenC {
void Dispatch_(const FloatImmObj* op, std::ostream& os) final;
void Dispatch_(const CallObj* op, std::ostream& os) final;
void Dispatch_(const prim::CastObj* op, std::ostream& os) final;
+ void Dispatch_(const prim::MinObj* op, std::ostream& os) final; // NOLINT(*)
+ void Dispatch_(const prim::MaxObj* op, std::ostream& os) final; // NOLINT(*)
+ template <typename T>
+ void PrintMinMaxNanPreservingImpl(const T* op, const char* opstr,
std::ostream& os);
+ void PrintVecBinaryOpNanPreserving(const std::string& op, const PrimType& t,
PrimExpr lhs,
+ PrimExpr rhs, const char* cmp,
std::ostream& os);
void Dispatch_(const EvaluateObj* op) final;
void Dispatch_(const ReturnObj* op) final;
void Dispatch_(const BindObj* op) final;
diff --git a/tests/python/backend/cuda/codegen/test_basic.py
b/tests/python/backend/cuda/codegen/test_basic.py
index 0b2ee6d0..f2596f0b 100644
--- a/tests/python/backend/cuda/codegen/test_basic.py
+++ b/tests/python/backend/cuda/codegen/test_basic.py
@@ -791,6 +791,139 @@ def test_vectorized_intrin1(*, backend_config):
run_test(*func, "float16")
+def _min_max_module(dtype, form, n=12):
+ # One kernel per (dtype, form) computes a 15-row matrix (2 rows for
int32), the
+ # shape of the host side's test (#31): rows 0-4 min of (A,B), (0,A),
(A,0), (nan,A),
+ # (A,nan); 5-9 the same for max; 10-11 max(A,B)+1 and max(0,A)+1 (the
ternary must
+ # stay parenthesized before "+"); 12 chained C=0; C=max(C,3); C=max(C,2)
(a stale
+ # cached read of C gives 2); 13-14 min(A,+inf), max(A,-inf). The constants
are
+ # FloatImms built outside the kernel (-inf as its own: a negated float16
FloatImm is
+ # not folded): scalar takes codegen's constant fast paths (NaN constant
inline,
+ # constant rhs as a reversed non-strict compare), vec4 the per-lane "||"
path.
+ is_float = dtype != "int32"
+ rows = 15 if is_float else 2
+ zero, one = T.cast(0, dtype), T.cast(1, dtype)
+ nan = T.cast(float("nan"), dtype) if is_float else None
+ inf = T.cast(float("inf"), dtype) if is_float else None
+ ninf = T.cast(float("-inf"), dtype) if is_float else None
+
+ @I.ir_module
+ class Module:
+ @T.function
+ def main(A: T.Tensor((n,), dtype), B: T.Tensor((n,), dtype), C:
T.Tensor((rows, n), dtype)):
+ T.func_attr({"tirx.noalias": True})
+ if T.constexpr(form == "vec4"):
+ with T.launch_thread("blockIdx.x", n // 4) as i0:
+ for i1 in T.vectorized(4):
+ if T.constexpr(is_float):
+ C[0, i0 * 4 + i1] = T.min(A[i0 * 4 + i1], B[i0 * 4
+ i1])
+ C[1, i0 * 4 + i1] = T.min(zero, A[i0 * 4 + i1])
+ C[2, i0 * 4 + i1] = T.min(A[i0 * 4 + i1], zero)
+ C[3, i0 * 4 + i1] = T.min(nan, A[i0 * 4 + i1])
+ C[4, i0 * 4 + i1] = T.min(A[i0 * 4 + i1], nan)
+ C[5, i0 * 4 + i1] = T.max(A[i0 * 4 + i1], B[i0 * 4
+ i1])
+ C[6, i0 * 4 + i1] = T.max(zero, A[i0 * 4 + i1])
+ C[7, i0 * 4 + i1] = T.max(A[i0 * 4 + i1], zero)
+ C[8, i0 * 4 + i1] = T.max(nan, A[i0 * 4 + i1])
+ C[9, i0 * 4 + i1] = T.max(A[i0 * 4 + i1], nan)
+ C[10, i0 * 4 + i1] = T.max(A[i0 * 4 + i1], B[i0 *
4 + i1]) + one
+ C[11, i0 * 4 + i1] = T.max(zero, A[i0 * 4 + i1]) +
one
+ C[12, i0 * 4 + i1] = zero
+ C[12, i0 * 4 + i1] = T.max(C[12, i0 * 4 + i1],
T.cast(3, dtype))
+ C[12, i0 * 4 + i1] = T.max(C[12, i0 * 4 + i1],
T.cast(2, dtype))
+ C[13, i0 * 4 + i1] = T.min(A[i0 * 4 + i1], inf)
+ C[14, i0 * 4 + i1] = T.max(A[i0 * 4 + i1], ninf)
+ else:
+ C[0, i0 * 4 + i1] = T.min(A[i0 * 4 + i1], B[i0 * 4
+ i1])
+ C[1, i0 * 4 + i1] = T.max(A[i0 * 4 + i1], B[i0 * 4
+ i1])
+ else:
+ with T.launch_thread("threadIdx.x", n) as i:
+ if T.constexpr(is_float):
+ C[0, i] = T.min(A[i], B[i])
+ C[1, i] = T.min(zero, A[i])
+ C[2, i] = T.min(A[i], zero)
+ C[3, i] = T.min(nan, A[i])
+ C[4, i] = T.min(A[i], nan)
+ C[5, i] = T.max(A[i], B[i])
+ C[6, i] = T.max(zero, A[i])
+ C[7, i] = T.max(A[i], zero)
+ C[8, i] = T.max(nan, A[i])
+ C[9, i] = T.max(A[i], nan)
+ C[10, i] = T.max(A[i], B[i]) + one
+ C[11, i] = T.max(zero, A[i]) + one
+ C[12, i] = zero
+ C[12, i] = T.max(C[12, i], T.cast(3, dtype))
+ C[12, i] = T.max(C[12, i], T.cast(2, dtype))
+ C[13, i] = T.min(A[i], inf)
+ C[14, i] = T.max(A[i], ninf)
+ else:
+ C[0, i] = T.min(A[i], B[i])
+ C[1, i] = T.max(A[i], B[i])
+
+ return Module
+
+
[email protected]
[email protected](not env.has_cuda(), reason="need cuda")
[email protected]("form", ["scalar", "vec4"])
[email protected]("dtype", ["float32", "float64", "float16",
"bfloat16", "int32"])
+def test_min_max_nan_preserving_cuda(dtype, form, *, backend_config):
+ n = 12
+ is_float = dtype != "int32"
+ mod = tvm_tirx.tirx.compile(
+ _min_max_module(dtype, form, n), target={"kind": "cuda",
"backend_config": backend_config}
+ )
+ # The CUDA header before the kernel has "||" in preprocessor lines, so the
text is
+ # checked from the kernel declaration on. float never falls back to bare
min()/max()
+ # (the general rows emit "||", the constant NaN rows CUDART_NAN_F, a
scalar constant
+ # rhs the reversed non-strict compare); int32 keeps the base min()/max()
path.
+ src = mod.mod.imports[0].inspect_source()
+ kernel = src[src.index("__global__") :]
+ if dtype == "float32":
+ assert "||" in kernel and "CUDART_NAN_F" in kernel, kernel
+ assert "min(" not in kernel and "max(" not in kernel, kernel
+ if form == "scalar":
+ assert " >= " in kernel or " <= " in kernel, kernel
+ elif not is_float:
+ assert "min(" in kernel and "max(" in kernel, kernel
+
+ if is_float:
+ inf = np.inf
+ a_np = np.array([np.nan, 1, np.nan, 0.0, -0.0, 3, 2, -5, inf, -inf, 1,
-1], dtype=dtype)
+ b_np = np.array([1, np.nan, np.nan, -0.0, 0.0, 2, 2, -4, 1, 1, inf,
-inf], dtype=dtype)
+ zero, nan = np.zeros(n, dtype=dtype), np.full(n, np.nan, dtype=dtype)
+ pairs = [(a_np, b_np), (zero, a_np), (a_np, zero), (nan, a_np), (a_np,
nan)]
+ rows = [
+ np.where(cmp(x, y) | np.isnan(x), x, y)
+ for cmp in (np.less, np.greater)
+ for x, y in pairs
+ ]
+ rows += [rows[5] + 1, rows[6] + 1, np.full(n, 3, dtype=dtype), a_np,
a_np]
+ else:
+ a_np = np.array([3, -5, 7, 0, -1, 9, 2, -8, 4, 4, -2, 0], dtype=dtype)
+ b_np = np.array([1, 9, -2, 4, -3, 9, 5, -8, 0, -4, -2, 1], dtype=dtype)
+ rows = [np.minimum(a_np, b_np), np.maximum(a_np, b_np)]
+ expected = np.stack(rows)
+ a = tvm_tirx.runtime.tensor(a_np, tvm_tirx.cuda(0))
+ b = tvm_tirx.runtime.tensor(b_np, tvm_tirx.cuda(0))
+ c = tvm_tirx.runtime.empty(expected.shape, dtype, tvm_tirx.cuda(0))
+
+ def run_and_check():
+ mod(a, b, c)
+ got = c.numpy()
+ if not is_float:
+ np.testing.assert_array_equal(got, expected)
+ return
+ # NaN lanes are checked with isnan (CUDART_NAN_F and a min/max
instruction
+ # normalize the payload); every finite lane, including the sign of
zero, bitwise.
+ nan_mask = np.isnan(expected)
+ np.testing.assert_array_equal(np.isnan(got), nan_mask)
+ bits = {2: "uint16", 4: "uint32", 8: "uint64"}[expected.dtype.itemsize]
+ np.testing.assert_array_equal(got.view(bits)[~nan_mask],
expected.view(bits)[~nan_mask])
+
+ tvm_tirx.testing.run_with_gpu_lock(run_and_check)
+
+
@pytest.mark.gpu
@pytest.mark.skipif(not env.has_cuda(), reason="need cuda")
def test_round_ties_to_even(*, backend_config):
