Disable zerocallusedregs in llvm-intel for CUDA

overlay.nix

@@ -19,6 +19,7 @@ let
     cudainfo = prev.callPackage ./pkgs/cudainfo/default.nix { };
     #extrae = callPackage ./pkgs/extrae/default.nix { }; # Broken and outdated
     gpi-2 = callPackage ./pkgs/gpi-2/default.nix { };
+    gromacs = callPackage ./pkgs/gromacs/default.nix { enableSYCL = true; };
     intel-apt = callPackage ./pkgs/intel-oneapi/packages.nix { };
     intelPackages_2023 = callPackage ./pkgs/intel-oneapi/2023.nix { };
     intelPackages_2024 = final.intel-apt.hpckit_2024;
@@ -33,6 +34,8 @@ let
     linuxPackages_latest = prev.linuxPackages_latest.extend(_final: _prev: {
       amd-uprof-driver = _prev.callPackage ./pkgs/amd-uprof/driver.nix { };
     });
+    llvm-intel-unwrapped = callPackage ./pkgs/llvm-intel/default.nix { };
+    llvm-intel = callPackage ./pkgs/llvm-intel/wrapped.nix { };
     lmbench = callPackage ./pkgs/lmbench/default.nix { };
     # Broken and unmantained
     # mcxx = callPackage ./pkgs/mcxx/default.nix { };
@@ -55,9 +58,12 @@ let
     ovni = callPackage ./pkgs/ovni/default.nix { };
     ovniGit = final.ovni.override { useGit = true; };
     paraverKernel = callPackage ./pkgs/paraver/kernel.nix { };
+    pocl = callPackage ./pkgs/pocl/default.nix { };
+    pocl-noicd = callPackage ./pkgs/pocl/default.nix { enableICD = false; };
     prometheus-slurm-exporter = prev.callPackage ./pkgs/slurm-exporter/default.nix { };
     #pscom = callPackage ./pkgs/parastation/pscom.nix { }; # Unmaintaned
     #psmpi = callPackage ./pkgs/parastation/psmpi.nix { }; # Unmaintaned
+    rodinia = callPackage ./pkgs/rodinia/default.nix { };
     slurm = import ./pkgs/slurm/default.nix { slurm = prev.slurm; };
     sonar = callPackage ./pkgs/sonar/default.nix { };
     stdenvClangOmpss2 = final.stdenv.override { cc = final.clangOmpss2; allowedRequisites = null; };
@@ -68,6 +74,7 @@ let
     tampi = callPackage ./pkgs/tampi/default.nix { };
     tasycl = callPackage ./pkgs/tasycl/default.nix { };
     tasycl-acpp = callPackage ./pkgs/tasycl/default.nix { useIntel = false; };
+    unified-runtime = callPackage ./pkgs/unified-runtime/default.nix { };
     upc-qaire-exporter = prev.callPackage ./pkgs/upc-qaire-exporter/default.nix { };
     wxparaver = callPackage ./pkgs/paraver/default.nix { };
   };
@@ -79,6 +86,11 @@ let
     hello-cpp = callPackage ./test/compilers/hello-cpp.nix { };
     hello-sycl = callPackage ./test/compilers/hello-sycl.nix { };
     hello-syclompss = callPackage ./test/compilers/icpx-ompss2.nix { };
+    hello-sycl-pocl = (callPackage ./test/compilers/hello-sycl.nix { }).withPocl;
+    hello-sycl-intel = (callPackage ./test/compilers/hello-sycl.nix { }).withIntel;
+    hello-syclompss-pocl = (callPackage ./test/compilers/icpx-ompss2.nix { }).withPocl;
+    hello-syclompss-intel = (callPackage ./test/compilers/icpx-ompss2.nix { }).withIntel;
+
     lto = callPackage ./test/compilers/lto.nix { };
     asan = callPackage ./test/compilers/asan.nix { };
     intel2023-icx-c   = hello-c.override   { stdenv = final.intelPackages_2023.stdenv; };
@@ -112,6 +124,15 @@ let
     clangNosvOmpv-ld = callPackage ./test/compilers/clang-openmp-ld.nix {
       stdenv = final.stdenvClangOmpss2NodesOmpv;
     };
+
+    ocl-build = callPackage ./test/compilers/opencl.nix { };
+    ocl-build-noicd = callPackage ./test/compilers/opencl.nix { ocl-icd = final.pocl-noicd; };
+
+    pocl = callPackage ./test/compilers/pocl.nix { };
+    ocl-run-pocl = ocl-build.test-icd;
+    ocl-run-pocl-noicd = ocl-build-noicd.test-noicd;
+
+    sycl-ls = callPackage ./test/compilers/sycl-ls.nix { };
   };
 
   # For now, only build toplevel packages in CI/Hydra

pkgs/gromacs/default.nix

@@ -0,0 +1,190 @@
+{
+  lib,
+  stdenv,
+  fetchurl,
+  cmake,
+  hwloc,
+  fftw,
+  perl,
+  blas,
+  lapack,
+  llvmPackages,
+  intelPackages ? null,
+  mpi,
+  cudaPackages,
+  plumed,
+  singlePrec ? true,
+  config,
+  enableSYCL ? false,
+  enableCuda ? config.cudaSupport,
+  enableMpi ? false,
+  enablePlumed ? false,
+  cpuAcceleration ? null,
+}:
+
+# CUDA is only implemented for single precission
+assert enableCuda -> singlePrec;
+
+assert !(enableSYCL && enableCuda);
+
+assert enableSYCL -> intelPackages != null;
+
+let
+  inherit (cudaPackages.flags) cmakeCudaArchitecturesString;
+
+  # Select reasonable defaults for all major platforms
+  # The possible values are defined in CMakeLists.txt:
+  # AUTO None SSE2 SSE4.1 AVX_128_FMA AVX_256 AVX2_256
+  # AVX2_128 AVX_512 AVX_512_KNL MIC ARM_NEON ARM_NEON_ASIMD
+  SIMD =
+    x:
+    if (cpuAcceleration != null) then
+      x
+    else if stdenv.hostPlatform.system == "i686-linux" then
+      "SSE2"
+    else if stdenv.hostPlatform.system == "x86_64-linux" then
+      "SSE4.1"
+    else if stdenv.hostPlatform.system == "x86_64-darwin" then
+      "SSE4.1"
+    else if stdenv.hostPlatform.system == "aarch64-linux" then
+      "ARM_NEON_ASIMD"
+    else
+      "None";
+
+  source =
+    if enablePlumed then
+      {
+        version = "2024.2";
+        hash = "sha256-gCp+M18uiVdw9XsVnk7DaOuw/yzm2sz3BsboAlw2hSs=";
+      }
+    else
+      {
+        version = "2025.3";
+        hash = "sha256-i9/KAmjz8Qp8o8BuWbYvc+oCQgxnIRwP85EvMteDPGU=";
+      };
+
+  stdenv' = if enableSYCL then intelPackages.stdenv else stdenv;
+
+in
+stdenv'.mkDerivation rec {
+  pname = "gromacs";
+  version = source.version;
+
+  src = fetchurl {
+    url = "ftp://ftp.gromacs.org/pub/gromacs/gromacs-${version}.tar.gz";
+    inherit (source) hash;
+  };
+
+  patches = [ (if enablePlumed then ./pkgconfig-2024.patch else ./pkgconfig-2025.patch) ];
+
+  postPatch = lib.optionalString enablePlumed ''
+    plumed patch -p -e gromacs-${source.version}
+  '';
+
+  outputs = [
+    "out"
+    "dev"
+    "man"
+  ];
+
+  nativeBuildInputs = [
+    cmake
+  ]
+  ++ lib.optional enablePlumed plumed
+  ++ lib.optionals enableCuda [ cudaPackages.cuda_nvcc ];
+
+  env.MKLROOT = intelPackages.mkl;
+
+  buildInputs = [
+    fftw
+    perl
+    hwloc
+    blas
+    lapack
+  ]
+  ++ lib.optional enableMpi mpi
+  ++ lib.optionals enableCuda [
+    cudaPackages.cuda_cccl
+    cudaPackages.cuda_cudart
+    cudaPackages.libcufft
+    cudaPackages.cuda_profiler_api
+  ]
+  ++ lib.optional stdenv.hostPlatform.isDarwin llvmPackages.openmp;
+
+  propagatedBuildInputs = lib.optional enableMpi mpi;
+  propagatedUserEnvPkgs = lib.optional enableMpi mpi;
+
+  cmakeFlags = [
+    (lib.cmakeBool "GMX_HWLOC" true)
+    "-DGMX_SIMD:STRING=${SIMD cpuAcceleration}"
+    "-DGMX_OPENMP:BOOL=TRUE"
+    "-DBUILD_SHARED_LIBS=ON"
+  ]
+  ++ (
+    if singlePrec then
+      [
+        "-DGMX_DOUBLE=OFF"
+      ]
+    else
+      [
+        "-DGMX_DOUBLE=ON"
+        "-DGMX_DEFAULT_SUFFIX=OFF"
+      ]
+  )
+  ++ (
+    if enableMpi then
+      [
+        "-DGMX_MPI:BOOL=TRUE"
+        "-DGMX_THREAD_MPI:BOOL=FALSE"
+      ]
+    else
+      [
+        "-DGMX_MPI:BOOL=FALSE"
+      ]
+  )
+  ++ (lib.optionals enableSYCL [
+    "-DGMX_GPU=SYCL"
+    "-DGMX_OPENMP=OFF" # TODO: enable OpenMP with SYCL?
+  ])
+  ++ lib.optionals enableCuda [
+    "-DGMX_GPU=CUDA"
+    (lib.cmakeFeature "CMAKE_CUDA_ARCHITECTURES" cmakeCudaArchitecturesString)
+
+    # Gromacs seems to ignore and override the normal variables, so we add this ad hoc:
+    (lib.cmakeFeature "GMX_CUDA_TARGET_COMPUTE" cmakeCudaArchitecturesString)
+  ];
+
+  postInstall = ''
+    moveToOutput share/cmake $dev
+  '';
+
+  meta = with lib; {
+    homepage = "https://www.gromacs.org";
+    license = licenses.lgpl21Plus;
+    description = "Molecular dynamics software package";
+    longDescription = ''
+      GROMACS is a versatile package to perform molecular dynamics,
+      i.e. simulate the Newtonian equations of motion for systems
+      with hundreds to millions of particles.
+
+      It is primarily designed for biochemical molecules like
+      proteins, lipids and nucleic acids that have a lot of
+      complicated bonded interactions, but since GROMACS is
+      extremely fast at calculating the nonbonded interactions (that
+      usually dominate simulations) many groups are also using it
+      for research on non-biological systems, e.g. polymers.
+
+      GROMACS supports all the usual algorithms you expect from a
+      modern molecular dynamics implementation, (check the online
+      reference or manual for details), but there are also quite a
+      few features that make it stand out from the competition.
+
+      See: https://www.gromacs.org/about.html for details.
+    '';
+    platforms = platforms.unix;
+    maintainers = with maintainers; [
+      sheepforce
+      markuskowa
+    ];
+  };
+}

pkgs/gromacs/pkgconfig-2024.patch

@@ -0,0 +1,24 @@
+diff --git a/src/external/muparser/muparser.pc.in b/src/external/muparser/muparser.pc.in
+index 646787cb53..9b97ad57f7 100644
+--- a/src/external/muparser/muparser.pc.in
++++ b/src/external/muparser/muparser.pc.in
+@@ -1,7 +1,5 @@
+-prefix=@CMAKE_INSTALL_PREFIX@
+-exec_prefix=${prefix}
+-libdir=${prefix}/@CMAKE_INSTALL_LIBDIR@
+-includedir=${prefix}/@CMAKE_INSTALL_INCLUDEDIR@
++libdir=@CMAKE_INSTALL_FULL_LIBDIR@
++includedir=@CMAKE_INSTALL_FULL_INCLUDEDIR@
+ 
+ Name: @PACKAGE_NAME@
+ Description: Mathematical expressions parser library
+diff --git a/src/gromacs/libgromacs.pc.cmakein b/src/gromacs/libgromacs.pc.cmakein
+index ec1ed6684e..ca1105474a 100644
+--- a/src/gromacs/libgromacs.pc.cmakein
++++ b/src/gromacs/libgromacs.pc.cmakein
+@@ -1,4 +1,4 @@
+-libdir=@CMAKE_INSTALL_PREFIX@/@CMAKE_INSTALL_LIBDIR@
++libdir=@CMAKE_INSTALL_FULL_LIBDIR@
+ 
+ Name: libgromacs@GMX_LIBS_SUFFIX@
+ Description: Gromacs library

pkgs/gromacs/pkgconfig-2025.patch

@@ -0,0 +1,36 @@
+diff --git a/src/external/muparser/muparser.pc.in b/src/external/muparser/muparser.pc.in
+index 646787cb53..d26e84de8f 100644
+--- a/src/external/muparser/muparser.pc.in
++++ b/src/external/muparser/muparser.pc.in
+@@ -1,11 +1,9 @@
+-prefix=@CMAKE_INSTALL_PREFIX@
+-exec_prefix=${prefix}
+-libdir=${prefix}/@CMAKE_INSTALL_LIBDIR@
+-includedir=${prefix}/@CMAKE_INSTALL_INCLUDEDIR@
+-
+-Name: @PACKAGE_NAME@
+-Description: Mathematical expressions parser library
+-Version: @MUPARSER_VERSION@
+-Requires:
+-Libs: -L${libdir} -lmuparser
+-Cflags: -I${includedir} @PKG_CONFIG_FLAGS@
++libdir=@CMAKE_INSTALL_FULL_LIBDIR@
++includedir=@CMAKE_INSTALL_FULL_INCLUDEDIR@
++
++Name: @PACKAGE_NAME@
++Description: Mathematical expressions parser library
++Version: @MUPARSER_VERSION@
++Requires:
++Libs: -L${libdir} -lmuparser
++Cflags: -I${includedir} @PKG_CONFIG_FLAGS@
+diff --git a/src/gromacs/libgromacs.pc.cmakein b/src/gromacs/libgromacs.pc.cmakein
+index af9b5a6dc0..5f58d549bf 100644
+--- a/src/gromacs/libgromacs.pc.cmakein
++++ b/src/gromacs/libgromacs.pc.cmakein
+@@ -1,5 +1,4 @@
+-prefix=@CMAKE_INSTALL_PREFIX@
+-libdir=${prefix}/@CMAKE_INSTALL_LIBDIR@
++libdir=@CMAKE_INSTALL_FULL_LIBDIR@
+
+ Name: libgromacs@GMX_LIBS_SUFFIX@
+ Description: Gromacs library

pkgs/intel-oneapi/patch_intel.nix

@@ -21,6 +21,7 @@
   rdma-core,
   ucx,
   zlib,
+  writeTextFile,
 }:
 
 lib.makeOverridable (
@@ -89,6 +90,9 @@ lib.makeOverridable (
           ln -s $out/bin/icpx $out/bin/c++
           ln -s $out/bin/icx  $out/bin/cc
 
+          ln -s $ccPath/sycl-ls $out/bin/sycl-ls
+          ln -s $ccPath/sycl-trace $out/bin/sycl-trace
+
           sed -i 's/.*isCxx=0/isCxx=1/' $out/bin/icpx
 
           # Use this to detect when a compiler subprocess is called
@@ -156,11 +160,17 @@ lib.makeOverridable (
           folder: version:
           let
             original = "${finalAttrs.finalPackage}/${folder}/${version}";
+
+            etc-vendors = writeTextFile {
+              name = "intel-ocl-icd";
+              text = "${original}/lib/libintelocl.so";
+              destination = "/etc/OpenCL/vendors/intel.icd";
+            };
           in
           symlinkJoin {
             pname = "intel-${folder}";
             inherit version;
-            paths = [ original ];
+            paths = [ original ] ++ lib.optionals (folder == "compiler") [ etc-vendors ];
             passthru = { inherit original; };
           }
         ) _components;

pkgs/llvm-intel/default.nix

@@ -0,0 +1,243 @@
+{
+  stdenv,
+  lib,
+  ninja,
+  autoAddDriverRunpath,
+  wrapCC,
+  cmake,
+  emhash,
+  fetchFromGitHub,
+  khronos-ocl-icd-loader,
+  libbacktrace,
+  opencl-headers,
+  parallel-hashmap,
+  perl,
+  pkg-config,
+  python3,
+  spirv-headers,
+  spirv-tools,
+  symlinkJoin,
+  zlib,
+
+  cudaPackages,
+  enableCuda ? true,
+}:
+
+let
+
+  llvmMajorVersion = "21";
+  version = "6.2.1";
+  src = fetchFromGitHub {
+    owner = "intel";
+    repo = "llvm";
+    tag = "v${version}";
+    sha256 = "sha256-j8+DmGKO0qDF5JjH+DlkLKs1kBz6dS7ukwySd/Crqv0=";
+  };
+
+  pinned = {
+    inherit parallel-hashmap;
+
+    vc-intrinsics.src = fetchFromGitHub {
+      owner = "intel";
+      repo = "vc-intrinsics";
+      rev = "4e51b2467104a257c22788e343dafbdde72e28bb";
+      sha256 = "sha256-AHXeKbih4bzmcuu+tx2TeS7Ixmk54uS1vKFVxI6ZP8g=";
+    };
+
+    unified-memory-framework.src = fetchFromGitHub {
+      owner = "oneapi-src";
+      repo = "unified-memory-framework";
+      tag = "v0.11.0";
+      sha256 = "sha256-k8QdP2u1QOoeCZ6ps4sM8+1iZq/H3q0lqCfVk8mz9KI=";
+    };
+
+    ocl-headers = opencl-headers.overrideAttrs {
+      version = "2024.10.24";
+      src = fetchFromGitHub {
+        owner = "KhronosGroup";
+        repo = "OpenCL-Headers";
+        tag = "v2024.10.24";
+        sha256 = "sha256-KDlruE0IG8d+lAChxYyc6dg5XOvqCMrMyO69sdAzejA=";
+      };
+    };
+
+    ocl-icd = khronos-ocl-icd-loader;
+
+    spirv-headers = spirv-headers.overrideAttrs {
+      version = "1.4.309.0";
+      src = fetchFromGitHub {
+        owner = "KhronosGroup";
+        repo = "SPIRV-Headers";
+        tag = "vulkan-sdk-1.4.309.0";
+        sha256 = "sha256-Q1i6i5XimULuGufP6mimwDW674anAETUiIEvDQwvg5Y=";
+      };
+    };
+  };
+
+  fetchcontent_cmakeFlags = lib.mapAttrsToList (
+    name: value: (lib.cmakeFeature "FETCHCONTENT_SOURCE_DIR_${lib.toUpper name}" (toString value.src))
+  ) pinned;
+
+  # Hack to fix #include <emhash/...>
+  emhash' = symlinkJoin {
+    name = "emhash";
+    paths = [ emhash ];
+    postBuild = ''
+      pushd $out/include
+      ln -s . emhash
+      popd
+    '';
+  };
+
+in
+stdenv.mkDerivation {
+  pname = "intel-llvm";
+  inherit src version;
+
+  enableParallelBuilding = true;
+  strictDeps = true;
+
+  nativeBuildInputs = [
+    perl
+    cmake
+    ninja
+    pkg-config
+    python3
+    spirv-tools
+    autoAddDriverRunpath
+  ]
+  ++ lib.optionals enableCuda [
+    cudaPackages.cuda_nvcc
+  ];
+
+  buildInputs = [
+    emhash'
+    libbacktrace
+    zlib
+
+    pinned.ocl-icd
+    pinned.ocl-headers
+    pinned.spirv-headers
+  ]
+  ++ lib.optionals enableCuda [
+    cudaPackages.cudatoolkit
+  ];
+
+  patches = [ ./fix-ur.patch ]; # Has been fixed in next release
+
+  preBuild = ''
+    chmod -R u+w /build/source
+  '';
+
+  cmakeDir = "../llvm";
+
+  cmakeFlags =
+    let
+      cmakeList = name: l: lib.cmakeOptionType "list" name (builtins.concatStringsSep ";" l);
+
+      LLVM_EXTERNAL_PROJECTS = [
+        "sycl"
+        "llvm-spirv"
+        "opencl"
+        "xpti"
+        "xptifw"
+        "compiler-rt"
+        "sycl-jit"
+        "libdevice"
+      ]
+      ++ lib.optionals enableCuda [
+        "libclc"
+      ];
+    in
+    [
+      (lib.cmakeBool "LLVM_ENABLE_ASSERTIONS" true)
+
+      (cmakeList "LLVM_TARGETS_TO_BUILD" (
+        [
+          "host"
+          "SPIRV"
+        ]
+        ++ lib.optionals enableCuda [ "NVPTX" ]
+      ))
+      (cmakeList "LLVM_EXTERNAL_PROJECTS" LLVM_EXTERNAL_PROJECTS)
+      (cmakeList "LLVM_ENABLE_PROJECTS" ([ "clang" ] ++ LLVM_EXTERNAL_PROJECTS))
+
+      (lib.cmakeBool "UR_USE_EXTERNAL_UMF" false)
+      (lib.cmakeOptionType "path" "UR_OPENCL_INCLUDE_DIR" (toString pinned.ocl-headers))
+      (lib.cmakeOptionType "path" "UR_OPENCL_ICD_LOADER_LIBRARY" (toString pinned.ocl-icd))
+      (lib.cmakeFeature "UR_OPENCL_ICD_LOADER_LIBRARY" (toString pinned.ocl-icd))
+
+      (lib.cmakeBool "BUILD_SHARED_LIBS" false)
+      (lib.cmakeBool "LLVM_BUILD_TOOLS" true)
+      (lib.cmakeBool "LLVM_ENABLE_ZSTD" true)
+      (lib.cmakeBool "LLVM_USE_STATIC_ZSTD" true)
+      (lib.cmakeBool "LLVM_ENABLE_DOXYGEN" false)
+      (lib.cmakeBool "LLVM_ENABLE_LLD" false)
+      (lib.cmakeBool "LLVM_ENABLE_SPHINX" false)
+      (lib.cmakeBool "LLVM_SPIRV_ENABLE_LIBSPIRV_DIS" false)
+      (lib.cmakeBool "SYCL_ENABLE_WERROR" false)
+      (lib.cmakeBool "SYCL_INCLUDE_TESTS" false)
+      (lib.cmakeBool "SYCL_ENABLE_EXTENSION_JIT" true)
+      (lib.cmakeBool "SYCL_ENABLE_MAJOR_RELEASE_PREVIEW_LIB" false)
+      (lib.cmakeBool "SYCL_ENABLE_XPTI_TRACING" true)
+      (lib.cmakeBool "XPTI_ENABLE_WERROR" false)
+
+      (cmakeList "SYCL_ENABLE_BACKENDS" ([ "opencl" ] ++ lib.optionals enableCuda [ "cuda" ]))
+    ]
+    ++ fetchcontent_cmakeFlags
+    ++ lib.optionals enableCuda [
+      (lib.cmakeOptionType "list" "LIBCLC_TARGETS_TO_BUILD" "nvptx64--nvidiacl")
+      (lib.cmakeBool "LIBCLC_GENERATE_REMANGLED_VARIANTS" true)
+      (lib.cmakeBool "LIBCLC_NATIVECPU_HOST_TARGET" false)
+    ];
+
+  hardeningDisable = lib.optionals enableCuda [ "zerocallusedregs" ]; # not supported for nvidia
+
+  postPatch =
+
+    let
+      # See the postPatch phase for details on why this is used
+      ccWrapperStub = wrapCC (
+        stdenv.mkDerivation {
+          name = "ccWrapperStub";
+          dontUnpack = true;
+          installPhase = ''
+            mkdir -p $out/bin
+            cat > $out/bin/clang-${llvmMajorVersion} <<'EOF'
+            #!/bin/sh
+            exec "$NIX_BUILD_TOP/source/build/bin/clang-${llvmMajorVersion}" "$@"
+            EOF
+            chmod +x $out/bin/clang-${llvmMajorVersion}
+            cp $out/bin/clang-${llvmMajorVersion} $out/bin/clang
+            cp $out/bin/clang-${llvmMajorVersion} $out/bin/clang++
+          '';
+          passthru.isClang = true;
+        }
+      );
+    in
+    ''
+      # Parts of libdevice are built using the freshly-built compiler.
+      # As it tries to link to system libraries, we need to wrap it with the
+      # usual nix cc-wrapper.
+      # Since the compiler to be wrapped is not available at this point,
+      # we use a stub that points to where it will be later on
+      # in `$NIX_BUILD_TOP/source/build/bin/clang-${llvmMajorVersion}`
+      substituteInPlace libdevice/cmake/modules/SYCLLibdevice.cmake \
+        --replace-fail "\''${clang_exe}" "${ccWrapperStub}/bin/clang++"
+    '';
+
+  passthru = {
+    inherit llvmMajorVersion;
+  };
+
+  meta = {
+    homepage = "https://github.com/intel/llvm";
+    description = "Intel staging area for llvm.org contribution";
+    maintainers = with lib.maintainers.bsc; [ abonerib ];
+    platforms = lib.platforms.linux;
+    license = [
+      lib.licenses.asl20
+      lib.licenses.llvm-exception
+    ];
+  };
+}

pkgs/llvm-intel/fix-ur.patch

@@ -0,0 +1,23 @@
+diff --git a/unified-runtime/source/adapters/cuda/CMakeLists.txt b/unified-runtime/source/adapters/cuda/CMakeLists.txt
+index a8401c71a..2ab363926 100644
+--- a/unified-runtime/source/adapters/cuda/CMakeLists.txt
++++ b/unified-runtime/source/adapters/cuda/CMakeLists.txt
+@@ -62,14 +62,14 @@ add_library(cudadrv SHARED IMPORTED GLOBAL)
+ if (WIN32)
+   set_target_properties(
+     cudadrv PROPERTIES 
+-      IMPORTED_IMPLIB               ${CUDA_cuda_driver_LIBRARY}
+-      INTERFACE_INCLUDE_DIRECTORIES ${CUDAToolkit_INCLUDE_DIRS}
++      IMPORTED_IMPLIB               "${CUDA_cuda_driver_LIBRARY}"
++      INTERFACE_INCLUDE_DIRECTORIES "${CUDAToolkit_INCLUDE_DIRS}"
+   )
+ else()
+   set_target_properties(
+     cudadrv PROPERTIES 
+-      IMPORTED_LOCATION             ${CUDA_cuda_driver_LIBRARY}
+-      INTERFACE_INCLUDE_DIRECTORIES ${CUDAToolkit_INCLUDE_DIRS}
++      IMPORTED_LOCATION             "${CUDA_cuda_driver_LIBRARY}"
++      INTERFACE_INCLUDE_DIRECTORIES "${CUDAToolkit_INCLUDE_DIRS}"
+   )
+ endif()
+ 

pkgs/llvm-intel/wrapped.nix

@@ -0,0 +1,60 @@
+{
+  stdenv,
+  wrapCCWith,
+  llvm-intel-unwrapped,
+  gcc,
+  libgcc,
+}:
+
+let
+  wrapIntel =
+    cc:
+    let
+      targetConfig = stdenv.targetPlatform.config;
+    in
+    wrapCCWith {
+      inherit cc;
+
+      nixSupport = {
+        cc-ldflags = [
+          "-L${gcc.cc}/lib/gcc/${targetConfig}/${gcc.version}"
+          "-L${gcc.cc.lib}/lib"
+          "-L${libgcc.out}/lib"
+          "-L${cc}/lib"
+        ];
+        cc-cflags = [
+          "--gcc-toolchain=${gcc.cc}"
+          "-isystem \"${cc}/lib/clang/${cc.llvmMajorVersion}/include\""
+          "-isystem ${cc}/include"
+          "-isystem ${cc}/include/intel64"
+          "-isystem ${gcc.cc}/lib/gcc/${targetConfig}/${gcc.version}/include"
+        ];
+        libcxx-cxxflags = [
+          # "--gcc-toolchain=${gcc.cc}"
+          "-isystem ${gcc.cc}/include/c++/${gcc.version}"
+          "-isystem ${gcc.cc}/include/c++/${gcc.version}/${targetConfig}"
+        ];
+      };
+
+      extraBuildCommands = ''
+        # FIXME: We should find a better way to modify the PATH instead of using
+        # this ugly hack. See https://jungle.bsc.es/git/rarias/bscpkgs/issues/9
+        echo 'path_backup="${gcc.cc}/bin:$path_backup"' >>$out/nix-support/cc-wrapper-hook
+
+        # Disable hardening by default
+        echo "" > $out/nix-support/add-hardening.sh
+
+        ln -s $ccPath/sycl-ls $out/bin/sycl-ls
+        ln -s $ccPath/sycl-trace $out/bin/sycl-trace
+
+        # Use this to detect when a compiler subprocess is called
+        # from icpx (--fsycl-host-compiler)
+        echo 'export NIX_CC_WRAPPER_INTEL=1' >>$out/nix-support/cc-wrapper-hook
+
+        # oneMath looks for sycl libraries in bin/../lib
+        ln -s ${cc}/lib $out/lib
+        ln -s ${cc}/include $out/include
+      '';
+    };
+in
+wrapIntel llvm-intel-unwrapped

pkgs/pocl/0001-cmake-do-not-use-suffix.patch

@@ -0,0 +1,26 @@
+From ccf301659caac9b5e973ba1f2d32352acf617a98 Mon Sep 17 00:00:00 2001
+From: =?UTF-8?q?Aleix=20Bon=C3=A9?= <aleix.boneribo@bsc.es>
+Date: Tue, 2 Jul 2024 16:45:05 +0200
+Subject: [PATCH] cmake: do not use suffix
+
+---
+ cmake/LLVM.cmake                              |  3 ++
+ 1 files changed, 3 insertions(+)
+
+diff --git a/cmake/LLVM.cmake b/cmake/LLVM.cmake
+index f4dbda065..e29144dce 100644
+--- a/cmake/LLVM.cmake
++++ b/cmake/LLVM.cmake
+@@ -65,6 +65,9 @@ else()
+     message(WARNING "Cannot determine llvm binary suffix from ${LLVM_CONFIG}")
+   endif()
+   message(STATUS "LLVM binaries suffix : ${LLVM_BINARY_SUFFIX}")
++
++  # We don't want suffixes in nix
++  set(LLVM_BINARY_SUFFIX "")
+ endif()
+ 
+ get_filename_component(LLVM_CONFIG_LOCATION "${LLVM_CONFIG}" DIRECTORY)
+-- 
+2.45.1
+

pkgs/pocl/0001-cmake-native-build-tools.patch

@@ -0,0 +1,33 @@
+From f24b456c50f587b05cc8f2699c94d8cdefc5b13e Mon Sep 17 00:00:00 2001
+From: =?UTF-8?q?Aleix=20Bon=C3=A9?= <aleix.boneribo@bsc.es>
+Date: Tue, 2 Jul 2024 17:11:11 +0200
+Subject: [PATCH] cmake: native build tools
+
+---
+ cmake/LLVM.cmake                    |   7 +-
+ 1 files changed, 6 insertions(+), 1 deletion(-)
+
+diff --git a/cmake/LLVM.cmake b/cmake/LLVM.cmake
+index e29144dce..b9f14ce6a 100644
+--- a/cmake/LLVM.cmake
++++ b/cmake/LLVM.cmake
+@@ -269,10 +269,15 @@ endforeach()
+ 
+ ####################################################################
+ 
++if(DEFINED LLVM_NATIVE_TOOL_DIR)
++  set(TOOL_DIR "${LLVM_NATIVE_TOOL_DIR}")
++  message(STATUS "TOOL_DIR=${TOOL_DIR}")
++endif()
++
+ macro(find_program_or_die OUTPUT_VAR PROG_NAME DOCSTRING)
+   find_program(${OUTPUT_VAR}
+     NAMES "${PROG_NAME}${LLVM_BINARY_SUFFIX}${CMAKE_EXECUTABLE_SUFFIX}"
+-    HINTS "${LLVM_BINDIR}" "${LLVM_CONFIG_LOCATION}"
++    HINTS "${TOOL_DIR}" "${LLVM_BINDIR}" "${LLVM_CONFIG_LOCATION}"
+     DOC "${DOCSTRING}"
+     NO_CMAKE_PATH
+     NO_CMAKE_ENVIRONMENT_PATH
+-- 
+2.45.1
+

pkgs/pocl/0001-cmake-use-clang-from-cmake-variable.patch

@@ -0,0 +1,48 @@
+From 4688b5ce895761c884ae15fc0234ed8d623b988b Mon Sep 17 00:00:00 2001
+From: =?UTF-8?q?Aleix=20Bon=C3=A9?= <aleix.boneribo@bsc.es>
+Date: Tue, 2 Jul 2024 15:41:47 +0200
+Subject: [PATCH] cmake: use clang from cmake variable
+
+---
+ cmake/LLVM.cmake | 22 ++++++++++++++++++++--
+ 1 file changed, 20 insertions(+), 2 deletions(-)
+
+diff --git a/cmake/LLVM.cmake b/cmake/LLVM.cmake
+index 71b786dc7..f4dbda065 100644
+--- a/cmake/LLVM.cmake
++++ b/cmake/LLVM.cmake
+@@ -281,11 +281,29 @@ macro(find_program_or_die OUTPUT_VAR PROG_NAME DOCSTRING)
+   endif()
+ endmacro()
+ 
+-find_program_or_die( CLANG "clang" "clang binary")
++# -DCMAKE_C_COMPILER=riscv64-unknown-linux-gnu-clang -DCMAKE_CXX_COMPILER=riscv64-unknown-linux-gnu-clang++
++
++# if(DEFINED CLANG_EXECUTABLE)
++if(DEFINED CMAKE_C_COMPILER)
++  # set(CLANG "${CLANG_EXECUTABLE}")
++  set(CLANG "${CMAKE_C_COMPILER}")
++  message(STATUS "Using CLANG_EXECUTABLE: ${CLANG}")
++else()
++  find_program_or_die( CLANG "clang" "clang binary")
++endif()
+ execute_process(COMMAND "${CLANG}" "--version" OUTPUT_VARIABLE LLVM_CLANG_VERSION RESULT_VARIABLE CLANG_RES)
++
+ # TODO this should be optional
+-find_program_or_die( CLANGXX "clang++" "clang++ binary")
++# if(DEFINED CLANGXX_EXECUTABLE)
++if(DEFINED CMAKE_CXX_COMPILER)
++  # set(CLANGXX "${CLANGXX_EXECUTABLE}")
++  set(CLANGXX "${CMAKE_CXX_COMPILER}")
++  message(STATUS "Using CLANGXX_EXECUTABLE: ${CLANGXX}")
++else()
++  find_program_or_die( CLANGXX "clang++" "clang++ binary")
++endif()
+ execute_process(COMMAND "${CLANGXX}" "--version" OUTPUT_VARIABLE LLVM_CLANGXX_VERSION RESULT_VARIABLE CLANGXX_RES)
++
+ if(CLANGXX_RES OR CLANG_RES)
+   message(FATAL_ERROR "Failed running clang/clang++ --version")
+ endif()
+-- 
+2.45.1
+

pkgs/pocl/default.nix

@@ -0,0 +1,255 @@
+{
+  lib,
+  llvmPackages,
+  ninja,
+  cmake,
+  hwloc,
+  ocl-icd,
+  opencl-headers,
+  getconf,
+  pkg-config,
+  spirv-llvm-translator,
+  spirv-tools,
+  libxml2, # required for statically linked llvm
+  python3,
+  writableTmpDirAsHomeHook,
+  writeText,
+  runCommand,
+  opencl-cts,
+
+  fetchFromGitHub,
+
+  useGit ? false,
+  gitUrl ? "git@gitlab-internal.bsc.es:task-awareness/pocl/pocl.git",
+  gitBranch ? "feat/nosv",
+  gitCommit ? "556542309fc3c5900cf05dae5b41d200cd3de6a0",
+
+  staticLLVM ? true,
+  enableICD ? true,
+  enableSPIRV ? true,
+  enableHWLOC ? true,
+  enableRemoteServer ? false,
+  enableRemoteClient ? false,
+  enableDistroVariants ? false,
+
+  lttng-ust,
+  enableLTTNG ? false,
+
+  onetbb,
+  enableTBB ? false,
+
+  nosv ? null,
+  ovni ? null,
+  enableNOSV ? useGit,
+  enableOVNI ? useGit,
+
+  buildPackages,
+  targetPackages,
+  nix-update-script,
+}:
+
+assert (enableNOSV || enableOVNI) -> useGit;
+
+let
+
+  release = rec {
+    version = "7.1";
+    src = fetchFromGitHub {
+      owner = "pocl";
+      repo = "pocl";
+      tag = "v${version}";
+      hash = "sha256-bS6vTIjLO7YLs7qYLKW0cYYbEJ/hRS/+IjjAKbkj8ac=";
+    };
+  };
+
+  git = rec {
+    version = src.shortRev;
+    src = builtins.fetchGit {
+      url = gitUrl;
+      ref = gitBranch;
+      rev = gitCommit;
+    };
+  };
+
+  source = if (useGit) then git else release;
+
+  stdenv = llvmPackages.stdenv;
+
+  buildLlvmPackages = buildPackages.llvmPackages;
+
+  isCross = stdenv.hostPlatform != stdenv.buildPlatform;
+
+  featureList = lib.concatStringsSep "+" (
+    (lib.optionals enableICD [ "icd" ])
+    ++ (lib.optionals enableNOSV [ "nosv" ])
+    ++ (lib.optionals enableOVNI [ "ovni" ])
+    ++ (lib.optionals enableSPIRV [ "spirv" ])
+    ++ (lib.optionals enableTBB [ "tbb" ])
+    ++ (lib.optionals enableRemoteClient [ "server" ])
+    ++ (lib.optionals enableRemoteServer [ "client" ])
+  );
+in
+
+stdenv.mkDerivation (finalAttrs: {
+  pname = "pocl";
+  version = featureList + "-" + source.version;
+  inherit (source) src;
+
+  patches = lib.optionals useGit [
+    ./0001-cmake-use-clang-from-cmake-variable.patch
+    ./0001-cmake-do-not-use-suffix.patch
+    ./0001-cmake-native-build-tools.patch
+  ];
+
+  cmakeFlags = [
+    # TODO: all these are broken when cross compiling. Upstream has refactored
+    # all the cmake infra for cross compilation, but it's not in a release yet
+    (lib.cmakeOptionType "filepath" "WITH_LLVM_CONFIG" (
+      lib.getExe' buildLlvmPackages.llvm.dev "llvm-config"
+    ))
+    (lib.cmakeOptionType "filepath" "CLANG" (lib.getExe' buildLlvmPackages.clangUseLLVM "clang"))
+    (lib.cmakeOptionType "filepath" "CLANGXX" (lib.getExe' buildLlvmPackages.clangUseLLVM "clang++"))
+    (lib.cmakeOptionType "path" "CLANG_RESOURCE_DIR" "${llvmPackages.stdenv.cc}/resource-root")
+
+    (lib.cmakeBool "ENABLE_TESTS" finalAttrs.finalPackage.doCheck)
+
+    (lib.cmakeBool "ENABLE_RELOCATION" true)
+    (lib.cmakeBool "ENABLE_POCL_BUILDING" true)
+    (lib.cmakeBool "POCL_ICD_ABSOLUTE_PATH" true)
+
+    (lib.cmakeBool "ENABLE_TBB_DEVICE" enableTBB)
+    (lib.cmakeBool "ENABLE_ICD" enableICD)
+    (lib.cmakeBool "ENABLE_HWLOC" enableHWLOC)
+    (lib.cmakeBool "ENABLE_LTTNG" enableLTTNG)
+
+    (lib.cmakeBool "ENABLE_REMOTE_CLIENT" enableRemoteClient)
+    (lib.cmakeBool "ENABLE_REMOTE_SERVER" enableRemoteServer)
+
+    # avoid the runtime linker pulling in a different llvm e.g. from graphics drivers
+    (lib.cmakeBool "STATIC_LLVM" staticLLVM)
+  ]
+  ++ lib.optionals (enableNOSV || enableOVNI) [
+    (lib.cmakeBool "ENABLE_NOSV" enableNOSV)
+    (lib.cmakeBool "ENABLE_OVNI" enableOVNI)
+  ]
+  ++ lib.optionals (!isCross && enableDistroVariants) [
+    (lib.cmakeFeature "KERNELLIB_HOST_CPU_VARIANTS" "distro") # TODO: check if we can do it when cross compiling
+  ]
+  ++ lib.optionals isCross [
+    (lib.cmakeFeature "CMAKE_SYSTEM_NAME" "Linux")
+    (lib.cmakeFeature "CMAKE_FIND_ROOT_PATH_MODE_PROGRAM" "NEVER")
+    (lib.cmakeFeature "CMAKE_FIND_ROOT_PATH_MODE_LIBRARY" "ONLY")
+    (lib.cmakeFeature "CMAKE_FIND_ROOT_PATH_MODE_INCLUDE" "ONLY")
+    (lib.cmakeFeature "CMAKE_FIND_ROOT_PATH_MODE_PACKAGE" "ONLY")
+  ]
+  ++ lib.optionals stdenv.hostPlatform.isRiscV [
+    (lib.cmakeFeature "LLC_TRIPLE" "riscv64-unknown-linux-gnu")
+    (lib.cmakeFeature "LLC_HOST_CPU" "rv64g")
+    (lib.cmakeFeature "CLANG_MARCH_FLAG" "-march=")
+  ];
+
+  # Fixes error: ld: cannot find -lgcc / -lgcc_s when linking from libpocl on
+  # host
+  preConfigure = ''
+    cmakeFlagsArray+=(
+      -DEXTRA_HOST_LD_FLAGS="${lib.escapeShellArg (builtins.readFile "${targetPackages.stdenv.cc}/nix-support/cc-ldflags")}"
+      -DEXTRA_KERNEL_FLAGS="${lib.escapeShellArg (builtins.readFile "${targetPackages.stdenv.cc}/nix-support/cc-ldflags")}"
+    )
+  '';
+
+  nativeBuildInputs = [
+    cmake
+    getconf
+    ninja
+    pkg-config
+    python3
+  ];
+
+  buildInputs = [
+    opencl-headers
+    llvmPackages.libclang
+    llvmPackages.llvm
+  ]
+  ++ lib.optionals staticLLVM [ libxml2 ]
+  ++ lib.optionals enableNOSV [ nosv ]
+  ++ lib.optionals enableOVNI [ ovni ]
+  ++ lib.optionals enableHWLOC [ hwloc ]
+  ++ lib.optionals enableTBB [ onetbb ]
+  ++ lib.optionals enableICD [ ocl-icd ]
+  ++ lib.optionals enableLTTNG [ lttng-ust ]
+  ++ lib.optionals enableSPIRV [
+    (spirv-llvm-translator.override { inherit (llvmPackages) llvm; })
+    spirv-tools
+  ];
+
+  nativeInstallCheckInputs = [
+    writableTmpDirAsHomeHook # needed for POCL_CACHE_DIR
+  ];
+
+  doInstallCheck = true;
+
+  installCheckPhase = lib.optionalString enableICD ''
+    runHook preInstallCheck
+
+    export OCL_ICD_VENDORS=$out/etc/OpenCL/vendors
+    $out/bin/poclcc -o poclcc.cl.pocl $src/examples/poclcc/poclcc.cl
+
+    runHook postInstallCheck
+  '';
+
+  setupHook = lib.optionalDrvAttr enableICD (
+    writeText "setup-hook" ''
+      addToSearchPath OCL_ICD_VENDORS @out@/etc/OpenCL/vendors
+    ''
+  );
+
+  passthru =
+    let
+      self = finalAttrs.finalPackage;
+    in
+    {
+      updateScript = nix-update-script { };
+      git = self.override { useGit = true; };
+
+      withICD = enableICD;
+
+      test = {
+        # Run builtin tests (causes a rebuild)
+        self = self.overrideAttrs { doCheck = true; };
+
+        # Run basic test from Khronos OpenCL Conformance Test Suite
+        # WARN: despite its name, test_basic is very exhaustive, and can take more
+        # than 1 hour in a marenostrum5 node.
+        cts =
+          assert enableICD;
+          lib.genAttrs [ "api" "basic" "c11_atomics" "printf" "svm" ] (
+            name:
+            runCommand "pocl-cts-test_${name}" {
+              nativeBuildInputs = [ self ];
+              requiredSystemFeatures = [ "sys-devices" ];
+              env = {
+                POCL_DEBUG = "error,warn";
+                POCL_CACHE_DIR = "/build/pocl_cache";
+              };
+            } "${opencl-cts}/bin/Release/test_${name} | tee $out"
+          );
+      };
+    };
+
+  env.HWLOC_SYNTHETIC = lib.optionalDrvAttr enableHWLOC "node:1 core:1 pu:1";
+
+  propagatedBuildInputs = [ stdenv.cc.cc ];
+
+  meta = with lib; {
+    description = "A portable open source (MIT-licensed) implementation of the OpenCL standard";
+    homepage = "http://portablecl.org";
+    license = licenses.mit;
+    maintainers = with maintainers; [
+      leixb
+      jansol
+      xddxdd
+      bsc.maintainers.abonerib
+    ];
+    platforms = platforms.linux ++ platforms.darwin;
+  };
+})

pkgs/rodinia/default.nix

@@ -0,0 +1,58 @@
+{
+  stdenv,
+  lib,
+  fetchurl,
+  ocl-icd,
+  opencl-headers,
+}:
+
+stdenv.mkDerivation {
+  pname = "rodinia-opencl";
+  version = "3.1";
+  src = fetchurl {
+    url = "http://www.cs.virginia.edu/~skadron/lava/Rodinia/Packages/rodinia_3.1.tar.bz2";
+    hash = "sha256-+uusfBHtj4/Pa/LX6FwwhvwtEfciBNbfwo3FsujyrP0=";
+  };
+
+  patches = [
+    ./rodinia-fixes.patch
+    ./rodinia-rm.patch
+  ];
+  buildInputs = [
+    opencl-headers
+    ocl-icd
+  ];
+  buildFlags = [
+    "OPENCL"
+    "--always-make"
+  ];
+
+  preBuild = ''
+    mkdir -p bin/linux/{cuda,omp,opencl}
+  '';
+
+  installPhase = ''
+    runHook preInstall
+
+    mkdir -p $out/share
+    cp -r bin data opencl $out
+    cp LICENSE README $out/share
+
+    runHook postInstall
+  '';
+
+  meta = {
+    homepage = "https://rodinia.cs.virginia.edu/";
+    description = "Test Suite of Compute-Intensive Applications with Accelerators";
+    maintainers = with lib.maintainers.bsc; [ abonerib ];
+    platforms = lib.platforms.linux;
+    license = {
+      deprecated = false;
+      free = false;
+      fullName = "Rodinia License";
+      redistributable = true;
+      shortName = "rodinia";
+      url = "https://www.cs.virginia.edu/~skadron/lava/rodinia/license.htm";
+    };
+  };
+}

pkgs/unified-runtime/default.nix

@@ -0,0 +1,63 @@
+{
+  lib,
+  stdenv,
+  fetchFromGitHub,
+  cmake,
+  python3,
+  cudatoolkit,
+  libbacktrace,
+  intel-compute-runtime,
+  level-zero,
+  unified-memory-framework,
+  autoAddDriverRunpath,
+
+  enableL0 ? false,
+}:
+
+stdenv.mkDerivation rec {
+  pname = "unified-runtime";
+  version = "2025-WW45";
+
+  src = fetchFromGitHub {
+    owner = "intel";
+    repo = "llvm";
+    tag = version;
+    sha256 = "sha256-yf8pFEeNEWf9LpEhAy4vMdFDlA08x5XaCegnq+5aXRw=";
+  };
+
+  sourceRoot = "${src.name}/unified-runtime";
+
+  nativeBuildInputs = [
+    cmake
+    python3
+    libbacktrace
+    (python3.withPackages (
+      python-pkgs: with python-pkgs; [
+        lit
+        filecheck
+      ]
+    ))
+    autoAddDriverRunpath
+  ];
+
+  buildInputs = [
+    unified-memory-framework
+    cudatoolkit
+  ]
+  ++ lib.optionals enableL0 [
+    intel-compute-runtime
+    level-zero
+  ];
+
+  cmakeFlags = [
+    (lib.cmakeBool "UR_USE_EXTERNAL_UMF" true)
+    (lib.cmakeBool "UR_BUILD_TESTS" false)
+
+    (lib.cmakeBool "UR_BUILD_ADAPTER_CUDA" true)
+    (lib.cmakeBool "UR_BUILD_ADAPTER_L0" enableL0)
+
+    (lib.cmakeBool "UR_BUILD_ADAPTER_OPENCL" false)
+    (lib.cmakeBool "UR_BUILD_ADAPTER_HIP" false)
+    (lib.cmakeBool "UR_BUILD_ADAPTER_NATIVE_CPU" false)
+  ];
+}

test/compilers/hello-sycl.nix

@@ -2,6 +2,7 @@
   intelPackages,
   writeText,
   strace,
+  pocl,
 }:
 
 let
@@ -35,10 +36,10 @@ let
   '';
 in
 
-stdenv.mkDerivation {
+stdenv.mkDerivation (finalAttrs: {
   version = "0.0.1";
   name = "hello-sycl";
-  buildInputs = [
+  nativeBuildInputs = [
     stdenv
     strace
   ];
@@ -52,11 +53,39 @@ stdenv.mkDerivation {
     echo CXX=$CXX
     command -v $CXX
     $CXX -fsycl hello.cpp -o hello
-    ./hello
+    ./hello | tee test-output
     set +x
   '';
 
   installPhase = ''
-    touch $out
+    cp test-output $out
   '';
-}
+
+  passthru = {
+    withPocl = finalAttrs.finalPackage.overrideAttrs (old: {
+      nativeBuildInputs = old.nativeBuildInputs ++ [ pocl ];
+      env = (old.env or { }) // {
+        POCL_DEBUG = "error,warn";
+        POCL_CACHE_DIR = "/build/pocl_cache";
+
+        # Make PoCL report Intel vendor id so oneapi SYCL works
+        POCL_DRIVER_VERSION_OVERRIDE = "2024.18.6.0.02_160000";
+        POCL_CPU_VENDOR_ID_OVERRIDE = 32902;
+      };
+      doInstallCheck = true;
+      installCheckPhase = ''
+        grep "Hello World! (on device)" $out
+      '';
+    });
+    withIntel = finalAttrs.finalPackage.overrideAttrs (old: {
+      env = (old.env or { }) // {
+        OCL_ICD_VENDORS = intelPackages.compiler + "/etc/OpenCL/vendors";
+      };
+      doInstallCheck = true;
+      installCheckPhase = ''
+        grep "Hello World! (on device)" $out
+      '';
+
+    });
+  };
+})

test/compilers/icpx-ompss2.nix

@@ -5,6 +5,7 @@
   nosv,
   clangOmpss2Nodes,
   strace,
+  pocl,
 }:
 
 let
@@ -36,7 +37,7 @@ let
   '';
 in
 
-intelPackages.stdenv.mkDerivation {
+intelPackages.stdenv.mkDerivation (finalAttrs: {
   version = "0.0.1";
   name = "hello-syclompss";
 
@@ -48,6 +49,34 @@ intelPackages.stdenv.mkDerivation {
     nosv
   ];
 
+  passthru = {
+    withPocl = finalAttrs.finalPackage.overrideAttrs (old: {
+      nativeBuildInputs = old.nativeBuildInputs ++ [ pocl ];
+      env = (old.env or { }) // {
+        POCL_DEBUG = "error,warn";
+        POCL_CACHE_DIR = "/build/pocl_cache";
+
+        # Make PoCL report Intel vendor id so oneapi SYCL works
+        POCL_DRIVER_VERSION_OVERRIDE = "2024.18.6.0.02_160000";
+        POCL_CPU_VENDOR_ID_OVERRIDE = 32902;
+      };
+      doInstallCheck = true;
+      installCheckPhase = ''
+        grep "Hello World! (on device)" $out
+      '';
+    });
+
+    withIntel = finalAttrs.finalPackage.overrideAttrs (old: {
+      env = (old.env or { }) // {
+        OCL_ICD_VENDORS = intelPackages.compiler + "/etc/OpenCL/vendors";
+      };
+      doInstallCheck = true;
+      installCheckPhase = ''
+        grep "Hello World! (on device)" $out
+      '';
+    });
+  };
+
   dontUnpack = true;
   dontConfigure = true;
 
@@ -71,11 +100,11 @@ intelPackages.stdenv.mkDerivation {
         $NODES_HOME/lib/nodes-main-wrapper.o \
         hello.cpp -o hello
 
-    ./hello
+    ./hello | tee test-output
     set +x
   '';
 
   installPhase = ''
-    touch $out
+    cp test-output $out
   '';
-}
+})

test/compilers/opencl.c

@@ -0,0 +1,127 @@
+#include <CL/cl.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+#define BUF_LEN 128
+
+cl_int exit_err = CL_SUCCESS;
+
+#define CHECK(cmd)                                                             \
+  do {                                                                         \
+    cl_int err = cmd;                                                          \
+    if (err != CL_SUCCESS) {                                                   \
+      printf("[ERROR] " #cmd " (Error code: %d)\n"                             \
+             "@ "__FILE__                                                      \
+             ":%d",                                                            \
+             err, __LINE__);                                                   \
+      exit_err = err;                                                          \
+      goto cleanup;                                                            \
+    }                                                                          \
+  } while (0)
+
+const char *kernel_source = "__kernel void vector_add(__global const float *a, "
+                            "__global const float *b, __global float *c) {\n"
+                            "    int gid = get_global_id(0);\n"
+                            "    c[gid] = a[gid] + b[gid];\n"
+                            "}\n";
+
+cl_int test_kernel_compilation(cl_platform_id platform) {
+  cl_device_id device = NULL;
+  cl_context context = NULL;
+  cl_program program = NULL;
+  cl_int err = CL_SUCCESS;
+  char device_name[BUF_LEN];
+
+  // Get first device for this platform
+  CHECK(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &device, NULL));
+
+  // Get device name
+  CHECK(clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(device_name),
+                        device_name, NULL));
+  printf("    Device: %s\n", device_name);
+
+  context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
+  CHECK(err);
+
+  program = clCreateProgramWithSource(context, 1, &kernel_source, NULL, &err);
+  CHECK(err);
+
+  err = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
+  if (err != CL_SUCCESS) {
+    printf("      [ERROR] Kernel compilation failed (Error code: %d)\n", err);
+
+    // Build log
+    size_t log_size;
+    CHECK(clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL,
+                                &log_size));
+    char *log = (char *)malloc(log_size);
+    CHECK(clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size,
+                                log, NULL));
+    printf("      Build log:\n%s\n", log);
+    free(log);
+    goto cleanup;
+  }
+
+  printf("      [OK] Kernel compiled successfully!\n");
+
+cleanup:
+  if (program)
+    clReleaseProgram(program);
+  if (context)
+    clReleaseContext(context);
+  printf("\n");
+
+  return exit_err;
+}
+
+int main() {
+  cl_uint num_platforms;
+  cl_platform_id *platforms = NULL;
+
+  // Get number of platforms
+  CHECK(clGetPlatformIDs(0, NULL, &num_platforms));
+
+  printf("Found %d OpenCL platform(s)\n\n", num_platforms);
+
+  if (num_platforms == 0) {
+    printf("No OpenCL platforms found!\n");
+    return EXIT_FAILURE;
+  }
+
+  // Allocate memory for platforms
+  platforms = (cl_platform_id *)malloc(sizeof(cl_platform_id) * num_platforms);
+
+  // Get platform IDs
+  CHECK(clGetPlatformIDs(num_platforms, platforms, NULL));
+
+  // Query each platform and test kernel compilation
+  for (cl_uint i = 0; i < num_platforms; i++) {
+    char platform_name[BUF_LEN];
+    char platform_vendor[BUF_LEN];
+    char platform_version[BUF_LEN];
+
+    CHECK(clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME,
+                            sizeof(platform_name), platform_name, NULL));
+    CHECK(clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR,
+                            sizeof(platform_vendor), platform_vendor, NULL));
+    CHECK(clGetPlatformInfo(platforms[i], CL_PLATFORM_VERSION,
+                            sizeof(platform_version), platform_version, NULL));
+
+    printf("Platform %d: %s\n", i, platform_name);
+    printf("  Vendor: %s\n", platform_vendor);
+    printf("  Version: %s\n", platform_version);
+
+    // Test kernel compilation
+    CHECK(test_kernel_compilation(platforms[i]));
+  }
+
+cleanup:
+  if (platforms)
+    free(platforms);
+
+  if (exit_err == CL_SUCCESS)
+    return EXIT_SUCCESS;
+
+  return EXIT_FAILURE;
+}

test/compilers/opencl.nix

@@ -0,0 +1,104 @@
+{
+  stdenv,
+  ocl-icd,
+  opencl-headers,
+  pocl,
+  clinfo,
+  runCommand,
+  lib,
+}:
+
+stdenv.mkDerivation (finalAttrs: {
+  pname = "opencl-pocl-test";
+  version = "1.0.0";
+
+  src = ./opencl.c;
+
+  dontUnpack = true;
+  dontConfigure = true;
+
+  buildInputs = [
+    ocl-icd
+    opencl-headers
+  ];
+
+  buildPhase = ''
+    runHook preBuild
+
+    $CC -Wall -Wextra -D CL_TARGET_OPENCL_VERSION=100 $src -o ocl-hello -lOpenCL
+
+    runHook postBuild
+  '';
+
+  installPhase = ''
+    runHook preInstall
+
+    mkdir -p $out/bin
+    cp ocl-hello $out/bin
+
+    runHook postInstall
+  '';
+
+  passthru = {
+    test-icd =
+      runCommand "custom-clinfo"
+        {
+          nativeBuildInputs =
+            assert !(lib.hasPrefix "pocl" ocl-icd.pname);
+            [
+              clinfo
+              finalAttrs.finalPackage
+              pocl
+            ];
+
+          env = {
+            POCL_DEBUG = "error,warn";
+            POCL_CACHE_DIR = "/build/pocl_cache";
+          };
+
+          requiredSystemFeatures = [ "sys-devices" ];
+        }
+
+        ''
+          set -x
+
+          mkdir $out
+
+          clinfo -l >$out/clinfo
+          clinfo >$out/clinfo-full
+
+          ocl-hello >$out/hello
+
+          for i in $out/*; do
+            grep "Portable Computing Language" $i
+          done
+
+          echo $OCL_ICD_VENDORS >$out/env
+
+          grep "pocl.*/etc/OpenCL/vendors" $out/env
+
+          set +x
+        '';
+    # this needs: ocl-icd = pocl-noicd
+    test-noicd =
+      runCommand "custom-clinfo"
+        {
+          nativeBuildInputs = [
+            finalAttrs.finalPackage
+          ];
+
+          env = {
+            POCL_DEBUG = "error,warn";
+            POCL_CACHE_DIR = "/build/pocl_cache";
+          };
+
+          requiredSystemFeatures = [ "sys-devices" ];
+        }
+
+        ''
+          ocl-hello >$out
+
+          grep "Portable Computing Language" $out
+        '';
+  };
+})

test/compilers/pocl.nix

@@ -0,0 +1,27 @@
+{
+  runCommand,
+  pocl,
+  clinfo,
+}:
+
+runCommand "clinfo-pocl"
+  {
+    nativeBuildInputs = [
+      clinfo
+      pocl
+    ];
+
+    requiredSystemFeatures = [ "sys-devices" ];
+
+    env = {
+      POCL_DEBUG = "error,warn";
+      POCL_CACHE_DIR = "/build/pocl_cache";
+    };
+  }
+  ''
+    echo $OCL_ICD_VENDORS
+    clinfo >$out
+
+    # check that we have PoCL:
+    grep "Portable Computing Language" $out
+  ''

test/compilers/sycl-ls.nix

@@ -0,0 +1,40 @@
+{
+  intelPackages,
+  pocl,
+  runCommandWith,
+  symlinkJoin,
+  clinfo,
+}:
+
+runCommandWith
+  {
+    name = "sycl-ls";
+    inherit (intelPackages) stdenv;
+    runLocal = true;
+    derivationArgs = {
+      nativeBuildInputs = [ clinfo ];
+      env = {
+        OCL_ICD_VENDORS = symlinkJoin {
+          name = "vendors";
+          paths = builtins.map (p: p + "/etc/OpenCL/vendors") [
+            intelPackages.compiler
+            pocl
+          ];
+        };
+
+        POCL_DRIVER_VERSION_OVERRIDE = "2025.20.6.0.04_224945";
+        POCL_CPU_VENDOR_ID_OVERRIDE = 32902;
+        POCL_DEBUG = "error,warn";
+        POCL_CACHE_DIR = "/build/pocl_cache";
+
+        HWLOC_SYNTHETIC = "node:1 core:1 pu:1";
+      };
+    };
+  }
+  # bash
+  ''
+    clinfo -l | tee -a $out
+    sycl-ls | tee -a $out
+
+    test x6 = x"$(cat $out | wc -l)" # 2 lines per clinfo entry and 1 per sycl-ls
+  ''