[SYCLomatic] Add migration support for cuMemGetAllocationPropertiesFromHandle (#2623)

Author: Deepak Raj H R

clang/lib/DPCT/RulesLang/APINamesMemory.inc

@@ -717,7 +717,21 @@ CONDITIONAL_FACTORY_ENTRY(
     UNSUPPORT_FACTORY_ENTRY("cuMemCreate",
                             Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cuMemCreate"),
-                            ARG("--use-experimental-features=virtual_memory")))
+                            ARG("--use-experimental-features=virtual_mem")))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExpVirtualMemory,
+    ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        ASSIGN_FACTORY_ENTRY(
+            "cuMemGetAllocationPropertiesFromHandle",
+            MEMBER_EXPR(DEREF(0), false, LITERAL("location.id")),
+            CALL(MapNames::getDpctNamespace() + "get_device_id",
+                 MEMBER_CALL(ARG(1), true, "get_device"))))),
+    UNSUPPORT_FACTORY_ENTRY("cuMemGetAllocationPropertiesFromHandle",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cuMemGetAllocationPropertiesFromHandle"),
+                            ARG("--use-experimental-features=virtual_mem")))
 
 CONDITIONAL_FACTORY_ENTRY(
     UseExpVirtualMemory,
@@ -734,7 +748,7 @@ CONDITIONAL_FACTORY_ENTRY(
     UNSUPPORT_FACTORY_ENTRY("cuMemAddressReserve",
                             Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cuMemAddressReserve"),
-                            ARG("--use-experimental-features=virtual_memory")))
+                            ARG("--use-experimental-features=virtual_mem")))
 
 CONDITIONAL_FACTORY_ENTRY(
     UseExpVirtualMemory,
@@ -750,7 +764,7 @@ CONDITIONAL_FACTORY_ENTRY(
     UNSUPPORT_FACTORY_ENTRY("cuMemAddressFree",
                             Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cuMemAddressFree"),
-                            ARG("--use-experimental-features=virtual_memory")))
+                            ARG("--use-experimental-features=virtual_mem")))
 
 CONDITIONAL_FACTORY_ENTRY(
     UseExpVirtualMemory,
@@ -770,7 +784,7 @@ CONDITIONAL_FACTORY_ENTRY(
     UNSUPPORT_FACTORY_ENTRY("cuMemGetAllocationGranularity",
                             Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cuMemGetAllocationGranularity"),
-                            ARG("--use-experimental-features=virtual_memory")))
+                            ARG("--use-experimental-features=virtual_mem")))
 
 CONDITIONAL_FACTORY_ENTRY(
     UseExpVirtualMemory,
@@ -780,7 +794,7 @@ CONDITIONAL_FACTORY_ENTRY(
     UNSUPPORT_FACTORY_ENTRY("cuMemRelease",
                             Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cuMemRelease"),
-                            ARG("--use-experimental-features=virtual_memory")))
+                            ARG("--use-experimental-features=virtual_mem")))
 
 CONDITIONAL_FACTORY_ENTRY(
     UseExpVirtualMemory,
@@ -795,7 +809,7 @@ CONDITIONAL_FACTORY_ENTRY(
             ARG(2)))),
     UNSUPPORT_FACTORY_ENTRY("cuMemMap", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cuMemMap"),
-                            ARG("--use-experimental-features=virtual_memory")))
+                            ARG("--use-experimental-features=virtual_mem")))
 
 CONDITIONAL_FACTORY_ENTRY(
     UseExpVirtualMemory,
@@ -810,7 +824,7 @@ CONDITIONAL_FACTORY_ENTRY(
                          "get_current_device().get_context()"))))),
     UNSUPPORT_FACTORY_ENTRY("cuMemUnmap", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cuMemUnmap"),
-                            ARG("--use-experimental-features=virtual_memory")))
+                            ARG("--use-experimental-features=virtual_mem")))
 
 CONDITIONAL_FACTORY_ENTRY(
     UseExpVirtualMemory,
@@ -828,7 +842,7 @@ CONDITIONAL_FACTORY_ENTRY(
     UNSUPPORT_FACTORY_ENTRY("cuMemSetAccess",
                             Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cuMemSetAccess"),
-                            ARG("--use-experimental-features=virtual_memory")))
+                            ARG("--use-experimental-features=virtual_mem")))
 
 CONDITIONAL_FACTORY_ENTRY(
     UsePeerAccess(),

clang/lib/DPCT/RulesLang/RulesLang.cpp

@@ -8055,7 +8055,8 @@ void VirtualMemRule::registerMatcher(ast_matchers::MatchFinder &MF) {
   auto virtualmemoryAPI = [&]() {
     return hasAnyName("cuMemCreate", "cuMemAddressReserve", "cuMemMap",
                       "cuMemUnmap", "cuMemAddressFree", "cuMemRelease",
-                      "cuMemSetAccess", "cuMemGetAllocationGranularity");
+                      "cuMemSetAccess", "cuMemGetAllocationGranularity",
+                      "cuMemGetAllocationPropertiesFromHandle");
   };
   auto virtualmemoryType = [&]() {
     return hasAnyName("CUmemAllocationProp", "CUmemGenericAllocationHandle",

clang/lib/DPCT/SrcAPI/APINames.inc

@@ -1725,7 +1725,7 @@ ENTRY(cuMemCreate, cuMemCreate, true, NO_FLAG, P4, "comment")
 ENTRY(cuMemExportToShareableHandle, cuMemExportToShareableHandle, false, NO_FLAG, P4, "comment")
 ENTRY(cuMemGetAccess, cuMemGetAccess, false, NO_FLAG, P4, "comment")
 ENTRY(cuMemGetAllocationGranularity, cuMemGetAllocationGranularity, true, NO_FLAG, P4, "comment")
-ENTRY(cuMemGetAllocationPropertiesFromHandle, cuMemGetAllocationPropertiesFromHandle, false, NO_FLAG, P4, "comment")
+ENTRY(cuMemGetAllocationPropertiesFromHandle, cuMemGetAllocationPropertiesFromHandle, true, NO_FLAG, P4, "Successful")
 ENTRY(cuMemImportFromShareableHandle, cuMemImportFromShareableHandle, false, NO_FLAG, P4, "comment")
 ENTRY(cuMemMap, cuMemMap, true, NO_FLAG, P4, "comment")
 ENTRY(cuMemMapArrayAsync, cuMemMapArrayAsync, false, NO_FLAG, P7, "comment")

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

@@ -677,14 +677,14 @@ namespace experimental {
 typedef sycl::ext::oneapi::experimental::physical_mem *physical_mem_ptr;
 
 struct mem_location {
-  int id;
-  int type; // Location type. Value 1 means device location, and thus, id is a
-            // device id.
+  int id = 0;
+  // Location type. Value 1 means device location, and thus, id is a device id.
+  int type = 1;
 };
 
 struct mem_prop {
   mem_location location;
-  int type; // Memory type. Value 1 means default device memory.
+  int type = 1; // Memory type. Value 1 means default device memory.
 };
 
 struct mem_access_desc {

clang/test/dpct/virtual_memory.cu

@@ -38,14 +38,17 @@ int main() {
     cuMemCreate(&allocHandle, POOL_SIZE, &prop, 0);
     cuMemMap(reserved_addr, POOL_SIZE, 0, allocHandle, 0);
 
+    // CHECK: prop.location.id = dpct::get_device_id(allocHandle->get_device());
+    cuMemGetAllocationPropertiesFromHandle(&prop, allocHandle);
+
 // CHECK:  dpct::experimental::mem_access_desc accessDesc = {};
 // CHECK:  accessDesc.location.type = 1;
-// CHECK:  accessDesc.location.id = device;
+// CHECK:  accessDesc.location.id = prop.location.id;
 // CHECK:  accessDesc.flags = sycl::ext::oneapi::experimental::address_access_mode::read_write;
 // CHECK:  sycl::ext::oneapi::experimental::set_access_mode(reserved_addr, POOL_SIZE, accessDesc.flags, dpct::get_device(accessDesc.location.id).get_context());
     CUmemAccessDesc accessDesc = {};
     accessDesc.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
-    accessDesc.location.id = device;
+    accessDesc.location.id = prop.location.id;
     accessDesc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
     cuMemSetAccess(reserved_addr, POOL_SIZE, &accessDesc, 1);
     int* host_data = new int[SIZE];