# [Incorrect result / possible VMFault] Hygon FlagGems `aten::scatter_add_.default` ignores stride-zero expanded index

[wzx226226]

Describe the bug

Summary

On the Hygon DCU mixed FlagGems route, aten::scatter_add_.default produces
incorrect results when index is a valid view(...).expand(...) tensor with a
zero broadcast stride. DAS boxing returns the correct PyTorch result for the
identical input on the same DCU card.

expand does not materialize the broadcast dimension. The logical index has
shape (340, 12) and stride (1, 0). The FlagGems 2D fast path treats it as a
contiguous (12, 1)-strided tensor instead.

This issue uses controlled extra storage so that the functional error is
deterministic and does not depend on whether an out-of-bounds access happens to
hit a protected device page. In a real DiffDock run using the normal, unpadded
expand storage, the same bug was also observed as:

>>>>>>>> KERNEL VMFault !!!! <<<<<<
HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION
SIGABRT / core dump

Environment

• Hardware: Hygon DCU / BW1000, device flagos:0
• Torch-FL backend: backends_dcu_flaggems.conf
• FlagGems vendor: hygon
• PyTorch: 2.10.0

Minimal reproducer

Save the following complete program as repro_scatter_add_expanded_index.py.
The route is selected before import torch_fl, so each invocation starts a
fresh process with the intended backend configuration.

import argparse
import os

parser = argparse.ArgumentParser()
parser.add_argument("--route", choices=("boxing", "flaggems"), required=True)
args = parser.parse_args()
if args.route == "flaggems":
    os.environ["FLAGOS_USE_FLAGGEMS"] = "1"
else:
    os.environ.pop("FLAGOS_USE_FLAGGEMS", None)

import torch
import torch_fl

rows, columns, output_rows = 340, 12, 10
device = "flagos:0"
output = torch.zeros((output_rows, columns), dtype=torch.float32, device=device)
source = torch.ones((rows, columns), dtype=torch.float32, device=device)

# Preserve the logical expanded layout, stride=(1, 0). The controlled tail
# makes a stride-ignoring implementation fail deterministically without a
# VMFault being required for reproduction.
storage = torch.zeros(rows * columns, dtype=torch.int64, device=device)
storage[:rows] = torch.arange(rows, dtype=torch.int64, device=device) % output_rows
index = storage[:rows].view(rows, 1).expand(rows, columns)
assert index.stride() == (1, 0)

actual = output.scatter_add_(0, index, source)
torch.flagos.synchronize()
reference = torch.zeros((output_rows, columns), dtype=torch.float32)
reference.scatter_add_(0, index.cpu(), source.cpu())
actual_cpu = actual.cpu()
different = actual_cpu != reference
print("route:", args.route)
print("actual[0, :]   :", actual_cpu[0].tolist())
print("reference[0, :]:", reference[0].tolist())
print("mismatched elements:", different.sum().item(), "/", different.numel())
if different.any():
    row, column = different.nonzero()[0].tolist()
    print("first mismatch:", (row, column), actual_cpu[row, column].item(), reference[row, column].item())
torch.testing.assert_close(actual_cpu, reference)

Run the same program twice:

export FLAGOS_LOG_DISPATCH=1
python repro_scatter_add_expanded_index.py --route boxing
python repro_scatter_add_expanded_index.py --route flaggems

DAS boxing versus mixed FlagGems

Both commands used the same node, flagos:0, tensor shapes, dtypes, values,
and program. The only difference was the selected backend before importing
Torch-FL.

Route	Dispatch log	Comparison with CPU reference
DAS boxing	scatter_add_ -> cuda	0 / 120 mismatched elements; assertion passes
mixed FlagGems	scatter_add_ -> flagos_python	120 / 120 mismatched elements; assertion fails

DAS boxing output

[flagos dispatch] scatter_add_ -> cuda
actual[0, :]   : [34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0]
reference[0, :]: [34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0]
mismatched elements: 0 / 120

Mixed FlagGems output

[flagos] loading backend config from backends_dcu_flaggems.conf
[flagos dispatch] scatter_add_ -> flagos_python
actual[0, :]   : [317.0, 311.0, 316.0, 311.0, 317.0, 312.0, 318.0, 312.0, 318.0, 312.0, 318.0, 312.0]
reference[0, :]: [34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0, 34.0]
mismatched elements: 120 / 120
first mismatch: (0, 0) 317.0 34.0
AssertionError: Tensor-likes are not close!
Greatest absolute difference: 284.0 at index (0, 6)

Root cause

File: src/flag_gems/ops/scatter_add.py, scatter_add_2d_kernel.

The current kernel computes the index address as if index were contiguous:

row = offsets // idx_ncols
col = offsets % idx_ncols
idx_offsets = row * idx_ncols + col
idx = tl.load(index_ptr + idx_offsets, mask=mask, other=0)

For this reproducer, the correct address is:

index_ptr + row * 1 + col * 0

The current calculation instead reads row * 12 + col. Its final logical
element accesses offset 339 * 12 + 11 = 4079, although an ordinary expanded
index has only 340 physical int64 elements. The value read is then used as a
destination row for tl.atomic_add, which explains both the wrong output and
the possible device-memory fault in the unpadded case.

Expected behavior and suggested fix

scatter_add_ must accept valid non-contiguous indices, including zero-stride
views created by expand, and match DAS boxing/PyTorch.

Please either pass index_stride0 and index_stride1 into the 2D kernel and
address the index as row * index_stride0 + col * index_stride1, or restrict
this fast path to contiguous indices and dispatch other layouts to a
stride-aware implementation. This reproducer should be added as a regression
test.