HIP backend: add temporary parallel_for impl
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -56,3 +56,4 @@ endfunction()
# ****************************************************************************
add_hipthrust_test("thrust.hip.memory" test_memory.cpp)
+add_hipthrust_test("thrust.hip.for_each" test_for_each.cpp)
\ No newline at end of file

ADD
test/test_for_each.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/for_each.h>
+#include <thrust/device_ptr.h>
+#include <thrust/iterator/counting_iterator.h>
+
+// HIP API
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <hip/hip_runtime_api.h>
+
+#define HIP_CHECK(condition) ASSERT_EQ(condition, hipSuccess)
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+template<
+ class InputType
+>
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class ForEachTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<int>,
+ Params<unsigned short>
+> ForEachTestsParams;
+
+TYPED_TEST_CASE(ForEachTests, ForEachTestsParams);
+
+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 <typename T>
+struct mark_processed_functor
+{
+ T * ptr;
+ __host__ __device__ void operator()(T x){ ptr[static_cast<int>(x)] = 1; }
+};
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+TYPED_TEST(ForEachTests, SimpleTest)
+{
+ thrust::device_system_tag tag;
+ using T = typename TestFixture::input_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ auto ptr = thrust::malloc<T>(tag, sizeof(T) * size);
+ auto raw_ptr = thrust::raw_pointer_cast(ptr);
+ if(size > 0) ASSERT_NE(raw_ptr, nullptr);
+
+ // Zero input memory
+ if(size > 0) HIP_CHECK(hipMemset(raw_ptr, 0, sizeof(T) * size));
+
+ // Create unary function
+ mark_processed_functor<T> func;
+ func.ptr = raw_ptr;
+
+ // Run for_each in [0; end] range
+ auto end = size < 2 ? size : size/2;
+ auto result = thrust::for_each(
+ thrust::make_counting_iterator<size_t>(0),
+ thrust::make_counting_iterator<size_t>(end),
+ func
+ );
+ ASSERT_EQ(result, thrust::make_counting_iterator<size_t>(end));
+
+ std::vector<T> output(size);
+ HIP_CHECK(
+ hipMemcpy(
+ output.data(), raw_ptr,
+ size * sizeof(T),
+ hipMemcpyDeviceToHost
+ )
+ );
+
+ for(size_t i = 0; i < size; i++)
+ {
+ if(i < end)
+ {
+ ASSERT_EQ(output[i], T(1)) << "where index = " << i;
+ }
+ else
+ {
+ ASSERT_EQ(output[i], T(0)) << "where index = " << i;
+ }
+ }
+
+ // Free
+ thrust::free(tag, ptr);
+ }
+}
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

MODIFY
thrust/system/hip/detail/for_each.h
thrust/system/hip/detail/for_each.h
#pragma once
-#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
#include <iterator>
#include <thrust/system/hip/config.h>
@@ -52,7 +52,8 @@ namespace hip_rocprim {
: input(input), op(op) {}
template <class Size>
- THRUST_DEVICE_FUNCTION void operator()(Size idx)
+ THRUST_FUNCTION
+ void operator()(Size idx)
{
op(raw_reference_cast(input[idx]));
}

MODIFY
thrust/system/hip/detail/parallel_for.h
thrust/system/hip/detail/parallel_for.h
@@ -40,8 +40,51 @@ BEGIN_NS_THRUST
namespace hip_rocprim {
+// STREAMHPC TODO implement parallel_for (for_each) in rocPRIM
namespace __parallel_for {
+ template<unsigned int BlockSize, unsigned int ItemsPerThread>
+ struct kernel_config
+ {
+ static constexpr unsigned int block_size = BlockSize;
+ static constexpr unsigned int items_per_thread = ItemsPerThread;
+ };
+
+ template<unsigned int BlockSize,
+ unsigned int ItemsPerThread,
+ class F,
+ class Size>
+ __global__ void
+ kernel(F f, Size num_items)
+ {
+ constexpr auto items_per_block = BlockSize * ItemsPerThread;
+ Size tile_base = blockIdx.x * items_per_block;
+ Size num_remaining = num_items - tile_base;
+ const unsigned int items_in_tile = static_cast<unsigned int>(
+ num_remaining < items_per_block ? num_remaining : items_per_block
+ );
+
+ if(items_in_tile == items_per_block)
+ {
+ #pragma unroll
+ for(unsigned int i = 0; i < ItemsPerThread; i++)
+ {
+ unsigned int idx = BlockSize * i + threadIdx.x;
+ f(tile_base + idx);
+ }
+ }
+ else
+ {
+ #pragma unroll
+ for(unsigned int i = 0; i < ItemsPerThread; i++)
+ {
+ unsigned int idx = BlockSize * i + threadIdx.x;
+ if(idx < items_in_tile)
+ f(tile_base + idx);
+ }
+ }
+ }
+
template <class F,
class Size>
THRUST_RUNTIME_FUNCTION hipError_t
@@ -49,10 +92,27 @@ namespace __parallel_for {
F f,
hipStream_t stream)
{
+ using config = kernel_config<256, 1>;
+
bool debug_sync = THRUST_DEBUG_SYNC_FLAG;
- // STREAMHPC TODO implement parallel_for
+ // STREAMHPC Use debug_sync
(void) debug_sync;
- (void) num_items; (void) f; (void) stream;
+
+ constexpr unsigned int block_size = config::block_size;
+ constexpr unsigned int items_per_thread = config::items_per_thread;
+ constexpr auto items_per_block = block_size * items_per_thread;
+ const auto number_of_blocks = (num_items + items_per_block - 1)/items_per_block;
+
+ hipLaunchKernelGGL(
+ HIP_KERNEL_NAME(kernel<
+ block_size, items_per_thread, F, Size
+ >),
+ dim3(number_of_blocks), dim3(block_size), 0, stream,
+ f, num_items
+ );
+
+ auto error = hipPeekAtLastError();
+ if(error != hipSuccess) return error;
return hipSuccess;
}
} // __parallel_for
@@ -69,19 +129,20 @@ parallel_for(execution_policy<Derived> &policy,
if (count == 0)
return;
- if(__THRUST_HAS_HIPRT__)
+#if __THRUST_HAS_HIPRT__
{
+ std::cout << "HOST PATH" << std::endl;
hipStream_t stream = hip_rocprim::stream(policy);
hipError_t status = __parallel_for::parallel_for(count, f, stream);
hip_rocprim::throw_on_error(status, "parallel_for failed");
}
- else
+#else
{
-#if !__THRUST_HAS_HIPRT__
+ (void) policy;
for (Size idx = 0; idx != count; ++idx)
f(idx);
-#endif
}
+#endif
}
} // namespace hip_rocprim

HIP backend: add temporary parallel_for impl