Add CI workflow to build and test dpnp with the open source LLVM SYCL compiler - #3048
Open
antonwolfy wants to merge 18 commits into
Open
Add CI workflow to build and test dpnp with the open source LLVM SYCL compiler#3048antonwolfy wants to merge 18 commits into
antonwolfy wants to merge 18 commits into
Conversation
Adapt dpctl's os-llvm-sycl-build workflow to dpnp: download the nightly open source DPC++ toolchain, build dpnp against it via build_locally.py, and run the test suite. Triggered manually only for now.
The open source DPC++ nightly ships an IntelSYCLConfig.cmake that does not define add_sycl_to_target(), which dpnp's CMake relies on. Point IntelSYCL_DIR at dpnp's vendored Modules directory so the module that provides the function is used instead.
The vendored IntelSYCLConfig.cmake only recognized the Intel compiler
version string ("Intel(R) ... Compiler X.Y.Z"). The open source DPC++
nightly reports "clang version X.Y.Zgit", so parse_compiler_version()
returned empty and the module bailed out with an early return() before
defining add_sycl_to_target(), leaving dpnp's CMake unable to use it.
Add an else branch to parse the "clang version X.Y.Z" string so the
module completes and defines add_sycl_to_target() when building with the
open source LLVM SYCL compiler. The Intel compiler path is unchanged.
Configuration succeeded but compilation failed with "'oneapi/dpl/numeric' file not found": find_package(oneDPL) loaded the vendored oneDPLConfig.cmake (via PATHS), which resolves headers relative to the source tree where oneapi/dpl does not exist. The pip onedpl-devel package ships a proper oneDPLConfig.cmake under <prefix>/lib/cmake/oneDPL that resolves headers to <prefix>/include. Export oneDPL_ROOT pointing at the Python prefix so that config wins (oneDPL_ROOT has higher search priority than PATHS in CONFIG mode).
Setting oneDPL_ROOT only affects the CMake oneDPL target's include dirs, but the dpnp extensions include <oneapi/dpl/numeric> etc. without linking that target, relying on the compiler's default include path. The Intel oneAPI compiler bundles these headers; the open source DPC++ nightly does not, so compilation failed with "'oneapi/dpl/numeric' file not found". Export CPATH with the pip package's include directory so the oneDPL, MKL and TBB headers are on the default compiler search path.
The final "pip install -e ." failed resolving dpnp's own runtime pin dpctl>=0.23.0dev0, which is published only on the dppy dev channel and not on PyPI (max 0.22.1). Install dpctl from environments/dpctl_pkg.txt (which points at the dppy dev index) instead of PyPI so the pin is satisfied.
The dependencies step now reads environments/dpctl_pkg.txt, which only exists after checkout. Move the checkout step ahead of the dependency install so the requirements file is present.
CPATH adds directories as -I, so headers from the pip packages were treated as non-system and their deprecation warnings were shown, despite dpnp's CMake marking the compiler's SYCL include dir as SYSTEM. Switch to C_INCLUDE_PATH/CPLUS_INCLUDE_PATH, which are treated as -isystem, so the oneDPL/MKL/SYCL headers are found and their warnings stay suppressed.
The dpnp build linked against the open source nightly libsycl, but the pip dpctl (dppy dev channel) is built against Intel's released compiler and its libsycl.so.9 lacks symbols the nightly provides (e.g. sycl::_V1::detail::defaultAsyncHandler), so importing dpnp failed with an undefined symbol error at runtime. Check out IntelPython/dpctl and build it from source with the same nightly compiler (mirroring dpctl's own os-llvm-sycl-build workflow) so dpctl and dpnp share one SYCL runtime. Drop the pip dpctl install.
Building dpctl editable compiled the extensions but skipped CMake's install step, so dpctl/resources/cmake (where `python -m dpctl --cmakedir` points) had no DpctlConfig.cmake and dpnp's find_package(Dpctl) failed. Install dpctl with a plain "pip install . --no-build-isolation --no-deps" so scikit-build runs the install step and lays out include/ and resources/cmake properly. --no-deps keeps intel-sycl-rt off PyPI so the nightly libsycl on LD_LIBRARY_PATH remains the only SYCL runtime.
Pass --onemath so dpnp builds the oneMath interface from source (at the pinned tag, currently the latest v0.9) instead of linking the prebuilt oneMKL SYCL libraries, mirroring the tag-based job in check-onemath.yaml. This compiles the math layer with the same open source nightly compiler.
The MACRO_POPULATE_DISPATCH_* macros in populate.hpp defined the helper structs ContigFactory, StridedFactory and TypeMapFactory with the same name in the shared dpnp::extensions::ufunc::impl namespace across every per-ufunc translation unit. Each definition differs (it references that ufunc's own OutputType and *_impl kernel), so this is an ODR violation. The Intel oneAPI compiler inlines the factories' get() methods, so each translation unit uses its own definition. The open source intel/llvm DPC++ compiler instead emits them as weak COMDAT symbols and folds the identically-named instantiations across translation units, so several binary ufuncs get their contig/strided dispatch tables and result-type tables populated from the wrong (folded) translation unit. As a result fmin, fmod, heaviside, logaddexp2 and ldexp silently computed the elementwise maximum, and the integer result-type entries for fmin, fmax, fmod, gcd and lcm were dropped (int64,int64 resolved to no kernel). Give the factories internal, per-ufunc-unique names by pasting the ufunc name into them, so no two translation units share a symbol and there is nothing to fold.
Shorten the explanatory comments in the dpctl and dpnp build steps; no change to the build commands themselves.
Add a native concurrency group so only one build runs per ref, and drop the redundant styfle/cancel-workflow-action "Cancel Previous Runs" step that it supersedes.
Add a weekly (Saturday) workflow that opens a pull request when a newer intel/llvm oclcpuexp driver release or oneTBB release is available, keeping the pins in os-llvm-sycl-build.yml current. The nightly DPC++ compiler is discovered at run time and is not pinned.
antonwolfy
requested review from
ndgrigorian and
vlad-perevezentsev
as code owners
August 29, 2026 09:28
Bump peter-evans/create-pull-request to v8.1.1 (Node 24) to clear the Node 20 deprecation warning; pass find-latest-tag outputs to the resolve step via env instead of interpolating them into the shell body; and set sort-tags: false explicitly so the YYYY-WWNN driver tag is selected by release date rather than semver.
The pre-loop reset unset DEPLOY_NIGHTLY_TAG_SHA, a variable that is never used; the loop and downstream checks use DEPLOY_LLVM_TAG_SHA. Unset that name instead so stale loop state is actually cleared.
Give the ContigFactory/TypeMapFactory helpers unique per-op names in the vm populate macros and in the hand-written vm/erf_funcs, ufunc/erf_funcs and ufunc/float_power sources. Previously these shared one mangled name across many translation units in a single namespace, so their weak factory instantiations could be folded together by the open source LLVM SYCL compiler (the same ODR violation fixed in the ufunc populate.hpp macros). This matches the unique-naming convention lapack/blas/statistics already follow.
Collaborator
|
coverage: 78.494% (+0.01%) from 78.483% — antonwolfy:llvm-sycl-workflow into IntelPython:master |
Contributor
Author
|
Both new workflows are passing: |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
This PR adds a manually-triggered CI workflow that builds and tests dpnp against the open source
intel/llvmDPC++ (LLVM SYCL) nightly toolchain, instead of the released Intel oneAPI compiler. This gives early warning when an upstream compiler change breaks dpnp, and exercises the source against a second, independent SYCL implementation. Building against the open source nightly surfaced one real correctness bug in dpnp (fixed here) and required a few build-system adjustments, all included.What is added
.github/workflows/os-llvm-sycl-build.yml—workflow_dispatchjob that downloads the latestintel/llvmnightly compiler, the OpenCL CPU runtime (oclcpuexp) and oneTBB, builds dpctl and dpnp from source against the nightly (dpnp built against the oneMath interface), and runs the fulldpnp/testssuite..github/workflows/bump-sycl-deps.yml— weekly (Saturday) workflow that opens a PR when a newer oclcpuexp driver release or oneTBB release is available, keeping the pins inos-llvm-sycl-build.ymlcurrent. The nightly compiler itself is discovered at run time and is not pinned.dpnp/backend/cmake/Modules/IntelSYCLConfig.cmake— parse the open source compiler version string (clang version X.Y.Z) in the vendored module, which is forced for this build because the nightly does not shipadd_sycl_to_target().Correctness fix
Building the ufunc extension with the open source nightly exposed an ODR violation in
dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp: theMACRO_POPULATE_DISPATCH_*macros defined the helper structsContigFactory,StridedFactoryandTypeMapFactorywith the same name in the shareddpnp::extensions::ufunc::implnamespace across every per-ufunc translation unit, each with a different body. The Intel compiler inlines these away, but the open source compiler emits them as weak COMDAT symbols and folds the identically-named instantiations across translation units, so several binary ufuncs got their dispatch and result-type tables populated from the wrong translation unit. Observable effect:fmin,fmod,heaviside,logaddexp2andldexpsilently computed the elementwise maximum, and the integer result-type entries forfmin,fmax,fmod,gcdandlcmwere dropped. The fix gives the factories per-ufunc-unique names so no two translation units share a symbol.