From 5d412b049d6443045c270612dad2f8761aaa0c5e Mon Sep 17 00:00:00 2001
From: "Cai, Justin" <justin.cai@intel.com>
Date: Fri, 21 Oct 2022 11:28:16 -0700
Subject: [PATCH] [SYCLomatic] Add migration for the thrust::device_reference
 type

Signed-off-by: Cai, Justin <justin.cai@intel.com>
---
 clang/lib/DPCT/ASTTraversal.cpp               |  3 ++-
 clang/lib/DPCT/MapNames.cpp                   |  3 +++
 .../DplExtrasMemory/api_test11.cu             | 19 +++++++++++++++++
 clang/test/dpct/types002.cu                   | 21 +++++++++++++++++++
 4 files changed, 45 insertions(+), 1 deletion(-)
 create mode 100644 clang/test/dpct/test_api_level/DplExtrasMemory/api_test11.cu

diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp
index 9f7ffee1083e..bdd8cff7fff6 100644
--- a/clang/lib/DPCT/ASTTraversal.cpp
+++ b/clang/lib/DPCT/ASTTraversal.cpp
@@ -2428,7 +2428,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
index e339b96a3989..3674f74afd5e 100644
--- 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
new file mode 100644
index 000000000000..d4478f5673cb
--- /dev/null
+++ 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
index 1c1b34780921..9a66ec9f5cb8 100644
--- 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
 {