Fix warnings for DPC++ (AMReX-Codes#1214)

## Summary

- Add more warnings flags to dpcpp.
- Fix warnings for DPC++.
- Fix a bug in DPC++ version of ReduceLogicalOr.

## Checklist

The proposed changes:
- [x] fix a bug or incorrect behavior in AMReX
- [ ] add new capabilities to AMReX
- [ ] changes answers in the test suite to more than roundoff level
- [ ] are likely to significantly affect the results of downstream AMReX users
- [ ] are described in the proposed changes to the AMReX documentation, if appropriate

Src/Base/AMReX_BlockMutex.H

@@ -32,6 +32,7 @@ struct BlockMutex
     void lock (int i) noexcept {
 #ifdef AMREX_USE_DPCPP
 // xxxxx DPCPP todo
+        amrex::ignore_unused(i);
 #else
         int blockid = blockIdx.z*blockDim.x*blockDim.y + blockIdx.y*blockDim.x + blockIdx.x;
         state_t old = m_state[i];
@@ -57,6 +58,7 @@ struct BlockMutex
     void unlock (int i) noexcept {
 #ifdef AMREX_USE_DPCPP
 // xxxxx DPCPP todo
+        amrex::ignore_unused(i);
 #else
         state_t old = m_state[i];
         state_t assumed;

Src/Base/AMReX_BlockMutex.cpp

@@ -6,6 +6,7 @@ namespace amrex {
 
 void BlockMutex::init_states (state_t* state, int N) noexcept {
 #ifdef AMREX_USE_DPCPP
+    amrex::ignore_unused(state,N);
     amrex::Abort("xxxxx DPCPP todo");
 #else
     amrex::launch((N+255)/256, 256, Gpu::nullStream(),

Src/Base/AMReX_FabArrayUtility.H

@@ -82,6 +82,7 @@ template <class FAB, class F>
 amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, typename FAB::value_type>
 ReduceSum_host_wrapper (FabArray<FAB> const& fa, IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa,nghost,f);
     amrex::Abort("ReduceSum: Launch Region is off. Device lambda cannot be called by host.");
     return 0;
 }
@@ -189,6 +190,7 @@ amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, typename FAB1::val
 ReduceSum_host_wrapper (FabArray<FAB1> const& fa1, FabArray<FAB2> const& fa2,
                         IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa1,fa2,nghost,f);
     amrex::Abort("ReduceSum: Launch Region is off. Device lambda cannot be called by host.");
     return 0;
 }
@@ -301,6 +303,7 @@ amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, typename FAB1::val
 ReduceSum_host_wrapper (FabArray<FAB1> const& fa1, FabArray<FAB2> const& fa2,
                         FabArray<FAB3> const& fa3, IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa1,fa2,fa3,nghost,f);
     amrex::Abort("ReduceSum: Launch Region is off. Device lambda cannot be called by host.");
     return 0;
 }
@@ -405,6 +408,7 @@ template <class FAB, class F>
 amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, typename FAB::value_type>
 ReduceMin_host_wrapper (FabArray<FAB> const& fa, IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa,nghost,f);
     amrex::Abort("ReduceMin: Launch Region is off. Device lambda cannot be called by host.");
     return 0;
 }
@@ -514,6 +518,7 @@ amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, typename FAB1::val
 ReduceMin_host_wrapper (FabArray<FAB1> const& fa1, FabArray<FAB2> const& fa2,
                         IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa1,fa2,nghost,f);
     amrex::Abort("ReduceMin: Launch Region is off. Device lambda cannot be called by host.");
     return 0;
 }
@@ -628,6 +633,7 @@ amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, typename FAB1::val
 ReduceMin_host_wrapper (FabArray<FAB1> const& fa1, FabArray<FAB2> const& fa2,
                         FabArray<FAB3> const& fa3, IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa1,fa2,fa3,nghost,f);
     amrex::Abort("ReduceMin: Launch Region is off. Device lambda lambda cannot be called by host.");
     return 0;
 }
@@ -733,6 +739,7 @@ template <class FAB, class F>
 amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, typename FAB::value_type>
 ReduceMax_host_wrapper (FabArray<FAB> const& fa, IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa,nghost,f);
     amrex::Abort("ReduceMax: Launch Region is off. Device lambda cannot be called by host.");
     return 0;
 }
@@ -842,6 +849,7 @@ amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, typename FAB1::val
 ReduceMax_host_wrapper (FabArray<FAB1> const& fa1, FabArray<FAB2> const& fa2,
                         IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa1,fa2,nghost,f);
     amrex::Abort("ReduceMax: Launch Region is off. Device lambda cannot be called by host.");
     return 0;
 }
@@ -956,6 +964,7 @@ amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, typename FAB1::val
 ReduceMax_host_wrapper (FabArray<FAB1> const& fa1, FabArray<FAB2> const& fa2,
                         FabArray<FAB3> const& fa3, IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa1,fa2,fa3,nghost,f);
     amrex::Abort("ReduceMax: Launch Region is off. Device lambda lambda cannot be called by host.");
     return 0;
 }
@@ -1058,6 +1067,7 @@ template <class FAB, class F>
 amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, bool>
 ReduceLogicalAnd_host_wrapper (FabArray<FAB> const& fa, IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa,nghost,f);
     amrex::Abort("ReduceLogicalAnd: Launch Region is off. Device lambda cannot be called by host.");
     return false;
 }
@@ -1165,6 +1175,7 @@ amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, bool>
 ReduceLogicalAnd_host_wrapper (FabArray<FAB1> const& fa1, FabArray<FAB2> const& fa2,
                                IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa1,fa2,nghost,f);
     amrex::Abort("ReduceLogicalAnd: Luanch Region is off. Device lambda cannot be called by host.");
     return false;
 }
@@ -1267,6 +1278,7 @@ template <class FAB, class F>
 amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, bool>
 ReduceLogicalOr_host (FabArray<FAB> const& fa, IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa,nghost,f);
     amrex::Abort("ReduceLogicalOr: Launch Region is off. Device lambda cannot be called by host.");
     return 0;
 }
@@ -1350,6 +1362,7 @@ ReduceLogicalOr_device (FabArray<FAB1> const& fa1, FabArray<FAB2> const& fa2,
             [=] AMREX_GPU_DEVICE (Box const& b) -> ReduceTuple
             {
                 int tr = f(b, arr1, arr2);
+                return {tr};
             });
         }
 
@@ -1373,6 +1386,7 @@ amrex::EnableIf_t<amrex::DefinitelyNotHostRunnable<F>::value, bool>
 ReduceLogicalOr_host_wrapper (FabArray<FAB1> const& fa1, FabArray<FAB2> const& fa2,
                               IntVect const& nghost, F&& f)
 {
+    amrex::ignore_unused(fa1,fa2,nghost,f);
     amrex::Abort("ReeuceLogicalOr: Launch Region is off. Device lambda cannot be called by host.");
     return false;
 }

Src/Base/AMReX_GpuAtomic.H

@@ -82,6 +82,7 @@ namespace detail {
         sycl::atomic<T,as> a{sycl::multi_ptr<T,as>(sum)};
         return sycl::atomic_fetch_add(a, value, mo);
 #else
+        amrex::ignore_unused(sum, value);
         return T(); // should never get here, but have to return something
 #endif
     }
@@ -143,6 +144,7 @@ namespace detail {
         sycl::atomic<T,as> a{sycl::multi_ptr<T,as>(m)};
         return sycl::atomic_fetch_min(a, value, mo);
 #else
+        amrex::ignore_unused(m,value);
         return T(); // should never get here, but have to return something
 #endif
     }
@@ -202,6 +204,7 @@ namespace detail {
         sycl::atomic<T,as> a{sycl::multi_ptr<T,as>(m)};
         return sycl::atomic_fetch_max(a, value, mo);
 #else
+        amrex::ignore_unused(m,value);
         return T(); // should never get here, but have to return something
 #endif
     }

Src/Base/AMReX_GpuDevice.cpp

@@ -1046,7 +1046,8 @@ Device::freeMemAvailable ()
     std::size_t f, t;
     AMREX_HIP_OR_CUDA_OR_DPCPP( AMREX_HIP_SAFE_CALL(hipMemGetInfo(&f,&t));,
                                 AMREX_CUDA_SAFE_CALL(cudaMemGetInfo(&f,&t));,
-                                f = device_prop.totalGlobalMem; ); // xxxxx DPCPP tod
+                                f = device_prop.totalGlobalMem; ); // xxxxx DPCPP todo
+    amrex::ignore_unused(t);
     return f;
 #else
     return 0;

Src/Base/AMReX_GpuUtility.H

@@ -49,6 +49,7 @@ namespace Gpu {
         CUresult r = cuPointerGetAttributes(1, &attrib, data, (CUdeviceptr)p);
         return r == CUDA_SUCCESS && is_managed;
 #elif defined(AMREX_USE_DPCPP)
+        amrex::ignore_unused(p);
         // xxxxx DPCPP todo: get_pointer_type
 //        auto const info = sycl::get_pointer_info(p);
 //        auto type = sycl::get_pointer_type(p,Device::syclContext());
@@ -72,6 +73,7 @@ namespace Gpu {
         CUresult r = cuPointerGetAttributes(1, &attrib, data, (CUdeviceptr)p);
         return r == CUDA_SUCCESS && mem_type == CU_MEMORYTYPE_DEVICE;
 #elif defined(AMREX_USE_DPCPP)
+        amrex::ignore_unused(p);
         // xxxxx DPCPP todo: get_pointer_type
         return false;
 #else
@@ -93,6 +95,7 @@ namespace Gpu {
         CUresult r = cuPointerGetAttributes(1, &attrib, data, (CUdeviceptr)p);
         return r == CUDA_SUCCESS && mem_type == CU_MEMORYTYPE_HOST;
 #elif defined(AMREX_USE_DPCPP)
+        amrex::ignore_unused(p);
         // xxxxx DPCPP todo: get_pointer_type
         return false;
 #else
@@ -124,6 +127,7 @@ namespace Gpu {
              mem_type == CU_MEMORYTYPE_ARRAY  ||
              mem_type == CU_MEMORYTYPE_UNIFIED);
 #elif defined(AMREX_USE_DPCPP)
+        amrex::ignore_unused(p);
         // xxxxx DPCPP todo: get_pointer_type
         return false;
 #else

Src/Base/AMReX_Random.cpp

@@ -33,18 +33,25 @@ namespace
     constexpr int gpu_nstates_default = 1e5;
 
     int gpu_nstates_h = 0;
+#ifndef AMREX_USE_DPCPP
     AMREX_GPU_DEVICE int gpu_nstates_d = 0;
+#endif
 
     randState_t* d_states_h_ptr = nullptr;
+#ifndef AMREX_USE_DPCPP
     AMREX_GPU_DEVICE randState_t* d_states_d_ptr;
+#endif
 
     amrex::BlockMutex* h_mutex_h_ptr = nullptr;
     amrex::BlockMutex* d_mutex_h_ptr = nullptr;
 
+#ifndef AMREX_USE_DPCPP
     AMREX_GPU_DEVICE
     amrex::BlockMutex* d_mutex_d_ptr = nullptr;
 #endif
 
+#endif
+
 }
 
 void
@@ -74,6 +81,7 @@ int amrex::get_state (int tid)
 {
 #ifdef AMREX_USE_DPCPP
 // xxxxx DPCPP todo
+    amrex::ignore_unused(tid);
     return 0;
 #else
     // block size must evenly divide # of RNG states so we cut off the excess states
@@ -91,6 +99,7 @@ AMREX_GPU_DEVICE
 void amrex::free_state (int tid)
 {
 #ifdef AMREX_USE_DPCPP
+    amrex::ignore_unused(tid);
 // xxxxx DPCPP todo
 #else
     int bsize = blockDim.x * blockDim.y * blockDim.z;
@@ -128,6 +137,7 @@ amrex::RandomNormal (amrex::Real mean, amrex::Real stddev)
 
 #elif defined(__SYCL_DEVICE_ONLY__)
 
+    amrex::ignore_unused(mean,stddev);
     assert(0);
     rand = 0.0_rt;
     return rand;
@@ -206,8 +216,9 @@ amrex::RandomPoisson (amrex::Real lambda)
 
 #elif defined(__SYCL_DEVICE_ONLY__)
 
+    amrex::ignore_unused(lambda);
     assert(0);
-    rand = 0.0_rt;
+    rand = 0;
     return rand;
 
 #else
@@ -245,6 +256,7 @@ amrex::Random_int (unsigned int n)
 
 #elif defined(__SYCL_DEVICE_ONLY__)
 
+    amrex::ignore_unused(n);
     assert(0);
     return 0;
 

Src/Base/AMReX_Reduce.H

@@ -596,9 +596,6 @@ bool AnyOf (Box const& box, P&& pred)
     Gpu::LaunchSafeGuard lsg(true);
     Gpu::DeviceScalar<int> ds(0);
     int* dp = ds.dataPtr();
-    int ncells = box.numPts();
-    const auto lo  = amrex::lbound(box);
-    const auto len = amrex::length(box);
 #ifdef AMREX_USE_DPCPP
     // xxxxx DPCPP todo: Anyof Box: better version
     amrex::ParallelFor(box, [=] (int i, int j, int k) noexcept
@@ -607,6 +604,9 @@ bool AnyOf (Box const& box, P&& pred)
         Gpu::Atomic::LogicalOr(dp, r);
     });
 #else
+    int ncells = box.numPts();
+    const auto lo  = amrex::lbound(box);
+    const auto len = amrex::length(box);
     auto ec = Gpu::ExecutionConfig(ncells);
     ec.numBlocks.x = std::min(ec.numBlocks.x, static_cast<unsigned int>(Gpu::Device::maxBlocksPerLaunch()));
     AMREX_LAUNCH_KERNEL(ec.numBlocks, ec.numThreads, 0, 0,

Src/Base/AMReX_parstream.cpp

@@ -97,6 +97,7 @@ namespace amrex
 // in serial, this does absolutely nothing
   static void openFile()
   {
+    amrex::ignore_unused(s_pout);
   }
 #endif
 

Tools/GNUMake/comps/dpcpp.mak

@@ -27,8 +27,8 @@ F90FLAGS =
 
 ifeq ($(DEBUG),TRUE)
 
-  CXXFLAGS += -g -O0 -Wall -Wextra -Wno-sign-compare -Wno-unused-parameter -Wno-unused-variable #-ftrapv
-  CFLAGS   += -g -O0 -Wall -Wextra -Wno-sign-compare -Wno-unused-parameter -Wno-unused-variable #-ftrapv
+  CXXFLAGS += -g -O0 #-ftrapv
+  CFLAGS   += -g -O0 #-ftrapv
 
   FFLAGS   += -g -O0 -ggdb -fbounds-check -fbacktrace -Wuninitialized -Wunused -ffpe-trap=invalid,zero -finit-real=snan -finit-integer=2147483647 #-ftrapv
   F90FLAGS += -g -O0 -ggdb -fbounds-check -fbacktrace -Wuninitialized -Wunused -ffpe-trap=invalid,zero -finit-real=snan -finit-integer=2147483647 #-ftrapv
@@ -44,6 +44,29 @@ else
 
 endif
 
+CXXFLAGS += -Wno-pass-failed  # disable this warning
+
+ifeq ($(WARN_ALL),TRUE)
+  warning_flags = -Wall -Wextra -Wno-sign-compare -Wunreachable-code -Wnull-dereference
+  warning_flags += -Wfloat-conversion -Wextra-semi
+
+  ifneq ($(USE_CUDA),TRUE)
+    warning_flags += -Wpedantic
+  endif
+
+  ifneq ($(WARN_SHADOW),FALSE)
+    warning_flags += -Wshadow
+  endif
+
+  CXXFLAGS += $(warning_flags) -Woverloaded-virtual
+  CFLAGS += $(warning_flags)
+endif
+
+ifeq ($(WARN_ERROR),TRUE)
+  CXXFLAGS += -Werror
+  CFLAGS += -Werror
+endif
+
 ########################################################################
 
 ifdef CXXSTD

Tools/GNUMake/comps/llvm.mak

@@ -38,10 +38,11 @@ else
 
 endif
 
+CXXFLAGS += -Wno-pass-failed  # disable this warning
+
 ifeq ($(WARN_ALL),TRUE)
   warning_flags = -Wall -Wextra -Wno-sign-compare -Wunreachable-code -Wnull-dereference
   warning_flags += -Wfloat-conversion -Wextra-semi
-  warning_flags += -Wno-pass-failed  # disable this warning
 
   ifneq ($(USE_CUDA),TRUE)
     warning_flags += -Wpedantic