rocThrust for thrust functionality in HIP (#560)

* Add rocThrust and basic test

* Add rocThrust library

* Thrust -> rocThrust

* Add HIP workflow

* Add HIP compiler options: only -fPIE at the moment

* Switch to new rocm clang container

* Move to unversioned rocm_clang version until we have a new machines tag

* Add -fPIE to compilation

* CMake updates, copyright

* Tweaking the rocThrust setup according to my own taste.

So that it could be used with our existing ROCm/HIP Docker
images as well.

---------

Co-authored-by: Stewart Martin-Haugh [email protected] <[email protected]>
Co-authored-by: Attila Krasznahorkay <[email protected]>

.github/workflows/builds.yml

@@ -30,6 +30,9 @@ jobs:
           - name: CPU
             container: ghcr.io/acts-project/ubuntu2004:v30
             options: -DTRACCC_USE_ROOT=FALSE
+          - name: HIP
+            container: ghcr.io/acts-project/ubuntu2004_rocm:v42
+            options: -DTRACCC_BUILD_HIP=TRUE -DTRACCC_SETUP_ROCTHRUST=TRUE
           - name: CUDA
             container: ghcr.io/acts-project/ubuntu2004_cuda:v30
             options: -DTRACCC_BUILD_CUDA=TRUE

CMakeLists.txt

@@ -51,6 +51,7 @@ endif()
 # Flags controlling which parts of traccc to build.
 option( TRACCC_BUILD_CUDA "Build the CUDA sources included in traccc"
    ${TRACCC_BUILD_CUDA_DEFAULT} )
+option( TRACCC_BUILD_HIP "Build the HIP sources included in traccc" FALSE)
 option( TRACCC_BUILD_SYCL "Build the SYCL sources included in traccc" FALSE )
 option( TRACCC_BUILD_FUTHARK "Build the Futhark sources included in traccc"
    FALSE )
@@ -118,6 +119,26 @@ set( TRACCC_THRUST_OPTIONS "" CACHE STRING
 mark_as_advanced( TRACCC_THRUST_OPTIONS )
 thrust_create_target( traccc::Thrust ${TRACCC_THRUST_OPTIONS} )
 
+# Set up rocThrust.
+option( TRACCC_SETUP_ROCTHRUST
+   "Set up the rocThrust target(s) explicitly" FALSE )
+option( TRACCC_USE_SYSTEM_ROCTHRUST
+   "Pick up an existing installation of rocThrust from the build environment"
+   ${TRACCC_USE_SYSTEM_LIBS} )
+if( TRACCC_SETUP_ROCTHRUST )
+   set( ROCM_WARN_TOOLCHAIN_VAR FALSE CACHE BOOL "Don't print ROCm warnings" )
+   set( ROCM_ERROR_TOOLCHAIN_VAR FALSE CACHE BOOL "Don't print ROCm errors" )
+   mark_as_advanced( ROCM_WARN_TOOLCHAIN_VAR ROCM_ERROR_TOOLCHAIN_VAR )
+   if( TRACCC_USE_SYSTEM_ROCTHRUST )
+      find_package( rocThrust REQUIRED )
+   else()
+      add_subdirectory( extern/rocThrust )
+   endif()
+   # Dress up the rocthrust target a little.
+   target_compile_definitions( rocthrust INTERFACE
+      THRUST_IGNORE_CUB_VERSION_CHECK )
+endif()
+
 # Set up TBB.
 option( TRACCC_SETUP_TBB
    "Set up the TBB target(s) explicitly" TRUE )

cmake/traccc-config.cmake.in

@@ -1,6 +1,6 @@
 # TRACCC library, part of the ACTS project (R&D line)
 #
-# (c) 2022-2023 CERN for the benefit of the ACTS project
+# (c) 2022-2024 CERN for the benefit of the ACTS project
 #
 # Mozilla Public License Version 2.0
 
@@ -9,6 +9,7 @@
 
 # Remember the options that traccc was built with.
 set( TRACCC_BUILD_CUDA     @TRACCC_BUILD_CUDA@ )
+set( TRACCC_BUILD_HIP      @TRACCC_BUILD_HIP@ )
 set( TRACCC_BUILD_SYCL     @TRACCC_BUILD_SYCL@ )
 set( TRACCC_BUILD_FUTHARK  @TRACCC_BUILD_FUTHARK@ )
 set( TRACCC_BUILD_KOKKOS   @TRACCC_BUILD_KOKKOS@ )

extern/rocThrust/CMakeLists.txt

@@ -0,0 +1,49 @@
+# TRACCC library, part of the ACTS project (R&D line)
+#
+# (c) 2024 CERN for the benefit of the ACTS project
+#
+# Mozilla Public License Version 2.0
+
+# CMake include(s).
+cmake_minimum_required( VERSION 3.14 )
+include( FetchContent )
+
+# Silence FetchContent warnings with CMake >=3.24.
+if( POLICY CMP0135 )
+   cmake_policy( SET CMP0135 NEW )
+endif()
+
+# Tell the user what's happening.
+message( STATUS "Building rocThrust as part of the TRACCC project" )
+
+# Declare where to get rocThrust from.
+set( TRACCC_ROCTHRUST_SOURCE
+   "URL;https://github.com/ROCm/rocThrust/archive/refs/tags/rocm-6.1.1.tar.gz;URL_MD5;038abf313688c00555fe1efc51e1307b"
+   CACHE STRING "Source for rocThrust, when built as part of this project" )
+set( TRACCC_ROCTHRUST_PATCH
+   "PATCH_COMMAND;patch;-p1;<;${CMAKE_CURRENT_SOURCE_DIR}/rocm-6.1.1.patch"
+   CACHE STRING "Patch for rocThrust, when built as part of this project" )
+mark_as_advanced( TRACCC_ROCTHRUST_SOURCE )
+FetchContent_Declare( rocThrust
+   ${TRACCC_ROCTHRUST_SOURCE}
+   ${TRACCC_ROCTHRUST_PATCH} )
+
+# Settings for the rocThrust build.
+if( DEFINED CACHE{BUILD_TESTING} )
+   set( _buildTestingValue ${BUILD_TESTING} )
+endif()
+set( BUILD_TESTING FALSE CACHE INTERNAL "Forceful setting of BUILD_TESTING" )
+set( HIP_COMPILER "clang" )
+set( HIP_CXX_COMPILER "hipcc" )
+
+# Get it into the current directory.
+FetchContent_MakeAvailable( rocThrust )
+
+# Reset the BUILD_TESTING variable.
+if( DEFINED _buildTestingValue )
+   set( BUILD_TESTING ${_buildTestingValue} CACHE BOOL "Turn tests on/off"
+      FORCE )
+   unset( _buildTestingValue )
+else()
+   unset( BUILD_TESTING CACHE )
+endif()

extern/rocThrust/rocm-6.1.1.patch

@@ -0,0 +1,26 @@
+diff -ur rocThrust-rocm-6.1.1-orig/cmake/Dependencies.cmake rocThrust-rocm-6.1.1-fixed/cmake/Dependencies.cmake
+--- rocThrust-rocm-6.1.1-orig/cmake/Dependencies.cmake	2024-02-07 00:12:37.000000000 +0100
++++ rocThrust-rocm-6.1.1-fixed/cmake/Dependencies.cmake	2024-05-17 09:33:52.423910609 +0200
+@@ -23,7 +23,7 @@
+     GIT_REPOSITORY      https://github.com/ROCmSoftwarePlatform/rocPRIM.git
+     GIT_TAG             develop
+     INSTALL_DIR         ${CMAKE_CURRENT_BINARY_DIR}/deps/rocprim
+-    CMAKE_ARGS          -DBUILD_TEST=OFF -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR> -DCMAKE_PREFIX_PATH=/opt/rocm
++    CMAKE_ARGS          -DBUILD_TEST=OFF -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR> -DCMAKE_PREFIX_PATH=/opt/rocm -DHIP_COMPILER=${HIP_COMPILER} -DHIP_CXX_COMPILER=${HIP_CXX_COMPILER}
+     LOG_DOWNLOAD        TRUE
+     LOG_CONFIGURE       TRUE
+     LOG_BUILD           TRUE
+diff -ur rocThrust-rocm-6.1.1-orig/CMakeLists.txt rocThrust-rocm-6.1.1-fixed/CMakeLists.txt
+--- rocThrust-rocm-6.1.1-orig/CMakeLists.txt	2024-02-07 00:12:37.000000000 +0100
++++ rocThrust-rocm-6.1.1-fixed/CMakeLists.txt	2024-05-16 16:52:07.600882025 +0200
+@@ -63,10 +63,6 @@
+ # Get dependencies
+ include(cmake/Dependencies.cmake)
+
+-# Verify that supported compilers are used
+-if (NOT WIN32)
+-  include(cmake/VerifyCompiler.cmake)
+-endif()
+ # Build options
+ # Disable -Werror
+ option(DISABLE_WERROR "Disable building with Werror" ON)

tests/CMakeLists.txt

@@ -47,6 +47,10 @@ if( TRACCC_BUILD_CUDA )
     add_subdirectory( cuda )
 endif()
 
+if( TRACCC_BUILD_HIP )
+    add_subdirectory( hip )
+endif()
+
 if( TRACCC_BUILD_SYCL )
     add_subdirectory( sycl )
 endif()

tests/alpaka/CMakeLists.txt

@@ -7,6 +7,7 @@
 if(alpaka_ACC_GPU_CUDA_ENABLE)
   enable_language(CUDA)
   set_source_files_properties(alpaka_basic.cpp PROPERTIES LANGUAGE CUDA)
+  include( traccc-compiler-options-cuda )
   list(APPEND DEVICE_LIBRARIES vecmem::cuda)
 elseif(alpaka_ACC_GPU_HIP_ENABLE)
   enable_language(HIP)
@@ -22,3 +23,9 @@ traccc_add_test( alpaka
    vecmem::core
    ${DEVICE_LIBRARIES}
 )
+
+#Can only do this once target is defined, so need another if here
+if(alpaka_ACC_GPU_HIP_ENABLE)
+  set_target_properties( traccc_test_alpaka PROPERTIES
+     POSITION_INDEPENDENT_CODE TRUE )
+endif()

tests/hip/CMakeLists.txt

@@ -0,0 +1,20 @@
+# TRACCC library, part of the ACTS project (R&D line)
+#
+# (c) 2024 CERN for the benefit of the ACTS project
+#
+# Mozilla Public License Version 2.0
+
+enable_language(HIP)
+traccc_add_test(
+    hip 
+    # Define the sources for the test.
+    test_thrust.hip
+    LINK_LIBRARIES
+    rocthrust
+    GTest::gtest_main
+    vecmem::core
+    vecmem::hip
+)
+
+set_target_properties( traccc_test_hip PROPERTIES
+   POSITION_INDEPENDENT_CODE TRUE )

tests/hip/test_thrust.hip

@@ -0,0 +1,101 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2024 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+// VecMem include(s).
+#include <vecmem/containers/data/vector_buffer.hpp>
+#include <vecmem/containers/device_vector.hpp>
+#include <vecmem/containers/vector.hpp>
+#include <vecmem/memory/hip/device_memory_resource.hpp>
+#include <vecmem/memory/memory_resource.hpp>
+#include <vecmem/utils/copy.hpp>
+#include <vecmem/utils/hip/copy.hpp>
+
+// Thrust include(s).
+#include <thrust/execution_policy.h>
+#include <thrust/fill.h>
+#include <thrust/scan.h>
+#include <thrust/sort.h>
+
+// GTest include(s).
+#include <gtest/gtest.h>
+
+// This defines the local frame test suite
+
+namespace {
+vecmem::hip::copy copy;
+vecmem::host_memory_resource host_resource;
+vecmem::hip::device_memory_resource device_resource;
+
+}  // namespace
+
+TEST(thrust, sort) {
+
+    vecmem::vector<unsigned int> host_vector{{3, 2, 1, 8, 4}, &host_resource};
+
+    auto host_buffer = vecmem::get_data(host_vector);
+    auto device_buffer = copy.to(vecmem::get_data(host_vector), device_resource,
+                                 vecmem::copy::type::host_to_device);
+
+    vecmem::device_vector<unsigned int> device_vector(device_buffer);
+
+    thrust::sort(thrust::device, device_vector.begin(), device_vector.end());
+
+    copy(device_buffer, host_buffer, vecmem::copy::type::device_to_host);
+
+    ASSERT_EQ(host_vector[0], 1);
+    ASSERT_EQ(host_vector[1], 2);
+    ASSERT_EQ(host_vector[2], 3);
+    ASSERT_EQ(host_vector[3], 4);
+    ASSERT_EQ(host_vector[4], 8);
+}
+
+TEST(thrust, scan) {
+
+    vecmem::vector<unsigned int> host_vector{{3, 2, 1, 8, 4}, &host_resource};
+
+    auto host_buffer = vecmem::get_data(host_vector);
+    auto device_buffer = copy.to(vecmem::get_data(host_vector), device_resource,
+                                 vecmem::copy::type::host_to_device);
+
+    vecmem::device_vector<unsigned int> device_vector(device_buffer);
+
+    thrust::inclusive_scan(thrust::device, device_vector.begin(),
+                           device_vector.end(), device_vector.begin());
+
+    copy(device_buffer, host_buffer, vecmem::copy::type::device_to_host);
+
+    ASSERT_EQ(host_vector[0], 3);
+    ASSERT_EQ(host_vector[1], 5);
+    ASSERT_EQ(host_vector[2], 6);
+    ASSERT_EQ(host_vector[3], 14);
+    ASSERT_EQ(host_vector[4], 18);
+}
+
+TEST(thrust, fill) {
+
+    vecmem::vector<unsigned int> host_vector{{1, 1, 1, 1, 1, 1, 1},
+                                             &host_resource};
+
+    auto host_buffer = vecmem::get_data(host_vector);
+    auto device_buffer = copy.to(vecmem::get_data(host_vector), device_resource,
+                                 vecmem::copy::type::host_to_device);
+
+    vecmem::device_vector<unsigned int> device_vector(device_buffer);
+
+    thrust::fill(thrust::device, device_vector.begin(), device_vector.end(),
+                 112);
+
+    copy(device_buffer, host_buffer, vecmem::copy::type::device_to_host);
+
+    ASSERT_EQ(host_vector[0], 112);
+    ASSERT_EQ(host_vector[1], 112);
+    ASSERT_EQ(host_vector[2], 112);
+    ASSERT_EQ(host_vector[3], 112);
+    ASSERT_EQ(host_vector[4], 112);
+    ASSERT_EQ(host_vector[5], 112);
+    ASSERT_EQ(host_vector[6], 112);
+}