compiler: Avoid int32 overflow in linearized host-device transfer size#2939
compiler: Avoid int32 overflow in linearized host-device transfer size#2939gaoflow wants to merge 4 commits into
Conversation
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.
|
Good point — dropped the sections = ''.join([f'[{ccode(i)}:{ccode(as_long(j))}]'
for i, j in self.sections])The generated code is unchanged: the start bound |
|
Gentle ping — the review feedback was addressed back in 454434b (dropped the |
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## main #2939 +/- ##
==========================================
- Coverage 82.85% 79.23% -3.63%
==========================================
Files 249 249
Lines 52279 52288 +9
Branches 4503 4505 +2
==========================================
- Hits 43318 41431 -1887
- Misses 8188 10042 +1854
- Partials 773 815 +42
Flags with carried forward coverage won't be shown. Click here to find out more. ☔ View full report in Codecov by Harness. 🚀 New features to boost your workflow:
|
Description
Fixes #2777.
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 Cints, so the productsize[0]*size[1]*size[2]*size[3]is evaluated in 32-bit arithmetic. For aFunctionwith more than~2**31elements (e.g. the reporter's1295**3 ≈ 2.17e9points, ~24.5 GB) the product overflowsintbefore it is used as the transfer bound, producing a bogus size (the reporter saw18446744065653020036) and a corrupt / failed device transfer — independent ofindex-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 uselongfor all of those". Non-transfer expressions (e.g. free-space guards, TMA descriptors) are unaffected.Reproduction
Before:
After:
Verification
openacccopyin/copyout/deleteandopenmpmap(to:/release:)) and to 2D/3D Functions.[0:s0][0:s1]...) is unchanged — there is no product there, hence no overflow.TestPassesOptional::test_linearize_transfer_no_overflowasserting eachsize[i]factor of a linearized transfer is cast tolongand that no bare 32-bit product remains.test_gpu_openmp.pyexpectations (test_basic,test_multiple_eqns) whose OpenMP transfers use the flattened product form.linearize=True, build and run unchanged (no device transfers emitted).flake8clean on the changed files.