Skip to content

[SYCL] Minor cleanup to improve build times and also fix failed tests - #52

Open
abagusetty wants to merge 13 commits into
wavefunction91:masterfrom
abagusetty:cleanup-sycl
Open

[SYCL] Minor cleanup to improve build times and also fix failed tests #52
abagusetty wants to merge 13 commits into
wavefunction91:masterfrom
abagusetty:cleanup-sycl

Conversation

@abagusetty

@abagusetty abagusetty commented Mar 12, 2026

Copy link
Copy Markdown
Collaborator
  1. Getting rid of legacy SYCL code and keep the implementations on par with CUDA
  2. Remove explicit nd_item arg to be passed around and replace with SYCL free-function for cleaner equivalents like CUDA
  3. Rename sycl::queue vars to stream to align with much of the other backends
  4. Added SYCL opencl:cpu CI for git runners

@abagusetty
abagusetty marked this pull request as ready for review March 13, 2026 00:20
@abagusetty abagusetty added the SYCL SYCL backend label Mar 13, 2026
abagusetty and others added 9 commits March 29, 2026 11:49
Add spir64_x86_64 to the accepted EXCHCXX_SYCL_TARGET values so the CPU
can be targeted ahead-of-time through opencl-aot, alongside the existing
GPU aliases. No new user-facing variable is introduced.

Two motivations:

- The OpenCL CPU runtime currently JIT-compiles the SPIR-V at first
  kernel launch and makes its own vectorization/FP choices at that
  point, outside the reach of the -Xsycl-target-frontend -fp-model=precise
  flag applied further down this file. Compiling ahead of time keeps
  device codegen under the SYCL toolchain. This is a candidate
  explanation for the XC_KERNEL signed-zero/vrho mismatches seen only on
  the OpenCL CPU backend, while the same tests pass under Level Zero on
  PVC and under GNU on CPU.
- It removes the one-time JIT warm-up from test execution.

No -march is passed, so opencl-aot leaves CL_CONFIG_CPU_TARGET_ARCH
unset and the output stays portable across x86_64 CI hosts.

Also add a guard rejecting any allowlisted target that has no flag
mapping. Previously nvidia_gpu_* and amd_gpu_* passed the allowlist but
matched no branch, yielding empty option lists: a silent JIT build that
still printed "AoT enabled".

Enable the new target in the CPU CI workflow to exercise the path.
-Xsycl-target-frontend -fp-model=precise, applied just above, is
forwarded only to the device frontend. icpx compiles the host side with
-fp-model=fast by default, which implies -fno-signed-zeros and permits
folding -0.0 and +0.0 together.

This shows up in TEST_CASE("Builtin Correctness Test"), which evaluates
the builtin kernels on host pointers with no device dispatch: R2SCANL_C
produces vlapl = -0.0 where libxc yields +0.0, and Catch2's Approx
distinguishes the two. The same test passes when ExchCXX is built with
GCC, which preserves signed zeros by default -- so the discrepancy
tracks the host compiler, not the SYCL backend.

Ask for -fp-model=precise on host code as well so builtin and libxc
results stay comparable across compilers. Guarded by
check_cxx_compiler_flag, so GCC/Clang builds are unaffected.
Move the host-side -fp-model=precise out of src/sycl/exchcxx_sycl.cmake,
where it only reached the exchcxx target, and into the top-level
CMakeLists.txt as a directory-scope flag set before dependencies are
configured. Everything in the build now inherits IEEE FP semantics:

  * libxc, FetchContent'd below and compiled with icx. It provides the
    reference values the builtin kernels are validated against, so it
    has to agree on signed-zero and reassociation rules as well --
    otherwise the comparison is between two differently-relaxed builds.
  * exchcxx itself.
  * the unit tests and benchmarks, which perform their own arithmetic on
    values returned by both backends.

C and C++ are probed separately, since libxc is C and exchcxx is C++,
and both checks are guarded so non-Intel compilers driving a SYCL build
are unaffected.

The device-side -Xsycl-target-frontend -fp-model=precise stays in
exchcxx_sycl.cmake; a bare -fp-model flag is not forwarded to the SYCL
device frontend, so the two are genuinely separate knobs.
@abagusetty

Copy link
Copy Markdown
Collaborator Author

@wavefunction91 Hi David, Its been a while I looked into this PR. Got it cleaned up with SYCL CI finally!

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

SYCL SYCL backend

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant