[SYCLomatic][Driver] Enabled migration of cuPointerGetAttributes

clang/examples/DPCT/Driver/cuPointerGetAttributes.cu

@@ -0,0 +1,8 @@
+void test(unsigned int numAttr, CUpointer_attribute *attr,
+          void **data, CUdeviceptr ptr) {
+  // Start
+  cuPointerGetAttributes(numAttr /*unsigned int*/,
+                        attr /*CUpointer_attribute **/, data /*void ***/,
+                        ptr /*CUdeviceptr*/);
+  // End
+}

clang/lib/DPCT/APINames.inc

@@ -1764,7 +1764,7 @@ ENTRY(cuMemPrefetchAsync_v2, cuMemPrefetchAsync_v2, false, NO_FLAG, P4, "Partial
 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(cuPointerGetAttributes, cuPointerGetAttributes, false, NO_FLAG, P4, "comment")
+ENTRY(cuPointerGetAttributes, cuPointerGetAttributes, true, NO_FLAG, P4, "Partial: Only 5 attributes are supported")
 ENTRY(cuPointerSetAttribute, cuPointerSetAttribute, false, NO_FLAG, P4, "comment")
 
 // Stream Management

clang/lib/DPCT/APINamesMemory.inc

@@ -398,6 +398,13 @@ ASSIGNABLE_FACTORY(
         MEMBER_CALL_FACTORY_ENTRY(
             "cudaPointerGetAttributes", DEREF(0), false, "init", ARG(1))))
 
+ASSIGNABLE_FACTORY(
+    FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        CALL_FACTORY_ENTRY(
+            "cuPointerGetAttributes", CALL(MapNames::getDpctNamespace() + "pointer_attributes::get",
+            ARG(0), ARG(1), ARG(2), ARG(3)))))
+
 ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
     HelperFeatureEnum::device_ext,
     CONDITIONAL_FACTORY_ENTRY(

clang/lib/DPCT/ASTTraversal.cpp

@@ -1728,9 +1728,11 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
               "__nv_bfloat16", "cooperative_groups::__v1::thread_group",
               "cooperative_groups::__v1::thread_block", "libraryPropertyType_t",
               "libraryPropertyType", "cudaDataType_t", "cudaDataType",
-              "cublasComputeType_t", "cublasAtomicsMode_t", "CUmem_advise_enum",
-              "CUmem_advise", "thrust::tuple_element", "thrust::tuple_size",
-              "cublasMath_t", "cudaPointerAttributes", "thrust::zip_iterator",
+              "cublasComputeType_t", "cublasAtomicsMode_t", "cublasMath_t",
+              "CUmem_advise_enum", "CUmem_advise", "CUmemorytype",
+              "CUmemorytype_enum", "thrust::tuple_element",
+              "thrust::tuple_size", "thrust::zip_iterator",
+              "cudaPointerAttributes", "CUpointer_attribute",
               "cusolverEigRange_t", "cudaUUID_t", "cusolverDnFunction_t",
               "cusolverAlgMode_t", "cusparseIndexType_t", "cusparseFormat_t",
               "cusparseDnMatDescr_t", "cusparseOrder_t", "cusparseDnVecDescr_t",
@@ -3232,6 +3234,7 @@ void EnumConstantRule::registerMatcher(MatchFinder &MF) {
                   "libraryPropertyType_t", "cudaDataType_t",
                   "CUmem_advise_enum", "cufftType_t",
                   "cufftType", "cudaMemoryType", "CUctx_flags_enum",
+                  "CUpointer_attribute_enum", "CUmemorytype_enum",
                   "cudaGraphicsMapFlags", "cudaGraphicsRegisterFlags"))),
               matchesName("CUDNN_.*"), matchesName("CUSOLVER_.*")))))
           .bind("EnumConstant"),
@@ -6141,10 +6144,11 @@ void FunctionCallRule::registerMatcher(MatchFinder &MF) {
         "cuDeviceCanAccessPeer", "cudaFuncSetAttribute",
         "cudaRuntimeGetVersion", "clock64", "__nanosleep",
         "cudaFuncSetSharedMemConfig", "cuFuncSetCacheConfig",
-        "cudaPointerGetAttributes", "cuCtxSetCacheConfig", "cuCtxSetLimit",
-        "cudaCtxResetPersistingL2Cache", "cuCtxResetPersistingL2Cache",
-        "cudaStreamSetAttribute", "cudaStreamGetAttribute", "cudaProfilerStart",
-        "cudaProfilerStop", "__trap", "cuCtxEnablePeerAccess");
+        "cudaPointerGetAttributes", "cuPointerGetAttributes",
+        "cuCtxSetCacheConfig", "cuCtxSetLimit", "cudaCtxResetPersistingL2Cache",
+        "cuCtxResetPersistingL2Cache", "cudaStreamSetAttribute",
+        "cudaStreamGetAttribute", "cudaProfilerStart", "cudaProfilerStop",
+        "__trap", "cuCtxEnablePeerAccess");
   };
 
   MF.addMatcher(

clang/lib/DPCT/MapNames.cpp

@@ -189,6 +189,8 @@ void MapNames::setExplicitNamespaceMap(
       {"CUfunction",
        std::make_shared<TypeNameRule>(getDpctNamespace() + "kernel_function",
                                       HelperFeatureEnum::device_ext)},
+      {"CUpointer_attribute",
+       std::make_shared<TypeNameRule>(getDpctNamespace() + "pointer_attributes::type")},
       {"cudaPointerAttributes",
        std::make_shared<TypeNameRule>(getDpctNamespace() + "pointer_attributes",
                                       HelperFeatureEnum::device_ext)},
@@ -451,6 +453,10 @@ void MapNames::setExplicitNamespaceMap(
                ? getClNamespace() + "ext::oneapi::experimental::queue_state"
                : "cudaStreamCaptureStatus")},
       {"CUmem_advise", std::make_shared<TypeNameRule>("int")},
+      {"CUmemorytype",
+       std::make_shared<TypeNameRule>(getClNamespace() + "usm::alloc")},
+      {"CUmemorytype_enum",
+       std::make_shared<TypeNameRule>(getClNamespace() + "usm::alloc")},
       {"cudaPos", std::make_shared<TypeNameRule>(getClNamespace() + "id<3>")},
       {"cudaExtent",
        std::make_shared<TypeNameRule>(getClNamespace() + "range<3>")},
@@ -1258,6 +1264,61 @@ void MapNames::setExplicitNamespaceMap(
        std::make_shared<EnumNameRule>("get_device_info().get_local_mem_size",
                                       HelperFeatureEnum::device_ext)},
 
+      // enum CUpointer_attribute
+      {"CU_POINTER_ATTRIBUTE_CONTEXT",
+       std::make_shared<EnumNameRule>(getDpctNamespace() +
+                                      "pointer_attributes::type::unsupported")},
+      {"CU_POINTER_ATTRIBUTE_MEMORY_TYPE",
+       std::make_shared<EnumNameRule>(getDpctNamespace() +
+                                      "pointer_attributes::type::memory_type")},
+      {"CU_POINTER_ATTRIBUTE_DEVICE_POINTER",
+       std::make_shared<EnumNameRule>(getDpctNamespace() +
+                                      "pointer_attributes::type::device_pointer")},
+      {"CU_POINTER_ATTRIBUTE_HOST_POINTER",
+       std::make_shared<EnumNameRule>(getDpctNamespace() +
+                                      "pointer_attributes::type::host_pointer")},
+      {"CU_POINTER_ATTRIBUTE_P2P_TOKENS",
+       std::make_shared<EnumNameRule>(getDpctNamespace() +
+                                      "pointer_attributes::type::unsupported")},
+      {"CU_POINTER_ATTRIBUTE_SYNC_MEMOPS",
+       std::make_shared<EnumNameRule>(getDpctNamespace() +
+                                      "pointer_attributes::type::unsupported")},
+      {"CU_POINTER_ATTRIBUTE_BUFFER_ID",
+       std::make_shared<EnumNameRule>(getDpctNamespace() +
+                                      "pointer_attributes::type::unsupported")},
+      {"CU_POINTER_ATTRIBUTE_IS_MANAGED",
+       std::make_shared<EnumNameRule>(getDpctNamespace() +
+                                      "pointer_attributes::type::is_managed")},
+      {"CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL",
+       std::make_shared<EnumNameRule>(getDpctNamespace() +
+                                      "pointer_attributes::type::device_id")},
+      {"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")},
+      {"CU_POINTER_ATTRIBUTE_MAPPED",
+       std::make_shared<EnumNameRule>(getDpctNamespace() +
+                                      "pointer_attributes::type::unsupported")},
+      {"CU_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES",
+       std::make_shared<EnumNameRule>(getDpctNamespace() +
+                                      "pointer_attributes::type::unsupported")},
+
+      // enum CUmemorytype Type
+      {"CU_MEMORYTYPE_HOST",
+       std::make_shared<EnumNameRule>(getClNamespace() + "usm::alloc::host",
+                                      HelperFeatureEnum::device_ext)},
+      {"CU_MEMORYTYPE_DEVICE",
+       std::make_shared<EnumNameRule>(getClNamespace() + "usm::alloc::device",
+                                      HelperFeatureEnum::device_ext)},
+      {"CU_MEMORYTYPE_UNIFIED",
+       std::make_shared<EnumNameRule>(getClNamespace() + "usm::alloc::shared",
+                                      HelperFeatureEnum::device_ext)},
+
       // enum CUlimit
       {"CU_LIMIT_PRINTF_FIFO_SIZE", std::make_shared<EnumNameRule>("INT_MAX")},
 

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

@@ -1663,6 +1663,15 @@ using usm_device_allocator = detail::deprecated::usm_allocator<T, sycl::usm::all
 
 class pointer_attributes {
 public:
+  enum class type {
+    memory_type,
+    device_pointer,
+    host_pointer,
+    is_managed,
+    device_id,
+    unsupported
+  };
+
   void init(const void *ptr,
               sycl::queue &q = dpct::get_default_queue()) {
 #ifdef DPCT_USM_LEVEL_NONE
@@ -1684,6 +1693,41 @@ class pointer_attributes {
 #endif
   }
 
+  // Query pointer propreties listed in attributes and store the results in data array
+  static void get(unsigned int numAttributes, type *attributes,
+                  void **data, device_ptr ptr) {
+    pointer_attributes sycl_attributes;
+
+    sycl_attributes.init(ptr);
+
+    for (int i = 0; i < numAttributes; i++) {
+      switch (attributes[i]) {
+      case type::memory_type:
+        *static_cast<int *>(data[i]) =
+            static_cast<int>(sycl_attributes.get_memory_type());
+        break;
+      case type::device_pointer:
+        *(reinterpret_cast<void **>(data[i])) =
+            const_cast<void *>(sycl_attributes.get_device_pointer());
+        break;
+      case type::host_pointer:
+        *(reinterpret_cast<void **>(data[i])) =
+            const_cast<void *>(sycl_attributes.get_host_pointer());
+        break;
+      case type::is_managed:
+        *static_cast<unsigned int *>(data[i]) =
+            sycl_attributes.is_memory_shared();
+        break;
+      case type::device_id:
+        *static_cast<unsigned int *>(data[i]) = sycl_attributes.get_device_id();
+        break;
+      default:
+        data[i] = nullptr;
+        break;
+      }
+    }
+  }
+
   sycl::usm::alloc get_memory_type() {
     return memory_type;
   }
@@ -1710,5 +1754,6 @@ class pointer_attributes {
   const void *host_pointer = nullptr;
   unsigned int device_id = -1;
 };
+
 } // namespace dpct
 #endif // __DPCT_MEMORY_HPP__

clang/test/dpct/cuPointerGetAttributes.cu

@@ -0,0 +1,132 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2
+// RUN: dpct --format-range=none -out-root %T/Out/cuPointerGetAttributes %s --cuda-include-path="%cuda-path/include"
+// RUN: FileCheck %s --match-full-lines --input-file %T/Out/cuPointerGetAttributes/cuPointerGetAttributes.dp.cpp
+// RUN: %if build_lit %{icpx -c -fsycl %T/Out/cuPointerGetAttributes/cuPointerGetAttributes.dp.cpp -o %T/Out/cuPointerGetAttributes/cuPointerGetAttributes.dp.o %}
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <stdlib.h>
+#include <iostream>
+
+int main() {
+  int N = 2048;
+  size_t size = N * sizeof(float);
+
+  float *h_A = (float *)malloc(size);
+
+  float *d_A;
+  cudaMalloc((void **)&d_A, size);
+
+  cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
+
+  unsigned int numAttributes = 5;
+
+  // CHECK: dpct::pointer_attributes::type attributes[] = {
+  // CHECK:   dpct::pointer_attributes::type::memory_type,
+  // CHECK:   dpct::pointer_attributes::type::device_pointer,
+  // CHECK:   dpct::pointer_attributes::type::host_pointer,
+  // CHECK:   dpct::pointer_attributes::type::is_managed,
+  // CHECK:   dpct::pointer_attributes::type::device_id
+  CUpointer_attribute attributes[] = {
+    CU_POINTER_ATTRIBUTE_MEMORY_TYPE,
+    CU_POINTER_ATTRIBUTE_DEVICE_POINTER,
+    CU_POINTER_ATTRIBUTE_HOST_POINTER,
+    CU_POINTER_ATTRIBUTE_IS_MANAGED,
+    CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL
+  };
+
+  // CHECK: dpct::pointer_attributes::type attributes_unsupported[] = {
+  // CHECK:   dpct::pointer_attributes::type::unsupported,
+  // CHECK:   dpct::pointer_attributes::type::unsupported
+  CUpointer_attribute attributes_unsupported[] = {
+    CU_POINTER_ATTRIBUTE_CONTEXT,
+    CU_POINTER_ATTRIBUTE_IS_LEGACY_CUDA_IPC_CAPABLE
+  };
+
+  // CHECK: sycl::usm::alloc memType;
+  CUmemorytype memType;
+  void* hostPtr;
+  unsigned int isManaged;
+  int deviceID;
+  // CHECK: dpct::device_ptr devPtr;
+  CUdeviceptr devPtr;
+  CUcontext cuCtx;
+  bool is_legacy_cuda_ipc_capable;
+
+  void* attributeValues[] = {
+    &memType,
+    &devPtr,
+    &hostPtr,
+    &isManaged,
+    &deviceID
+  };
+
+  void* attributeValues_unsupported[] = {
+    &cuCtx,
+    &is_legacy_cuda_ipc_capable
+  };
+
+  // CHECK: dpct::pointer_attributes::get(numAttributes, attributes, attributeValues, (dpct::device_ptr) h_A);
+  cuPointerGetAttributes(
+    numAttributes,
+    attributes,
+    attributeValues,
+    (CUdeviceptr) h_A
+  );
+
+  std::cout << "====== Host Attributes =======" << std::endl;
+  std::cout << deviceID << std::endl;
+  std::cout << static_cast<int>(memType) << std::endl;
+  std::cout << hostPtr << std::endl;
+  std::cout << devPtr << std::endl;
+  std::cout << isManaged << std::endl;
+
+  void * malloc_host;
+  cudaMallocHost((void **)&malloc_host, size);
+  // CHECK: dpct::pointer_attributes::get(numAttributes, attributes, attributeValues, (dpct::device_ptr) malloc_host);
+  cuPointerGetAttributes(
+    numAttributes,
+    attributes,
+    attributeValues,
+    (CUdeviceptr) malloc_host
+  );
+  std::cout << "====== Malloc Host Attributes =======" << std::endl;
+  std::cout << "malloc host " << malloc_host << std::endl;
+  std::cout << deviceID << std::endl;
+  std::cout << static_cast<int>(memType) << std::endl;
+  std::cout << hostPtr << std::endl;
+  std::cout << devPtr << std::endl;
+  std::cout << isManaged << std::endl;
+
+  // CHECK: dpct::pointer_attributes::get(numAttributes, attributes, attributeValues, (dpct::device_ptr) d_A);
+  cuPointerGetAttributes(
+    numAttributes,
+    attributes,
+    attributeValues,
+    (CUdeviceptr) d_A
+  );
+  std::cout << "====== Device Attributes =======" << std::endl;
+  std::cout << *static_cast<int *>(attributeValues[0]) << std::endl;
+  std::cout << attributeValues[1] << std::endl;
+  std::cout << attributeValues[2] << std::endl;
+  std::cout << *static_cast<unsigned int *>(attributeValues[3]) << std::endl;
+  std::cout << *static_cast<int *>(attributeValues[4]) << std::endl;
+
+  // CHECK: if (memType == sycl::usm::alloc::host) {
+  if (memType == CU_MEMORYTYPE_HOST) {
+    return 0;
+  // CHECK: } else if (memType == sycl::usm::alloc::device) {
+  } else if (memType == CU_MEMORYTYPE_DEVICE) {
+    return 1;
+  } else if (isManaged) {
+    return 2;
+  }
+
+  // CHECK: dpct::pointer_attributes::get(2, attributes_unsupported, attributeValues_unsupported, (dpct::device_ptr) d_A);
+  cuPointerGetAttributes(
+    2,
+    attributes_unsupported,
+    attributeValues_unsupported,
+    (CUdeviceptr) d_A
+  );
+}

clang/test/dpct/query_api_mapping/Driver/test.cu

@@ -265,6 +265,14 @@
 // CUMEMPREFETCHASYNC-NEXT:   dpct::queue_ptr cs;
 // CUMEMPREFETCHASYNC-NEXT:   cs->prefetch(pd, s);
 
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuPointerGetAttributes | FileCheck %s -check-prefix=CUPOINTERGETATTRIBUTES
+// CUPOINTERGETATTRIBUTES: CUDA API:
+// CUPOINTERGETATTRIBUTES-NEXT:   cuPointerGetAttributes(numAttr /*unsigned int*/,
+// CUPOINTERGETATTRIBUTES-NEXT:                            attr /*CUpointer_attribute **/, data /*void ***/,
+// CUPOINTERGETATTRIBUTES-NEXT:                            ptr /*CUdeviceptr*/);
+// CUPOINTERGETATTRIBUTES-NEXT: Is migrated to:
+// CUPOINTERGETATTRIBUTES-NEXT:   dpct::pointer_attributes::get(numAttr, attr, data, ptr);
+
 /// Stream Management
 
 // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuStreamAddCallback | FileCheck %s -check-prefix=CUSTREAMADDCALLBACK

clang/test/dpct/query_api_mapping/test_all.cu

@@ -695,6 +695,7 @@
 // CHECK-NEXT: cuModuleLoadDataEx
 // CHECK-NEXT: cuModuleUnload
 // CHECK-NEXT: cuOccupancyMaxActiveBlocksPerMultiprocessor
+// CHECK-NEXT: cuPointerGetAttributes
 // CHECK-NEXT: cuStreamAddCallback
 // CHECK-NEXT: cuStreamAttachMemAsync
 // CHECK-NEXT: cuStreamCreate

docs/dev_guide/api-mapping-status/Runtime_and_Driver_API_migration_status.csv

@@ -1577,7 +1577,7 @@ cuMemPrefetchAsync,YES, Partial
 cuMemRangeGetAttribute,NO,
 cuMemRangeGetAttributes,NO,
 cuPointerGetAttribute,NO,
-cuPointerGetAttributes,NO,
+cuPointerGetAttributes,YES, Partial
 cuPointerSetAttribute,NO,
 cuGraphicsVDPAURegisterOutputSurface,NO,
 cuGraphicsVDPAURegisterVideoSurface,NO,