[SYCLomatic] Add migration for the thrust::device_reference type (#319)

Signed-off-by: Cai, Justin <justin.cai@intel.com>
oneapi-src · Nov 9, 2022 · 579be32 · 579be32
1 parent 1f0aa19
commit 579be32
Show file tree

Hide file tree

Showing 4 changed files with 45 additions and 1 deletion.
diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp
@@ -2444,7 +2444,8 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
               hasAnyName(
                   "cudaError", "curandStatus", "cublasStatus", "CUstream",
                   "CUstream_st", "thrust::complex", "thrust::device_vector",
-                  "thrust::device_ptr", "thrust::host_vector", "cublasHandle_t",
+                  "thrust::device_ptr", "thrust::device_reference",
+		  "thrust::host_vector", "cublasHandle_t",
                   "CUevent_st", "__half", "half", "__half2", "half2",
                   "cudaMemoryAdvise", "cudaError_enum", "cudaDeviceProp",
                   "cudaPitchedPtr", "thrust::counting_iterator",

diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp
@@ -228,6 +228,9 @@ void MapNames::setExplicitNamespaceMap() {
        std::make_shared<TypeNameRule>(
            getDpctNamespace() + "device_pointer",
            HelperFeatureEnum::DplExtrasMemory_device_pointer_forward_decl)},
+      {"thrust::device_reference",
+       std::make_shared<TypeNameRule>(getDpctNamespace() + "device_reference",
+				      HelperFeatureEnum::DplExtrasMemory_device_reference)},
       {"thrust::device_vector",
        std::make_shared<TypeNameRule>(
            getDpctNamespace() + "device_vector",

diff --git a/clang/test/dpct/test_api_level/DplExtrasMemory/api_test11.cu b/clang/test/dpct/test_api_level/DplExtrasMemory/api_test11.cu
@@ -0,0 +1,19 @@
+// UNSUPPORTED: cuda-8.0
+// UNSUPPORTED: v8.0
+// RUN: dpct --format-range=none  --usm-level=none  --use-custom-helper=api -out-root %T/DplExtrasMemory/api_test11_out %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only -std=c++17
+// RUN: grep "IsCalled" %T/DplExtrasMemory/api_test11_out/MainSourceFiles.yaml | wc -l > %T/DplExtrasMemory/api_test11_out/count.txt
+// RUN: FileCheck --input-file %T/DplExtrasMemory/api_test11_out/count.txt --match-full-lines %s
+// RUN: rm -rf %T/DplExtrasMemory/api_test11_out
+
+// CHECK: 37
+// TEST_FEATURE: DplExtrasMemory_device_reference
+
+
+#include <thrust/device_vector.h>
+
+int main() {
+  thrust::device_vector<int> dvec(1, 13);
+  thrust::device_reference<int> r1 = dvec[0];
+  thrust::device_reference<int> r2 = dvec[0];
+  r2 = r1 + r2;
+}
diff --git a/clang/test/dpct/types002.cu b/clang/test/dpct/types002.cu
@@ -27,6 +27,15 @@ int main(int argc, char **argv) {
   a = sizeof(device_p);
   a = sizeof device_p;
 
+  //CHECK:dpct::device_reference<int> device_ref = device_vec[0];
+  //CHECK-NEXT:a = sizeof(dpct::device_reference<int>);
+  //CHECK-NEXT:a = sizeof(device_ref);
+  //CHECK-NEXT:a = sizeof device_ref;
+  thrust::device_reference<int> device_ref = device_vec[0];
+  a = sizeof(thrust::device_reference<int>);
+  a = sizeof(device_ref);
+  a = sizeof device_ref;
+
   //CHECK:std::vector<int> host_vec;
   //CHECK-NEXT:a = sizeof(std::vector<int>);
   //CHECK-NEXT:a = sizeof(host_vec);
@@ -37,6 +46,18 @@ int main(int argc, char **argv) {
   a = sizeof host_vec;
 }
 
+template <typename T>
+struct alloc {
+  // CHECK:      typedef dpct::device_pointer<T> pointer;
+  // CHECK-NEXT: typedef dpct::device_pointer<const T> const_pointer;
+  // CHECK-NEXT: typedef dpct::device_reference<T> reference;
+  // CHECK-NEXT: typedef dpct::device_reference<const T> const_reference;
+  typedef thrust::device_ptr<T> pointer;
+  typedef thrust::device_ptr<const T> const_pointer;
+  typedef thrust::device_reference<T> reference;
+  typedef thrust::device_reference<const T> const_reference;
+};
+
 template <typename type>
 struct bar
 {