Implement tabulate
MODIFY
testing/CMakeLists.txt
testing/CMakeLists.txt
@@ -160,7 +160,7 @@ add_thrust_test("stable_sort_by_key_large")
add_thrust_test("stable_sort")
add_thrust_test("stable_sort_large")
# add_thrust_test("swap_ranges")
-# add_thrust_test("tabulate")
+add_thrust_test("tabulate")
add_thrust_test("transform")
# add_thrust_test("transform_iterator")
# add_thrust_test("transform_output_iterator")

MODIFY
testing/tabulate.cu
testing/tabulate.cu
template<typename ForwardIterator, typename UnaryOperation>
-void tabulate(my_system &system, ForwardIterator first, ForwardIterator, UnaryOperation unary_op)
+void tabulate(my_system &system, ForwardIterator, ForwardIterator, UnaryOperation)
{
system.validate_dispatch();
}
@@ -24,7 +24,7 @@ DECLARE_UNITTEST(TestTabulateDispatchExplicit);
template<typename ForwardIterator, typename UnaryOperation>
-void tabulate(my_tag, ForwardIterator first, ForwardIterator, UnaryOperation unary_op)
+void tabulate(my_tag, ForwardIterator first, ForwardIterator, UnaryOperation)
{
*first = 13;
}
@@ -108,4 +108,3 @@ void TestTabulateToDiscardIterator(size_t n)
// nothing to check -- just make sure it compiles
}
DECLARE_VARIABLE_UNITTEST(TestTabulateToDiscardIterator);
-

ADD
thrust/system/hip/detail/sequence.h
+/*
+ * Copyright 2008-2013 NVIDIA Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+
+// this system has no special version of this algorithm
+

ADD
thrust/system/hip/detail/tabulate.h
+/******************************************************************************
+ * 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_HCC
+#include <thrust/distance.h>
+#include <thrust/system/hip/config.h>
+#include <thrust/system/hip/execution_policy.h>
+#include <thrust/system/hip/detail/transform.h>
+
+BEGIN_NS_THRUST
+namespace hip_rocprim {
+
+template <class Derived,
+ class Iterator,
+ class TabulateOp>
+void __host__ __device__
+tabulate(execution_policy<Derived>& policy,
+ Iterator first,
+ Iterator last,
+ TabulateOp tabulate_op)
+{
+ typedef typename iterator_traits<Iterator>::difference_type size_type;
+
+ hip_rocprim::transform(policy,
+ thrust::make_counting_iterator<size_type>(0),
+ thrust::make_counting_iterator<size_type>(thrust::distance(first, last)),
+ first,
+ [tabulate_op](size_type i) { return tabulate_op(i); });
+}
+
+} // namespace hip_rocprim
+END_NS_THRUST
+#endif

MODIFY
thrust/system/hip/execution_policy.h
thrust/system/hip/execution_policy.h
// #include <thrust/system/hip/detail/reverse.h>
// #include <thrust/system/hip/detail/scatter.h>
#include <thrust/system/hip/detail/swap_ranges.h>
-// #include <thrust/system/hip/detail/tabulate.h>
+#include <thrust/system/hip/detail/tabulate.h>
#include <thrust/system/hip/detail/transform.h>
// #include <thrust/system/hip/detail/transform_reduce.h>
// #include <thrust/system/hip/detail/transform_scan.h>

Implement tabulate