intel
Device code not split in modules when using aspect::atomic64
https://github.com/intel/llvm/pull/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 https://github.com/intel/llvm/pull/12673#issuecomment-1946701751, adding the flag shouldn't be necessary. https://github.com/intel/llvm/pull/12673#issuecomment-1946771946 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 https://github.com/intel/llvm/commit/73d34739bff5848cca075fbec415e06f67848efe
