Compile Errors saying "reference to __host__ function while calling function on device" in the following algorithms:
Adjacent Difference
Merge
Merge by Key
Partition
Find
MODIFY
test/test_adjacent_difference.cpp
test/test_adjacent_difference.cpp
@@ -238,4 +238,44 @@ TEST(AdjacentDifferenceTests, TestAdjacentDifferenceDispatchImplicit)
ASSERT_EQ(13, d_input.front());
}
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void AdjacentDifferenceKernel(int const N, int* array)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> begin(array);
+ thrust::device_ptr<int> end(array + N);
+ thrust::advance(begin,2);
+ thrust::adjacent_difference(thrust::hip::par, begin,end,begin);
+ }
+}
+
+TEST(AdjacentDifferenceTests, TestAdjacentDifferenceDevice)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ for(auto size: {0, 1, 2, 4, 6, 12, 16, 24, 32, 64, 84, 128, 160, 256} )
+ {
+ SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+ for(auto seed : get_seeds())
+ {
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+ thrust::host_vector<int> h_data = get_random_data<int>(size, 0, size, seed);
+ thrust::device_vector<int> d_data = h_data;
+ thrust::adjacent_difference(h_data.begin(),h_data.end(),h_data.begin());
+ hipLaunchKernelGGL(AdjacentDifferenceKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_data[0]));
+
+ ASSERT_EQ(h_data, d_data);
+ }
+ }
+}
+
#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HIP

MODIFY
test/test_fill.cpp
test/test_fill.cpp
@@ -535,3 +535,50 @@ TEST(FillTests, TestFillNDispatchImplicit)
ASSERT_EQ(13, vec.front());
}
+
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void FillKernel(int const N, int* array, int fill_value)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> begin(array);
+ thrust::device_ptr<int> end(array + N);
+ thrust::fill(thrust::hip::par, begin, end,fill_value);
+ }
+}
+
+TEST(FillTests, TestFillDevice)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+ for(auto size: {0, 1, 2, 4, 6, 12, 16, 24, 32, 64, 84, 128, 160, 256} )
+ {
+ SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+ for(auto seed : get_seeds())
+ {
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+ thrust::host_vector<int> h_data = get_random_data<int>(size, 0, size, seed);
+ thrust::device_vector<int> d_data = h_data;
+
+ int fill_value = get_random_data<int>(1,0,size,seed)[0];
+ SCOPED_TRACE(testing::Message() << "with fill_value= " <<fill_value);
+
+ thrust::fill(h_data.begin(), h_data.end(),fill_value);
+ hipLaunchKernelGGL(FillKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_data[0]),
+ fill_value);
+
+ ASSERT_EQ(h_data, d_data);
+ }
+ }
+}

MODIFY
test/test_find.cpp
test/test_find.cpp
@@ -385,3 +385,54 @@ TYPED_TEST(FindTests, TestFindIfNot)
}
}
}
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void FindKernel(int const N, int* in_array, int value, int *out_array)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> in_begin(in_array);
+ thrust::device_ptr<int> in_end(in_array + N);
+
+ auto x = thrust::find(thrust::hip::par, in_begin, in_end,value);
+ }
+}
+TEST(FindTests, TestFindDevice)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ for(auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size= " << size);
+ for(auto seed : get_seeds())
+ {
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+ thrust::host_vector<int> h_data = get_random_data<int>(size, std::numeric_limits<int>::min(), std::numeric_limits<int>::max(), seed);
+ thrust::device_vector<int> d_data = h_data;
+ thrust::device_vector<int> d_output(1);
+
+
+ int h_index = thrust::find(h_data.begin(), h_data.end(), 0) - h_data.begin();
+
+ hipLaunchKernelGGL(FindKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_data[0]),
+ 0,
+ thrust::raw_pointer_cast(&d_output[0]));
+ ASSERT_EQ(h_index,d_output[0]);
+
+ // for(size_t i = 1; i < size; i *= 2)
+ // {
+ // T sample = h_data[i];
+ // h_iter = thrust::find(h_data.begin(), h_data.end(), sample);
+ // d_iter = thrust::find(d_data.begin(), d_data.end(), sample);
+ // ASSERT_EQ(h_iter - h_data.begin(), d_iter - d_data.begin());
+ // }
+ }
+ }
+}

MODIFY
test/test_merge.cpp
test/test_merge.cpp
@@ -276,3 +276,78 @@ TYPED_TEST(PrimitiveMergeTests, MergeDescending)
}
}
}
+
+template<class T>
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void MergeKernel(int const N, T* inA_array, T* inB_array, T *out_array)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> inA_begin(inA_array);
+ thrust::device_ptr<int> inA_end(inA_array + N);
+ thrust::device_ptr<int> inB_begin(inB_array);
+ thrust::device_ptr<int> inB_end(inB_array + N);
+ thrust::device_ptr<int> out_begin(out_array);
+
+ thrust::merge(thrust::hip::par, inA_begin, inA_end, inB_begin, inB_end,out_begin);
+ thrust::merge_by_key(thrust::hip::par, inA_begin, inA_end, inB_begin, inB_end,inA_begin,inB_begin,out_begin,out_begin);
+
+ }
+}
+TEST(PrimitiveMergeTests, TestMergeDevice)
+{
+ using T = int;
+
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+ for(auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+ for(auto seed : get_seeds())
+ {
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+ thrust::host_vector<T> h_a = get_random_data<T>(
+ size, std::numeric_limits<T>::min(), std::numeric_limits<T>::max(), seed);
+ thrust::host_vector<T> h_b = get_random_data<T>(
+ size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max(),
+ seed + seed_value_addition
+ );
+
+ thrust::stable_sort(h_a.begin(), h_a.end(), thrust::greater<T>());
+ thrust::stable_sort(h_b.begin(), h_b.end(), thrust::greater<T>());
+
+ thrust::device_vector<T> d_a = h_a;
+ thrust::device_vector<T> d_b = h_b;
+
+ thrust::host_vector<T> h_result(h_a.size() + h_b.size());
+ thrust::device_vector<T> d_result(d_a.size() + d_b.size());
+
+ typename thrust::host_vector<T>::iterator h_end;
+ typename thrust::device_vector<T>::iterator d_end;
+
+ h_end = thrust::merge(h_a.begin(),
+ h_a.end(),
+ h_b.begin(),
+ h_b.end(),
+ h_result.begin(),
+ thrust::greater<T>());
+
+ hipLaunchKernelGGL(MergeKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_a[0]),
+ thrust::raw_pointer_cast(&d_b[0]),
+ thrust::raw_pointer_cast(&d_result[0]));
+
+ ASSERT_EQ(h_result, d_result);
+ }
+ }
+}

MODIFY
test/test_partition.cpp
test/test_partition.cpp
@@ -1796,3 +1796,47 @@ TEST(PartitionTests, TestStablePartitionCopyStencilDispatchImplicit)
ASSERT_EQ(13, vec.front());
}
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void PartitionKernel(int const N, int* array)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> begin(array);
+ thrust::device_ptr<int> end(array + N);
+ thrust::partition(thrust::hip::par, begin, end,is_even<int>());
+
+ }
+}
+
+TEST(PartitionTests,TestPartitionDevice)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+ for(auto size: {0, 1, 2, 4, 6, 12, 16, 24, 32, 64, 84, 128, 160, 256} )
+ {
+ SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+ for(auto seed : get_seeds())
+ {
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+ thrust::host_vector<int> h_data = get_random_data<int>(size, 0, size, seed);
+
+ thrust::device_vector<int> d_data = h_data;
+
+ thrust::partition(h_data.begin(), h_data.end(),is_even<int>());
+
+ hipLaunchKernelGGL(PartitionKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_data[0]));
+
+ ASSERT_EQ(h_data, d_data);
+ }
+ }
+}

Compile Errors saying "reference to __host__ function while calling function on device" in the following algorithms:
- Adjacent Difference
- Merge
- Merge by Key
- Partition
- Find