Updated cmake script to include rocPRIM headers
MODIFY
cmake/Dependencies.cmake
cmake/Dependencies.cmake
@@ -43,6 +43,24 @@ if (NOT Git_FOUND)
message(FATAL_ERROR "Please ensure Git is installed on the system")
endif()
+# rocPRIM (https://github.com/ROCmSoftwarePlatform/rocPRIM)
+message(STATUS "Downloading and building rocPRIM.")
+set(ROCPRIM_ROOT ${CMAKE_CURRENT_BINARY_DIR}/rocPRIM CACHE PATH "")
+download_project(
+ PROJ rocPRIM
+ GIT_REPOSITORY https://github.com/ROCmSoftwarePlatform/rocPRIM.git
+ GIT_TAG master
+ INSTALL_DIR ${ROCPRIM_ROOT}
+ CMAKE_ARGS -DBUILD_TEST=OFF -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR>
+ LOG_DOWNLOAD TRUE
+ LOG_CONFIGURE TRUE
+ LOG_BUILD TRUE
+ LOG_INSTALL TRUE
+ BUILD_PROJECT TRUE
+ ${UPDATE_DISCONNECTED_IF_AVAILABLE}
+)
+find_package(rocprim REQUIRED CONFIG PATHS ${ROCPRIM_ROOT})
+
# Test dependencies
if(BUILD_TEST)
# Google Test (https://github.com/google/googletest)

MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -31,6 +31,7 @@ function(add_hipthrust_test TEST_NAME TEST_SOURCES)
target_link_libraries(${TEST_TARGET}
PRIVATE
hipthrust
+ roc::rocprim_hip
${GTEST_BOTH_LIBRARIES}
)
foreach(amdgpu_target ${AMDGPU_TARGETS})

MODIFY
thrust/system/hip/detail/transform.h
thrust/system/hip/detail/transform.h
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
#include <thrust/system/hip/config.h>
+// This include is needed to circumvent issue with thrust::pointer
+// not having a value_type member when passing to a rocPRIM function
#include <thrust/system/hip/pointer.h>
#include <thrust/system/hip/detail/util.h>
#include <thrust/detail/type_traits/iterator/is_output_iterator.h>
#include <thrust/detail/type_traits/result_of_adaptable_function.h>
#include <thrust/system/hip/detail/parallel_for.h>
-#include <thrust/system/hip/detail/rocprim/functional.hpp>
-#include <thrust/system/hip/detail/rocprim/device/device_transform_hip.hpp>
#include <thrust/distance.h>
+// rocPRIM includes
+#include <rocprim/functional.hpp>
+#include <rocprim/device/device_transform_hip.hpp>
+
BEGIN_NS_THRUST
namespace hip_rocprim {
/// STREAMHPC TODO Use rocPRIM transform
+/// - rocPRIM transform for stencil unary transform
+/// - rocPRIM transform for binary transform
+/// - rocPRIM transform for stencil binary transform
namespace __transform {
struct no_stencil_tag
@@ -237,32 +244,6 @@ namespace __transform {
return result + num_items;
}
- template <class InputIt,
- class OutputIt,
- class Size,
- class TransformOp>
- hipError_t THRUST_HIP_FUNCTION
- rocprim_unary(InputIt items,
- OutputIt result,
- Size num_items,
- TransformOp transform_op,
- hipStream_t stream)
- {
- typedef typename thrust::iterator_value<InputIt>::type InputTy;
-
- InputTy * first_ptr = thrust::raw_pointer_cast(&items[0]);
- InputTy * result_ptr = thrust::raw_pointer_cast(&result[0]);
-
- return ::rocprim::transform(
- first_ptr,
- result_ptr,
- num_items,
- transform_op,
- stream,
- false
- );
- }
-
template <class Policy,
class InputIt1,
class InputIt2,
@@ -378,8 +359,8 @@ transform(execution_policy<Derived> &policy,
if (num_items == 0)
return result;
- // Workaround, so kernel called by __transform::rocprim_unary is not lost,
- // Implicit instantiation of __transform::rocprim_unary function template
+ // Workaround, so kernel(s) called by ::rocprim::transform is/are not lost,
+ // Implicit instantiation of ::rocprim::transform function template
// that will be used in #if __THRUST_HAS_HIPRT__ block.
{
auto ptr = ::rocprim::transform<::rocprim::default_config, InputIt, OutputIt, TransformOp>;

MODIFY
thrust/system/hip/pointer.h
thrust/system/hip/pointer.h
@@ -36,6 +36,6 @@ public:
typedef typename ptr::iterator_category iterator_category;
typedef typename ptr::value_type value_type;
typedef typename ptr::difference_type difference_type;
- typedef ptr pointer;
+ typedef Element * pointer;
typedef typename ptr::reference reference;
};

Updated cmake script to include rocPRIM headers