[hipCUB]Mergeback 6.4 hotfixes (#444)

* Include device_copy header (#433) The device/device_copy.hpp header was missing from the main hipcub.hpp header file. This change just adds it back. * [windows][hipCUB] Removed usage of std::unary_function and std::binary_function to prevent syntax error (#435) * removed ussage of std::unary/binary_function * updated changelog * Enable CMake HIP language (#434) If the HIP language can be used, the HIP_USECXX variable sets rocThrust to use the CMake HIP language rather than CXX. Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> * Add hipcub::AliasTemporaries and some macros (#438) * Resolve "Add missing macros in HIPCUB API" * Resolve "Add AliasTemporaries function" * Added macro names to CHANGELOG --------- Co-authored-by: Wayne Franz <wayfranz@amd.com> Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Nick Breed <78807921+NB4444@users.noreply.github.com>
ROCm · Dec 9, 2024 · d781413 · d781413
1 parent 91c65e3
commit d781413
Show file tree

Hide file tree

Showing 20 changed files with 744 additions and 46 deletions.
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -12,8 +12,11 @@ Full documentation for hipCUB is available at [https://rocm.docs.amd.com/project
 * Added `--emulation` option for `rtest.py`
   * Unit tests can be run with `[--emulation|-e|--test|-t]=<test_name>`
 * Added `DeviceSelect::FlaggedIf` and its inplace overload.
+* Added CUB macros missing from hipCUB: `HIPCUB_MAX`, `HIPCUB_MIN`, `HIPCUB_QUOTIENT_FLOOR`, `HIPCUB_QUOTIENT_CEILING`, `HIPCUB_ROUND_UP_NEAREST` and `HIPCUB_ROUND_DOWN_NEAREST`.
+* Added `hipcub::AliasTemporaries` function for CUB parity.
 
 ### Changed
+* Removed usage of `std::unary_function` and `std::binary_function` in `test_hipcub_device_adjacent_difference.cpp`
 * Changed the subset of tests that are run for smoke tests such that the smoke test will complete with faster run-time and to never exceed 2GB of vram usage. Use `python rtest.py [--emulation|-e|--test|-t]=smoke` to run these tests.
 * The `rtest.py` options have changed. `rtest.py` is now run with at least either `--test|-t` or `--emulation|-e`, but not both options.
 * The NVIDIA backend now requires CUB, Thrust and libcu++ 2.5.0. If it is not found it will be downloaded from the NVIDIA CCCL repository.

diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -29,13 +29,26 @@ set(CMAKE_INSTALL_PREFIX "/opt/rocm" CACHE PATH "Install path prefix, prepended
 # hipCUB project
 project(hipcub LANGUAGES CXX)
 
-# Set the ROCM install directory.
-if(WIN32)
-  set(ROCM_ROOT "$ENV{HIP_PATH}" CACHE PATH "Root directory of the ROCm installation")
-else()
-  set(ROCM_ROOT "/opt/rocm" CACHE PATH "Root directory of the ROCm installation")
+# Set CXX flags
+if (NOT DEFINED CMAKE_CXX_STANDARD)
+  set(CMAKE_CXX_STANDARD 17)
+endif()
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_CXX_EXTENSIONS OFF)
+
+if (CMAKE_CXX_STANDARD EQUAL 14)
+  message(WARNING "C++14 will be deprecated in the next major release")
+elseif(NOT CMAKE_CXX_STANDARD EQUAL 17)
+  message(FATAL_ERROR "Only C++14 and C++17 are supported")
 endif()
 
+# Set HIP flags
+set(CMAKE_HIP_STANDARD 14)
+set(CMAKE_HIP_STANDARD_REQUIRED ON)
+set(CMAKE_HIP_EXTENSIONS OFF)
+
+include(CheckLanguage)
+include(CMakeDependentOption)
 
 # Build options
 option(BUILD_TEST "Build tests (requires googletest)" OFF)
@@ -45,6 +58,17 @@ option(BUILD_BENCHMARK "Build benchmarks" OFF)
 option(BUILD_EXAMPLE "Build Examples" OFF)
 option(BUILD_ADDRESS_SANITIZER "Build with address sanitizer enabled" OFF)
 
+check_language(HIP)
+cmake_dependent_option(USE_HIPCXX "Use CMake HIP language support" OFF CMAKE_HIP_COMPILER OFF)
+
+# Set the ROCM install directory.
+if(WIN32)
+  set(ROCM_ROOT "$ENV{HIP_PATH}" CACHE PATH "Root directory of the ROCm installation")
+else()
+  set(ROCM_ROOT "/opt/rocm" CACHE PATH "Root directory of the ROCm installation")
+endif()
+
+
 # Set the header wrapper ON by default.
 option(BUILD_FILE_REORG_BACKWARD_COMPATIBILITY "Build with file/folder reorg with backward compatibility enabled" OFF)
 
@@ -60,45 +84,36 @@ endif()
 
 set(CMAKE_INSTALL_RPATH_USE_LINK_PATH TRUE CACHE BOOL "Add paths to linker search and installed rpath")
 
-# Set CXX flags
-if (NOT DEFINED CMAKE_CXX_STANDARD)
-  set(CMAKE_CXX_STANDARD 17)
-endif()
-set(CMAKE_CXX_STANDARD_REQUIRED ON)
-set(CMAKE_CXX_EXTENSIONS OFF)
-
-
-if (CMAKE_CXX_STANDARD EQUAL 14)
-  message(WARNING "C++14 will be deprecated in the next major release")
-elseif(NOT CMAKE_CXX_STANDARD EQUAL 17)
-  message(FATAL_ERROR "Only C++14 and C++17 are supported")
-endif()
-
 # rocm-cmake has to be included early so that it's available to set GPU_TARGETS
 # If hip is included prior to setting that then it defaults to building only for the current architecture
 include(ROCmCMakeBuildToolsDependency)
 
 # Setup GPU targets for rocm platform
-if(NOT (CMAKE_CXX_COMPILER MATCHES ".*nvcc$" OR "${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU"))
-  if(NOT DEFINED AMDGPU_TARGETS)
-    set(GPU_TARGETS "all" CACHE STRING "GPU architectures to compile for")
-  else()
-    set(GPU_TARGETS "${AMDGPU_TARGETS}" CACHE STRING "GPU architectures to compile for")
-  endif()
-  set_property(CACHE GPU_TARGETS PROPERTY STRINGS "all")
-
-  if(GPU_TARGETS STREQUAL "all")
-    if(BUILD_ADDRESS_SANITIZER)
-      # ASAN builds require xnack
-      rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
-        TARGETS "gfx908:xnack+;gfx90a:xnack+;gfx942:xnack+"
-      )
+if(USE_HIPCXX)
+  enable_language(HIP)
+else()
+  # Setup GPU targets for rocm platform
+  if(NOT (CMAKE_CXX_COMPILER MATCHES ".*nvcc$" OR "${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU"))
+    if(NOT DEFINED AMDGPU_TARGETS)
+      set(GPU_TARGETS "all" CACHE STRING "GPU architectures to compile for")
     else()
-      rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
-        TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201"
-      )
+      set(GPU_TARGETS "${AMDGPU_TARGETS}" CACHE STRING "GPU architectures to compile for")
+    endif()
+    set_property(CACHE GPU_TARGETS PROPERTY STRINGS "all")
+
+    if(GPU_TARGETS STREQUAL "all")
+      if(BUILD_ADDRESS_SANITIZER)
+        # ASAN builds require xnack
+        rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
+          TARGETS "gfx908:xnack+;gfx90a:xnack+;gfx942:xnack+"
+        )
+      else()
+        rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
+          TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201"
+        )
+      endif()
+      set(GPU_TARGETS "${DEFAULT_AMDGPU_TARGETS}" CACHE STRING "GPU architectures to compile for" FORCE)
     endif()
-    set(GPU_TARGETS "${DEFAULT_AMDGPU_TARGETS}" CACHE STRING "GPU architectures to compile for" FORCE)
   endif()
 endif()
 
@@ -109,8 +124,7 @@ include(VerifyCompiler)
 include(Dependencies)
 
 if(BUILD_ADDRESS_SANITIZER)
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fsanitize=address -shared-libasan")
-  set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -fsanitize=address -shared-libasan")
+  add_compile_options(-fsanitize=address -shared-libasan)
   add_link_options(-fuse-ld=lld)
 endif()
 

diff --git a/benchmark/CMakeLists.txt b/benchmark/CMakeLists.txt
@@ -22,6 +22,9 @@
 
 function(add_hipcub_benchmark BENCHMARK_SOURCE)
   get_filename_component(BENCHMARK_TARGET ${BENCHMARK_SOURCE} NAME_WE)
+  if(USE_HIPCXX)
+    set_source_files_properties(${BENCHMARK_SOURCE} PROPERTIES LANGUAGE HIP)
+  endif()
   add_executable(${BENCHMARK_TARGET} ${BENCHMARK_SOURCE})
   target_include_directories(${BENCHMARK_TARGET} SYSTEM BEFORE
     PUBLIC

diff --git a/cmake/Summary.cmake b/cmake/Summary.cmake
@@ -26,10 +26,17 @@ function(print_configuration_summary)
     message(STATUS "General:")
     message(STATUS "  System                : ${CMAKE_SYSTEM_NAME}")
     message(STATUS "  HIP ROOT              : ${HIP_ROOT_DIR}")
+if(USE_HIPCXX)
+    message(STATUS "  HIP compiler          : ${CMAKE_HIP_COMPILER}")
+    message(STATUS "  HIP compiler version  : ${CMAKE_HIP_COMPILER_VERSION}")
+    string(STRIP "${CMAKE_HIP_FLAGS}" CMAKE_HIP_FLAGS_STRIP)
+    message(STATUS "  HIP flags             : ${CMAKE_HIP_FLAGS_STRIP}")
+else()
     message(STATUS "  C++ compiler          : ${CMAKE_CXX_COMPILER}")
     message(STATUS "  C++ compiler version  : ${CMAKE_CXX_COMPILER_VERSION}")
     string(STRIP "${CMAKE_CXX_FLAGS}" CMAKE_CXX_FLAGS_STRIP)
     message(STATUS "  CXX flags             : ${CMAKE_CXX_FLAGS_STRIP}")
+endif()
 if(HIP_COMPILER STREQUAL "nvcc")
     string(REPLACE ";" " " HIP_NVCC_FLAGS_STRIP "${HIP_NVCC_FLAGS}")
     string(STRIP "${HIP_NVCC_FLAGS_STRIP}" HIP_NVCC_FLAGS_STRIP)
@@ -40,8 +47,10 @@ if(HIP_COMPILER STREQUAL "nvcc")
 endif()
     message(STATUS "  Build type            : ${CMAKE_BUILD_TYPE}")
     message(STATUS "  Install prefix        : ${CMAKE_INSTALL_PREFIX}")
-if(HIP_COMPILER STREQUAL "clang")
-    message(STATUS "  Device targets        : ${GPU_TARGETS}")
+if(USE_HIPCXX)
+    message(STATUS "  Device targets        : ${CMAKE_HIP_ARCHITECTURES}")
+elseif(HIP_COMPILER STREQUAL "clang")
+    message(STATUS "  Device targets        : ${AMDGPU_TARGETS}")
 else()
     message(STATUS "  Device targets        : ${NVGPU_TARGETS}")
 endif()

diff --git a/cmake/VerifyCompiler.cmake b/cmake/VerifyCompiler.cmake
@@ -30,9 +30,15 @@ if(HIP_COMPILER STREQUAL "nvcc")
         message(WARNING "On CUDA platform 'g++' is recommended C++ compiler.")
     endif()
 elseif(HIP_COMPILER STREQUAL "clang")
-    if(NOT (HIP_CXX_COMPILER MATCHES ".*hipcc" OR HIP_CXX_COMPILER MATCHES ".*clang\\+\\+"))
-        message(FATAL_ERROR "On ROCm platform 'hipcc' or HIP-aware Clang must be used as C++ compiler.")
+    if(USE_HIPCXX)
+        if(NOT (CMAKE_HIP_COMPILER MATCHES ".*hipcc$" OR CMAKE_HIP_COMPILER MATCHES ".*clang\\+\\+"))
+            message(FATAL_ERROR "On ROCm platform 'hipcc' or HIP-aware Clang must be used as HIP compiler.")
+        endif()
+    else()
+        if(NOT (CMAKE_CXX_COMPILER MATCHES ".*hipcc$" OR CMAKE_CXX_COMPILER MATCHES ".*clang\\+\\+"))
+            message(FATAL_ERROR "On ROCm platform 'hipcc' or HIP-aware Clang must be used as C++ compiler.")
+        endif()
     endif()
 else()
-    message(FATAL_ERROR "HIP_COMPILER must be 'clang' (AMD ROCm platform) or `nvcc` (NVIDIA CUDA platform).")
+    message(FATAL_ERROR "HIP_COMPILER must be `clang` (AMD ROCm platform)")
 endif()
diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
@@ -43,6 +43,10 @@ endforeach()
 function(add_hipcub_example EXAMPLE_NAME EXAMPLE_SOURCES)
   list(GET EXAMPLE_SOURCES 0 EXAMPLE_MAIN_SOURCE)
   get_filename_component(EXAMPLE_TARGET ${EXAMPLE_MAIN_SOURCE} NAME_WE)
+  if(USE_HIPCXX)
+    set_source_files_properties(${EXAMPLE_SOURCES} PROPERTIES LANGUAGE HIP)
+  endif()
+
   add_executable(${EXAMPLE_TARGET} ${EXAMPLE_SOURCES})
   target_include_directories(${EXAMPLE_TARGET} SYSTEM BEFORE
     PUBLIC

diff --git a/hipcub/include/hipcub/backend/cub/hipcub.hpp b/hipcub/include/hipcub/backend/cub/hipcub.hpp
@@ -112,5 +112,7 @@
 // These functions must be wrapped so they return
 // hipError_t instead of cudaError_t
 #include "util_allocator.hpp"
+#include "util_device.hpp"
+#include "util_temporary_storage.hpp"
 
 #endif // HIPCUB_CUB_HIPCUB_HPP_
diff --git a/hipcub/include/hipcub/backend/cub/util_device.hpp b/hipcub/include/hipcub/backend/cub/util_device.hpp
@@ -0,0 +1,39 @@
+/******************************************************************************
+ * Copyright (c) 2011, Duane Merrill.  All rights reserved.
+ * Copyright (c) 2011-2024, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024, Advanced Micro Devices, Inc.  All rights reserved.
+ * 
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ * 
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+
+#ifndef HIPCUB_CUB_UTIL_DEVICE_HPP_
+#define HIPCUB_CUB_UTIL_DEVICE_HPP_
+
+#include "../../config.hpp"
+
+#include "util_temporary_storage.hpp"
+
+#include "cub/util_device.cuh"
+
+#endif // HIPCUB_CUB_UTIL_DEVICE_HPP_
diff --git a/hipcub/include/hipcub/backend/cub/util_macro.hpp b/hipcub/include/hipcub/backend/cub/util_macro.hpp
@@ -0,0 +1,71 @@
+/******************************************************************************
+ * Copyright (c) 2011, Duane Merrill.  All rights reserved.
+ * Copyright (c) 2011-2018, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024, Advanced Micro Devices, Inc.  All rights reserved.
+ * 
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ * 
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+
+#ifndef HIPCUB_CUB_MACRO_HPP_
+#define HIPCUB_CUB_MACRO_HPP_
+
+#include "../../config.hpp"
+
+#include "cub/util_macro.cuh"
+
+BEGIN_HIPCUB_NAMESPACE
+
+#ifndef HIPCUB_MAX
+    /// Select maximum(a, b)
+    #define HIPCUB_MAX(a, b) CUB_MAX(a, b)
+#endif
+
+#ifndef HIPCUB_MIN
+    /// Select minimum(a, b)
+    #define HIPCUB_MIN(a, b) CUB_MIN(a, b)
+#endif
+
+#ifndef HIPCUB_QUOTIENT_FLOOR
+    /// Quotient of x/y rounded down to nearest integer
+    #define HIPCUB_QUOTIENT_FLOOR(x, y) CUB_QUOTIENT_FLOOR(x, y)
+#endif
+
+#ifndef HIPCUB_QUOTIENT_CEILING
+    /// Quotient of x/y rounded up to nearest integer
+    #define HIPCUB_QUOTIENT_CEILING(x, y) CUB_QUOTIENT_CEILING(x, y)
+#endif
+
+#ifndef HIPCUB_ROUND_UP_NEAREST
+    /// x rounded up to the nearest multiple of y
+    #define HIPCUB_ROUND_UP_NEAREST(x, y) CUB_ROUND_UP_NEAREST(x, y)
+#endif
+
+#ifndef HIPCUB_ROUND_DOWN_NEAREST
+    /// x rounded down to the nearest multiple of y
+    #define HIPCUB_ROUND_DOWN_NEAREST(x, y) CUB_ROUND_DOWN_NEAREST(x, y)
+#endif
+
+END_HIPCUB_NAMESPACE
+
+#endif // HIPCUB_CUB_MACRO_HPP_