Implement count and count_if functions. Add test_count test file.
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -58,6 +58,7 @@ endfunction()
add_hipthrust_test("thrust.hip.allocator" test_allocator.cpp)
add_hipthrust_test("thrust.hip.copy" test_copy.cpp)
+add_hipthrust_test("thrust.hip.count" test_count.cpp)
add_hipthrust_test("thrust.hip.counting_iterator" test_counting_iterator.cpp)
add_hipthrust_test("thrust.hip.device_delete" test_device_delete.cpp)
add_hipthrust_test("thrust.hip.device_ptr" test_device_ptr.cpp)

ADD
test/test_count.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.
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/count.h>
+#include <thrust/tabulate.h>
+#include <thrust/functional.h>
+#include <thrust/iterator/discard_iterator.h>
+#include <thrust/iterator/retag.h>
+#include <thrust/device_vector.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
+
+#include "test_utils.hpp"
+
+template<
+ class InputType
+>
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class CountTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<thrust::host_vector<short>>,
+ Params<thrust::host_vector<int>>,
+ Params<thrust::host_vector<long long>>,
+ Params<thrust::host_vector<unsigned short>>,
+ Params<thrust::host_vector<unsigned int>>,
+ Params<thrust::host_vector<unsigned long long>>,
+ Params<thrust::host_vector<float>>,
+ Params<thrust::host_vector<double>>,
+ Params<thrust::device_vector<short>>,
+ Params<thrust::device_vector<int>>,
+ Params<thrust::device_vector<long long>>,
+ Params<thrust::device_vector<unsigned short>>,
+ Params<thrust::device_vector<unsigned int>>,
+ Params<thrust::device_vector<unsigned long long>>,
+ Params<thrust::device_vector<float>>,
+ Params<thrust::device_vector<double>>
+> CountTestsParams;
+
+TYPED_TEST_CASE(CountTests, CountTestsParams);
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+TYPED_TEST(CountTests, TestCountSimple)
+{
+ using Vector = typename TestFixture::input_type;
+
+ Vector data(5);
+ data[0] = 1; data[1] = 1; data[2] = 0; data[3] = 0; data[4] = 1;
+
+ ASSERT_EQ(thrust::count(data.begin(), data.end(), 0), 2);
+ ASSERT_EQ(thrust::count(data.begin(), data.end(), 1), 3);
+ ASSERT_EQ(thrust::count(data.begin(), data.end(), 2), 0);
+}
+
+TYPED_TEST(CountTests, TestCount)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<T> h_data = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+ thrust::device_vector<T> d_data = h_data;
+
+ size_t cpu_result = thrust::count(h_data.begin(), h_data.end(), T(5));
+ size_t gpu_result = thrust::count(d_data.begin(), d_data.end(), T(5));
+
+ ASSERT_EQ(cpu_result, gpu_result);
+ }
+}
+
+template <typename T>
+struct greater_than_five
+{
+ __host__ __device__ bool operator()(const T &x) const {return x > 5;}
+};
+
+TYPED_TEST(CountTests, TestCountIfSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector data(5);
+ data[0] = 1; data[1] = 6; data[2] = 1; data[3] = 9; data[4] = 2;
+
+ ASSERT_EQ(thrust::count_if(data.begin(), data.end(), greater_than_five<T>()), 2);
+
+}
+
+TYPED_TEST(CountTests, TestCountIf)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<T> h_data = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+ thrust::device_vector<T> d_data = h_data;
+ size_t cpu_result = thrust::count_if(h_data.begin(), h_data.end(), greater_than_five<T>());
+ size_t gpu_result = thrust::count_if(d_data.begin(), d_data.end(), greater_than_five<T>());
+
+ ASSERT_EQ(cpu_result, gpu_result);
+ }
+}
+
+TYPED_TEST(CountTests, TestCountFromConstIteratorSimple)
+{
+ using Vector = typename TestFixture::input_type;
+
+ Vector data(5);
+ data[0] = 1; data[1] = 1; data[2] = 0; data[3] = 0; data[4] = 1;
+
+ ASSERT_EQ(thrust::count(data.cbegin(), data.cend(), 0), 2);
+ ASSERT_EQ(thrust::count(data.cbegin(), data.cend(), 1), 3);
+ ASSERT_EQ(thrust::count(data.cbegin(), data.cend(), 2), 0);
+}
+
+template<typename InputIterator, typename EqualityComparable>
+int count(my_system &system, InputIterator, InputIterator, EqualityComparable x)
+{
+ system.validate_dispatch();
+ return x;
+}
+
+TEST(CountTests, TestCountDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::count(sys,
+ vec.begin(),
+ vec.end(),
+ 13);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename InputIterator, typename EqualityComparable>
+int count(my_tag, InputIterator, InputIterator, EqualityComparable x)
+{
+ return x;
+}
+
+TEST(EqualTests, TestEqualDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ int result = thrust::count(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.end()),
+ 13);
+
+ ASSERT_EQ(13, result);
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

MODIFY
thrust/system/detail/adl/count.h
thrust/system/detail/adl/count.h
#if 0
#include <thrust/system/cpp/detail/count.h>
#include <thrust/system/cuda/detail/count.h>
+#include <thrust/system/hip/detail/count.h>
#include <thrust/system/omp/detail/count.h>
#include <thrust/system/tbb/detail/count.h>
#endif

MODIFY
thrust/system/hip/detail/count.h
thrust/system/hip/detail/count.h
-// STREAMHPC TODO implement
+/******************************************************************************
+ * Copyright (c) 2016, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ * * Neither the name of the NVIDIA CORPORATION nor the
+ * names of its contributors may be used to endorse or promote products
+ * derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+#pragma once
+
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
+#include <thrust/system/hip/config.h>
+
+#include <thrust/system/hip/detail/util.h>
+#include <thrust/system/hip/detail/reduce.h>
+#include <thrust/distance.h>
+
+BEGIN_NS_THRUST
+namespace hip_rocprim {
+
+ template <class Derived,
+ class InputIt,
+ class UnaryPred>
+ typename iterator_traits<InputIt>::difference_type __host__ __device__
+ count_if(execution_policy<Derived> &policy,
+ InputIt first,
+ InputIt last,
+ UnaryPred unary_pred)
+ {
+ typedef typename iterator_traits<InputIt>::difference_type size_type;
+ typedef transform_input_iterator_t<size_type,
+ InputIt,
+ UnaryPred>
+ flag_iterator_t;
+
+ return hip_rocprim::reduce_n(policy,
+ flag_iterator_t(first, unary_pred),
+ thrust::distance(first, last),
+ size_type(0),
+ plus<size_type>());
+ }
+
+ template <class Derived,
+ class InputIt,
+ class Value>
+ typename iterator_traits<InputIt>::difference_type __host__ __device__
+ count(execution_policy<Derived> &policy,
+ InputIt first,
+ InputIt last,
+ Value const & value)
+ {
+ return hip_rocprim::count_if(policy,
+ first,
+ last,
+ detail::equal_to_value<Value>(value));
+ }
+
+} // namespace cuda_cub
+END_NS_THRUST
+#endif

Implement count and count_if functions. Add test_count test file.