compiler: Avoid int32 overflow in linearized host-device transfer size by gaoflow · Pull Request #2939 · devitocodes/devito · GitHub

Description

When a host↔device data transfer is linearized, its array section size is emitted as a product of the Function's per-dimension sizes, for example:

#pragma acc enter data copyin(u[0:u_vec->size[0]*u_vec->size[1]*u_vec->size[2]*u_vec->size[3]])

The u_vec->size[i] fields are 32-bit C ints, so the product size[0]*size[1]*size[2]*size[3] is evaluated in 32-bit arithmetic. For a Function with more than ~2**31 elements (e.g. the reporter's 1295**3 ≈ 2.17e9 points, ~24.5 GB) the product overflows int before it is used as the transfer bound, producing a bogus size (the reporter saw 18446744065653020036) and a corrupt / failed device transfer — independent of index-mode=int64/linearize=True, because those control the kernel index type, not the type of the transfer-clause arithmetic.

As @mloubout noted on the issue, the fix is to perform the size multiplication in 64-bit. This casts each factor of a product section bound to a 64-bit integer:

#pragma acc enter data copyin(u[0:(long)(u_vec->size[0])*(long)(u_vec->size[1])*(long)(u_vec->size[2])*(long)(u_vec->size[3])])

Casting the whole product ((long)(a*b*c)) would be too late — the overflow would already have happened in 32-bit — so each factor is cast individually, which forces every multiplication to be 64-bit regardless of operand ordering.

The change lives in PragmaTransfer._generate, so it is scoped to host-device transfer clauses only. Non-product bounds (a single dimension size, an offset, a constant) cannot overflow and are left untouched, addressing the concern that there is "no reason to use long for all of those". Non-transfer expressions (e.g. free-space guards, TMA descriptors) are unaffected.

Reproduction

from devito import Eq, Grid, Operator, TimeFunction

grid = Grid(shape=(4, 5, 6))
u = TimeFunction(name='u', grid=grid)
op = Operator(Eq(u.forward, u + 1), platform='nvidiaX', language='openacc',
              opt=('advanced', {'linearize': True}))
print(op.body.maps[0].ccode.value)

Before:

acc enter data copyin(u[0:u_vec->size[0]*u_vec->size[1]*u_vec->size[2]*u_vec->size[3]])

After:

acc enter data copyin(u[0:(long)(u_vec->size[0])*(long)(u_vec->size[1])*(long)(u_vec->size[2])*(long)(u_vec->size[3])])

Verification

The fix applies to both backends (openacc copyin/copyout/delete and openmp map(to:/release:)) and to 2D/3D Functions.
The non-linearized transfer path (separate per-dimension sections [0:s0][0:s1]...) is unchanged — there is no product there, hence no overflow.
Added TestPassesOptional::test_linearize_transfer_no_overflow asserting each size[i] factor of a linearized transfer is cast to long and that no bare 32-bit product remains.
Updated the existing test_gpu_openmp.py expectations (test_basic, test_multiple_eqns) whose OpenMP transfers use the flattened product form.
Host (CPU) operators, including linearize=True, build and run unchanged (no device transfers emitted). flake8 clean on the changed files.

Note: the GPU test modules are skipif(['nodevice']), so the codegen assertions run on the GPU CI runners. They were validated locally by forcing platform='nvidiaX'.

Good point — dropped the cast helper and the is_Mul check in 454434b, applying as_long to the section extent directly:

sections = ''.join([f'[{ccode(i)}:{ccode(as_long(j))}]'
                    for i, j in self.sections])

The generated code is unchanged: the start bound i is always 0/an offset (left as-is, can't overflow) and the extent j is the size product that as_long promotes to 64-bit, e.g. (long)(size[3])*(long)(size[2])*.... I left i uncast rather than wrapping both bounds, since the overflow site (#2777) is the extent product.

Gentle ping — the review feedback was addressed back in 454434b (dropped the cast helper and the is_Mul check, switched to as_list). The branch is behind main but conflict-free. Happy to do anything further; otherwise this should be ready for another look when you have a moment.

Codecov Report

❌ Patch coverage is 43.75000% with 9 lines in your changes missing coverage. Please review.
✅ Project coverage is 79.20%. Comparing base (f335bb1) to head (b83821a).

Files with missing lines	Patch %	Lines
tests/test_gpu_common.py	0.00%	9 Missing ⚠️

Additional details and impacted files

@@            Coverage Diff             @@
##             main    #2939      +/-   ##
==========================================
- Coverage   83.57%   79.20%   -4.38%     
==========================================
  Files         257      257              
  Lines       53841    53852      +11     
  Branches     4609     4611       +2     
==========================================
- Hits        44999    42654    -2345     
- Misses       8044    10348    +2304     
- Partials      798      850      +52

Flag	Coverage Δ
pytest-gpu-aomp-amdgpuX	?
pytest-gpu-gcc-	?
pytest-gpu-icx-	?
pytest-gpu-nvc-nvidiaX	?

Flags with carried forward coverage won't be shown. Click here to find out more.

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

Just to keep this moving, this PR will need a rebase and failing tests fixing before it can be merged

Rebased onto current main (flattened the old merge commit into the three commits) and fixed the isort failure in devito/symbolics/manipulation.py that was breaking the lint job. CI should be running fresh now.

Check out this pull request on

See visual diffs & provide feedback on Jupyter Notebooks.

Powered by ReviewNB

Rebased onto current main and fixed the real CI failures: applying as_long unconditionally also rewrote the plain per-dimension extents (they are IndexedPointer leaves too), which churned every non-linearized transfer pragma — that's what broke the openacc codegen tests and the 01_gpu notebook. Restored the product-only guard at the call site (only a product of the 32-bit sizes can overflow) and updated the notebook's reference output for the linearized clauses that legitimately gain the casts. The arm64/MPI failures were infra — apt mirror timeouts while building the docker images.

Heads-up: the CI runs for the rebase are sitting at "awaiting approval" (fork workflow gating), so nothing has actually run on the new head. Could someone approve the runs when convenient?

When a host-device data transfer is linearized, its array section size is emitted as a product of the Function's per-dimension sizes, e.g. `copyin(u[0:u_vec->size[0]*u_vec->size[1]*u_vec->size[2]*u_vec->size[3]])`. The `size[i]` fields are 32-bit C ints, so for a Function with more than ~2**31 elements the product overflows `int` before it is used as the transfer bound, yielding a bogus size and a corrupt/failed device transfer. Cast each factor of the product to a 64-bit integer so the multiplication is carried out in 64-bit arithmetic. Casting the whole product would be too late (the overflow would already have occurred), so each factor is cast individually. Non-product bounds (a single size, an offset, a constant) cannot overflow and are left untouched, as are non-transfer expressions. Fixes devitocodes#2777

Address review: replace the ad-hoc _avoid_overflow helper with the existing as_long. as_long only substituted plain Symbols (retrieve_symbols), so it was a no-op on the IndexedPointer size factors (vec->size[i]) of a linearized transfer bound; extend it to retrieve_terminals so Indexed/IndexedPointer leaves are cast too. Keep the cast scoped to Mul products in PragmaTransfer so non-linearized multi-dimensional sections are not needlessly upcast.

…nt directly Per review: as_long already walks the expression args, so the cast() helper and its is_Mul check are unnecessary. Apply as_long to the section extent directly. Output is unchanged: the start bound is always 0/an offset (left as-is) and the extent is the size product that as_long promotes to 64-bit.

…tebook Applying as_long unconditionally also rewrote plain per-dimension extents (IndexedPointer leaves), churning every non-linearized transfer pragma and breaking the openacc codegen tests and the 01_gpu notebook. Only a product of 32-bit sizes can overflow int, so restrict the promotion to Mul bounds and update the notebook's reference output for the linearized transfer clauses that legitimately gained the casts.

Rebased onto current main and force-pushed (b83821a). The transfer-size change is unchanged; CI is re-running. test_w_data also passes locally on the rebased tree.

The fresh runs after the rebase are sitting at "awaiting approval" (fork workflow gate) — could someone approve them when convenient? Thanks.

FabioLuporini requested changes May 29, 2026

View reviewed changes

Comment thread devito/passes/iet/parpragma.py Outdated

mloubout added the compiler label May 29, 2026

mloubout reviewed Jun 1, 2026

View reviewed changes

Comment thread devito/passes/iet/parpragma.py

mloubout approved these changes Jul 6, 2026

View reviewed changes

gaoflow force-pushed the fix-2777-transfer-size-overflow branch from 43311ac to 3e83986 Compare July 15, 2026 07:56

gaoflow force-pushed the fix-2777-transfer-size-overflow branch from 3e83986 to 01bd794 Compare July 16, 2026 07:14

gaoflow added 4 commits July 31, 2026 21:36

gaoflow force-pushed the fix-2777-transfer-size-overflow branch from 01bd794 to b83821a Compare July 31, 2026 19:36

mloubout merged commit de79615 into devitocodes:main Aug 1, 2026
38 of 39 checks passed

compiler: Avoid int32 overflow in linearized host-device transfer size by gaoflow · Pull Request #2939 · devitocodes/devito · GitHub

Uh oh!

compiler: Avoid int32 overflow in linearized host-device transfer size - #2939

compiler: Avoid int32 overflow in linearized host-device transfer size#2939
mloubout merged 4 commits into
devitocodes:mainfrom
gaoflow:fix-2777-transfer-size-overflow

gaoflow commented May 29, 2026

Uh oh!

Uh oh!

Uh oh!

gaoflow commented Jun 1, 2026

Uh oh!

gaoflow commented Jun 29, 2026

Uh oh!

codecov Bot commented Jul 6, 2026 •
edited
Loading

Uh oh!

Uh oh!

EdCaunt commented Jul 14, 2026

Uh oh!

gaoflow commented Jul 15, 2026

Uh oh!

review-notebook-app Bot commented Jul 16, 2026

Uh oh!

gaoflow commented Jul 16, 2026

Uh oh!

gaoflow commented Jul 17, 2026

Uh oh!

gaoflow commented Jul 31, 2026

Uh oh!

gaoflow commented Jul 31, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

compiler: Avoid int32 overflow in linearized host-device transfer size by gaoflow · Pull Request #2939 · devitocodes/devito · GitHub

Uh oh!

compiler: Avoid int32 overflow in linearized host-device transfer size - #2939

compiler: Avoid int32 overflow in linearized host-device transfer size#2939mloubout merged 4 commits intodevitocodes:maindevitocodes/devito:mainfrom gaoflow:fix-2777-transfer-size-overflowgaoflow/devito:fix-2777-transfer-size-overflowCopy head branch name to clipboard

Conversation

gaoflow commented May 29, 2026

Description

Reproduction

Verification

Uh oh!

Uh oh!

Uh oh!

gaoflow commented Jun 1, 2026

Uh oh!

gaoflow commented Jun 29, 2026

Uh oh!

codecov Bot commented Jul 6, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

EdCaunt commented Jul 14, 2026

Uh oh!

gaoflow commented Jul 15, 2026

Uh oh!

review-notebook-app Bot commented Jul 16, 2026

Uh oh!

gaoflow commented Jul 16, 2026

Uh oh!

gaoflow commented Jul 17, 2026

Uh oh!

gaoflow commented Jul 31, 2026

Uh oh!

gaoflow commented Jul 31, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

compiler: Avoid int32 overflow in linearized host-device transfer size#2939
mloubout merged 4 commits into
devitocodes:mainfrom
gaoflow:fix-2777-transfer-size-overflow

codecov Bot commented Jul 6, 2026 •
edited
Loading