Use rocprim::transform in unary transform
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -58,3 +58,4 @@ endfunction()
add_hipthrust_test("thrust.hip.copy" test_copy.cpp)
add_hipthrust_test("thrust.hip.for_each" test_for_each.cpp)
add_hipthrust_test("thrust.hip.memory" test_memory.cpp)
+add_hipthrust_test("thrust.hip.transform" test_transform.cpp)

ADD
test/test_transform.cpp
+// MIT License
+//
+// Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#include <iostream>
+#include <type_traits>
+#include <cstdlib>
+#include <algorithm>
+#include <vector>
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/memory.h>
+#include <thrust/transform.h>
+// STREAMHPC TODO replace <thrust/detail/seq.h> with <thrust/execution_policy.h>
+#include <thrust/detail/seq.h>
+#include <thrust/device_ptr.h>
+#include <thrust/device_malloc.h>
+#include <thrust/device_free.h>
+#include <thrust/iterator/counting_iterator.h>
+
+// HIP API
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <hip/hip_runtime_api.h>
+#include <hip/hip_runtime.h>
+
+#define HIP_CHECK(condition) ASSERT_EQ(condition, hipSuccess)
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+template<
+ class InputType,
+ class OutputType = InputType
+>
+struct Params
+{
+ using input_type = InputType;
+ using output_type = OutputType;
+};
+
+template<class Params>
+class TransformTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+ using output_type = typename Params::output_type;
+};
+
+typedef ::testing::Types<
+ Params<int>,
+ Params<unsigned short>
+> TransformTestsParams;
+
+TYPED_TEST_CASE(TransformTests, TransformTestsParams);
+
+std::vector<size_t> get_sizes()
+{
+ std::vector<size_t> sizes = {
+ 0, 1, 2, 12, 63, 64, 211, 256, 344,
+ 1024, 2048, 5096, 34567, (1 << 17) - 1220
+ };
+ return sizes;
+}
+
+template<class T>
+struct transform
+{
+ __device__ __host__ inline
+ constexpr T operator()(const T& a) const
+ {
+ return a + 5;
+ }
+};
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+TYPED_TEST(TransformTests, UnaryTransform)
+{
+ using T = typename TestFixture::input_type;
+ using U = typename TestFixture::output_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ std::vector<T> input(size, 0);
+ std::vector<U> output(size, 0);
+
+ thrust::device_ptr<T> input_ptr = thrust::device_malloc<T>(size);
+ thrust::device_ptr<U> output_ptr = thrust::device_malloc<U>(size);
+
+ for(size_t i = 0; i < size; i++)
+ {
+ input[i] = i;
+ //input_ptr[i] = i;
+ }
+
+ hipMemcpy(
+ thrust::raw_pointer_cast(input_ptr), input.data(),
+ input.size() * sizeof(T),
+ hipMemcpyHostToDevice
+ );
+ //thrust::copy(input, input + size, input_ptr);
+
+ // Calculate expected results on host
+ std::vector<U> expected(input.size());
+ std::transform(input.begin(), input.end(), expected.begin(), transform<U>());
+
+ thrust::transform(input_ptr, input_ptr + size, output_ptr, transform<U>());
+
+ hipMemcpy(
+ output.data(), thrust::raw_pointer_cast(output_ptr),
+ size * sizeof(U),
+ hipMemcpyDeviceToHost
+ );
+ //thrust::copy(output_ptr, output_ptr + size, output);
+
+ for(size_t i = 0; i < size; i++)
+ {
+ ASSERT_EQ(output[i], expected[i]);
+ }
+
+ // Free
+ thrust::device_free(input_ptr);
+ thrust::device_free(output_ptr);
+ }
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

MODIFY
thrust/system/hip/detail/transform.h
thrust/system/hip/detail/transform.h
#include <thrust/system/hip/config.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>
BEGIN_NS_THRUST
@@ -233,6 +236,32 @@ 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,
@@ -342,12 +371,44 @@ transform(execution_policy<Derived> &policy,
OutputIt result,
TransformOp transform_op)
{
- return hip_rocprim::transform_if(policy,
+ typedef typename iterator_traits<InputIt>::difference_type size_type;
+ size_type num_items = static_cast<size_type>(thrust::distance(first, last));
+
+ 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
+ // that will be used in #if __THRUST_HAS_HIPRT__ block.
+ {
+ auto ptr = ::rocprim::transform<::rocprim::default_config, InputIt, OutputIt, TransformOp>;
+ (void) ptr;
+ }
+ #if __THRUST_HAS_HIPRT__
+ {
+ hipStream_t stream = hip_rocprim::stream(policy);
+ hipError_t status = ::rocprim::transform(
first,
- last,
result,
+ num_items,
transform_op,
- __transform::always_true_predicate());
+ stream,
+ false
+ );
+ hip_rocprim::throw_on_error(status, "transform failed");
+
+ return result + num_items;
+ }
+ #else
+ {
+ (void) policy;
+ while(first != last)
+ {
+ *result++ = transform_op(*first++);
+ }
+ return result;
+ }
+ #endif
} // func transform
//-------------------------

Use rocprim::transform in unary transform