diff --git a/.github/workflows/ubuntu-latest-oneapi.yml b/.github/workflows/ubuntu-latest-oneapi.yml new file mode 100644 index 0000000000..bebd04efad --- /dev/null +++ b/.github/workflows/ubuntu-latest-oneapi.yml @@ -0,0 +1,45 @@ +name: Build - Ubuntu/dpcpp + +on: + pull_request: + merge_group: + workflow_dispatch: + +concurrency: + group: ${{ github.workflow }}-${{ github.head_ref || github.ref_name }} + cancel-in-progress: true + +env: + # Customize the CMake build type here (Release, Debug, RelWithDebInfo, etc.) + BUILD_TYPE: Release + +jobs: + build: + strategy: + matrix: + ONEAPI_VERSION: [ + '2024.2.1-0-devel-ubuntu22.04', + 'latest' # 2025.0.0-0-devel-ubuntu24.04 at the time of commit + ] + + runs-on: ubuntu-latest + container: intel/oneapi-basekit:${{ matrix.ONEAPI_VERSION }} + + steps: + - uses: actions/checkout@v4 + + - name: Configure CMake + run: | + cmake \ + -B ${{github.workspace}}/build \ + -D CMAKE_BUILD_TYPE=${{env.BUILD_TYPE}} \ + -D CMAKE_C_COMPILER=$(which icx) \ + -D CMAKE_CXX_COMPILER=$(which icpx) \ + -D CMAKE_CXX_FLAGS="-fsycl" \ + -D SUNDIALS_BUILD_WITH_PROFILING=ON \ + -D ENABLE_ALL_WARNINGS=ON \ + -D ENABLE_WARNINGS_AS_ERRORS=ON \ + -D ENABLE_SYCL=ON + + - name: Build + run: cmake --build ${{github.workspace}}/build --config ${{env.BUILD_TYPE}} diff --git a/CHANGELOG.md b/CHANGELOG.md index 74d31b0d91..f9e1929344 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -72,6 +72,9 @@ inner stepper object, `ARKodeCreateMRIStepInnerStepper`. ### Bug Fixes +Fixed a build failure with the SYCL NVector when using Intel oneAPI 2025.0 +compilers. See GitHub Issue [#596](https://github.com/LLNL/sundials/issues/596). + Fixed a bug where `CVodeSetProjFailEta` would ignore the `eta` parameter. Fixed a bug in the SPTFQMR linear solver where recoverable preconditioner errors diff --git a/doc/shared/RecentChanges.rst b/doc/shared/RecentChanges.rst index 09e022695a..5e61e9f216 100644 --- a/doc/shared/RecentChanges.rst +++ b/doc/shared/RecentChanges.rst @@ -81,6 +81,9 @@ inner stepper object, :c:func:`ARKodeCreateMRIStepInnerStepper`. **Bug Fixes** +Fixed a build failure with the SYCL NVector when using Intel oneAPI 2025.0 +compilers. See GitHub Issue `#596 `__. + Fixed a bug where :c:func:`CVodeSetProjFailEta` would ignore the `eta` parameter. diff --git a/examples/cvode/CXX_sycl/cvAdvDiff_kry_sycl.cpp b/examples/cvode/CXX_sycl/cvAdvDiff_kry_sycl.cpp index 9997ad52d9..84718e9e54 100644 --- a/examples/cvode/CXX_sycl/cvAdvDiff_kry_sycl.cpp +++ b/examples/cvode/CXX_sycl/cvAdvDiff_kry_sycl.cpp @@ -248,9 +248,9 @@ static int f(sunrealtype t, N_Vector u, N_Vector udot, void* user_data) h.parallel_for(sycl::range{MX, MY}, [=](sycl::id<2> idx) { - sunindextype i = idx[0]; - sunindextype j = idx[1]; - sunindextype tid = i * MY + j; + size_t i = idx[0]; + size_t j = idx[1]; + size_t tid = i * MY + j; sunrealtype uij = udata[tid]; sunrealtype udn = (j == 0) ? ZERO : udata[tid - 1]; @@ -293,9 +293,9 @@ static int jtv(N_Vector v, N_Vector Jv, sunrealtype t, N_Vector u, N_Vector fu, h.parallel_for(sycl::range{MX, MY}, [=](sycl::id<2> idx) { - sunindextype i = idx[0]; - sunindextype j = idx[1]; - sunindextype tid = i * MY + j; + size_t i = idx[0]; + size_t j = idx[1]; + size_t tid = i * MY + j; // set the tid-th element of Jv Jvdata[tid] = -TWO * (verdc + hordc) * vdata[tid]; diff --git a/include/nvector/nvector_sycl.h b/include/nvector/nvector_sycl.h index 82935ccca6..11613cabfb 100644 --- a/include/nvector/nvector_sycl.h +++ b/include/nvector/nvector_sycl.h @@ -105,7 +105,7 @@ static inline sunrealtype* N_VGetDeviceArrayPointer_Sycl(N_Vector x) * NVECTOR API functions * ----------------------------------------------------------------- */ -static inline N_Vector_ID N_VGetVectorID_Sycl(N_Vector v) +static inline N_Vector_ID N_VGetVectorID_Sycl(N_Vector) { return SUNDIALS_NVEC_SYCL; } diff --git a/include/sundials/sundials_sycl_policies.hpp b/include/sundials/sundials_sycl_policies.hpp index db0a887bfa..cf9950f4f9 100644 --- a/include/sundials/sundials_sycl_policies.hpp +++ b/include/sundials/sundials_sycl_policies.hpp @@ -50,13 +50,14 @@ class ThreadDirectExecPolicy : public ExecPolicy : blockDim_(ex.blockDim_) {} - virtual size_t gridSize(size_t numWorkUnits = 0, size_t blockDim = 0) const + virtual size_t gridSize(size_t numWorkUnits = 0, size_t /* blockDim */ = 0) const { /* ceil(n/m) = floor((n + m - 1) / m) */ return (numWorkUnits + blockSize() - 1) / blockSize(); } - virtual size_t blockSize(size_t numWorkUnits = 0, size_t gridDim = 0) const + virtual size_t blockSize(size_t /* numWorkUnits */ = 0, + size_t /* gridDim */ = 0) const { return blockDim_; } @@ -86,12 +87,14 @@ class GridStrideExecPolicy : public ExecPolicy : blockDim_(ex.blockDim_), gridDim_(ex.gridDim_) {} - virtual size_t gridSize(size_t numWorkUnits = 0, size_t blockDim = 0) const + virtual size_t gridSize(size_t /* numWorkUnits */ = 0, + size_t /* blockDim */ = 0) const { return gridDim_; } - virtual size_t blockSize(size_t numWorkUnits = 0, size_t gridDim = 0) const + virtual size_t blockSize(size_t /* numWorkUnits */ = 0, + size_t /* gridDim */ = 0) const { return blockDim_; } @@ -124,7 +127,7 @@ class BlockReduceExecPolicy : public ExecPolicy : blockDim_(ex.blockDim_), gridDim_(ex.gridDim_) {} - virtual size_t gridSize(size_t numWorkUnits = 0, size_t blockDim = 0) const + virtual size_t gridSize(size_t numWorkUnits = 0, size_t /* blockDim */ = 0) const { if (gridDim_ == 0) { @@ -133,7 +136,8 @@ class BlockReduceExecPolicy : public ExecPolicy return gridDim_; } - virtual size_t blockSize(size_t numWorkUnits = 0, size_t gridDim = 0) const + virtual size_t blockSize(size_t /* numWorkUnits */ = 0, + size_t /* gridDim */ = 0) const { return blockDim_; } diff --git a/src/nvector/sycl/nvector_sycl.cpp b/src/nvector/sycl/nvector_sycl.cpp index 2f5b6a0a61..13d0155d55 100644 --- a/src/nvector/sycl/nvector_sycl.cpp +++ b/src/nvector/sycl/nvector_sycl.cpp @@ -1146,7 +1146,7 @@ sunrealtype N_VWrmsNorm_Sycl(N_Vector x, N_Vector w) { const sunindextype N = NVEC_SYCL_LENGTH(x); const sunrealtype sum = N_VWSqrSumLocal_Sycl(x, w); - return std::sqrt(sum / N); + return ::sycl::sqrt(sum / N); } sunrealtype N_VWSqrSumMaskLocal_Sycl(N_Vector x, N_Vector w, N_Vector id) @@ -1192,7 +1192,7 @@ sunrealtype N_VWrmsNormMask_Sycl(N_Vector x, N_Vector w, N_Vector id) { const sunindextype N = NVEC_SYCL_LENGTH(x); const sunrealtype sum = N_VWSqrSumMaskLocal_Sycl(x, w, id); - return std::sqrt(sum / N); + return ::sycl::sqrt(sum / N); } sunrealtype N_VMin_Sycl(N_Vector x) @@ -1233,7 +1233,7 @@ sunrealtype N_VMin_Sycl(N_Vector x) sunrealtype N_VWL2Norm_Sycl(N_Vector x, N_Vector w) { - return std::sqrt(N_VWSqrSumLocal_Sycl(x, w)); + return ::sycl::sqrt(N_VWSqrSumLocal_Sycl(x, w)); } sunrealtype N_VL1Norm_Sycl(N_Vector x) @@ -2285,14 +2285,14 @@ static int FusedBuffer_CopyRealArray(N_Vector v, sunrealtype* rdata, int nval, return SUN_ERR_GENERIC; } - sunrealtype* h_buffer = (sunrealtype*)((char*)(vcp->fused_buffer_host->ptr) + - vcp->fused_buffer_offset); + sunrealtype* h_buffer = reinterpret_cast( + (char*)(vcp->fused_buffer_host->ptr) + vcp->fused_buffer_offset); for (int j = 0; j < nval; j++) { h_buffer[j] = rdata[j]; } /* Set shortcut to the device buffer and update offset*/ - *shortcut = (sunrealtype*)((char*)(vcp->fused_buffer_dev->ptr) + - vcp->fused_buffer_offset); + *shortcut = reinterpret_cast( + (char*)(vcp->fused_buffer_dev->ptr) + vcp->fused_buffer_offset); vcp->fused_buffer_offset += nval * sizeof(sunrealtype); @@ -2314,14 +2314,14 @@ static int FusedBuffer_CopyPtrArray1D(N_Vector v, N_Vector* X, int nvec, return SUN_ERR_GENERIC; } - sunrealtype** h_buffer = (sunrealtype**)((char*)(vcp->fused_buffer_host->ptr) + - vcp->fused_buffer_offset); + sunrealtype** h_buffer = reinterpret_cast( + (char*)(vcp->fused_buffer_host->ptr) + vcp->fused_buffer_offset); for (int j = 0; j < nvec; j++) { h_buffer[j] = NVEC_SYCL_DDATAp(X[j]); } /* Set shortcut to the device buffer and update offset*/ - *shortcut = (sunrealtype**)((char*)(vcp->fused_buffer_dev->ptr) + - vcp->fused_buffer_offset); + *shortcut = reinterpret_cast( + (char*)(vcp->fused_buffer_dev->ptr) + vcp->fused_buffer_offset); vcp->fused_buffer_offset += nvec * sizeof(sunrealtype*); @@ -2342,8 +2342,8 @@ static int FusedBuffer_CopyPtrArray2D(N_Vector v, N_Vector** X, int nvec, return SUN_ERR_GENERIC; } - sunrealtype** h_buffer = (sunrealtype**)((char*)(vcp->fused_buffer_host->ptr) + - vcp->fused_buffer_offset); + sunrealtype** h_buffer = reinterpret_cast( + (char*)(vcp->fused_buffer_host->ptr) + vcp->fused_buffer_offset); for (int j = 0; j < nvec; j++) { @@ -2354,8 +2354,8 @@ static int FusedBuffer_CopyPtrArray2D(N_Vector v, N_Vector** X, int nvec, } /* Set shortcut to the device buffer and update offset*/ - *shortcut = (sunrealtype**)((char*)(vcp->fused_buffer_dev->ptr) + - vcp->fused_buffer_offset); + *shortcut = reinterpret_cast( + (char*)(vcp->fused_buffer_dev->ptr) + vcp->fused_buffer_offset); /* Update the offset */ vcp->fused_buffer_offset += nvec * nsum * sizeof(sunrealtype*);