[rocPRIM] merge back 6.4 changes

CHANGELOG.md

@@ -72,6 +72,7 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec
 * Fixed incorrect 128-bit signed and unsigned integers type traits.
 * Fixed compilation issue when `rocprim::radix_key_codec<...>` is specialized with a 128-bit integer.
 * Fixed the warp-level reduction `rocprim::warp_reduce.reduce` DPP implementation to avoid undefined intermediate values during the reduction.
+* Fixed an issue that caused a segmentation fault when `hipStreamLegacy` was passed to some API functions.
 
 ### Upcoming changes
 * Using the initialisation constructor of `rocprim::reverse_iterator` will throw a deprecation warning. It will be marked as explicit in the next major release.

rocprim/include/rocprim/block/detail/block_histogram_sort.hpp

@@ -103,14 +103,10 @@ class block_histogram_sort
         ::rocprim::syncthreads(); // Fix race condition that appeared on Vega10 hardware, storage LDS is reused below.
 
         ROCPRIM_UNROLL
-        for(unsigned int offset = 0; offset < Bins; offset += BlockSize)
+        for(unsigned int offset = 0; offset + flat_tid < Bins; offset += BlockSize)
         {
-            const unsigned int offset_tid = offset + flat_tid;
-            if(offset_tid < Bins)
-            {
-                storage_.start[offset_tid] = tile_size;
-                storage_.end[offset_tid] = tile_size;
-            }
+            storage_.start[offset + flat_tid] = tile_size;
+            storage_.end[offset + flat_tid] = tile_size;
         }
         ::rocprim::syncthreads();
 

rocprim/include/rocprim/config.hpp

@@ -127,6 +127,39 @@
     #define ROCPRIM_TARGET_GCN3 1
 #endif
 
+#undef ROCPRIM_TARGET_GCN3
+#undef ROCPRIM_TARGET_GCN5
+#undef ROCPRIM_TARGET_RDNA1
+#undef ROCPRIM_TARGET_RDNA2
+#undef ROCPRIM_TARGET_RDNA3
+#undef ROCPRIM_TARGET_RDNA4
+#undef ROCPRIM_TARGET_CDNA1
+#undef ROCPRIM_TARGET_CDNA2
+#undef ROCPRIM_TARGET_CDNA3
+
+// See https://llvm.org/docs/AMDGPUUsage.html#instructions
+#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+    #define ROCPRIM_TARGET_CDNA3 1
+#elif defined(__gfx90a__)
+    #define ROCPRIM_TARGET_CDNA2 1
+#elif defined(__gfx908__)
+    #define ROCPRIM_TARGET_CDNA1 1
+#elif defined(__gfx900__) || defined(__gfx902__) || defined(__gfx904__) || defined(__gfx906__) \
+    || defined(__gfx90c__)
+    #define ROCPRIM_TARGET_GCN5 1
+#elif defined(__GFX12__)
+    #define ROCPRIM_TARGET_RDNA4 1
+#elif defined(__GFX11__)
+    #define ROCPRIM_TARGET_RDNA3 1
+#elif defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || defined(__gfx1033__) \
+    || defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__)
+    #define ROCPRIM_TARGET_RDNA2 1
+#elif defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__) || defined(__gfx1013__)
+    #define ROCPRIM_TARGET_RDNA1 1
+#elif defined(__GFX8__)
+    #define ROCPRIM_TARGET_GCN3 1
+#endif
+
 // DPP is supported only after Volcanic Islands (GFX8+)
 // Only defined when support is present, in contrast to ROCPRIM_DETAIL_USE_DPP, which should be
 // always defined

rocprim/include/rocprim/device/config_types.hpp

@@ -154,6 +154,8 @@ enum class target_arch : unsigned int
     gfx1030 = 1030,
     gfx1100 = 1100,
     gfx1102 = 1102,
+    gfx1200 = 1200,
+    gfx1201 = 1201,
     unknown = std::numeric_limits<unsigned int>::max(),
 };
 #endif // DOXYGEN_SHOULD_SKIP_THIS
@@ -193,7 +195,9 @@ constexpr target_arch get_target_arch_from_name(const char* const arch_name, con
                                                     "gfx942",
                                                     "gfx1030",
                                                     "gfx1100",
-                                                    "gfx1102"};
+                                                    "gfx1102",
+                                                    "gfx1200",
+                                                    "gfx1201"};
     constexpr target_arch target_architectures[] = {
         target_arch::gfx803,
         target_arch::gfx900,
@@ -204,6 +208,8 @@ constexpr target_arch get_target_arch_from_name(const char* const arch_name, con
         target_arch::gfx1030,
         target_arch::gfx1100,
         target_arch::gfx1102,
+        target_arch::gfx1200,
+        target_arch::gfx1201,
     };
     static_assert(sizeof(target_names) / sizeof(target_names[0])
                       == sizeof(target_architectures) / sizeof(target_architectures[0]),
@@ -264,6 +270,10 @@ auto dispatch_target_arch(const target_arch target_arch)
             return Config::template architecture_config<target_arch::gfx1100>::params;
         case target_arch::gfx1102:
             return Config::template architecture_config<target_arch::gfx1102>::params;
+        case target_arch::gfx1200:
+            return Config::template architecture_config<target_arch::gfx1200>::params;
+        case target_arch::gfx1201:
+            return Config::template architecture_config<target_arch::gfx1201>::params;
         case target_arch::invalid:
             assert(false && "Invalid target architecture selected at runtime.");
     }
@@ -321,7 +331,15 @@ inline hipError_t get_device_arch(int device_id, target_arch& arch)
 inline hipError_t get_device_from_stream(const hipStream_t stream, int& device_id)
 {
     static constexpr hipStream_t default_stream = 0;
-    if(stream == default_stream || stream == hipStreamPerThread || stream == hipStreamLegacy)
+
+    // hipStreamLegacy is supported in HIP >= 6.1.0
+#if (HIP_VERSION_MAJOR >= 6 && HIP_VERSION_MINOR >= 1)
+    const bool is_legacy_stream = (stream == hipStreamLegacy);
+#else
+    const bool is_legacy_stream = false;
+#endif
+
+    if (stream == default_stream || stream == hipStreamPerThread || is_legacy_stream);
     {
         const hipError_t result = hipGetDevice(&device_id);
         if(result != hipSuccess)

rocprim/include/rocprim/device/detail/config/device_adjacent_find.hpp

@@ -201,6 +201,142 @@ struct default_adjacent_find_config<
     : adjacent_find_config<1024, 4>
 {};
 
+// Based on input_type = double
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1200),
+    input_type,
+    std::enable_if_t<(bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 8) && (sizeof(input_type) > 4))>>
+    : adjacent_find_config<64, 2>
+{};
+
+// Based on input_type = float
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1200),
+    input_type,
+    std::enable_if_t<(bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 4) && (sizeof(input_type) > 2))>>
+    : adjacent_find_config<128, 4>
+{};
+
+// Based on input_type = rocprim::half
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1200),
+    input_type,
+    std::enable_if_t<(bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 2))>> : adjacent_find_config<64, 64>
+{};
+
+// Based on input_type = int64_t
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1200),
+    input_type,
+    std::enable_if_t<(!bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 8) && (sizeof(input_type) > 4))>>
+    : adjacent_find_config<64, 4>
+{};
+
+// Based on input_type = int
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1200),
+    input_type,
+    std::enable_if_t<(!bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 4) && (sizeof(input_type) > 2))>>
+    : adjacent_find_config<128, 32>
+{};
+
+// Based on input_type = short
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1200),
+    input_type,
+    std::enable_if_t<(!bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 2) && (sizeof(input_type) > 1))>>
+    : adjacent_find_config<256, 32>
+{};
+
+// Based on input_type = int8_t
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1200),
+    input_type,
+    std::enable_if_t<(!bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 1))>> : adjacent_find_config<256, 64>
+{};
+
+// Based on input_type = double
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1201),
+    input_type,
+    std::enable_if_t<(bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 8) && (sizeof(input_type) > 4))>>
+    : adjacent_find_config<64, 2>
+{};
+
+// Based on input_type = float
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1201),
+    input_type,
+    std::enable_if_t<(bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 4) && (sizeof(input_type) > 2))>>
+    : adjacent_find_config<64, 32>
+{};
+
+// Based on input_type = rocprim::half
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1201),
+    input_type,
+    std::enable_if_t<(bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 2))>> : adjacent_find_config<256, 16>
+{};
+
+// Based on input_type = int64_t
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1201),
+    input_type,
+    std::enable_if_t<(!bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 8) && (sizeof(input_type) > 4))>>
+    : adjacent_find_config<512, 2>
+{};
+
+// Based on input_type = int
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1201),
+    input_type,
+    std::enable_if_t<(!bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 4) && (sizeof(input_type) > 2))>>
+    : adjacent_find_config<64, 16>
+{};
+
+// Based on input_type = short
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1201),
+    input_type,
+    std::enable_if_t<(!bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 2) && (sizeof(input_type) > 1))>>
+    : adjacent_find_config<256, 2>
+{};
+
+// Based on input_type = int8_t
+template<class input_type>
+struct default_adjacent_find_config<
+    static_cast<unsigned int>(target_arch::gfx1201),
+    input_type,
+    std::enable_if_t<(!bool(rocprim::is_floating_point<input_type>::value)
+                      && (sizeof(input_type) <= 1))>> : adjacent_find_config<128, 64>
+{};
+
 // Based on input_type = double
 template<class input_type>
 struct default_adjacent_find_config<