[Bug][Hygon/DCU] remainder returns NaN for finite dividends with infinite divisors #7755

Describe the bug

Summary

aten.remainder.Tensor on the Hygon FlagGems backend returns NaN for cases such as remainder(2,+Inf), remainder(0,+Inf), and remainder(-2,-Inf), where CPU and DAS return the finite dividend. This changes exceptional-value masks and can propagate NaNs downstream.

Environment and verified scope

• Platform: Hygon DCU, Torch-FL PrivateUse1 (flagos), eager execution; one accelerator.
• Device properties report BW, gfx936:sramecc+:xnack-, 80 CUs, warp size 64, 68702699520 bytes memory. This is a device-API report, not a separate PCI-ID audit.
• OS: Linux 5.10.134-17.1.3.sga8.x86_64, glibc 2.28.
• Python: 3.10.20, GCC 14.3.0.
• PyTorch frontend: 2.10.0+cpu, source 449b1768410104d3ed79d3bcfe4ba1d65c7f22c0; effective HIP metadata 6.3.26113. The frontend suffix does not mean the tested tensors executed on CPU.
• Torch-FL: 2.10.0+das.dtk2604.glibc228.g6d84d12b, source 6d84d12b715c307e5b660fc33e0f51e660602f6a.
• FlagGems: 5.4.0rc2.post1+ge7b4a865f, source e7b4a865fce6d85861ee91a6aca56564ef9acf7d.
• FlagTree: 0.6.0+hcu3.6.glibc228.compat1.git8a1dda94; exposed Triton version 3.6.0.
• Loaded hycu driver-module version: 6.3.31-V1.5.0a. A native HIP shared-library path/PCI query was not captured; no stronger driver identity is inferred.
• Tested release: r5-20260928-tfl6d84d12b-fge7b4a865. Prebuilt deployed environment; build commands were not run or recovered in this reproduction. No package install or shared source/configuration modification.

The report was prompted by a 2026-10-21 test export, but is supported by new deterministic three-way controls, not by the CSV verdict alone. On 2026-10-26, a full supplemental diagnostic ran 53 cases per route; a separate reduced-reproducer job then executed the scripts below in three fresh processes (CPU, DAS, FlagGems). The reduced job completed successfully in 133 seconds; success means the harness matched the expected good controls and expected FlagGems failure, not that FlagGems passed correctness. Actual ATen calls and Python implementation entries were witnessed, source/configuration hashes were retained and loaded source hashes stayed unchanged. No performance or fix claim is made.

A preceding reduced job failed because its witness filter omitted the Hygon backend override; that failed result was retained, the diagnostic filter alone was corrected in a separate revision, and the independent successful run is used here. These are synthetic public inputs, not a replay of the unavailable original Lilith harness/inputs. Current FlagGems mainline was inspected at 415ea38c63141382b0051eac5ef0fa22cbbbec69; current mainline and proposed fixes were not hardware-tested.

Reproduction

Save the code as repro.py. Run each route in a fresh process in the pinned environment:

ROUTE=cpu python repro.py
ROUTE=das python repro.py
ROUTE=flaggems python repro.py

The script sets the relevant FLAGOS_OP_* overrides before importing Torch-FL (cuda is the DAS-boxing route selector here; it is not an NVIDIA device claim). No private file/configuration is needed by the reproducer. The observed route was independently checked with ATen tracing and actual FlagGems entry calls; CPU and DAS controls had no FlagGems operator-entry calls. Targeted overrides were used rather than changing the installed default configuration.

OPS = ['remainder__Tensor']
# Standalone route-controlled correctness reproduction.
import os
route = os.environ.get("ROUTE", "flaggems")
assert route in ("cpu", "das", "flaggems")
if route != "cpu":
    for op in OPS:
        os.environ["FLAGOS_OP_" + op] = "cuda" if route == "das" else "flaggems"
    import torch_fl
import torch
if route != "cpu":
    torch.cuda.set_device(0)
device = "cpu" if route == "cpu" else "flagos:0"
base = torch.tensor([-float("inf"), -2., -0., 0., 2., float("inf"), float("nan"), 1e-45, -1e-45, 1e30, -1e30])
base = base.repeat(1490)[:16384].reshape(128, 128)
den = torch.tensor([-2., 2., -float("inf"), float("inf"), 0., float("nan"), 0.5]).repeat(2341)[:16384].reshape(128, 128)
ref = torch.remainder(base, den)
out = torch.remainder(base.to(device), den.to(device)).cpu()
for i in (2, 3, 23, 59):
    print(route, i, base.flatten()[i].item(), den.flatten()[i].item(), ref.flatten()[i].item(), out.flatten()[i].item())
torch.testing.assert_close(out, ref, rtol=0, atol=0, equal_nan=True)

Expected vs actual

CPU/DAS return 2 for remainder(2,+Inf), +0 for remainder(+0,+Inf), -2 for remainder(-2,-Inf), and -0 for remainder(-0,-Inf). FlagGems returns NaN in these cases. The full deterministic special-value matrix has 3406 mismatches in the FP32 supplemental test; the public reduced run also reproduces the exceptional-value failures.

Actual route and diagnostics

Target FlagGems entry witness(es): {"flag_gems/runtime/backend/_hygon/ops/div.py:remainder": 1}. Target ATen overloads are present in the captured dispatcher trace. Route controls and implementation-entry witnesses were checked independently; a configuration setting alone was not treated as proof.

Observed stdout:

cpu 2 -0.0 -inf -0.0 -0.0
cpu 3 0.0 inf 0.0 0.0
cpu 23 -2.0 -inf -2.0 -2.0
cpu 59 2.0 inf 2.0 2.0
das 2 -0.0 -inf -0.0 -0.0
das 3 0.0 inf 0.0 0.0
das 23 -2.0 -inf -2.0 -2.0
das 59 2.0 inf 2.0 2.0
flaggems 2 -0.0 -inf -0.0 nan
flaggems 3 0.0 inf 0.0 nan
flaggems 23 -2.0 -inf -2.0 nan
flaggems 59 2.0 inf 2.0 nan

Full caught failure traceback from the instrumented reduced run (only absolute environment/repro roots redacted; diagnostic frames retained):

Traceback (most recent call last):
  File "<repro>/repro_runner.py", line 65, in <module>
    with contextlib.redirect_stdout(output),torch.inference_mode(),Trace():runpy.run_path(str(R/'repros'/name),run_name='__main__')
  File "<python>/runpy.py", line 289, in run_path
    return _run_module_code(code, init_globals, run_name,
  File "<python>/runpy.py", line 96, in _run_module_code
    _run_code(code, mod_globals, init_globals,
  File "<python>/runpy.py", line 86, in _run_code
    exec(code, run_globals)
  File "<repro>/repros/03-remainder-inf.py", line 21, in <module>
    torch.testing.assert_close(out, ref, rtol=0, atol=0, equal_nan=True)
  File "<python>/site-packages/torch/testing/_comparison.py", line 1600, in assert_close
    raise error_metas[0].to_error(msg)
AssertionError: Tensor-likes are not equal!

Mismatched elements: 3406 / 16384 (20.8%)
Greatest absolute difference: nan at index (0, 2)
Greatest relative difference: nan at index (0, 2)

Loaded source SHA256 (stable before/after run):

{
  "flag_gems/ops/remainder.py": "d629f7c1c7880966c3c0ce2be43dd92779826e59a1fffac1ec92d8fccc13968d",
  "flag_gems/runtime/backend/_hygon/ops/div.py": "708457acfae99a9e2922afae86bb75bad3232e62c28510278f8399229d440343",
  "triton/language/core.py": "b8c33dc0501719f2bf88c15242e3e5157faf2256fd887fa160489c1d5984f167"
}

Execution chain and root-cause boundary

CPU-created deterministic data (or per-backend seeded RNG domain tests) -> requested dtype/device -> eager torch API / exact ATen overload -> Torch-FL flagos routing -> witnessed FlagGems implementation -> Triton/FlagTree backend kernel (except round, which rejects before launch) -> CPU result comparison. DAS controls use the corresponding vendor boxing path. No torch.compile/fusion path or performance measurement was involved.

The executed entry was flag_gems/runtime/backend/_hygon/ops/div.py:remainder, not the generic remainder.py entry. It dispatches tensor/tensor inputs to rem_tt, then _remainder: r = x % y; c1 = r != 0; c2 = (x < 0) ^ (y < 0); return tl.where(c1 & c2, r + y, r). This path has no explicit treatment for finite dividend/infinite divisor. The tested same-sign examples already fail without needing the sign-correction branch to explain them, so attributing everything to r+y would be unsupported. The evidence localizes the operator/backend arithmetic path; the precise remainder lowering still needs kernel/IR inspection before blaming a compiler or runtime. Primary ownership is the FlagGems Hygon override, with FlagTree involvement only if isolated lowering proves necessary.

Related code locations

• FlagGems/src/flag_gems/runtime/backend/_hygon/ops/div.py

  Lines 340 to 376 in e7b4a86

  	@triton.jit
  	def _remainder(x, y):
  	r = x % y
  	c1 = r != 0
  	c2 = (x < 0) ^ (y < 0)
  	return tl.where(c1 & c2, r + y, r)
  	@pointwise_dynamic(promotion_methods=[(0, 1, "DEFAULT")])
  	@triton.jit
  	def rem_tt(x, y):
  	return _remainder(x, y)
  	@pointwise_dynamic(is_tensor=[True, False], promotion_methods=[(0, 1, "DEFAULT")])
  	@triton.jit
  	def rem_ts(x, y):
  	return _remainder(x, y)
  	@pointwise_dynamic(is_tensor=[False, True], promotion_methods=[(0, 1, "DEFAULT")])
  	@triton.jit
  	def rem_st(x, y):
  	return _remainder(x, y)
  	def remainder(A, B):
  	logger.debug("GEMS_HYGON FLOOR_DIVIDE")
  	if isinstance(A, torch.Tensor) and isinstance(B, torch.Tensor):
  	return rem_tt(A, B)
  	elif isinstance(A, torch.Tensor):
  	return rem_ts(A, B)
  	elif isinstance(B, torch.Tensor):
  	return rem_st(A, B)
  	else:
  	# Both scalar
  	return torch.tensor(A % B)

• FlagGems/src/flag_gems/ops/remainder.py

  Lines 1 to 62 in e7b4a86

  	# Copyright 2026 FlagOS Contributors
  	#
  	# Licensed under the Apache License, Version 2.0 (the "License");
  	# you may not use this file except in compliance with the License.
  	# You may obtain a copy of the License at
  	#
  	# http://www.apache.org/licenses/LICENSE-2.0
  	#
  	# Unless required by applicable law or agreed to in writing, software
  	# distributed under the License is distributed on an "AS IS" BASIS,
  	# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  	# See the License for the specific language governing permissions and
  	# limitations under the License.
  	# Generated by KernelGen: https://github.com/flagos-ai/KernelGen
  	import logging
  	import torch
  	import triton
  	import triton.language as tl
  	from flag_gems.utils import pointwise_dynamic
  	logger = logging.getLogger(__name__)
  	@triton.jit
  	def _remainder(x, y):
  	r = x % y
  	c1 = r != 0
  	c2 = (x < 0) ^ (y < 0)
  	return tl.where(c1 & c2, r + y, r)
  	@pointwise_dynamic(promotion_methods=[(0, 1, "DEFAULT")])
  	@triton.jit
  	def rem_tt(x, y):
  	return _remainder(x, y)
  	@pointwise_dynamic(is_tensor=[True, False], promotion_methods=[(0, 1, "DEFAULT")])
  	@triton.jit
  	def rem_ts(x, y):
  	return _remainder(x, y)
  	@pointwise_dynamic(is_tensor=[False, True], promotion_methods=[(0, 1, "DEFAULT")])
  	@triton.jit
  	def rem_st(x, y):
  	return _remainder(x, y)
  	def remainder(A, B):
  	logger.debug("GEMS REMAINDER")
  	if isinstance(A, torch.Tensor) and isinstance(B, torch.Tensor):
  	return rem_tt(A, B)
  	elif isinstance(A, torch.Tensor):
  	return rem_ts(A, B)
  	elif isinstance(B, torch.Tensor):
  	return rem_st(A, B)
  	else:
  	# Both scalar

Suggested fix and acceptance

Handle finite/Inf and signed-zero semantics explicitly in the Hygon override (and audit the corresponding generic arithmetic pattern), matching PyTorch remainder rather than only ordinary finite modulo. Isolate and inspect the % lowering before assigning a compiler defect. Do not use nan_to_num as a semantic fix.

Exercise all sign combinations for finite dividend and +/-Inf divisor, signed zero, NaN/Inf inputs, zero divisor and normal finite cases; compare masks, finite values and signed zero separately against CPU/DAS. Extend scalar/tensor, broadcast, dtype and layout variants after the tensor/tensor reproducer.

Duplicate check and limitations

Open and closed issues/PRs were searched again before filing on 2026-10-26. No same-root issue was found in the inspected results. This is scoped to the pinned deployed R5 pair; static inspection of current source is not latest-version hardware validation. The original CSV exporter and other reduction-error findings are not attributed to this issue. No fix was applied or validated.

Prepared by a Codex coding agent; exact model identifier was not exposed/recorded. Submitted with the user's authorization.