Port back ordered block id feature to rocPRIM rocm 7.0.x

projects/rocprim/CHANGELOG.md

@@ -2,6 +2,13 @@
 
 Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projects/rocPRIM/en/latest/](https://rocm.docs.amd.com/projects/rocPRIM/en/latest/).
 
+## rocPRIM 4.0.2 for ROCm 7.0.2
+
+### Added
+
+* Atomic Ordered Block Id alternate implementation for lookback scan for optimal performance on MI3xx.
+* Added a new cmake option, `BUILD_OFFLOAD_COMPRESS`. When rocPRIM is build with this option enabled, the `--offload-compress` switch is passed to the compiler. This causes the compiler to compress the binary that it generates. Compression can be useful in cases where you are compiling for a large number of targets, since this often results in a large binary. Without compression, in some cases, the generated binary may become so large symbols are placed out of range, resulting in linking errors. The new `BUILD_OFFLOAD_COMPRESS` option is set to `ON` by default.
+
 ## rocPRIM 4.0.1 for ROCm 7.0.1
 
 ### Resolved issues

projects/rocprim/CMakeLists.txt

@@ -73,10 +73,22 @@ cmake_dependent_option(BUILD_DOCS "Build documentation (requires sphinx)" OFF "N
 option(BUILD_CODE_COVERAGE "Build with code coverage enabled" OFF)
 option(ROCPRIM_INSTALL "Enable installation of rocPRIM (projects embedding rocPRIM may want to turn this OFF)" ON)
 option(ROCPRIM_ENABLE_ASSERTS "Enable asserts in release build)" OFF)
+option(BUILD_OFFLOAD_COMPRESS "Build rocPRIM with offload compression" ON)
 
 check_language(HIP)
 cmake_dependent_option(USE_HIPCXX "Use CMake HIP language support" OFF CMAKE_HIP_COMPILER OFF)
 
+include(CheckCXXCompilerFlag)
+
+if(BUILD_OFFLOAD_COMPRESS)
+  check_cxx_compiler_flag("--offload-compress" CXX_COMPILER_SUPPORTS_OFFLOAD_COMPRESS)
+  if(CXX_COMPILER_SUPPORTS_OFFLOAD_COMPRESS)
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --offload-compress")
+  else()
+    message(STATUS "Warning: BUILD_OFFLOAD_COMPRESS=ON but flag not supported by compiler. Ignoring option.")
+  endif()
+endif()
+
 if (CMAKE_CURRENT_SOURCE_DIR STREQUAL CMAKE_SOURCE_DIR)
   set(ROCPRIM_PROJECT_IS_TOP_LEVEL TRUE)
 else()
@@ -149,13 +161,20 @@ else()
       )
     else()
       rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
-        TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201"
+        TARGETS "gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201"
       )
     endif()
     set(GPU_TARGETS "${DEFAULT_AMDGPU_TARGETS}" CACHE STRING "GPU architectures to compile for" FORCE)
   endif()
 endif()
 
+# Compressed offload binaries are currently not working with the SPIR-V target
+if("amdgcnspirv" IN_LIST GPU_TARGETS)
+  if(BUILD_OFFLOAD_COMPRESS)
+    message(FATAL_ERROR "Cannot combine SPIR-V and BUILD_OFFLOAD_COMPRESS")
+  endif()
+endif()
+
 # TODO: Fix VerifyCompiler for HIP on Windows
 if (NOT WIN32)
   include(cmake/VerifyCompiler.cmake)

projects/rocprim/cmake/Summary.cmake

@@ -86,13 +86,14 @@ function(print_configuration_summary)
     message(STATUS "  Device targets        : ${GPU_TARGETS}")
   endif()
   message(STATUS "")
-  message(STATUS "  ONLY_INSTALL          : ${ONLY_INSTALL}")
-  message(STATUS "  BUILD_TEST            : ${BUILD_TEST}")
-  message(STATUS "  WITH_ROCRAND          : ${WITH_ROCRAND}")
-  message(STATUS "  BUILD_BENCHMARK       : ${BUILD_BENCHMARK}")
-  message(STATUS "  BUILD_NAIVE_BENCHMARK : ${BUILD_NAIVE_BENCHMARK}")
-  message(STATUS "  BUILD_EXAMPLE         : ${BUILD_EXAMPLE}")
-  message(STATUS "  BUILD_DOCS            : ${BUILD_DOCS}")
+  message(STATUS "  ONLY_INSTALL           : ${ONLY_INSTALL}")
+  message(STATUS "  BUILD_TEST             : ${BUILD_TEST}")
+  message(STATUS "  WITH_ROCRAND           : ${WITH_ROCRAND}")
+  message(STATUS "  BUILD_BENCHMARK        : ${BUILD_BENCHMARK}")
+  message(STATUS "  BUILD_NAIVE_BENCHMARK  : ${BUILD_NAIVE_BENCHMARK}")
+  message(STATUS "  BUILD_EXAMPLE          : ${BUILD_EXAMPLE}")
+  message(STATUS "  BUILD_DOCS             : ${BUILD_DOCS}")
+  message(STATUS "  BUILD_OFFLOAD_COMPRESS : ${BUILD_OFFLOAD_COMPRESS}")
   message(STATUS "")
   message(STATUS "Detailed:")
   message(STATUS "  C++ compiler details  : \n${CMAKE_CXX_COMPILER_VERBOSE_DETAILS}")

projects/rocprim/rmake.py

@@ -41,9 +41,10 @@ def parse_args():
     parser.add_argument(      '--cmake-darg', required=False, dest='cmake_dargs', action='append', default=[],
                         help='List of additional cmake defines for builds (e.g. CMAKE_CXX_COMPILER_LAUNCHER=ccache)')
     parser.add_argument('-a', '--architecture', dest='gpu_architecture', required=False, default=default_gpus, #:sramecc+:xnack-" ) #gfx1030" ) #gfx906" ) # gfx1030" )
-                        help='Set GPU architectures, e.g. all, gfx000, gfx803, gfx906:xnack-;gfx1030;gfx1100 (optional, default: all)')
+                        help='Set GPU architectures, e.g. all, gfx000, gfx906:xnack-;gfx1030;gfx1100 (optional, default: all)')
     parser.add_argument('-v', '--verbose', required=False, default=False, action='store_true',
                         help='Verbose build (default: False)')
+    parser.add_argument('--no-offload-compress', required=False, default=False, action='store_true', help='Do not apply offload compression (deafult: False)')
     return parser.parse_args()
 
 def os_detect():
@@ -177,6 +178,9 @@ def config_cmd():
     if args.build_clients:
         cmake_options.append( f"-DBUILD_TEST=ON -DBUILD_BENCHMARK=ON -DBUILD_EXAMPLE=ON -DBUILD_DIR={build_dir}" )
 
+    if args.no_offload_compress:
+        cmake_options.append( f"-DBUILD_OFFLOAD_COMPRESS=OFF" )
+
     cmake_options.append( f"-DAMDGPU_TARGETS={args.gpu_architecture}" )
 
     if args.cmake_dargs:

projects/rocprim/rocprim/include/rocprim/config.hpp

@@ -139,8 +139,121 @@
     #define ROCPRIM_TARGET_UNKNOWN 1
 #endif
 
+#if defined(ROCPRIM_TARGET_SPIRV) && ROCPRIM_TARGET_SPIRV == 1
+    #define ROCPRIM_AMDGCN_CONSTEXPR
+#else
+    #define ROCPRIM_AMDGCN_CONSTEXPR constexpr
+#endif
+
+#if __has_builtin(__builtin_amdgcn_processor_is)
+    #if !defined(ROCPRIM_THREAD_LOAD_USE_CACHE_MODIFIERS)
+        #define ROCPRIM_THREAD_LOAD_USE_CACHE_MODIFIERS 1
+    #endif
+    #if !defined(ROCPRIM_THREAD_STORE_USE_CACHE_MODIFIERS)
+        #define ROCPRIM_THREAD_STORE_USE_CACHE_MODIFIERS 1
+    #endif
+    #define IS_CDNA3()                                                                     \
+        __builtin_amdgcn_processor_is("gfx942") || __builtin_amdgcn_processor_is("gfx950") \
+            || __builtin_amdgcn_processor_is("gfx9-4-generic")
+    #define IS_CDNA2() __builtin_amdgcn_processor_is("gfx90a")
+    #define IS_CDNA1() __builtin_amdgcn_processor_is("gfx908")
+    #define IS_GCN5()                                                                             \
+        __builtin_amdgcn_processor_is("gfx900") || __builtin_amdgcn_processor_is("gfx902")        \
+            || __builtin_amdgcn_processor_is("gfx904") || __builtin_amdgcn_processor_is("gfx906") \
+            || __builtin_amdgcn_processor_is("gfx90c")                                            \
+            || __builtin_amdgcn_processor_is("gfx9-generic")
+    #define IS_RDNA4()                                                                       \
+        __builtin_amdgcn_processor_is("gfx1200") || __builtin_amdgcn_processor_is("gfx1201") \
+        // TODO: enable when these structures are supported
+            /*|| __builtin_amdgcn_processor_is("gfx1250")                                      \
+            || __builtin_amdgcn_processor_is("gfx12-generic")*/
+    #define IS_RDNA3()                                                                       \
+        __builtin_amdgcn_processor_is("gfx1100") || __builtin_amdgcn_processor_is("gfx1101") \
+            || __builtin_amdgcn_processor_is("gfx1102")                                      \
+            || __builtin_amdgcn_processor_is("gfx1103")                                      \
+            || __builtin_amdgcn_processor_is("gfx11-generic")
+    #define IS_RDNA2()                                                                       \
+        __builtin_amdgcn_processor_is("gfx1030") || __builtin_amdgcn_processor_is("gfx1031") \
+            || __builtin_amdgcn_processor_is("gfx1032")                                      \
+            || __builtin_amdgcn_processor_is("gfx1033")                                      \
+            || __builtin_amdgcn_processor_is("gfx1034")                                      \
+            || __builtin_amdgcn_processor_is("gfx1035")                                      \
+            || __builtin_amdgcn_processor_is("gfx1036")                                      \
+            || __builtin_amdgcn_processor_is("gfx10-3-generic")
+    #define IS_RDNA1()                                                                       \
+        __builtin_amdgcn_processor_is("gfx1010") || __builtin_amdgcn_processor_is("gfx1011") \
+            || __builtin_amdgcn_processor_is("gfx1012")                                      \
+            || __builtin_amdgcn_processor_is("gfx1013")                                      \
+            || __builtin_amdgcn_processor_is("gfx10-1-generic")
+    #define IS_GCN3()                                                                             \
+        __builtin_amdgcn_processor_is("gfx801") || __builtin_amdgcn_processor_is("gfx802")        \
+            || __builtin_amdgcn_processor_is("gfx803") || __builtin_amdgcn_processor_is("gfx805") \
+            || __builtin_amdgcn_processor_is("gfx810")
+#else
+    #if defined(ROCPRIM_TARGET_CDNA3)
+        #define IS_CDNA3() 1
+    #else
+        #define IS_CDNA3() 0
+    #endif
+    #if defined(ROCPRIM_TARGET_CDNA2)
+        #define IS_CDNA2() 1
+    #else
+        #define IS_CDNA2() 0
+    #endif
+    #if defined(ROCPRIM_TARGET_CDNA1)
+        #define IS_CDNA1() 1
+    #else
+        #define IS_CDNA1() 0
+    #endif
+    #if defined(ROCPRIM_TARGET_GCN5)
+        #define IS_GCN5() 1
+    #else
+        #define IS_GCN5() 0
+    #endif
+    #if defined(ROCPRIM_TARGET_RDNA4)
+        #define IS_RDNA4() 1
+    #else
+        #define IS_RDNA4() 0
+    #endif
+    #if defined(ROCPRIM_TARGET_RDNA3)
+        #define IS_RDNA3() 1
+    #else
+        #define IS_RDNA3() 0
+    #endif
+    #if defined(ROCPRIM_TARGET_RDNA2)
+        #define IS_RDNA2() 1
+    #else
+        #define IS_RDNA2() 0
+    #endif
+    #if defined(ROCPRIM_TARGET_RDNA1)
+        #define IS_RDNA1() 1
+    #else
+        #define IS_RDNA1() 0
+    #endif
+    #if defined(ROCPRIM_TARGET_GCN3)
+        #define IS_GCN3() 1
+    #else
+        #define IS_GCN3() 0
+    #endif
+
+    #if !defined(ROCPRIM_THREAD_LOAD_USE_CACHE_MODIFIERS)
+        #if defined(ROCPRIM_TARGET_SPIRV)
+            #define ROCPRIM_THREAD_LOAD_USE_CACHE_MODIFIERS 0
+        #else
+            #define ROCPRIM_THREAD_LOAD_USE_CACHE_MODIFIERS 1
+        #endif
+    #endif
+    #if !defined(ROCPRIM_THREAD_STORE_USE_CACHE_MODIFIERS)
+        #if defined(ROCPRIM_TARGET_SPIRV)
+            #define ROCPRIM_THREAD_STORE_USE_CACHE_MODIFIERS 0
+        #else
+            #define ROCPRIM_THREAD_STORE_USE_CACHE_MODIFIERS 1
+        #endif
+    #endif
+#endif
+
 // SPIR-V and unknown targets do not support 128-bit atomics.
-#if defined(ROCPRIM_TARGET_UKNOWN) || defined(ROCPRIM_TARGET_SPIRV)
+#if defined(ROCPRIM_TARGET_UNKNOWN) || defined(ROCPRIM_TARGET_SPIRV)
     #define ROCPRIM_MAX_ATOMIC_SIZE 8
 #else
     #define ROCPRIM_MAX_ATOMIC_SIZE 16
@@ -150,34 +263,17 @@
 // Only defined when support is present, in contrast to ROCPRIM_DETAIL_USE_DPP, which should be
 // always defined
 #if defined(__HIP_DEVICE_COMPILE__) && defined(__AMDGCN__) \
-    && (!defined(__GFX6__) && !defined(__GFX7__))          \
-    && !(defined(ROCPRIM_TARGET_SPIRV) && ROCPRIM_TARGET_SPIRV == 1)
+    && (!defined(__GFX6__) && !defined(__GFX7__))
     #define ROCPRIM_DETAIL_HAS_DPP 1
 #endif
 
-#if(!defined(ROCPRIM_DISABLE_DPP) || ROCPRIM_DISABLE_DPP == 0) \
+#if (!defined(ROCPRIM_DISABLE_DPP) || ROCPRIM_DISABLE_DPP == 0) \
     && (defined(ROCPRIM_DETAIL_HAS_DPP) && ROCPRIM_DETAIL_HAS_DPP == 1)
     #define ROCPRIM_DETAIL_USE_DPP 1
 #else
     #define ROCPRIM_DETAIL_USE_DPP 0
 #endif
 
-#if defined(ROCPRIM_DETAIL_HAS_DPP) && (defined(__GFX8__) || defined(__GFX9__))
-    #define ROCPRIM_DETAIL_HAS_DPP_BROADCAST 1
-#endif
-
-#if defined(ROCPRIM_DETAIL_HAS_DPP) && (defined(__GFX8__) || defined(__GFX9__))
-    #define ROCPRIM_DETAIL_HAS_DPP_WF 1
-#endif
-
-#if !defined(ROCPRIM_THREAD_LOAD_USE_CACHE_MODIFIERS) && !defined(ROCPRIM_TARGET_SPIRV)
-    #define ROCPRIM_THREAD_LOAD_USE_CACHE_MODIFIERS 1
-#endif
-
-#if !defined(ROCPRIM_THREAD_STORE_USE_CACHE_MODIFIERS) && !defined(ROCPRIM_TARGET_SPIRV)
-    #define ROCPRIM_THREAD_STORE_USE_CACHE_MODIFIERS 1
-#endif
-
 #ifndef ROCPRIM_NAVI
     #if defined(__HIP_DEVICE_COMPILE__) \
         && (defined(__GFX10__) || defined(__GFX11__) || defined(__GFX12__))
@@ -196,7 +292,7 @@
 /// Quad size (group of 4 threads)
 #define ROCPRIM_QUAD_SIZE 4u
 
-#if(defined(_MSC_VER) && !defined(__clang__)) || (defined(__GNUC__) && !defined(__clang__))
+#if (defined(_MSC_VER) && !defined(__clang__)) || (defined(__GNUC__) && !defined(__clang__))
     #define ROCPRIM_UNROLL
     #define ROCPRIM_NO_UNROLL
 #else
@@ -275,4 +371,23 @@
     #define ROCPRIM_DETAIL_SUPPRESS_DEPRECATION_POP ROCPRIM_CLANG_SUPPRESS_WARNING_POP
 #endif
 
-#endif // ROCPRIM_CONFIG_HPP_
+#if __has_builtin(__builtin_amdgcn_is_invocable)
+    #define ROCPRIM_HAS_DPP() __builtin_amdgcn_is_invocable(__builtin_amdgcn_mov_dpp)
+    #define ROCPRIM_HAS_PERMLANE() __builtin_amdgcn_is_invocable(__builtin_amdgcn_permlane16)
+#elif defined(ROCPRIM_TARGET_SPIRV) && ROCPRIM_TARGET_SPIRV == 1
+    #define ROCPRIM_HAS_DPP() false
+    #define ROCPRIM_HAS_PERMLANE() false
+#else
+    #if defined(ROCPRIM_DETAIL_HAS_DPP) && ROCPRIM_DETAIL_HAS_DPP == 1
+        #define ROCPRIM_HAS_DPP() true
+    #else
+        #define ROCPRIM_HAS_DPP() false
+    #endif
+    #if defined(__GFX8__) || defined(__GFX9__)
+        #define ROCPRIM_HAS_PERMLANE() false
+    #else
+        #define ROCPRIM_HAS_PERMLANE() true
+    #endif
+#endif
+
+#endif // ROCPRIM_CONFIG_HPP_

projects/rocprim/rocprim/include/rocprim/detail/various.hpp

@@ -33,6 +33,11 @@
 
 #include <hip/hip_runtime.h>
 
+#include <chrono>
+#include <iostream>
+#include <type_traits>
+#include <variant>
+
 // Check for c++ standard library features, in a backwards compatible manner
 #ifndef __has_include
     #define __has_include(x) 0
@@ -498,6 +503,70 @@ inline float update_time_point(std::chrono::high_resolution_clock::time_point& t
     return delta_time;
 }
 
+template<typename T, T... Vs>
+struct constexpr_value_variant
+{
+    using variant = std::variant<std::integral_constant<T, Vs>...>;
+
+    static variant create(T value)
+    {
+        variant var{};
+        // Unfold over variadic enum values. For each value
+        // create and run a small lambda that sets our variant.
+        (
+            [&]
+            {
+                if(value == Vs)
+                {
+                    var = std::integral_constant<T, Vs>{};
+                }
+            }(),
+            ...);
+        return var;
+    }
+};
+
+template<typename... T>
+struct constexpr_type_variant
+{
+    using variant = std::variant<T...>;
+
+    static variant create(int index, T... values)
+    {
+        variant var{};
+        update_variant(index, var, std::make_tuple(values...), std::index_sequence_for<T...>{});
+        return var;
+    }
+
+    static variant create(int index)
+    {
+        return create(index, (T{}, ...));
+    }
+
+    template<typename F>
+    static variant create_with(int index, F f)
+    {
+        auto var = create(index, (T{}, ...));
+        std::visit(f, var);
+    };
+
+    static variant create(bool select, T... values)
+    {
+        return create(select ? 1 /* true */ : 0 /* false */, values...);
+    }
+
+private:
+    template<std::size_t... I>
+    static void update_variant(std::size_t      index,
+                               variant&         var,
+                               std::tuple<T...> t,
+                               std::index_sequence<I...> /* */)
+    {
+        ((index == I ? var = std::get<I>(t), 0 : 0), ...);
+    }
+};
+
+
 } // end namespace detail
 END_ROCPRIM_NAMESPACE