Add Hip-Cpu support initial

.gitlab-ci.yml

@@ -81,6 +81,25 @@ variables:
     paths:
       - $CMAKE_LATEST_PATH
 
+.deps-cpu:cmake-latest:
+  stage: build:cmake_latest
+  before_script:
+    - $SUDO_CMD apt update -qq
+    - $SUDO_CMD apt install -y -qq apt-transport-https software-properties-common
+    - $SUDO_CMD add-apt-repository ppa:ubuntu-toolchain-r/test
+    # | Used in the script    | Build tools | Fetch from https:// | rocminfo calls lsmod
+    - $SUDO_CMD apt install -y -qq wget tar xz-utils bzip2 libnuma-dev libunwind-dev git build-essential pkg-config ninja-build ca-certificates kmod g++-9
+    # Fetch CMake only if the cache has not been restored
+    - if [ ! -d $CMAKE_LATEST_PATH ]; then mkdir -p $CMAKE_LATEST_PATH; wget --no-check-certificate --quiet -O - $CMAKE_LATEST_URL | tar --strip-components=1 -xz -C $CMAKE_LATEST_PATH;
+    - fi;
+    - export PATH=$CMAKE_LATEST_PATH/bin:$PATH
+    # Debug printing of environment for context when errors occur
+    - hipconfig
+  cache:
+    key: $CMAKE_LATEST_VERSION
+    paths:
+      - $CMAKE_LATEST_PATH
+
 build:cmake-latest:
   extends:
     - .deps:cmake-latest
@@ -92,6 +111,7 @@ build:cmake-latest:
     - cmake
       -G Ninja
       -D CMAKE_CXX_COMPILER=hipcc
+      -D CMAKE_CXX_FLAGS="-Wall -Wextra -Werror"
       -D CMAKE_BUILD_TYPE=Release
       -D BUILD_TEST=ON
       -D BUILD_EXAMPLE=ON
@@ -112,6 +132,31 @@ build:cmake-latest:
     - $BUILD_LATEST_DIR/CTestTestfile.cmake
     expire_in: 2 weeks
 
+# TODO: Enable the hip-cpu CI step
+#build-cpu:cmake-latest:
+#  extends:
+#    - .deps-cpu:cmake-latest
+#  tags:
+#    - s9300
+#    - rocm
+#  script:
+#    - mkdir -p $BUILD_LATEST_DIR
+#    - cd $BUILD_LATEST_DIR
+#    - cmake
+#      -G Ninja
+#      -D CMAKE_CXX_COMPILER=g++-9
+#      -D CMAKE_CXX_FLAGS="-Wall -Wextra"
+#      -D CMAKE_BUILD_TYPE=Release
+#      -D BUILD_TEST=ON
+#      -D BUILD_EXAMPLE=OFF
+#      -D BUILD_BENCHMARK=OFF
+#      -D USE_HIP_CPU=ON
+#      -S $CI_PROJECT_DIR
+#      -B $BUILD_LATEST_DIR
+#    - cmake
+#      --build $BUILD_LATEST_DIR
+#      --parallel 3
+
 build:cmake-minimum:
   extends:
     - .deps:cmake-minimum
@@ -126,6 +171,7 @@ build:cmake-minimum:
     - cmake
       -G Ninja
       -D CMAKE_CXX_COMPILER=hipcc
+      -D CMAKE_CXX_FLAGS="-Wall -Wextra -Werror"
       -D CMAKE_BUILD_TYPE=Release
       -D BUILD_TEST=ON
       -D BUILD_EXAMPLE=ON

CHANGELOG.md

@@ -2,6 +2,10 @@
 
 Full documentation for rocPRIM is available at [https://codedocs.xyz/ROCmSoftwarePlatform/rocPRIM/](https://codedocs.xyz/ROCmSoftwarePlatform/rocPRIM/)
 
+## [Unreleased rocPRIM-Next]
+### Added
+- Experimental [HIP-CPU](https://github.com/ROCm-Developer-Tools/HIP-CPU) support; build using GCC/Clang/MSVC on Win/Linux. It is work in progress, many algorithms still known to fail.
+
 ## [Unreleased rocPRIM-2.10.11 for ROCm 4.4.0]
 ### Added
 - Code coverage tools build option

CMakeLists.txt

@@ -31,6 +31,14 @@ list( APPEND CMAKE_PREFIX_PATH /opt/rocm/llvm /opt/rocm )
 # rocPRIM project
 project(rocprim LANGUAGES CXX)
 
+# Build options
+option(BUILD_TEST "Build tests (requires googletest)" OFF)
+option(BUILD_BENCHMARK "Build benchmarks" OFF)
+option(BUILD_EXAMPLE "Build examples" OFF)
+option(USE_HIP_CPU "Prefer HIP-CPU runtime instead of HW acceleration" OFF)
+# Disables building tests, benchmarks, examples
+option(ONLY_INSTALL "Only install" OFF)
+
 # CMake modules
 list(APPEND CMAKE_MODULE_PATH
   ${CMAKE_CURRENT_SOURCE_DIR}/cmake
@@ -46,36 +54,35 @@ endif()
 
 set(CMAKE_INSTALL_RPATH_USE_LINK_PATH TRUE CACHE BOOL "Add paths to linker search and installed rpath")
 
-# Get dependencies
-include(cmake/Dependencies.cmake)
-
-# Set the AMDGPU_TARGETS with backward compatiblity
-# Use target ID syntax if supported for AMDGPU_TARGETS
-if(COMMAND rocm_check_target_ids)
-  rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
-    TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx1030"
-  )
-else()
-  # Detect compiler support for target ID
-  # This section is deprecated. Please use rocm_check_target_ids for future use.
-  if( CMAKE_CXX_COMPILER MATCHES ".*/hipcc$" )
-    execute_process(COMMAND ${CMAKE_CXX_COMPILER} "--help"
-      OUTPUT_VARIABLE CXX_OUTPUT
-      OUTPUT_STRIP_TRAILING_WHITESPACE
-      ERROR_STRIP_TRAILING_WHITESPACE)
-    string(REGEX MATCH ".mcode\-object\-version" TARGET_ID_SUPPORT ${CXX_OUTPUT})
-  endif()
-  if(TARGET_ID_SUPPORT)
-    set(DEFAULT_AMDGPU_TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-")
+if(NOT USE_HIP_CPU)
+  # Set the AMDGPU_TARGETS with backward compatiblity
+  # Use target ID syntax if supported for AMDGPU_TARGETS
+  if(COMMAND rocm_check_target_ids)
+    rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
+      TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx1030"
+    )
   else()
-    set(DEFAULT_AMDGPU_TARGETS "gfx803;gfx900;gfx906;gfx908")
+    # Detect compiler support for target ID
+    # This section is deprecated. Please use rocm_check_target_ids for future use.
+    if( CMAKE_CXX_COMPILER MATCHES ".*/hipcc$" )
+      execute_process(COMMAND ${CMAKE_CXX_COMPILER} "--help"
+        OUTPUT_VARIABLE CXX_OUTPUT
+        OUTPUT_STRIP_TRAILING_WHITESPACE
+        ERROR_STRIP_TRAILING_WHITESPACE)
+      string(REGEX MATCH ".mcode\-object\-version" TARGET_ID_SUPPORT ${CXX_OUTPUT})
+    endif()
+    if(TARGET_ID_SUPPORT)
+      set(DEFAULT_AMDGPU_TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-")
+    else()
+      set(DEFAULT_AMDGPU_TARGETS "gfx803;gfx900;gfx906;gfx908")
+    endif()
   endif()
-endif()
-set(AMDGPU_TARGETS "${DEFAULT_AMDGPU_TARGETS}" CACHE STRING "List of specific machine types for library to target")
-set(AMDGPU_TEST_TARGETS "" CACHE STRING "List of specific device types to test for") # Leave empty for default system device
+  set(AMDGPU_TARGETS "${DEFAULT_AMDGPU_TARGETS}" CACHE STRING "List of specific machine types for library to target")
+  set(AMDGPU_TEST_TARGETS "" CACHE STRING "List of specific device types to test for") # Leave empty for default system device
 
-# Verify that hcc compiler is used on ROCM platform
-include(cmake/VerifyCompiler.cmake)
+  # Verify that hcc compiler is used on ROCM platform
+  include(cmake/VerifyCompiler.cmake)
+endif()
 
 # Build options
 # Disable -Werror
@@ -98,11 +105,8 @@ set(CMAKE_CXX_STANDARD 14)
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
 set(CMAKE_CXX_EXTENSIONS OFF)
 
-if(DISABLE_WERROR)
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -Wextra")
-else()
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -Wextra -Werror")
-endif()
+# Get dependencies
+include(cmake/Dependencies.cmake)
 
 # Setup VERSION
 set(VERSION_STRING "2.10.9")

README.md

@@ -49,6 +49,11 @@ cd rocPRIM; mkdir build; cd build
 # before 'cmake' or setting cmake option 'CMAKE_CXX_COMPILER' to path to the compiler.
 # Using HIP-clang:
 [CXX=hipcc] cmake -DBUILD_BENCHMARK=ON ../.
+#
+# ! EXPERIMENTAL !
+# Alternatively one may build using the experimental (and highly incomplete) HIP-CPU back-end for host-side
+# execution using any C++17 conforming compiler (supported by HIP-CPU). AMDGPU_* options are unavailable in this case. 
+#   USE_HIP_CPU - OFF by default
 
 # Build
 make -j4

benchmark/CMakeLists.txt

@@ -24,18 +24,40 @@ option(BENCHMARK_CONFIG_TUNING "Benchmark device-level functions using various c
 
 function(add_rocprim_benchmark BENCHMARK_SOURCE)
   get_filename_component(BENCHMARK_TARGET ${BENCHMARK_SOURCE} NAME_WE)
+
   add_executable(${BENCHMARK_TARGET} ${BENCHMARK_SOURCE})
+
   target_link_libraries(${BENCHMARK_TARGET}
     PRIVATE
-      rocprim_hip
+      rocprim
       benchmark::benchmark
   )
-  foreach(amdgpu_target ${AMDGPU_TARGETS})
+  if(NOT USE_HIP_CPU)
+    target_link_libraries(${BENCHMARK_TARGET}
+      PRIVATE
+        rocprim_hip
+    )
+  else()
     target_link_libraries(${BENCHMARK_TARGET}
       PRIVATE
-        --amdgpu-target=${amdgpu_target}
+        Threads::Threads
+        hip_cpu_rt::hip_cpu_rt
     )
-  endforeach()
+    if(STL_DEPENDS_ON_TBB)
+      target_link_libraries(${BENCHMARK_TARGET}
+        PRIVATE
+          TBB::tbb
+      )
+    endif()
+  endif()
+
+  target_compile_options(${BENCHMARK_TARGET}
+    PRIVATE
+      $<$<CXX_COMPILER_ID:MSVC>:
+        /bigobj # number of sections exceeded object file format limit: compile with /bigobj
+      >
+  )
+
   set_target_properties(${BENCHMARK_TARGET}
     PROPERTIES
       RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/benchmark"

benchmark/benchmark_block_discontinuity.cpp

@@ -82,13 +82,13 @@ struct flag_heads
     __device__
     static void run(const T * d_input, T * d_output)
     {
-        const unsigned int lid = hipThreadIdx_x;
-        const unsigned int block_offset = hipBlockIdx_x * ItemsPerThread * BlockSize;
+        const unsigned int lid = threadIdx.x;
+        const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
         rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
-        #pragma nounroll
+        ROCPRIM_NO_UNROLL
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             rp::block_discontinuity<T, BlockSize> bdiscontinuity;
@@ -125,13 +125,13 @@ struct flag_tails
     __device__
     static void run(const T * d_input, T * d_output)
     {
-        const unsigned int lid = hipThreadIdx_x;
-        const unsigned int block_offset = hipBlockIdx_x * ItemsPerThread * BlockSize;
+        const unsigned int lid = threadIdx.x;
+        const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
         rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
-        #pragma nounroll
+        ROCPRIM_NO_UNROLL
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             rp::block_discontinuity<T, BlockSize> bdiscontinuity;
@@ -168,13 +168,13 @@ struct flag_heads_and_tails
     __device__
     static void run(const T * d_input, T * d_output)
     {
-        const unsigned int lid = hipThreadIdx_x;
-        const unsigned int block_offset = hipBlockIdx_x * ItemsPerThread * BlockSize;
+        const unsigned int lid = threadIdx.x;
+        const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
         rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
-        #pragma nounroll
+        ROCPRIM_NO_UNROLL
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             rp::block_discontinuity<T, BlockSize> bdiscontinuity;
@@ -217,8 +217,8 @@ void run_benchmark(benchmark::State& state, hipStream_t stream, size_t N)
     std::vector<T> input = get_random_data<T>(size, T(0), T(10));
     T * d_input;
     T * d_output;
-    HIP_CHECK(hipMalloc(&d_input, size * sizeof(T)));
-    HIP_CHECK(hipMalloc(&d_output, size * sizeof(T)));
+    HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_input), size * sizeof(T)));
+    HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_output), size * sizeof(T)));
     HIP_CHECK(
         hipMemcpy(
             d_input, input.data(),