HIP backend: Fix for_each host path
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -55,5 +55,5 @@ endfunction()
# Tests
# ****************************************************************************
-add_hipthrust_test("thrust.hip.memory" test_memory.cpp)
add_hipthrust_test("thrust.hip.for_each" test_for_each.cpp)
+add_hipthrust_test("thrust.hip.memory" test_memory.cpp)

MODIFY
test/test_for_each.cpp
test/test_for_each.cpp
// Thrust
#include <thrust/memory.h>
#include <thrust/for_each.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/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
@@ -81,7 +84,7 @@ struct mark_processed_functor
};
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
-TYPED_TEST(ForEachTests, SimpleTest)
+TYPED_TEST(ForEachTests, HostPathSimpleTest)
{
thrust::device_system_tag tag;
using T = typename TestFixture::input_type;
@@ -136,4 +139,68 @@ TYPED_TEST(ForEachTests, SimpleTest)
thrust::free(tag, ptr);
}
}
+
+template<class F>
+__global__
+void simple_test_kernel(F func, int size)
+{
+ // (void) func; (void) size;
+ thrust::for_each(
+ thrust::seq,
+ thrust::make_counting_iterator<int>(0),
+ thrust::make_counting_iterator<int>(size),
+ func
+ );
+}
+
+TYPED_TEST(ForEachTests, DevicePathSimpleTest)
+{
+ thrust::device_system_tag tag;
+ using T = typename TestFixture::input_type;
+ const size_t size = 1024;
+
+ auto ptr = thrust::malloc<T>(tag, sizeof(T) * size);
+ auto raw_ptr = thrust::raw_pointer_cast(ptr);
+ ASSERT_NE(raw_ptr, nullptr);
+
+ // Zero input memory
+ 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
+ size_t end = 375;
+ hipLaunchKernelGGL(
+ HIP_KERNEL_NAME(simple_test_kernel<mark_processed_functor<T>>),
+ dim3(1), dim3(1), 0, 0,
+ func, static_cast<int>(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/parallel_for.h
thrust/system/hip/detail/parallel_for.h
@@ -126,6 +126,13 @@ parallel_for(execution_policy<Derived> &policy,
if (count == 0)
return;
+ // Workaround, so kernel called by __parallel_for::parallel_for is not lost,
+ // Implicit instantiation of __parallel_for::parallel_for function template
+ // that will be used in #if __THRUST_HAS_HIPRT__ block.
+ {
+ auto ptr = __parallel_for::parallel_for<F, Size>;
+ (void) ptr;
+ }
#if __THRUST_HAS_HIPRT__
{
hipStream_t stream = hip_rocprim::stream(policy);

HIP backend: Fix for_each host path