Skip to content

[OpenCL][USM] urEnqueueUSMFill incorrectly assumes destination memory alignment and fails #2440

New issue

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

#

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

Already on GitHub? # to your account

Jump to bottom
Open
rafbiels opened this issue Dec 9, 2024 · 7 comments · Fixed by intel/llvm#18153 · May be fixed by intel/llvm#18423
Open

[OpenCL][USM] urEnqueueUSMFill incorrectly assumes destination memory alignment and fails #2440

rafbiels opened this issue Dec 9, 2024 · 7 comments · Fixed by intel/llvm#18153 · May be fixed by intel/llvm#18423
Assignees
@martygrant
Labels
bug Something isn't working opencl OpenCL adapter specific issues

Comments

@rafbiels
Copy link
Contributor

rafbiels commented Dec 9, 2024 •
edited
Loading

TLDR: OpenCL adapter implementation of urEnqueueUSMFill calls clEnqueueMemFillINTEL for power-of-2 pattern size without checking destination memory alignment required by clEnqueueMemFillINTEL.

Full Story:
I ran into this issue when trying to add sycl::queue::fill test in intel/llvm#15991 (specific CI run failure: https://github.com/intel/llvm/actions/runs/11971631535/job/33695502538?pr=15991)

terminate called after throwing an instance of 'sycl::_V1::exception'
  what():  Enqueue process failed.
Aborted (core dumped)

I will disable the OpenCL CPU backend in the e2e test, linking this issue in a comment, so it can be re-enabled when the problem is solved.

The minimal reproducer for the issue is:

#include <sycl/sycl.hpp>
#include <array>

constexpr size_t PatternSize{32}; // bytes
constexpr size_t NumElements{10};

int main() {
  sycl::queue q{};
  using T = std::array<uint8_t, PatternSize>;
  T value{};

  T *dptr{sycl::malloc_device<T>(NumElements, q)};
  q.fill(dptr, value, NumElements).wait();
  sycl::free(dptr, q);

  return 0;
}

compiled and ran with:

clang++ -fsycl -fsycl-targets=spir64_x86_64 -o mini minimal.cpp
ONEAPI_DEVICE_SELECTOR=opencl:cpu ./mini

Debugging further with this UR change:

diff --git a/source/adapters/opencl/usm.cpp b/source/adapters/opencl/usm.cpp
index dfcc1dfa..ed18659e 100644
--- a/source/adapters/opencl/usm.cpp
+++ b/source/adapters/opencl/usm.cpp
@@ -276,12 +276,15 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMFill(
         cl_ext::getExtFuncFromContext<clEnqueueMemFillINTEL_fn>(
             CLContext, cl_ext::ExtFuncPtrCache->clEnqueueMemFillINTELCache,
             cl_ext::EnqueueMemFillName, &EnqueueMemFill));
-
-    CL_RETURN_ON_FAILURE(
+    CLErr =
         EnqueueMemFill(cl_adapter::cast<cl_command_queue>(hQueue), ptr,
                        pPattern, patternSize, size, numEventsInWaitList,
                        cl_adapter::cast<const cl_event *>(phEventWaitList),
-                       cl_adapter::cast<cl_event *>(phEvent)));
+                       cl_adapter::cast<cl_event *>(phEvent));
+    std::cout << "EnqueueMemFillINTEL(patternSize=" << patternSize << ") return code " << CLErr << std::endl;
+    if (CLErr != CL_SUCCESS) {
+      return mapCLErrorToUR(CLErr);
+    }
     return UR_RESULT_SUCCESS;
   }

confirmed that it's the clEnqueueMemFillINTEL call which sometimes returns -30 (CL_INVALID_VALUE).

EnqueueMemFillINTEL(patternSize=32) return code -30
terminate called after throwing an instance of 'sycl::_V1::exception'
  what():  Enqueue process failed.
Aborted (core dumped)

I noticed that in my build of UR / DPC++ this happens when the binary name is shorter than 24 characters, but stops happening when it is longer. Simply renaming the file changes the behaviour. In another build I got the opposite behaviour where short-named binary succeeds but long-named fails. I assume what happens is that long file name causes heap allocation for argv[0] and shifts the memory layout, and thus alignment of the device allocation.

$ # short name fails
$ ./mini
EnqueueMemFillINTEL(patternSize=32) return code -30
terminate called after throwing an instance of 'sycl::_V1::exception'
  what():  Enqueue process failed.
Aborted (core dumped)

$ # rename the file
$ cp mini mini12341234123412341234

$ # long name succeeds
$ ./mini12341234123412341234
EnqueueMemFillINTEL(patternSize=32) return code 0

I note that clEnqueueMemFillINTEL as described here:
https://registry.khronos.org/OpenCL/extensions/intel/cl_intel_unified_shared_memory.html
may return:

CL_INVALID_VALUE if dst_ptr is NULL, or if dst_ptr is not aligned to pattern_size bytes

so I assume this is what happens as I checked the other conditions for returning CL_INVALID_VALUE are not met.

IIUC neither SYCL API nor UR API make any requirements about the destination memory alignment for their USM fill functions, therefore it is incorrect for the implementation to assume alignment. I think the solution here could be to check the alignment and take the other (slower) path which doesn't call clEnqueueMemFillINTEL when the alignment requirement is not met.

Side note: something seems to be lost in error handling here as the user is informed neither about the error code (INVALID_VALUE) nor its origin (USM fill). There is only a generic exception thrown by the SYCL runtime.
@rafbiels rafbiels mentioned this issue Dec 9, 2024
Merged
@rafbiels rafbiels added the opencl OpenCL adapter specific issues label Dec 9, 2024
@rafbiels rafbiels mentioned this issue Dec 20, 2024
Closed
@aelovikov-intel
Copy link
Contributor

I'm reverting the test addition in intel/llvm#16465, please ensure it's properly re-committed before closing this issue.

@rafbiels
Copy link
Contributor Author

rafbiels commented Jan 8, 2025

The test was added back with intel/llvm#16544 now properly marked as UNSUPPORTED: (opencl && cpu). The issue is still there and the test can be used to validate future fixes (in addition to the reproducer above).

@aelovikov-intel aelovikov-intel mentioned this issue Jan 10, 2025
Closed
@sarnex sarnex mentioned this issue Jan 10, 2025
Closed
@uditagarwal97 uditagarwal97 mentioned this issue Jan 10, 2025
Merged
@rafbiels
Copy link
Contributor Author

It's been shown to fail also on the opencl:fpga target and on a second thought there is no reason at UR level why it wouldn't affect GPU as well, even if it hasn't been observed. The UR code is the same for any device type, so I believe all can be affected depending on how the driver allocates memory. intel/llvm#16588 marked the test UNSUPPORTED for all OpenCL devices.

@kbenzie kbenzie added the bug Something isn't working label Apr 8, 2025
@martygrant martygrant self-assigned this Apr 10, 2025
@martygrant martygrant mentioned this issue Apr 23, 2025
Merged
@rafbiels
Copy link
Contributor Author

rafbiels commented May 7, 2025

This is not fixed, the reproducer still shows the issue with intel/llvm nightly-2025-05-07 and the following OpenCL CPU device:

[opencl:cpu][opencl:0] Intel(R) OpenCL, 12th Gen Intel(R) Core(TM) i9-12900K OpenCL 3.0 (Build 0) [2025.19.4.0.18_160000.xmain-hotfix]

$ clang++ --version
clang version 21.0.0git (https://github.com/intel/llvm 1e5d85da1c4e5a7bc62948a99f13885b565b0c14)
$ clang++ -fsycl -fsycl-targets=spir64_x86_64 -o mini minimal.cpp
$ ONEAPI_DEVICE_SELECTOR=opencl:cpu ./mini
terminate called after throwing an instance of 'sycl::_V1::exception'
  what():  Enqueue process failed.
Aborted (core dumped)

It looks like the CI configuration has changed and the affected test is no longer being run on the affected device. It was originally failing in Post Commit CI on OpenCL CPU and FPGA devices, but the Post Commit looks to run on GPU only now:
https://github.com/intel/llvm/actions/runs/14880949602

@rafbiels rafbiels reopened this May 7, 2025
@martygrant
Copy link
Contributor

@rafbiels I ran the actual failing e2e test and it passed on my machine with the same version of OpenCL as in your comment, is that something you can check locally too? I didn't realise the CI configuration has changed to not run on CPU I will investigate this. The reproducer you gave does look like it's failing though, it seems like there is a difference between that and the e2e test?

@rafbiels
Copy link
Contributor Author

rafbiels commented May 7, 2025

Hi @martygrant,
as the issue lies in memory alignment, any small change to host allocation prior to the to-be-filled memory may hide or bring up the crash. I found that just renaming the binary has this effect as the size of argv[0] influences how the subsequent allocations are aligned. I assume version differences of glibc, libstdc++ and other libraries may also have this influence, changing the behaviour from one system to another. You could try renaming the test binary locally to see if that triggers the issue for you (or just work with the reproducer).

@rafbiels rafbiels reopened this May 8, 2025
@martygrant
Copy link
Contributor

I have reverted my PR to re-enable the fill_any_size.cpp test (intel/llvm#18381) as it's passing was just being flaky and a proper fix needs to be investigated.

@martygrant martygrant linked a pull request May 12, 2025 that will close this issue
Draft
# for free to join this conversation on GitHub. Already have an account? # to comment
Assignees

@martygrant martygrant

Labels
bug Something isn't working opencl OpenCL adapter specific issues
Projects
None yet
Milestone
No milestone
4 participants
@kbenzie @martygrant @aelovikov-intel @rafbiels