Refactor Memory.h and simply the API#3742
Open
guangyey wants to merge 1 commit into
Open
Conversation
github-actions
Bot
added
disable_e2e
guangyey
force-pushed
the
guangyey/memcpy
branch
4 times, most recently
from
1b15397 to
529a550
Compare
Contributor
There was a problem hiding this comment.
Pull request overview
Skill file(s) read: .github/skills/xpu-ops-pr-review/SKILL.md.
This PR refactors XPU SYCL memory copy/memset helpers into src/comm/Memory.h, introducing a small set of wrappers intended to standardize synchronization and pinned-memory event recording, and migrates multiple callsites to use the new APIs.
Changes:
- Replaced scattered
queue.memcpy()/memset()usages (plus ad-hocrecord_event) with centralized wrappers (memcpyAndSync,memcpyAsync,memcpyHostToDeviceAsync,memcpyPinnedHostToDeviceAsync,memcpyDeviceToHostAsync,memsetAndSync,memsetAsync). - Migrated key callsites (copy path, foreach/multi-tensor metadata uploads, scalar extraction, sparse CSR add, resize) to use the wrappers.
Must-fix (blocking):
src/ATen/native/xpu/Copy.cppdevice-to-device “memcpy-eligible” path now enqueues the copy viaxpu::sycl::memcpyAsync, which uses the current stream internally rather than the providedcopy_stream. This can break the cross-device barrier/synchronization logic and cause races.
Reviewed changes
Copilot reviewed 7 out of 7 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
| src/comm/Memory.h | Adds centralized SYCL memcpy/memset wrappers with pinned-host staging and event recording. |
| src/ATen/native/xpu/XPUScalar.cpp | Uses memcpyAndSync for scalar extraction instead of raw queue memcpy+wait. |
| src/ATen/native/xpu/sycl/ResizeKernel.cpp | Uses centralized async memcpy for storage resize copy. |
| src/ATen/native/xpu/sycl/MultiTensorApply.h | Uses pinned-H2D wrapper for metadata uploads and encapsulated event recording. |
| src/ATen/native/xpu/sycl/ForeachReduceKernels.cpp | Uses pinned-H2D wrapper for metadata/count uploads and encapsulated event recording. |
| src/ATen/native/xpu/Copy.cpp | Refactors memcpy paths to use centralized wrappers (introduces stream/queue mismatch bug). |
| src/ATen/native/sparse/xpu/sycl/SparseCsrTensorAddKernels.cpp | Uses centralized memset/memcpy wrappers for async set and sync readback of nnz. |
| auto src = iter.data_ptr(1); | ||
| size_t size = iter.numel() * iter.element_size(0); | ||
| q.copy(src, dst, size); | ||
| xpu::sycl::memcpyAsync(dst, src, size); |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Centralize XPU memory copy/set operations into
comm/Memory.h. Previously, all APIs defined inMemory.hare unused. I refactor it and provide some APIs could be reused in code.Raw
queue.memcpy()+record_event()pairs were scattered across the codebase, making it easy to miss event recording and leak pinned memory lifetime bugs. This introduces a small set of typed wrappers that encapsulate the correct synchronization and event-recording protocol, then migrates all callsites to use them.Benefits: correctness by construction (can't forget record_event), single point of maintenance.