Skip to content

[SYCL] Disable HostPtr reuse when the pointer is read-only #10334

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Jul 13, 2023

Conversation

0x12CC
Copy link
Contributor

@0x12CC 0x12CC commented Jul 12, 2023

Mutable SYCL buffers can be initialized using a const T* hostData. This change ensures that these buffers allocate new memory so that their contents can be modified without changing the original host data. Fixes #10091.

Mutable SYCL buffers can be initialized using a `const T* hostData`.
This change ensures that these buffers allocate new memory so that
their contents can be modified without changing the original host data.
Fixes intel#10091.

Signed-off-by: Michael Aziz <[email protected]>
@0x12CC 0x12CC requested a review from a team as a code owner July 12, 2023 20:33
@0x12CC 0x12CC requested review from againull and steffenlarsen July 12, 2023 20:33
@0x12CC 0x12CC temporarily deployed to aws July 12, 2023 20:47 — with GitHub Actions Inactive
@0x12CC 0x12CC temporarily deployed to aws July 12, 2023 21:39 — with GitHub Actions Inactive
Copy link
Contributor

@steffenlarsen steffenlarsen left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Great find! 🚀

@againull againull merged commit d058186 into intel:sycl Jul 13, 2023
@bader
Copy link
Contributor

bader commented Jul 13, 2023

@0x12CC, @steffenlarsen, @againull, in my pre-commit the test fails on CUDA and Intel:

FAIL: SYCL :: Basic/host_write_back.cpp ([19](https://github.com/intel/llvm/actions/runs/5546033087/jobs/10129281860?pr=10363#step:16:20)0 of 1436)
******************** TEST 'SYCL :: Basic/host_write_back.cpp' FAILED ********************
Script:
--
: 'RUN: at line 1';    /__w/llvm/llvm/toolchain/bin//clang++   -fsycl -fsycl-targets=nvptx64-nvidia-cuda /__w/llvm/llvm/llvm/sycl/test-e2e/Basic/host_write_back.cpp -o /__w/llvm/llvm/build-e2e/Basic/Output/host_write_back.cpp.tmp.out
: 'RUN: at line 2';   env SYCL_PI_CUDA_ENABLE_IMAGE_SUPPORT=1 ONEAPI_DEVICE_SELECTOR=ext_oneapi_cuda:gpu  /__w/llvm/llvm/build-e2e/Basic/Output/host_write_back.cpp.tmp.out
--
Exit Code: -6

Command Output (stdout):
--
$ ":" "RUN: at line 1"
note: command had no output on stdout or stderr
$ "/__w/llvm/llvm/toolchain/bin//clang++" "-fsycl" "-fsycl-targets=nvptx64-nvidia-cuda" "/__w/llvm/llvm/llvm/sycl/test-e2e/Basic/host_write_back.cpp" "-o" "/__w/llvm/llvm/build-e2e/Basic/Output/host_write_back.cpp.tmp.out"
# command stderr:
clang++: warning: CUDA version 12.1 is only partially supported [-Wunknown-cuda-version]

$ ":" "RUN: at line 2"
note: command had no output on stdout or stderr
$ "env" "SYCL_PI_CUDA_ENABLE_IMAGE_SUPPORT=1" "ONEAPI_DEVICE_SELECTOR=ext_oneapi_cuda:gpu" "/__w/llvm/llvm/build-e2e/Basic/Output/host_write_back.cpp.tmp.out"
# command stderr:
host_write_back.cpp.tmp.out: /__w/llvm/llvm/llvm/sycl/test-e2e/Basic/host_write_back.cpp:[31](https://github.com/intel/llvm/actions/runs/5546033087/jobs/10129281860?pr=10363#step:16:32): void testHostAcessor(): Assertion `vec[0] == 1' failed.

error: command failed with exit status: -6

--

********************

Any ideas why?
https://github.com/intel/llvm/actions/runs/5546033087/jobs/10129281860?pr=10363 - CUDA log.
https://github.com/intel/llvm/actions/runs/5546033087/jobs/10129281286?pr=10363 - Intel log.

@bader
Copy link
Contributor

bader commented Jul 13, 2023

@aelovikov-intel told me that it's caused by misalignment between the compiler and tests, which is supposed to be fixed by 7c8932a. I.e. this is CI scripts issue.

@0x12CC 0x12CC deleted the buffer_write_back branch July 14, 2023 13:03
@0x12CC 0x12CC restored the buffer_write_back branch September 27, 2023 21:06
0x12CC added a commit to 0x12CC/llvm that referenced this pull request Sep 27, 2023
intel#10334 causes a performance regression since `HostPtr` cannot be reused
when it is read-only. This change fixes the regression by deferring the
copy operation to the creation of a writable accessor.

Signed-off-by: Michael Aziz <[email protected]>
@0x12CC 0x12CC deleted the buffer_write_back branch September 29, 2023 15:50
steffenlarsen pushed a commit that referenced this pull request Oct 11, 2023
#10334 causes a performance regression
since `HostPtr` can't be reused when it's read-only. This PR fixes the
regression by deferring the copy operation to the creation of a writable
accessor. It includes following the changes:

- A new `SYCLMemObjT::MCreateShadowCopy` to defer allocation. When the
`HostPtr` cannot be reused since it's read-only,
`SYCLMemObjT::handleHostData` sets this member to a function that will
allocate the shadow copy.
- A new `SYCLMemObjT::handleWriteAccessorCreation` member function. This
function calls `SYCLMemObjT::MCreateShadowCopy` and updates any existing
`MAllocaCommands` if `MUserPtr` changed.
- Whenever a writable host or device accessor is created,
`handleWriteAccessorCreation` gets called to ensure that any required
memory allocation occurs.

With this change, the allocation and copying overhead occurs during the
creation of the first writable accessor. There's no overhead if all of
the relevant accessors use `sycl::access_mode::read`.

---------

Signed-off-by: Michael Aziz <[email protected]>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

sycl::buffer write back
4 participants