Update for ROCm 6.1.0 (#3898)

* A few functions used in Scan have been deprecated in 6.1.0. * Remove -Wno-deprecated-declarations from HIP CIs because we no longer need it. This will help us catch deprecated functions earlier. The flag was added because of atomicNoAdd, which has been handled by `clang diagnostic ignore` in the source code. * No need for -Wno-gnu-zero-variadic-macro-arguments in HIP CIs anymore.
AMReX-Codes · Apr 25, 2024 · b029e4f · b029e4f
1 parent aba9891
commit b029e4f
Show file tree

Hide file tree

Showing 2 changed files with 22 additions and 22 deletions.
diff --git a/.github/workflows/hip.yml b/.github/workflows/hip.yml
@@ -24,21 +24,11 @@ jobs:
         restore-keys: |
              ccache-${{ github.workflow }}-${{ github.job }}-git-
     - name: Build & Install
-      # Have to have -Wno-deprecated-declarations due to deprecated atomicAddNoRet
-      # Have to have -Wno-gnu-zero-variadic-macro-arguments to avoid
-      #    amrex/Src/Base/AMReX_GpuLaunchGlobal.H:15:5: error: must specify at least one argument for '...' parameter of variadic macro [-Werror,-Wgnu-zero-variadic-macro-arguments]
-      #        __launch_bounds__(amrex_launch_bounds_max_threads)
-      #        ^
-      #    /opt/rocm-4.1.1/hip/include/hip/hcc_detail/hip_runtime.h:178:71: note: expanded from macro '__launch_bounds__'
-      #        select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__)
-      #                                                                          ^
-      #    /opt/rocm-4.1.1/hip/include/hip/hcc_detail/hip_runtime.h:176:9: note: macro 'select_impl_' defined here
-      #    #define select_impl_(_1, _2, impl_, ...) impl_
       # Have to remove "-fno-operator-names to avoid
       #    /opt/rocm-6.1.0/include/rocprim/device/detail/device_adjacent_difference.hpp:198:26: error: token is not a valid binary operator in a preprocessor subexpression
       #      198 | #if defined(__gfx1102__) or defined(__gfx1030__)
       #          |     ~~~~~~~~~~~~~~~~~~~~~^~
-      env: {CXXFLAGS: "-Werror -Wall -Wextra -Wpedantic -Wnull-dereference -Wfloat-conversion -Wshadow -Woverloaded-virtual -Wextra-semi -Wunreachable-code -Wnon-virtual-dtor -Wno-deprecated-declarations -Wno-gnu-zero-variadic-macro-arguments"}
+      env: {CXXFLAGS: "-Werror -Wall -Wextra -Wpedantic -Wnull-dereference -Wfloat-conversion -Wshadow -Woverloaded-virtual -Wextra-semi -Wunreachable-code -Wnon-virtual-dtor"}
       run: |
         export CCACHE_COMPRESS=1
         export CCACHE_COMPRESSLEVEL=10
@@ -92,21 +82,11 @@ jobs:
         restore-keys: |
              ccache-${{ github.workflow }}-${{ github.job }}-git-
     - name: Build & Install
-      # Have to have -Wno-deprecated-declarations due to deprecated atomicAddNoRet
-      # Have to have -Wno-gnu-zero-variadic-macro-arguments to avoid
-      #    amrex/Src/Base/AMReX_GpuLaunchGlobal.H:15:5: error: must specify at least one argument for '...' parameter of variadic macro [-Werror,-Wgnu-zero-variadic-macro-arguments]
-      #        __launch_bounds__(amrex_launch_bounds_max_threads)
-      #        ^
-      #    /opt/rocm-4.1.1/hip/include/hip/hcc_detail/hip_runtime.h:178:71: note: expanded from macro '__launch_bounds__'
-      #        select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__)
-      #                                                                          ^
-      #    /opt/rocm-4.1.1/hip/include/hip/hcc_detail/hip_runtime.h:176:9: note: macro 'select_impl_' defined here
-      #    #define select_impl_(_1, _2, impl_, ...) impl_
       # Have to remove "-fno-operator-names to avoid
       #    /opt/rocm-6.1.0/include/rocprim/device/detail/device_adjacent_difference.hpp:198:26: error: token is not a valid binary operator in a preprocessor subexpression
       #      198 | #if defined(__gfx1102__) or defined(__gfx1030__)
       #          |     ~~~~~~~~~~~~~~~~~~~~~^~
-      env: {CXXFLAGS: "-Werror -Wall -Wextra -Wpedantic -Wnull-dereference -Wfloat-conversion -Wshadow -Woverloaded-virtual -Wextra-semi -Wunreachable-code -Wnon-virtual-dtor -Wno-deprecated-declarations -Wno-gnu-zero-variadic-macro-arguments"}
+      env: {CXXFLAGS: "-Werror -Wall -Wextra -Wpedantic -Wnull-dereference -Wfloat-conversion -Wshadow -Woverloaded-virtual -Wextra-semi -Wunreachable-code -Wnon-virtual-dtor"}
       run: |
         export CCACHE_COMPRESS=1
         export CCACHE_COMPRESSLEVEL=10

diff --git a/Src/Base/AMReX_Scan.H b/Src/Base/AMReX_Scan.H
@@ -641,13 +641,33 @@ T PrefixSum (N n, FIN const& fin, FOUT const& fout, TYPE, RetSum a_ret_sum = ret
     using ScanTileState = rocprim::detail::lookback_scan_state<T>;
     using OrderedBlockId = rocprim::detail::ordered_block_id<unsigned int>;
 
+#if (defined(HIP_VERSION_MAJOR) && (HIP_VERSION_MAJOR < 6)) || \
+    (defined(HIP_VERSION_MAJOR) && (HIP_VERSION_MAJOR == 6) && \
+     defined(HIP_VERSION_MINOR) && (HIP_VERSION_MINOR == 0))
+
     std::size_t nbytes_tile_state = rocprim::detail::align_size
         (ScanTileState::get_storage_size(nblocks));
     std::size_t nbytes_block_id = OrderedBlockId::get_storage_size();
 
     auto dp = (char*)(The_Arena()->alloc(nbytes_tile_state+nbytes_block_id));
 
     ScanTileState tile_state = ScanTileState::create(dp, nblocks);
+
+#else
+
+    std::size_t nbytes_tile_state;
+    AMREX_HIP_SAFE_CALL(ScanTileState::get_storage_size(nblocks, stream, nbytes_tile_state));
+    nbytes_tile_state = rocprim::detail::align_size(nbytes_tile_state);
+
+    std::size_t nbytes_block_id = OrderedBlockId::get_storage_size();
+
+    auto dp = (char*)(The_Arena()->alloc(nbytes_tile_state+nbytes_block_id));
+
+    ScanTileState tile_state;
+    AMREX_HIP_SAFE_CALL(ScanTileState::create(tile_state, dp, nblocks, stream));
+
+#endif
+
     auto ordered_block_id = OrderedBlockId::create
         (reinterpret_cast<OrderedBlockId::id_type*>(dp + nbytes_tile_state));