Fix build errors and improve CI config (#42)

* 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.
* Do not push to registry in pull-request CI (not allowed in GitHub). Instead, push the image to a local registry and use that to build the code. Push to registry only in CI triggered by branch push.

Author: rafbiels

.github/workflows/ci.yml

@@ -4,6 +4,7 @@ on: [push, pull_request]
 
 permissions:
   contents: read
+  packages: write
 
 env:
   REGISTRY: ghcr.io
@@ -25,15 +26,20 @@ jobs:
           fallback-style: 'Google'
           include-regex: '^.*\.((((c|C)(c|pp|xx|\+\+)?$)|((h|H)h?(pp|xx|\+\+)?$))|(frag|vert))$'
 
-  build-container:
+  build-dpcpp-nightly:
     runs-on: ubuntu-latest
 
-    outputs:
-      image_tag: ${{ steps.set_image_tag.outputs.image_tag }}
+    services:
+      registry:
+        image: registry:2.8.3@sha256:a3d8aaa63ed8681a604f1dea0aa03f100d5895b6a58ace528858a7b332415373
+        ports:
+          - 5000:5000
 
     steps:
       - name: Checkout
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+        with:
+          submodules: 'recursive'
 
       - name: Path filter
         uses: dorny/paths-filter@de90cc6fb38fc0963ad72b210f1f284cd68cea36 # v3.0.2
@@ -46,75 +52,55 @@ jobs:
       - name: Set up Docker Buildx
         if: steps.changes.outputs.container == 'true'
         uses: docker/setup-buildx-action@6524bf65af31da8d45b59e8c27de4bd072b392f5 # v3.8.0
+        with:
+          driver-opts: network=host
 
       - name: Log in to the Container registry
-        if: steps.changes.outputs.container == 'true'
+        if: steps.changes.outputs.container == 'false' || github.event_name == 'push'
         uses: docker/login-action@9780b0c442fbb1117ed29e0efdff1e18412f7567 # v3.3.0
         with:
           registry: ${{ env.REGISTRY }}
           username: ${{ github.actor }}
           password: ${{ secrets.GITHUB_TOKEN }}
 
       - name: Extract metadata
+        if: steps.changes.outputs.container == 'false' || github.event_name == 'push'
         id: meta
         uses: docker/metadata-action@369eb591f429131d6889c46b94e711f089e6ca96 # v5.6.1
         with:
           images: ${{ env.REGISTRY }}/${{ env.IMAGE_NAME }}
           tags: |
-            # set latest tag for the main branch
-            type=raw,value=latest,enable=${{ github.ref == format('refs/heads/{0}', 'main') }},priority=5000
-            # use ref_name in all other cases
-            type=raw,value=${{github.ref_name}},priority=4000
-
-      - name: Build and push
-        if: steps.changes.outputs.container == 'true'
+            # push latest tag on main branch push
+            type=raw,value=latest,enable=${{ github.event_name == 'push' && github.ref == 'refs/heads/main' }},priority=5000
+            # push branch name as the tag for every other branch push
+            type=raw,value=${{github.ref_name}},enable=${{ github.event_name == 'push' && github.ref != 'refs/heads/main' }},priority=4000
+            # pull latest in PRs
+            type=raw,value=latest,enable=${{ github.event_name != 'push' }},priority=3000
+
+      - name: Build and push image
+        if: steps.changes.outputs.container == 'true' && github.event_name == 'push'
         uses: docker/build-push-action@ca877d9245402d1537745e0e356eab47c3520991 # v6.13.0
         with:
           context: '{{defaultContext}}:ci'
           push: true
           tags: ${{ steps.meta.outputs.tags }}
           labels: ${{ steps.meta.outputs.labels }}
 
-      - name: Set output image tag
-        id: set_image_tag
-        run: >
-          tags="${{ steps.meta.outputs.tags }}"
-          && tags=${tags%% *}
-          && image=${tags%:*}
-          && tag="${{ steps.changes.outputs.container == 'true' && '${tags#*:}' || 'latest' }}"
-          && echo "image_tag=${image}:${tag}"
-          && echo "image_tag=${image}:${tag}" >> "$GITHUB_OUTPUT"
-
-  cmake-build:
-    needs: build-container
-    runs-on: ubuntu-latest
-    container:
-      image: ${{ needs.build-container.outputs.image_tag }}
-      credentials:
-         username: ${{ github.actor }}
-         password: ${{ secrets.github_token }}
-    steps:
-      - name: Checkout
-        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+      - name: Build local image
+        if: steps.changes.outputs.container == 'true' && github.event_name != 'push'
+        uses: docker/build-push-action@ca877d9245402d1537745e0e356eab47c3520991 # v6.13.0
         with:
-          submodules: "recursive"
-
-      - name: Configure CMake
-        run: >
-          cmake -B ${{github.workspace}}/build
-          -DCMAKE_BUILD_TYPE=${{env.BUILD_TYPE}}
-          -DCMAKE_CXX_COMPILER=clang++
-          -DENABLE_GRAPHICS=ON
-          -DENABLE_SPIR=ON
-          -DENABLE_CUDA=ON -DCUDA_COMPUTE_CAPABILITY=80
-          -DENABLE_HIP=ON -DHIP_GFX_ARCH=gfx90a
-          -DCMAKE_CXX_FLAGS='-Wall -Wextra -Wpedantic -Werror'
-          -G Ninja
+          context: '{{defaultContext}}:ci'
+          push: true
+          tags: localhost:5000/sycl-samples:local
+          labels: ${{ steps.meta.outputs.labels }}
 
-      - name: Build
-        run: cmake --build ${{github.workspace}}/build -- -k 0
+      - name: CMake configure and build
+        run: |
+          if [[ "${{steps.changes.outputs.container}}" == "false"  || "${{github.event_name}}" == "push" ]]; then IMAGE_TAG=${{steps.meta.outputs.tags}}; else IMAGE_TAG=localhost:5000/sycl-samples:local; fi
+          docker run -v${{github.workspace}}:/work ${IMAGE_TAG} ${{env.BUILD_TYPE}}
 
-  cmake-build-nographics:
+  build-oneapi-basekit-nographics:
     runs-on: ubuntu-latest
     container:
       image: intel/oneapi-basekit:latest

ci/Dockerfile

@@ -1,44 +1,39 @@
 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
+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
 
-# Set up entry point
-ENTRYPOINT []
-CMD /bin/bash
+WORKDIR /work
+ENTRYPOINT ["/work/ci/build.sh"]

ci/build.sh

@@ -0,0 +1,18 @@
+#!/usr/bin/bash
+
+BUILD_TYPE=$1
+if [[ -z "$BUILD_TYPE" ]]; then
+  BUILD_TYPE=Release
+fi
+
+cmake -B build \
+  -DCMAKE_BUILD_TYPE=${BUILD_TYPE} \
+  -DCMAKE_CXX_COMPILER=clang++ \
+  -DENABLE_GRAPHICS=ON \
+  -DENABLE_SPIR=ON \
+  -DENABLE_CUDA=ON -DCUDA_COMPUTE_CAPABILITY=80 \
+  -DENABLE_HIP=ON -DHIP_GFX_ARCH=gfx90a \
+  -DCMAKE_CXX_FLAGS="-Wall -Wextra -Wpedantic -Werror --warning-suppression-mappings=${PWD}/ci/warning-suppresions.txt" \
+  -G Ninja
+
+cmake --build build -- -k 0

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);
   }
 };