[SYCLomatic] Support the cuPointerGetAttribute API migration when USM is none(#2737)

Signed-off-by: Chen, Sheng S <[email protected]>

Author: ShengchenJ

clang/lib/DPCT/RuleInfra/MapNames.cpp

@@ -1226,9 +1226,6 @@ void MapNames::setExplicitNamespaceMap(
       {"CU_POINTER_ATTRIBUTE_IS_LEGACY_CUDA_IPC_CAPABLE",
        std::make_shared<EnumNameRule>(getDpctNamespace() +
                                       "pointer_attributes::type::unsupported")},
-      {"CU_POINTER_ATTRIBUTE_RANGE_START_ADDR",
-       std::make_shared<EnumNameRule>(getDpctNamespace() +
-                                      "pointer_attributes::type::unsupported")},
       {"CU_POINTER_ATTRIBUTE_RANGE_SIZE",
        std::make_shared<EnumNameRule>(getDpctNamespace() +
                                       "pointer_attributes::type::unsupported")},

clang/lib/DPCT/RulesLang/APINamesMemory.inc

@@ -440,6 +440,20 @@ ASSIGNABLE_FACTORY(
             "cuPointerGetAttributes", CALL(MapNames::getDpctNamespace() + "pointer_attributes::get",
             ARG(0), ARG(1), ARG(2), ARG(3)))))
 
+ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
+    CheckEnumArgStr(1, "CU_POINTER_ATTRIBUTE_RANGE_START_ADDR"),
+    CONDITIONAL_FACTORY_ENTRY(
+        checkIsUSM(),
+        UNSUPPORT_FACTORY_ENTRY("cuPointerGetAttribute",
+                                Diagnostics::UNSUPPORTED_PARAM, ARG(1)),
+        FEATURE_REQUEST_FACTORY(
+            HelperFeatureEnum::device_ext,
+            ASSIGN_FACTORY_ENTRY(
+                "cuPointerGetAttribute", ARG(0),
+                CALL(MapNames::getDpctNamespace() + "get_base_addr", ARG(2))))),
+    UNSUPPORT_FACTORY_ENTRY("cuPointerGetAttribute",
+                            Diagnostics::UNSUPPORTED_PARAM, ARG(1))))
+
 ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
     HelperFeatureEnum::device_ext,
     CONDITIONAL_FACTORY_ENTRY(

clang/lib/DPCT/RulesLang/RewriterSYCLcompat.cpp

@@ -77,6 +77,7 @@ SYCLCOMPAT_UNSUPPORT("cub::StoreDirectWarpStriped")
 SYCLCOMPAT_UNSUPPORT("cub::ShuffleDown")
 SYCLCOMPAT_UNSUPPORT("cub::ShuffleUp")
 SYCLCOMPAT_UNSUPPORT("cuPointerGetAttributes")
+SYCLCOMPAT_UNSUPPORT("cuPointerGetAttribute")
   });
   // clang-format on
   initRewriterMethodMapCooperativeGroupsSYCLcompat(RewriterMap);

clang/lib/DPCT/RulesLang/RulesLang.cpp

@@ -2183,10 +2183,10 @@ void FunctionCallRule::registerMatcher(MatchFinder &MF) {
         "cudaRuntimeGetVersion", "clock64", "__nanosleep",
         "cudaFuncSetSharedMemConfig", "cuFuncSetCacheConfig",
         "cudaPointerGetAttributes", "cuPointerGetAttributes",
-        "cuCtxSetCacheConfig", "cuCtxSetLimit", "cudaCtxResetPersistingL2Cache",
-        "cuCtxResetPersistingL2Cache", "cudaStreamSetAttribute",
-        "cudaStreamGetAttribute", "cudaProfilerStart", "cudaProfilerStop",
-        "__trap", "cuCtxEnablePeerAccess");
+        "cuPointerGetAttribute", "cuCtxSetCacheConfig", "cuCtxSetLimit",
+        "cudaCtxResetPersistingL2Cache", "cuCtxResetPersistingL2Cache",
+        "cudaStreamSetAttribute", "cudaStreamGetAttribute", "cudaProfilerStart",
+        "cudaProfilerStop", "__trap", "cuCtxEnablePeerAccess");
   };
 
   MF.addMatcher(

clang/lib/DPCT/SrcAPI/APINames.inc

@@ -1765,7 +1765,7 @@ ENTRY(cuMemPrefetchAsync, cuMemPrefetchAsync, true, NO_FLAG, P4, "Partial: USM o
 ENTRY(cuMemPrefetchAsync_v2, cuMemPrefetchAsync_v2, false, NO_FLAG, P4, "Partial: USM only, DPCT1007")
 ENTRY(cuMemRangeGetAttribute, cuMemRangeGetAttribute, false, NO_FLAG, P4, "comment")
 ENTRY(cuMemRangeGetAttributes, cuMemRangeGetAttributes, false, NO_FLAG, P4, "comment")
-ENTRY(cuPointerGetAttribute, cuPointerGetAttribute, false, NO_FLAG, P4, "comment")
+ENTRY(cuPointerGetAttribute, cuPointerGetAttribute, true, NO_FLAG, P4, "Partial: none USM for start address API, DPCT1007.")
 ENTRY(cuPointerGetAttributes, cuPointerGetAttributes, true, NO_FLAG, P4, "Partial: Only 5 attributes are supported")
 ENTRY(cuPointerSetAttribute, cuPointerSetAttribute, false, NO_FLAG, P4, "comment")
 

clang/runtime/dpct-rt/include/dpct/detail/memory_detail.hpp

@@ -124,6 +124,11 @@ class mem_mgr {
     return it->second;
   }
 
+  void *get_base_addr(const void *ptr) {
+    allocation alloc = translate_ptr(ptr);
+    return alloc.alloc_ptr;
+  }
+
   /// Check if the pointer represents device pointer or not.
   bool is_device_ptr(const void *ptr) const {
     std::lock_guard<std::mutex> lock(m_mutex);

clang/runtime/dpct-rt/include/dpct/memory.hpp

@@ -124,6 +124,13 @@ static inline bool is_device_ptr(T ptr) {
   }
   return false;
 }
+/// Get the base address of the allocated memory for the given pointer \p ptr.
+template <class T> static inline void *get_base_addr(T ptr) {
+  if constexpr (std::is_pointer<T>::value) {
+    return detail::mem_mgr::instance().get_base_addr(ptr);
+  }
+  return nullptr;
+}
 #endif
 
 /// Get the buffer and the offset of a piece of memory pointed to by \p ptr.

clang/test/dpct/device_prop/cudaGetPointer.cu

@@ -0,0 +1,15 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2
+// RUN: dpct --format-range=none --usm-level=restricted -out-root %T/device_prop %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only -std=c++11
+// RUN: FileCheck --match-full-lines --input-file %T/device_prop/cudaGetPointer.dp.cpp %s
+
+
+#include <cuda.h>
+void test_attribute() {
+  void *base_ptr;
+  void *ptr;
+  // CHECK:/*
+  // CHECK-NEXT:DPCT1067:{{[0-9]+}}: The 'CU_POINTER_ATTRIBUTE_RANGE_START_ADDR' parameter could not be migrated. You may need to update the code manually.
+  // CHECK-NEXT:*/
+  if (cuPointerGetAttribute(base_ptr, CU_POINTER_ATTRIBUTE_RANGE_START_ADDR, (CUdeviceptr)ptr) != CUDA_SUCCESS);
+}

clang/test/dpct/device_prop/cudaGetPointer_none_usm.cu

@@ -0,0 +1,13 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2
+// RUN: dpct --format-range=none --usm-level=none -out-root %T/device_prop %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only -std=c++11
+// RUN: FileCheck --match-full-lines --input-file %T/device_prop/cudaGetPointer_none_usm.dp.cpp %s
+// RUN: %if build_lit %{icpx -c -fsycl %T/device_prop/cudaGetPointer_none_usm.dp.cpp -o %T/device_prop/cudaGetPointer_none_usm.dp.o %}
+
+#include <cuda.h>
+void test_attribute() {
+  void *base_ptr;
+  void *ptr;
+  //CHECK:  if (DPCT_CHECK_ERROR(base_ptr = dpct::get_base_addr((dpct::device_ptr)ptr)) != 0);
+  if (cuPointerGetAttribute(base_ptr, CU_POINTER_ATTRIBUTE_RANGE_START_ADDR, (CUdeviceptr)ptr) != CUDA_SUCCESS);
+}