Repository navigation

feat(hipccl): hipCUB's parity with CCCL 3.1.x #13083

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open

Beanavil wants to merge 66 commits into develop from users/beanavil/hipccl/hipcub-cccl-3-1-parity-upstream

+6,384 −2,546

projects/hipccl/hipccl3/hipcub/CHANGELOG.md

-Original file line number
+Diff line change
@@ Expand Up @@
     ## hipCUB-5.1.0 for ROCm 10.3.0
+    ### Added
+    * Added support for large number of segments and large number of items to `hipcub::DeviceSegmentedRadixSort::*`.
+    * Added new `hipcub::WarpReduce` overloadings `hipcub::WarpReduce::Sum`, `hipcub::WarpReduce::Max` and `hipcub::WarpReduce::Min`.
+    * Added fixed-size-segments variant for `hipcub::DeviceSegmentedReduce`.
+    * Added support for large number of segments to `hipcub::DeviceSegmentedReduce::Reduce`'s fixed-size-segments variant.
+    * Added support for large sizes to `hipcub::DeviceRunLengthEncode::Encode`.
+    * Extended fixed-size-segments variant of `hipcub::DeviceSegmentedReduce` to support all operators (`Sum`, `Min/Max`, and `ArgMin/Max`).
+    * Added support for large sizes to `hipcub::DeviceRunLengthEncode::NonTrivialRuns`.
+    * Added env-based API for `hipcub::DeviceReduce::Reduce`.
+    * Added env-based API for `hipcub::DeviceReduce::Sum`.
+    * Feature parity with CCCL/CUB 3.1.4.
     ### Resolved issues
     * Fixed a build failure when libhipcxx is not on the include path. `radix_rank_sort_operations.hpp` included the libhipcxx `bit` header unconditionally, which does not exist in the standard library fallback.
     ## hipCUB-5.0.0 for ROCm 10.2.0
+    ### Changed
+    * Changed `CCCL_MINIMUM_VERSION` to `3.1.4` to align with CUB.
+    ### Removed
+    * Removed the deprecated thread operators `Equality`, `Inequality`, `Sum`, `Difference`, `Division`, `Max` and `Min`.
+    ## hipCUB-5.0.0 for ROCm 10.2.0
     ### Added
     * Feature parity with CCCL/CUB 3.0.0.
@@ Expand Down @@

projects/hipccl/hipccl3/hipcub/CMakeLists.txt

            
                      Original file line number
                      Diff line number
                      Diff line change
                  
    @@ -50,8 +50,8 @@ endif()
  
    set(VERSION_STRING "5.0.0")

    # Set the CCCL-compatible version.

    set(HIPCUB_CCCL_VERSION_MAJOR 3)

    set(HIPCUB_CCCL_VERSION_MINOR 0)

    set(HIPCUB_CCCL_VERSION_PATCH 3)

    set(HIPCUB_CCCL_VERSION_MINOR 1)

    set(HIPCUB_CCCL_VERSION_PATCH 4)

    # Set the minimum required rocPRIM version

    set(MIN_ROCPRIM_PACKAGE_VERSION "5.0.0" CACHE STRING "Minimum version of rocPRIM to search for when ROCPRIM_FETCH_METHOD is set to PACKAGE.")

    # Set download branch for dependency rocPRIM

    @@ -136,13 +136,22 @@ set(ROCPRIM_FETCH_METHOD "PACKAGE" CACHE STRING "How to obtain the rocPRIM depen
  
    # This function checks to see if the fetch method variable it's passed is defined, and contains a valid value.

    # If it does not contain a valid value, it issues a fatal failure with an error message.

    function(check_fetch_method method)

      if (DEFINED ${method} AND NOT ${${method}} IN_LIST FETCH_METHOD_OPTIONS)

        message(FATAL_ERROR "Unrecognized ${method}: \"${${method}}\". Valid options are: ${FETCH_METHOD_OPTIONS}.")

    function(check_fetch_method method fetch_method_options)

      if (DEFINED ${method} AND NOT ${${method}} IN_LIST ${fetch_method_options})

        message(FATAL_ERROR "Unrecognized ${method}: \"${${method}}\". Valid options are: ${fetch_method_options}.")

      endif()

    endfunction()

    check_fetch_method(ROCPRIM_FETCH_METHOD)

    check_fetch_method(ROCPRIM_FETCH_METHOD FETCH_METHOD_OPTIONS)

    # TODO: update code below to use the above fetch options when that is possible

    # Set up options for obtaining dependency libhipcxx

    # PACKAGE: Search for an install package that contains the dependency.

    set(LIBHIPCXX_FETCH_METHOD_OPTIONS "PACKAGE")

    set(LIBHIPCXX_FETCH_METHOD "PACKAGE" CACHE STRING "How to obtain the libhipcxx dependency")

    check_fetch_method(LIBHIPCXX_FETCH_METHOD LIBHIPCXX_FETCH_METHOD_OPTIONS)

    # Add hipCUB's CMake modules

    list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")

projects/hipccl/hipccl3/hipcub/README.md

-Original file line number
+Diff line change
@@ Expand Up / @@ -69,6 +69,7 @@ mkdir build; cd build @@
     #   BUILD_TEST                   - OFF by default,
     #   BUILD_BENCHMARK              - OFF by default.
     #   ROCPRIM_FETCH_METHOD         - One of PACKAGE (default), DOWNLOAD, and MONOREPO. See below for a description of each.
+    #   LIBHIPCXX_FETCH_METHOD       - One of PACKAGE (default). See below for a description of each.
     #   EXTERNAL_DEPS_FORCE_DOWNLOAD - OFF by default, forces download for non-ROCm dependencies (e.g., Google Test).
     #   DOWNLOAD_CUB                 - OFF by default, (Nvidia CUB backend only) forces download of CUB instead of searching for an installed package.
     #   BUILD_OFFLOAD_COMPRESS       - ON by default, compresses device code to reduce the size of the generated binary.
@@ Expand Down Expand Up / @@ -105,6 +106,9 @@ make package @@
     * `DOWNLOAD` - Clones rocPRIM from the upstream repository. If git >= 2.25 is present, this option uses a sparse checkout that avoids downloading more than it needs to. If not, the whole monorepo is downloaded (this may take some time).
     * `MONOREPO` - This value is intended to be used if you are building hipCUB from within a copy of the rocm-libraries repository that you have cloned (and therefore already contains rocPRIM). When selected, the build will try find the dependency in the local repository tree. If it cannot be found, the build will attempt to use git to perform a sparse-checkout of rocPRIM. If that also fails, it will fall back to using the `DOWNLOAD` option described above.
+    `LIBHIPCXX_FETCH_METHOD` can be used to control how hipCUB obtains the libhipcxx dependency. It must be set to one of the following values:
+    * `PACKAGE` (default) - Searches for an installed package on the system that meets the minimum version requirement.
     ### HIP on Windows
     Initial support for HIP on Windows is available. You can install it using the provided `rmake.py` Python
@@ Expand Down @@

projects/hipccl/hipccl3/hipcub/benchmark/benchmark_block_scan.cpp

-Original file line number
+Diff line change
@@ Expand Up / @@ -43,9 +43,8 @@ struct inclusive_scan @@
         template<class T, unsigned int BlockSize, unsigned int ItemsPerThread>
         __device__
-        static void run(const T* input, T* output, const T init)
+        static void run(const T* input, T* output, [[maybe_unused]] const T init)
         {
-            (void)init;
             const unsigned int i = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
             T values[ItemsPerThread];
@@ Expand Down @@

projects/hipccl/hipccl3/hipcub/benchmark/benchmark_device_memory.cpp

-Original file line number
+Diff line change
@@ Expand Up @@
         using storage_type = empty_storage_type;
         HIPCUB_DEVICE
-        inline void
-            operator()(storage_type& storage,
-                       T (&input)[ItemsPerThread],
-                       T* global_mem_output = nullptr) const
+        inline void operator()([[maybe_unused]] storage_type& storage,
+                               T (&input)[ItemsPerThread],
+                               [[maybe_unused]] T* global_mem_output = nullptr) const
         {
-            (void)storage;
-            (void)global_mem_output;
             _CCCL_PRAGMA_UNROLL_FULL()
             for(unsigned int i = 0; i < ItemsPerThread; i++)
             {
@@ Expand All @@
         using storage_type = typename block_scan_type::TempStorage;
         HIPCUB_DEVICE
-        inline void
-            operator()(storage_type& storage,
-                       T (&input)[ItemsPerThread],
-                       T* global_mem_output = nullptr)
+        inline void operator()(storage_type& storage,
+                               T (&input)[ItemsPerThread],
+                               [[maybe_unused]] T* global_mem_output = nullptr)
         {
-            (void)global_mem_output;
             // sync before re-using shared memory from load
             __syncthreads();
             block_scan_type(storage).InclusiveScan(input, input, benchmark_utils::plus{});
@@ Expand All @@
         using storage_type = empty_storage_type;
         HIPCUB_DEVICE
-        inline void
-            operator()(storage_type& storage,
-                       T (&input)[ItemsPerThread],
-                       T* global_mem_output = nullptr)
+        inline void operator()([[maybe_unused]] storage_type& storage,
+                               [[maybe_unused]] T (&input)[ItemsPerThread],
+                               T* global_mem_output = nullptr)
         {
-            (void)storage;
-            (void)input;
             const unsigned int index
                 = threadIdx.x * ItemsPerThread + blockIdx.x * blockDim.x * ItemsPerThread;
             _CCCL_PRAGMA_UNROLL_FULL()
@@ Expand All @@
         using storage_type = empty_storage_type;
         HIPCUB_DEVICE
-        inline void
-            operator()(storage_type& storage,
-                       T (&input)[ItemsPerThread],
-                       T* global_mem_output = nullptr)
+        inline void operator()([[maybe_unused]] storage_type& storage,
+                               [[maybe_unused]] T (&input)[ItemsPerThread],
+                               T* global_mem_output = nullptr)
         {
-            (void)storage;
-            (void)input;
             const unsigned int index
                 = (threadIdx.x % warpSize) * ItemsPerThread + blockIdx.x * blockDim.x * ItemsPerThread;
             _CCCL_PRAGMA_UNROLL_FULL()
@@ Expand All @@
         using storage_type = empty_storage_type;
         HIPCUB_DEVICE
-        inline void
-            operator()(storage_type& storage,
-                       T (&input)[ItemsPerThread],
-                       T* global_mem_output = nullptr)
+        inline void operator()([[maybe_unused]] storage_type& storage,
+                               [[maybe_unused]] T (&input)[ItemsPerThread],
+                               T* global_mem_output = nullptr)
         {
-            (void)storage;
-            (void)input;
             const unsigned int index = threadIdx.x * ItemsPerThread;
             _CCCL_PRAGMA_UNROLL_FULL()
             for(unsigned int i = 0; i < ItemsPerThread; i++)
@@ Expand Down @@

projects/hipccl/hipccl3/hipcub/benchmark/benchmark_device_scan.cpp

            
                      Original file line number
                      Diff line number
                      Diff line change
                  
    @@ -55,17 +55,16 @@ auto run_device_scan(void*             temporary_storage,
  
    }

    template<bool Exclusive, class T, class BinaryFunction>

    auto run_device_scan(void*             temporary_storage,

                         size_t&           storage_size,

                         T*                input,

                         T*                output,

                         const T           initial_value,

                         const size_t      input_size,

                         BinaryFunction    scan_op,

                         const hipStream_t stream) ->

    auto run_device_scan(void*                    temporary_storage,

                         size_t&                  storage_size,

                         T*                       input,

                         T*                       output,

                         [[maybe_unused]] const T initial_value,

                         const size_t             input_size,

                         BinaryFunction           scan_op,

                         const hipStream_t        stream) ->

        typename std::enable_if<!Exclusive, hipError_t>::type

    {

        (void)initial_value;

        return hipcub::DeviceScan::InclusiveScan(temporary_storage,

                                                 storage_size,

                                                 input,

projects/hipccl/hipccl3/hipcub/benchmark/benchmark_device_segmented_radix_sort.cpp

            
                      Original file line number
                      Diff line number
                      Diff line change
                  
    @@ -48,8 +48,8 @@ class sort_keys_benchmark : public primbench::benchmark_interface
  
                                             size_t&,

                                             const Key*,

                                             Key*,

                                             int,

                                             int,

                                             _HIPCUB_STD::int64_t,

                                             _HIPCUB_STD::int64_t,

                                             offset_type*,

                                             offset_type*,

                                             int,

    @@ -164,8 +164,8 @@ class sort_pairs_benchmark : public primbench::benchmark_interface
  
                                             Key*,

                                             const Value*,

                                             Value*,

                                             int,

                                             int,

                                             _HIPCUB_STD::int64_t,

                                             _HIPCUB_STD::int64_t,

                                             offset_type*,

                                             offset_type*,

                                             int,

projects/hipccl/hipccl3/hipcub/benchmark/benchmark_warp_scan.cpp

-Original file line number
+Diff line change
@@ Expand Up / @@ -46,11 +46,9 @@ struct inclusive_scan @@
         template<class T, unsigned int WarpSize>
         __device__
-        static auto run(const T* input, T* output, const T init)
+        static auto run(const T* input, T* output, [[maybe_unused]] const T init)
             -> std::enable_if_t<benchmark_utils::device_test_enabled_for_warp_size_v<WarpSize>>
         {
-            (void)init;
             const unsigned int i     = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
             auto               value = input[i];
@@ Expand Down Expand Up / @@ -109,12 +107,10 @@ struct broadcast @@
         template<class T, unsigned int WarpSize>
         __device__
-        static auto run(const T* input, T* output, const T init)
+        static auto run(const T* input, T* output, [[maybe_unused]] const T init)
             -> std::enable_if_t<(benchmark_utils::device_test_enabled_for_warp_size_v<WarpSize>
                                  && benchmark_utils::is_power_of_two(WarpSize))>
         {
-            (void)init;
             const unsigned int i        = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
             const unsigned int warp_id  = i / WarpSize;
             const unsigned int src_lane = warp_id % WarpSize;
@@ Expand Down @@

projects/hipccl/hipccl3/hipcub/cmake/Dependencies.cmake

-Original file line number
+Diff line change
@@ Expand Up / @@ -426,7 +426,7 @@ endif(USER_BUILD_TEST) @@
     # CUB (only for CUDA platform)
     if(HIP_COMPILER STREQUAL "nvcc")
-      set(CCCL_MINIMUM_VERSION 3.0.0)
+      set(CCCL_MINIMUM_VERSION 3.1.4)
       if(NOT DOWNLOAD_CUB)
         find_package(CCCL ${CCCL_MINIMUM_VERSION} CONFIG)
       endif()
@@ Expand Down Expand Up / @@ -508,6 +508,20 @@ else() @@
           restore_cache_variable(BUILD_CODE_COVERAGE BOOL)
         endif()
       endif()
+      # libhipcxx (only for ROCm platform)
+      # TODO: use fetch_dep when it's possible
+      set(REPO_NAME "libhipcxx")
+      message(STATUS "Searching for ${REPO_NAME} package")
+      # Add default install location for WIN32 and non-WIN32 as hint
+      find_package(${REPO_NAME} CONFIG QUIET PATHS "${ROCM_ROOT}/lib/cmake/${REPO_NAME}")
+      if(NOT ${${REPO_NAME}_FOUND})
+        message(STATUS "No existing ${REPO_NAME} package was found.")
+      else()
+        message(STATUS "Package found (${${REPO_NAME}_DIR}) with version ${libhipcxx_VERSION}")
+      endif()
     endif()
     foreach(SHARED_OPTION BUILD_TEST BUILD_BENCHMARK BUILD_EXAMPLE)
@@ Expand Down @@

projects/hipccl/hipccl3/hipcub/cmake/Summary.cmake

-Original file line number
+Diff line change
@@ Expand Up / @@ -90,6 +90,7 @@ endif() @@
         message(STATUS "")
         message(STATUS "  EXTERNAL_DEPS_FORCE_DOWNLOAD: ${EXTERNAL_DEPS_FORCE_DOWNLOAD}")
         message(STATUS "  ROCPRIM_FETCH_METHOD        : ${ROCPRIM_FETCH_METHOD}")
+        message(STATUS "  LIBHIPCXX_FETCH_METHOD      : ${LIBHIPCXX_FETCH_METHOD}")
         message(STATUS "  BUILD_TEST                  : ${BUILD_TEST}")
     if(BUILD_TEST)
         message(STATUS "  CODE_COVERAGE               : ${CODE_COVERAGE}")
@@ Expand Down @@

projects/hipccl/hipccl3/hipcub/docs/install/build.rst

-Original file line number
+Diff line change
@@ Expand Up / @@ -98,6 +98,7 @@ The available build options are: @@
     * ``BUILD_OFFLOAD_COMPRESS``: Set this to ``OFF`` to prevent the ``--offload-compress`` switch from being passed to the compiler and compressing the binary. On by default.
     * ``USE_HIPCXX``: Set this to ``ON`` to build with CMake HIP language support. Setting this to ``ON`` eliminates the need to use ``CXX=hipcc``. Default is ``OFF``.
     * ``ROCPRIM_FETCH_METHOD``: Set this to the method to use to download rocPRIM. Can be set to ``PACKAGE``, ``DOWNLOAD``, or ``MONOREPO``. Set to ``MONOREPO`` if rocPRIM isn't already installed and you're building hipCUB from within a clone of the `rocm-libraries <https://github.com/ROCm/rocm-libraries/>`_ repository that also includes rocPRIM. Set to ``DOWNLOAD`` if rocPRIM isn't installed and you aren't in a clone of the ``rocm-libraries`` repository that includes rocPRIM. ``DOWNLOAD`` will clone the repository using sparse checkout so that only the necessary files are downloaded. Set to ``PACKAGE`` if rocPRIM is already installed. If you specify ``PACKAGE`` but rocPRIM isn't installed, the files will be downloaded using the same method as the ``DOWNLOAD`` option. The default method is ``PACKAGE``.
+    * ``LIBHIPCXX_FETCH_METHOD``: Set this to the method to use to download libhipcxx. Can be set to ``PACKAGE``.
     .. note::
@@ Expand Down @@

projects/hipccl/hipccl3/hipcub/examples/example_utils.hpp

-Original file line number
+Diff line change
@@ Expand Up / @@ -29,14 +29,17 @@ @@
     #ifndef EXAMPLES_EXAMPLE_UTILS_HPP
     #define EXAMPLES_EXAMPLE_UTILS_HPP
     #include "mersenne.h"
-    #include <vector>
-    #include <sstream>
-    #include <iostream>
+    #include <hipcub/config.hpp>
     #include <hipcub/util_allocator.hpp>
     #include <hipcub/util_type.hpp>
+    #include <iostream>
+    #include <sstream>
+    #include <vector>
     #include _HIPCUB_STD_INCLUDE(functional)
     #define AssertEquals(a, b) if ((a) != (b)) { std::cerr << "\n(" << __FILE__ << ": " << __LINE__ << ")\n"; exit(1);}
@@ Expand Down Expand Up / @@ -275,23 +278,38 @@ struct CommandLineArgs @@
                     exit(1);
                 }
-                device_giga_bandwidth = float(deviceProp.memoryBusWidth) * deviceProp.memoryClockRate * 2 / 8 / 1000 / 1000;
+                int memoryClockRate{};
+                error = HipcubDebug(
+                    hipDeviceGetAttribute(&memoryClockRate, hipDeviceAttributeClockRate, dev));
+                if(error)
+                {
+                    break;
+                }
+                int memoryBusWidth{};
+                error = HipcubDebug(
+                    hipDeviceGetAttribute(&memoryBusWidth, hipDeviceAttributeMemoryBusWidth, dev));
+                if(error)
+                {
+                    break;
+                }
+                device_giga_bandwidth = float(memoryBusWidth) * memoryClockRate * 2 / 8 / 1000 / 1000;
                 if (!CheckCmdLineFlag("quiet"))
                 {
-                    printf(
-                            "Using device %d: %s ( SM%d, %d SMs, "
-                            "%lld free / %lld total MB physmem, "
-                            "%.3f GB/s @ %d kHz mem clock, ECC %s)\n",
-                        dev,
-                        deviceProp.name,
-                        deviceProp.major * 100 + deviceProp.minor * 10,
-                        deviceProp.multiProcessorCount,
-                        (unsigned long long) device_free_physmem / 1024 / 1024,
-                        (unsigned long long) device_total_physmem / 1024 / 1024,
-                        device_giga_bandwidth,
-                        deviceProp.memoryClockRate,
-                        (deviceProp.ECCEnabled) ? "on" : "off");
+                    printf("Using device %d: %s ( SM%d, %d SMs, "
+                           "%lld free / %lld total MB physmem, "
+                           "%.3f GB/s @ %d kHz mem clock, ECC %s)\n",
+                           dev,
+                           deviceProp.name,
+                           deviceProp.major * 100 + deviceProp.minor * 10,
+                           deviceProp.multiProcessorCount,
+                           (unsigned long long)device_free_physmem / 1024 / 1024,
+                           (unsigned long long)device_total_physmem / 1024 / 1024,
+                           device_giga_bandwidth,
+                           memoryClockRate,
+                           (deviceProp.ECCEnabled) ? "on" : "off");
                     fflush(stdout);
                 }
@@ Expand Down @@

Navigation Menu

Repository navigation

feat(hipccl): hipCUB's parity with CCCL 3.1.x #13083

Diff view

Diff view

There are no files selected for viewing

Uh oh!

Uh oh!

feat(hipccl): hipCUB's parity with CCCL 3.1.x #13083

Are you sure you want to change the base?

feat(hipccl): hipCUB's parity with CCCL 3.1.x #13083

Uh oh!

Uh oh!

Diff view

Diff view

There are no files selected for viewing

Uh oh!

Uh oh!