[SYCLomatic] Migrate thrust::zip_iterator to dpct::zip_iterator (#563)

Introduce help type dpct::zip_iterator which only accept std::tuple as template argument to migrate thrust::zip_iterator.


Signed-off-by: Jiang, Zhiwei <[email protected]>

Author: zhiweij1

clang/lib/DPCT/ASTTraversal.cpp

@@ -2172,7 +2172,7 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
                   "cudaDataType_t", "cudaDataType", "cublasComputeType_t",
                   "cublasAtomicsMode_t", "CUmem_advise_enum", "CUmem_advise",
                   "thrust::tuple_element", "thrust::tuple_size", "cublasMath_t",
-                  "cudaPointerAttributes")
+                  "cudaPointerAttributes", "thrust::zip_iterator")
               )))))
           .bind("cudaTypeDef"),
       this);

clang/lib/DPCT/MapNames.cpp

@@ -282,6 +282,10 @@ void MapNames::setExplicitNamespaceMap() {
        std::make_shared<TypeNameRule>("std::tuple_size")},
        {"thrust::swap",
        std::make_shared<TypeNameRule>("std::swap")},
+       {"thrust::zip_iterator",
+       std::make_shared<TypeNameRule>(
+           getDpctNamespace() + "zip_iterator",
+           HelperFeatureEnum::DplExtrasIterators_zip_iterator)},
       {"cusolverDnHandle_t",
        std::make_shared<TypeNameRule>(getClNamespace() + "queue*")},
       {"cusolverEigType_t", std::make_shared<TypeNameRule>("int64_t")},

clang/runtime/dpct-rt/include/dpl_extras/iterators.h.inc

@@ -212,8 +212,29 @@ template <typename KeyTp, typename _ValueTp> struct make_key_value_pair {
 };
 // DPCT_LABEL_END
 
+// DPCT_LABEL_BEGIN|zip_iterator_impl|dpct::detail
+// DPCT_DEPENDENCY_EMPTY
+// DPCT_CODE
+template <class T> struct __zip_iterator_impl;
+template <class... Ts> struct __zip_iterator_impl<std::tuple<Ts...>> {
+  using type = oneapi::dpl::zip_iterator<Ts...>;
+};
+// DPCT_LABEL_END
+
 } // end namespace detail
 
+// DPCT_LABEL_BEGIN|zip_iterator|dpct
+// DPCT_DEPENDENCY_BEGIN
+// DplExtrasIterators|zip_iterator_impl
+// DPCT_DEPENDENCY_END
+// DPCT_CODE
+// dpct::zip_iterator can only accept std::tuple type as template argument for
+// compatibility purpose. Please use oneapi::dpl::zip_iterator if you want to
+// pass iterator's types directly.
+template <typename... Ts>
+using zip_iterator = typename detail::__zip_iterator_impl<Ts...>::type;
+// DPCT_LABEL_END
+
 // DPCT_LABEL_BEGIN|arg_index_input_iterator|dpct
 // DPCT_DEPENDENCY_BEGIN
 // DplExtrasIterators|make_key_value_pair

clang/test/dpct/helper_files_ref/include/dpl_extras/iterators.h

@@ -146,8 +146,19 @@ template <typename KeyTp, typename _ValueTp> struct make_key_value_pair {
   }
 };
 
+template <class T> struct __zip_iterator_impl;
+template <class... Ts> struct __zip_iterator_impl<std::tuple<Ts...>> {
+  using type = oneapi::dpl::zip_iterator<Ts...>;
+};
+
 } // end namespace detail
 
+// dpct::zip_iterator can only accept std::tuple type as template argument for
+// compatibility purpose. Please use oneapi::dpl::zip_iterator if you want to
+// pass iterator's types directly.
+template <typename... Ts>
+using zip_iterator = typename detail::__zip_iterator_impl<Ts...>::type;
+
 // arg_index_input_iterator is an iterator over a input iterator, with a index.
 // When dereferenced, it returns a key_value_pair, which can be interrogated for
 // the index key or the value from the input iterator

clang/test/dpct/test_api_level/DplExtrasIterators/api_test6.cu

@@ -0,0 +1,20 @@
+// RUN: dpct --format-range=none --use-custom-helper=api -out-root %T/DplExtrasIterators/api_test6_out %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only -std=c++17 -fno-delayed-template-parsing
+// RUN: grep "IsCalled" %T/DplExtrasIterators/api_test6_out/MainSourceFiles.yaml | wc -l > %T/DplExtrasIterators/api_test6_out/count.txt
+// RUN: FileCheck --input-file %T/DplExtrasIterators/api_test6_out/count.txt --match-full-lines %s
+// RUN: rm -rf %T/DplExtrasIterators/api_test6_out
+
+// CHECK: 4
+// TEST_FEATURE: DplExtrasIterators_zip_iterator
+
+#include <thrust/iterator/zip_iterator.h>
+#include <thrust/tuple.h>
+
+template<typename int_iterator>
+void foo() {
+  typedef thrust::tuple<int_iterator, int_iterator> iterator_tuple;
+  typedef thrust::zip_iterator<iterator_tuple> int_zip_iterator;
+}
+
+int main() {
+  return 0;
+}

clang/test/dpct/thrust-for-h2o4gpu.cu

@@ -225,14 +225,20 @@ void foo() {
  {
   //CHECK: dpct::device_vector<int> int_in(3);
   //CHECK-NEXT: dpct::device_vector<float> float_in(3);
-  //CHECK-NEXT: auto ret = oneapi::dpl::make_zip_iterator(std::make_tuple(int_in.begin(), float_in.begin()));
+  //CHECK-NEXT: typedef dpct::device_vector<int>::iterator int_iterator;
+  //CHECK-NEXT: typedef dpct::device_vector<float>::iterator float_iterator;
+  //CHECK-NEXT: typedef std::tuple<int_iterator, float_iterator> iterator_tuple;
+  //CHECK-NEXT: dpct::zip_iterator<iterator_tuple> ret = oneapi::dpl::make_zip_iterator(std::make_tuple(int_in.begin(), float_in.begin()));
   //CHECK-NEXT: auto arg = std::make_tuple(int_in.begin(), float_in.begin());
-  //CHECK-NEXT: auto ret_1 = oneapi::dpl::make_zip_iterator(arg);
+  //CHECK-NEXT: dpct::zip_iterator<iterator_tuple> ret_1 = oneapi::dpl::make_zip_iterator(arg);
   thrust::device_vector<int> int_in(3);
   thrust::device_vector<float> float_in(3);
-  auto ret = thrust::make_zip_iterator(thrust::make_tuple(int_in.begin(), float_in.begin()));
+  typedef thrust::device_vector<int>::iterator int_iterator;
+  typedef thrust::device_vector<float>::iterator float_iterator;
+  typedef thrust::tuple<int_iterator, float_iterator> iterator_tuple;
+  thrust::zip_iterator<iterator_tuple> ret = thrust::make_zip_iterator(thrust::make_tuple(int_in.begin(), float_in.begin()));
   auto arg = thrust::make_tuple(int_in.begin(), float_in.begin());
-  auto ret_1 = thrust::make_zip_iterator(arg);
+  thrust::zip_iterator<iterator_tuple> ret_1 = thrust::make_zip_iterator(arg);
  }
 
  {