Set AMD wavefront size according to AMReX_AMD_ARCH

.github/workflows/hip.yml

@@ -59,7 +59,7 @@ jobs:
             -DAMReX_FORTRAN=ON                            \
             -DAMReX_LINEAR_SOLVERS=ON                     \
             -DAMReX_GPU_BACKEND=HIP                       \
-            -DAMReX_AMD_ARCH=gfx908                       \
+            -DAMReX_AMD_ARCH=gfx1010                      \
             -DAMReX_ROCTX=ON                              \
             -DCMAKE_C_COMPILER=$(which clang)             \
             -DCMAKE_CXX_COMPILER=$(which clang++)         \

Src/Base/AMReX_GpuDevice.H

@@ -164,11 +164,8 @@ public:
 
     static int memoryPoolsSupported () noexcept { return memory_pools_supported; }
 
-#if defined(AMREX_USE_HIP) && (HIP_VERSION_MAJOR >= 4)
-    // definition: https://github.com/llvm/llvm-project/blob/62ec4ac90738a5f2d209ed28c822223e58aaaeb7/clang/lib/Basic/Targets/AMDGPU.cpp#L400
-    // overview wavefront size: https://github.com/llvm/llvm-project/blob/efc063b621ea0c4d1e452bcade62f7fc7e1cc937/clang/test/Driver/amdgpu-macros.cl#L70-L115
-    // gfx10XX has 32 threads per wavefront else 64
-    static AMREX_EXPORT constexpr int warp_size = __AMDGCN_WAVEFRONT_SIZE;
+#if defined(AMREX_USE_HIP)
+    static AMREX_EXPORT constexpr int warp_size = AMREX_AMDGCN_WAVEFRONT_SIZE;
 #elif defined(AMREX_USE_SYCL)
     static AMREX_EXPORT constexpr int warp_size = AMREX_SYCL_SUB_GROUP_SIZE;
 #else

Src/Base/AMReX_GpuDevice.cpp

@@ -53,6 +53,20 @@ namespace {
 }
 #endif
 
+#ifdef AMREX_USE_HIP
+namespace {
+    __host__ __device__ void amrex_check_wavefront_size () {
+#ifdef __HIP_DEVICE_COMPILE__
+        // https://github.com/AMReX-Codes/amrex/issues/3792
+        // __AMDGCN_WAVEFRONT_SIZE is valid in device code only.
+        // Thus we have to check it this way.
+        static_assert(__AMDGCN_WAVEFRONT_SIZE == AMREX_AMDGCN_WAVEFRONT_SIZE,
+                      "Please let the amrex team know if you encounter this");
+#endif
+    }
+}
+#endif
+
 namespace amrex::Gpu {
 
 int Device::device_id = 0;
@@ -308,6 +322,10 @@ Device::Initialize ()
     nvtxRangePop();
 #endif
 
+#if defined(AMREX_USE_HIP)
+    amrex::single_task(amrex_check_wavefront_size);
+#endif
+
     Device::profilerStart();
 
 #endif /* AMREX_USE_GPU */

Tools/CMake/AMReXSetDefines.cmake

@@ -59,8 +59,15 @@ if (AMReX_SYCL)
 endif()
 
 # HIP
-add_amrex_define( AMREX_USE_HIP NO_LEGACY IF AMReX_HIP )
-add_amrex_define( NDEBUG IF AMReX_HIP)  # This address a bug that causes slow build times
+if (AMReX_HIP)
+   add_amrex_define( AMREX_USE_HIP NO_LEGACY )
+   add_amrex_define( NDEBUG )  # This address a bug that causes slow build times
+   if (${AMReX_AMD_ARCH} MATCHES "gfx10.*")
+      add_amrex_define( AMREX_AMDGCN_WAVEFRONT_SIZE=32 NO_LEGACY )
+   else ()
+      add_amrex_define( AMREX_AMDGCN_WAVEFRONT_SIZE=64 NO_LEGACY )
+   endif ()
+endif()
 
 # Precision
 if (AMReX_PRECISION STREQUAL "SINGLE")

Tools/CMake/AMReX_Config_ND.H.in

@@ -46,6 +46,7 @@
 #cmakedefine AMREX_USE_EB
 #cmakedefine AMREX_USE_CUDA
 #cmakedefine AMREX_USE_HIP
+#cmakedefine AMREX_AMDGCN_WAVEFRONT_SIZE @AMREX_AMDGCN_WAVEFRONT_SIZE@
 #cmakedefine AMREX_USE_NVML
 #cmakedefine AMREX_GPU_MAX_THREADS @AMREX_GPU_MAX_THREADS@
 #ifndef AMREX_GPU_MAX_THREADS

Tools/GNUMake/Make.defs

@@ -1103,6 +1103,13 @@ else ifeq ($(USE_HIP),TRUE)
       AMD_ARCH = $(AMREX_AMD_ARCH)
     endif
 
+    # For AMD GPUs, the wavefront is 64 except for gfx10??.
+    ifeq ($(findstring gfx10,$(AMD_ARCH)),)
+       DEFINES += -DAMREX_AMDGCN_WAVEFRONT_SIZE=64
+    else
+       DEFINES += -DAMREX_AMDGCN_WAVEFRONT_SIZE=32
+    endif
+
     ifeq ($(HIP_SAVE_TEMPS),TRUE)
       ifeq ($(USE_GPU_RDC),TRUE)
         $(warning *** HIP_SAVE_TEMPS requires USE_GPU_RDC=FALSE to obtain the assembly files for AMD GPU kernels.)