[SYCLomatic] Refine the migration of thrust::device_malloc_allocator (#537)

tomflinda · web-flow · commit f7df54c40820 · 2023-04-23T09:26:35.000+08:00
Signed-off-by: chenwei.sun &lt;chenwei.sun@intel.com&gt;
diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp
@@ -2154,6 +2154,14 @@ bool TypeInDeclRule::replaceTemplateSpecialization(
     return false;
 
   const std::string RealTypeNameStr(Start, TyLen);
+
+  if (DpctGlobalInfo::getUsmLevel() == UsmLevel::UL_None &&
+      RealTypeNameStr.find("device_malloc_allocator") != std::string::npos) {
+    report(BeginLoc, Diagnostics::KNOWN_UNSUPPORTED_TYPE, false,
+            RealTypeNameStr);
+    return true;
+  }
+
   requestHelperFeatureForTypeNames(RealTypeNameStr, BeginLoc);
   std::string Replacement =
       MapNames::findReplacedName(MapNames::TypeNamesMap, RealTypeNameStr);
diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp
@@ -248,7 +248,8 @@ void MapNames::setExplicitNamespaceMap() {
            getDpctNamespace() + "device_vector",
            HelperFeatureEnum::DplExtrasVector_device_vector)},
       {"thrust::device_malloc_allocator",
-       std::make_shared<TypeNameRule>(getClNamespace() + "buffer_allocator")},
+       std::make_shared<TypeNameRule>(getDpctNamespace() + "deprecated::usm_device_allocator",
+       HelperFeatureEnum::Memory_usm_device_allocator_alias)},
       {"thrust::maximum",
        std::make_shared<TypeNameRule>("oneapi::dpl::maximum")},
       {"thrust::multiplies", std::make_shared<TypeNameRule>("std::multiplies")},
diff --git a/clang/runtime/dpct-rt/include/memory.hpp.inc b/clang/runtime/dpct-rt/include/memory.hpp.inc
@@ -1988,6 +1988,15 @@ template <typename T>
 using usm_host_allocator = detail::deprecated::usm_allocator<T, sycl::usm::alloc::host>;
 // DPCT_LABEL_END
 
+// DPCT_LABEL_BEGIN|usm_device_allocator_alias|dpct
+// DPCT_DEPENDENCY_BEGIN
+// Memory|dpct_usm_allocator
+// Device|get_default_queue
+// DPCT_DEPENDENCY_END
+// DPCT_CODE
+template <typename T>
+using usm_device_allocator = detail::deprecated::usm_allocator<T, sycl::usm::alloc::shared>;
+// DPCT_LABEL_END
 } // namespace deprecated
 
 // DPCT_LABEL_BEGIN|pointer_attributes|dpct
diff --git a/clang/test/dpct/helper_files_ref/include/memory.hpp b/clang/test/dpct/helper_files_ref/include/memory.hpp
@@ -1309,6 +1309,8 @@ namespace deprecated {
 template <typename T>
 using usm_host_allocator = detail::deprecated::usm_allocator<T, sycl::usm::alloc::host>;
 
+template <typename T>
+using usm_device_allocator = detail::deprecated::usm_allocator<T, sycl::usm::alloc::shared>;
 } // namespace deprecated
 
 class pointer_attributes {
diff --git a/clang/test/dpct/test_api_level/Memory/api_test46.cu b/clang/test/dpct/test_api_level/Memory/api_test46.cu
@@ -0,0 +1,23 @@
+// RUN: dpct --format-range=none   --use-custom-helper=api -out-root %T/Memory/api_test46_out %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
+// RUN: grep "IsCalled" %T/Memory/api_test46_out/MainSourceFiles.yaml | wc -l > %T/Memory/api_test46_out/count.txt
+// RUN: FileCheck --input-file %T/Memory/api_test46_out/count.txt --match-full-lines %s
+// RUN: rm -rf %T/Memory/api_test46_out
+
+// CHECK: 21
+// TEST_FEATURE: Memory_usm_device_allocator_alias
+
+#include <thrust/device_allocator.h>
+#include <thrust/device_malloc_allocator.h>
+#include <thrust/iterator/iterator_traits.h>
+
+namespace cusp
+{
+template <typename T, typename MemorySpace>
+struct default_memory_allocator
+    : thrust::device_malloc_allocator<T> 
+{};
+}
+
+int main(int argc, char *argv[]) {
+  return 0;
+}
diff --git a/clang/test/dpct/thrust-for-h2o4gpu-specific-case-noneusm.cu b/clang/test/dpct/thrust-for-h2o4gpu-specific-case-noneusm.cu
@@ -0,0 +1,47 @@
+// UNSUPPORTED: cuda-8.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-11.1, cuda-11.2, cuda-11.3
+// UNSUPPORTED: v8.0, v10.1, v10.2, v11.0, v11.1, v11.2, v11.3
+// RUN: dpct --usm-level=none --format-range=none -out-root %T/thrust-for-h2o4gpu-specific-case-noneusm %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
+// RUN: FileCheck --input-file %T/thrust-for-h2o4gpu-specific-case-noneusm/thrust-for-h2o4gpu-specific-case-noneusm.dp.cpp --match-full-lines %s
+
+// CHECK: #include <oneapi/dpl/execution>
+// CHECK-NEXT:#include <oneapi/dpl/algorithm>
+// CHECK-NEXT:#define DPCT_USM_LEVEL_NONE
+// CHECK-NEXT:#include <sycl/sycl.hpp>
+// CHECK-NEXT:#include <dpct/dpct.hpp>
+// CHECK-NEXT:#include <dpct/dpl_utils.hpp>
+#include <thrust/copy.h>
+#include <thrust/device_vector.h>
+#include <thrust/device_allocator.h>
+#include <thrust/device_malloc_allocator.h>
+#include <thrust/iterator/iterator_traits.h>
+
+// CHECK: /*
+// CHECK-NEXT:DPCT1082:{{[0-9]+}}: Migration of thrust::device_malloc_allocator type is not supported.
+// CHECK-NEXT:*/
+// CHECK-NEXT:template <typename T> void foo_cpy(dpct::device_vector<T, thrust::device_malloc_allocator<T>> &Do, dpct::device_vector<T, thrust::device_malloc_allocator<T>> &Di) {
+// CHECK-NEXT: return;
+// CHECK-NEXT:}
+template <typename T> void foo_cpy(thrust::device_vector<T, thrust::device_malloc_allocator<T>> &Do, thrust::device_vector<T, thrust::device_malloc_allocator<T>> &Di) {
+ return;
+}
+
+
+void foo_1() {
+// CHECK:  /*
+// CHECK-NEXT:  DPCT1082:{{[0-9]+}}: Migration of thrust::device_malloc_allocator type is not supported.
+// CHECK-NEXT:  */
+// CHECK-NEXT:  dpct::device_vector<int, thrust::device_malloc_allocator<int>> **tt=NULL;
+// CHECK-NEXT:  /*
+// CHECK-NEXT:  DPCT1082:{{[0-9]+}}: Migration of thrust::device_malloc_allocator type is not supported.
+// CHECK-NEXT:  */
+// CHECK-NEXT:  dpct::device_vector<int, thrust::device_malloc_allocator<int>> *dd[10];
+// CHECK-NEXT:  foo_cpy(*tt[0], *dd[0]);
+  thrust::device_vector<int, thrust::device_malloc_allocator<int>> **tt=NULL;
+  thrust::device_vector<int, thrust::device_malloc_allocator<int>> *dd[10];
+  foo_cpy(*tt[0], *dd[0]);
+}
+
+// CHECK:/*
+// CHECK-NEXT:DPCT1082:{{[0-9]+}}: Migration of thrust::device_malloc_allocator type is not supported.
+// CHECK-NEXT:*/
+namespace foo_struct { template <typename T, typename MemorySpace> struct default_memory_allocator : thrust::device_malloc_allocator<T>{};}
diff --git a/clang/test/dpct/thrust-for-h2o4gpu-specific-case.cu b/clang/test/dpct/thrust-for-h2o4gpu-specific-case.cu
@@ -1,5 +1,5 @@
-// UNSUPPORTED: cuda-8.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-11.1, cuda-11.2, cuda-11.3, cuda-11.4, cuda-11.5, cuda-11.6, cuda-11.7, cuda-11.8, cuda-12.0, cuda-12.1
-// UNSUPPORTED: v8.0, v10.1, v10.2, v11.0, v11.1, v11.2, v11.3, v11.4, v11.5, v11.6, v11.7, v11.8, v12.0, v12.1
+// UNSUPPORTED: cuda-8.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-11.1, cuda-11.2, cuda-11.3
+// UNSUPPORTED: v8.0, v10.1, v10.2, v11.0, v11.1, v11.2, v11.3
 // RUN: dpct --format-range=none -out-root %T/thrust-for-h2o4gpu-specific-case %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
 // RUN: FileCheck --input-file %T/thrust-for-h2o4gpu-specific-case/thrust-for-h2o4gpu-specific-case.dp.cpp --match-full-lines %s
 
@@ -11,24 +11,29 @@
 // CHECK-NEXT: #include <dpct/dpl_utils.hpp>
 #include <thrust/copy.h>
 #include <thrust/device_vector.h>
+#include <thrust/device_allocator.h>
+#include <thrust/device_malloc_allocator.h>
+#include <thrust/iterator/iterator_traits.h>
 
-//CHECK: template <typename T>
-//CHECK-NEXT:void foo_cpy(dpct::device_vector<T, sycl::buffer_allocator<T>> &Do, dpct::device_vector<T, sycl::buffer_allocator<T>> &Di) {
+//CHECK:template <typename T>
+//CHECK-NEXT:void foo_cpy(dpct::device_vector<T, dpct::deprecated::usm_device_allocator<T>> &Do, dpct::device_vector<T, dpct::deprecated::usm_device_allocator<T>> &Di) {
 //CHECK-NEXT: return;
 //CHECK-NEXT:}
 template <typename T>
 void foo_cpy(thrust::device_vector<T, thrust::device_malloc_allocator<T>> &Do, thrust::device_vector<T, thrust::device_malloc_allocator<T>> &Di) {
  return;
 }
 
-//CHECK: void foo_1() {
-//CHECK-NEXT: dpct::device_vector<int, sycl::buffer_allocator<int>> **tt=NULL;
-//CHECK-NEXT: dpct::device_vector<int, sycl::buffer_allocator<int>> *dd[10];
-//CHECK-NEXT: foo_cpy(*tt[0], *dd[0]);
+//CHECK:void foo_1() {
+//CHECK-NEXT:  dpct::device_vector<int, dpct::deprecated::usm_device_allocator<int>> **tt=NULL;
+//CHECK-NEXT:  dpct::device_vector<int, dpct::deprecated::usm_device_allocator<int>> *dd[10];
+//CHECK-NEXT:  foo_cpy(*tt[0], *dd[0]);
 //CHECK-NEXT:}
 void foo_1() {
   thrust::device_vector<int, thrust::device_malloc_allocator<int>> **tt=NULL;
   thrust::device_vector<int, thrust::device_malloc_allocator<int>> *dd[10];
   foo_cpy(*tt[0], *dd[0]);
 }
 
+//CHECK:namespace foo_struct { template <typename T, typename MemorySpace> struct default_memory_allocator : dpct::deprecated::usm_device_allocator<T>{};}
+namespace foo_struct { template <typename T, typename MemorySpace> struct default_memory_allocator : thrust::device_malloc_allocator<T>{};}
