Skip to content
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

Device code not split in modules when using aspect::atomic64 #12743

Closed
Alcpz opened this issue Feb 16, 2024 · 0 comments · Fixed by #14052
Closed

Device code not split in modules when using aspect::atomic64 #12743

Alcpz opened this issue Feb 16, 2024 · 0 comments · Fixed by #14052
Labels
bug Something isn't working confirmed

Comments

@Alcpz
Copy link
Contributor

Alcpz commented Feb 16, 2024

#12673 enabled tests for the opencl:fpga backend, which don't support 64bit atomics.
The following tests,

syclcompat/atomic/atomic_arith.cpp
syclcompat/atomic/atomic_comp_exchange.cpp

failed with the following output:

# .---command stderr------------
# | terminate called after throwing an instance of 'sycl::_V1::compile_program_error'
# |   what():  The program was built for 1 devices
# | Build program log for 'Intel(R) FPGA Emulation Device':
# | Compilation started
# | Unsupported SPIR-V module
# | int 64bit atomics are not supported on FPGA emulator.
# | Compilation failed
# |  -11 (PI_ERROR_BUILD_PROGRAM_FAILURE)
# `-----------------------------
# error: command failed with exit status: -6

The issue is avoided by adding per_kernel split to the tests (-fsycl-device-code-split=per_kernel), which indicates that kernels using atomic64 are not being split into a separate module.

According to #12673 (comment), adding the flag shouldn't be necessary.
#12673 (comment) provides additional context.

To Reproduce

Run syclcompat/atomic/atomic_arith.cpp or syclcompat/atomic/atomic_comp_exchange.cpp sycl-e2e tests using the opencl:fpga or opencl:acc backends.

Environment:

  • OS: Linux
  • Target device and vendor: OpenCL FPGA
  • DPC++ version: intel/llvm 73d3473
@Alcpz Alcpz added the bug Something isn't working label Feb 16, 2024
@Alcpz Alcpz changed the title Kernel modules not split when using aspect::atomic64 Device code not split in modules when using aspect::atomic64 Feb 16, 2024
dm-vodopyanov pushed a commit that referenced this issue Jun 13, 2024
atomic_ref<T *> uses 64-bit atomics and it should be decorated with the
corresponding aspect.

fixes: #12743
ianayl pushed a commit to ianayl/sycl that referenced this issue Jun 13, 2024
atomic_ref<T *> uses 64-bit atomics and it should be decorated with the
corresponding aspect.

fixes: intel#12743
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working confirmed
Projects
None yet
Development

Successfully merging a pull request may close this issue.

2 participants