Fix build errors and improve CI config

* Fix -Wmissing-template-arg-list-after-template-kw compilation errors in nbody and fluid demos.
* Fix sycl-ls regex in ConfigureSYCL.cmake following the removal of "ext_oneapi" from backend names.
* Update ConfigureSYCL.cmake default CUDA and HIP arch to the oldest supported by oneAPI 2025.0.
* Improve OpenMP CMake test in matrix_multiply_omp_compare sample. Previously it assumed all clang++ drivers support -fopenmp. This is only true in the oneAPI base toolkit version, but the open-source DPC++ nightly releases are built without libomp.so and linking with -fopenmp fails. Use try_compile to figure out if -fopenmp works.
* Add a warning suppresion file to suppress CI build failures in external code from submodules.
* Update the CI container with newer versions of CUDA and ROCm, and use a recent DPC++ open-source nightly instead of a oneAPI Toolkit release.
* Restore CI workflow's permission to push containers to registry.

Author: rafbiels

.github/workflows/ci.yml

@@ -4,6 +4,7 @@ on: [push, pull_request]
 
 permissions:
   contents: read
+  packages: write
 
 env:
   REGISTRY: ghcr.io
@@ -108,7 +109,7 @@ jobs:
           -DENABLE_SPIR=ON
           -DENABLE_CUDA=ON -DCUDA_COMPUTE_CAPABILITY=80
           -DENABLE_HIP=ON -DHIP_GFX_ARCH=gfx90a
-          -DCMAKE_CXX_FLAGS='-Wall -Wextra -Wpedantic -Werror'
+          -DCMAKE_CXX_FLAGS="-Wall -Wextra -Wpedantic -Werror --warning-suppression-mappings=${PWD}/ci/warning-suppresions.txt"
           -G Ninja
 
       - name: Build

ci/Dockerfile

@@ -1,44 +1,36 @@
 FROM ubuntu:24.04@sha256:3f85b7caad41a95462cf5b787d8a04604c8262cdcdf9a472b8c52ef83375fe15
 
 # Get basic dependencies from Ubuntu repositories
-RUN apt update && apt -y install wget gpg git cmake ninja-build \
-    gcc-12 libstdc++-12-dev libsdl2-dev \
+ARG DEBIAN_FRONTEND=noninteractive
+RUN apt update \
+    && apt -y install wget gpg git cmake ninja-build g++ libsdl2-dev \
     && apt clean
 
 # Install nvcc (dependency for compiling for a CUDA target)
-ARG CUDA_VERSION=12-4
+ARG CUDA_VERSION=12-8
 RUN wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.0-1_all.deb \
     && dpkg -i cuda-keyring_1.0-1_all.deb && rm cuda-keyring_1.0-1_all.deb \
     && apt update && apt -y install cuda-nvcc-${CUDA_VERSION} && apt clean
 
 # Install ROCm device libs (dependency for compiling for a HIP target)
-ARG ROCM_VERSION=5.4.3
+ARG ROCM_VERSION=6.2.4
 RUN wget https://repo.radeon.com/rocm/rocm.gpg.key -O - \
     | gpg --dearmor | tee /etc/apt/keyrings/rocm.gpg > /dev/null \
     && echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/rocm.gpg] https://repo.radeon.com/rocm/apt/${ROCM_VERSION} jammy main" \
     | tee /etc/apt/sources.list.d/rocm.list \
-    && apt update && apt -y install rocm-device-libs && apt clean
+    && apt update && apt -y install rocm-device-libs${ROCM_VERSION} && apt clean
 
-# Install DPC++ and remove parts we don't need to reduce the container size
-ARG ONEAPI_VERSION=2024.1
-RUN wget -O- https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB \
-    | gpg --dearmor | tee /usr/share/keyrings/oneapi-archive-keyring.gpg > /dev/null \
-    && echo "deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main" \
-    | tee /etc/apt/sources.list.d/oneAPI.list \
-    && apt update && apt -y install intel-oneapi-compiler-dpcpp-cpp-${ONEAPI_VERSION} \
-    && apt clean \
-    && cd /opt/intel/oneapi \
-    && rm -rf conda_channel debugger dev-utilities dpl compiler/latest/linux/lib/oclfpga
+# Download DPC++ nightly release
+ARG DPCPP_NIGHTLY=2025-03-06
+RUN mkdir /opt/dpcpp \
+    && wget -q -P /opt/dpcpp https://github.com/intel/llvm/releases/download/nightly-${DPCPP_NIGHTLY}/sycl_linux.tar.gz \
+    && tar -C /opt/dpcpp -xzf /opt/dpcpp/sycl_linux.tar.gz \
+    && rm /opt/dpcpp/sycl_linux.tar.gz
 
 # Set up the environment
-ENV ONEAPI_ROOT=/opt/intel/oneapi
-ENV CMPLR_ROOT=${ONEAPI_ROOT}/compiler/latest
-ENV PATH=${CMPLR_ROOT}/bin:${CMPLR_ROOT}/bin/compiler:${PATH}
-ENV CPATH=${CMPLR_ROOT}/include:${CPATH}
-ENV LIBRARY_PATH=${CMPLR_ROOT}/lib:${LIBRARY_PATH}
-ENV LD_LIBRARY_PATH=${CMPLR_ROOT}/lib:${LD_LIBRARY_PATH}
-ENV HIP_DEVICE_LIB_PATH=/usr/lib/x86_64-linux-gnu/amdgcn/bitcode
-
-# Set up entry point
-ENTRYPOINT []
-CMD /bin/bash
+ENV DPCPP_ROOT=/opt/dpcpp
+ENV PATH=${DPCPP_ROOT}/bin:${PATH}
+ENV CPATH=${DPCPP_ROOT}/include:${CPATH}
+ENV LIBRARY_PATH=${DPCPP_ROOT}/lib:${LIBRARY_PATH}
+ENV LD_LIBRARY_PATH=${DPCPP_ROOT}/lib:${LD_LIBRARY_PATH}
+ENV HIP_DEVICE_LIB_PATH=/opt/rocm/amdgcn/bitcode

ci/warning-suppresions.txt

@@ -0,0 +1,8 @@
+# Suppress warnings in external code from submodules
+
+[nontrivial-memcall]
+src:*/modules/imgui/*
+
+[deprecated-literal-operator]
+src:*/modules/corrade/*
+src:*/modules/magnum/*

cmake/ConfigureSYCL.cmake

@@ -22,10 +22,10 @@
 # Detect available backends
 # ------------------------------------------------
 execute_process(
-    COMMAND bash -c "! sycl-ls | grep -q ext_oneapi_cuda"
+    COMMAND bash -c "! sycl-ls | grep -q cuda"
     RESULT_VARIABLE CUDA_BACKEND_AVAILABLE)
 execute_process(
-    COMMAND bash -c "! sycl-ls | grep -q ext_oneapi_hip"
+    COMMAND bash -c "! sycl-ls | grep -q hip"
     RESULT_VARIABLE HIP_BACKEND_AVAILABLE)
 execute_process(
     COMMAND bash -c "! sycl-ls | grep -q 'opencl\\|level_zero'"
@@ -41,7 +41,7 @@ set(SYCL_TARGETS "")
 # ------------------------------------------------
 if(${ENABLE_CUDA})
     string(JOIN "," SYCL_TARGETS "${SYCL_TARGETS}" "nvptx64-nvidia-cuda")
-    set(DEFAULT_CUDA_COMPUTE_CAPABILITY "50")
+    set(DEFAULT_CUDA_COMPUTE_CAPABILITY "60")
     set(CUDA_COMPUTE_CAPABILITY "" CACHE BOOL
         "CUDA architecture (compute capability), e.g. sm_80. Default value is auto-configured using nvidia-smi.")
     # Auto-configure if not specified by user
@@ -65,7 +65,7 @@ endif()
 # ------------------------------------------------
 if(${ENABLE_HIP})
     string(JOIN "," SYCL_TARGETS "${SYCL_TARGETS}" "amdgcn-amd-amdhsa")
-    set(DEFAULT_HIP_GFX_ARCH "gfx906")
+    set(DEFAULT_HIP_GFX_ARCH "gfx908")
     set(HIP_GFX_ARCH "" CACHE BOOL
         "HIP architecture tag, e.g. gfx90a. Default value is auto-configured using rocminfo.")
     # Auto-configure if not specified by user

src/fluid/fluid.h

@@ -80,7 +80,7 @@ class SYCLFluidContainer {
   // Returns a pointer to the pixel data buffer.
   template <typename Func>
   void WithData(Func&& func) {
-    auto acc{img.template get_host_access(sycl::read_only)};
+    auto acc{img.template get_host_access<>(sycl::read_only)};
     func(acc.get_pointer());
   }
 
@@ -207,8 +207,8 @@ class SYCLFluidContainer {
     // Update the image pixel data with the appropriate color for a given
     // density.
     queue.submit([&](sycl::handler& cgh) {
-      auto img_acc{img.template get_access(cgh, sycl::write_only)};
-      auto density_a{density_b.template get_access(cgh, sycl::read_write)};
+      auto img_acc{img.template get_access<>(cgh, sycl::write_only)};
+      auto density_a{density_b.template get_access<>(cgh, sycl::read_write)};
       cgh.parallel_for<image_kernal>(
           sycl::range<1>(size * size), [=](sycl::item<1> item) {
             auto index{item.get_id(0)};
@@ -240,7 +240,7 @@ class SYCLFluidContainer {
   // Creates read_write accessors from a buffer.
   template <typename T>
   static read_write_accessor CreateAccessor(sycl::handler& cgh, T buffer) {
-    return buffer.template get_access(cgh, sycl::read_write);
+    return buffer.template get_access<>(cgh, sycl::read_write);
   }
 
   // Get clamped index based off of coordinates.

src/matrix_multiply_omp_compare/CMakeLists.txt

@@ -4,10 +4,25 @@ set(MATRIX_MULTIPLY_FLAGS ${SYCL_FLAGS})
 find_package(OpenMP QUIET)
 if(OpenMP_CXX_FOUND)
     target_link_libraries(matrix_multiply_omp_compare PUBLIC OpenMP::OpenMP_CXX)
+    message(STATUS "Found OpenMP: ${OpenMP_omp_LIBRARY}")
 elseif(CMAKE_CXX_COMPILER MATCHES "clang")
-    # CMake's FindOpenMP doesn't recognise LLVM internal implementation, but
-    # it's still possible to parallelise OpenMP loops with clang++ -fopenmp
-    list(APPEND MATRIX_MULTIPLY_FLAGS -fopenmp)
+    # CMake's FindOpenMP doesn't recognise that oneAPI Base Toolkit clang++
+    # compiler driver supports the -fopenmp flag. However, other clang++ builds
+    # may not support it. Test whether clang++ -fopenmp works to determine if we
+    # can use it.
+    file(WRITE ${CMAKE_CURRENT_BINARY_DIR}/omp-test.cpp "int main(){}")
+    try_compile(CLANG_SUPPORTS_FOPENMP
+        ${CMAKE_CURRENT_BINARY_DIR}
+        ${CMAKE_CURRENT_BINARY_DIR}/omp-test.cpp
+        COMPILE_DEFINITIONS -fopenmp
+        LINK_OPTIONS -fopenmp)
+    file(REMOVE ${CMAKE_CURRENT_BINARY_DIR}/omp-test.cpp)
+    if (CLANG_SUPPORTS_FOPENMP)
+        message(STATUS "Found OpenMP: supported with clang++ -fopenmp flag")
+        list(APPEND MATRIX_MULTIPLY_FLAGS -fopenmp)
+    else()
+        message(STATUS "OpenMP not found, matrix_multiply_omp_compare will compare SYCL to serial CPU execution")
+    endif()
 else()
     message(STATUS "OpenMP not found, matrix_multiply_omp_compare will compare SYCL to serial CPU execution")
 endif()

src/nbody/sycl_bufs.hpp

@@ -31,7 +31,7 @@ struct BufToReadAccFunc {
   template <typename In>
   AUTO_FUNC(
       // pair of (buffer, handler)
-      operator()(In && in), std::forward<In>(in).first.template get_access(
+      operator()(In && in), std::forward<In>(in).first.template get_access<>(
                                 *std::forward<In>(in).second, sycl::read_only))
 };
 
@@ -40,27 +40,27 @@ struct BufToDcdWriteAccFunc {
   // pair of (buffer, handler)
   template <typename In>
   AUTO_FUNC(operator()(In && in),
-            std::forward<In>(in).first.template get_access(
+            std::forward<In>(in).first.template get_access<>(
                 *std::forward<In>(in).second, sycl::write_only))
 };
 
 // Template function object which transforms buffers to host read accessors
 struct BufToHostReadAccFunc {
   template <typename In>
   auto operator()(In&& in)
-      -> decltype(std::forward<In>(in).template get_host_access(
+      -> decltype(std::forward<In>(in).template get_host_access<>(
           sycl::read_only)) {
-    return std::forward<In>(in).template get_host_access(sycl::read_only);
+    return std::forward<In>(in).template get_host_access<>(sycl::read_only);
   }
 };
 
 // Template function object which transforms buffers to host write accessors
 struct BufToHostDcdWriteAccFunc {
   template <typename In>
   auto operator()(In&& in)
-      -> decltype(std::forward<In>(in).template get_host_access(
+      -> decltype(std::forward<In>(in).template get_host_access<>(
           sycl::write_only)) {
-    return std::forward<In>(in).template get_host_access(sycl::write_only);
+    return std::forward<In>(in).template get_host_access<>(sycl::write_only);
   }
 };