[SYCLomatic] Enable migration of 8 CUDA driver virtual memory management API

clang/include/clang/DPCT/DPCTOptions.inc

@@ -787,6 +787,11 @@ DPCT_ENUM_OPTION(
             "be accessed within a kernel using syntax similar to C++ global "
             "variables.\n",
             false),
+        DPCT_OPTION_ENUM_VALUE(
+            "virtual_memory", int(ExperimentalFeatures::Exp_VirtualMemory),
+            "Experimental extension that allows map an address range onto "
+            "multiple allocations of physical memory.",
+            false),
         DPCT_OPTION_ENUM_VALUE(
             "non-stdandard-sycl-builtins",
             int(ExperimentalFeatures::Exp_NonStandardSYCLBuiltins),

clang/lib/DPCT/APINames.inc

@@ -32,6 +32,8 @@
 *
 ****************************************************************************/
 
+// clang-format off
+
 // API Names refer to "v12.2"
 // CUDA runtime API
 // Device management functions of runtime API
@@ -1717,20 +1719,20 @@ ENTRY(cuMipmappedArrayGetMemoryRequirements, cuMipmappedArrayGetMemoryRequiremen
 ENTRY(cuMipmappedArrayGetSparseProperties, cuMipmappedArrayGetSparseProperties, false, NO_FLAG, P7, "comment")
 
 // Virtual Memory Management
-ENTRY(cuMemAddressFree, cuMemAddressFree, false, NO_FLAG, P4, "comment")
-ENTRY(cuMemAddressReserve, cuMemAddressReserve, false, NO_FLAG, P4, "comment")
-ENTRY(cuMemCreate, cuMemCreate, false, NO_FLAG, P4, "comment")
+ENTRY(cuMemAddressFree, cuMemAddressFree, true, NO_FLAG, P4, "comment")
+ENTRY(cuMemAddressReserve, cuMemAddressReserve, true, NO_FLAG, P4, "comment")
+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, false, NO_FLAG, P4, "comment")
+ENTRY(cuMemGetAllocationGranularity, cuMemGetAllocationGranularity, true, NO_FLAG, P4, "comment")
 ENTRY(cuMemGetAllocationPropertiesFromHandle, cuMemGetAllocationPropertiesFromHandle, false, NO_FLAG, P4, "comment")
 ENTRY(cuMemImportFromShareableHandle, cuMemImportFromShareableHandle, false, NO_FLAG, P4, "comment")
-ENTRY(cuMemMap, cuMemMap, false, NO_FLAG, P4, "comment")
+ENTRY(cuMemMap, cuMemMap, true, NO_FLAG, P4, "comment")
 ENTRY(cuMemMapArrayAsync, cuMemMapArrayAsync, false, NO_FLAG, P7, "comment")
-ENTRY(cuMemRelease, cuMemRelease, false, NO_FLAG, P4, "comment")
+ENTRY(cuMemRelease, cuMemRelease, true, NO_FLAG, P4, "comment")
 ENTRY(cuMemRetainAllocationHandle, cuMemRetainAllocationHandle, false, NO_FLAG, P7, "comment")
-ENTRY(cuMemSetAccess, cuMemSetAccess, false, NO_FLAG, P4, "comment")
-ENTRY(cuMemUnmap, cuMemUnmap, false, NO_FLAG, P4, "comment")
+ENTRY(cuMemSetAccess, cuMemSetAccess, true, NO_FLAG, P4, "comment")
+ENTRY(cuMemUnmap, cuMemUnmap, true, NO_FLAG, P4, "comment")
 
 // Stream Ordered Memory Allocator
 ENTRY(cuMemAllocAsync, cuMemAllocAsync, false, NO_FLAG, P7, "comment")
@@ -2307,3 +2309,5 @@ ENTRY(__assert_fail, __assert_fail, true, NO_FLAG, P4, "Successful")
 ENTRY(__assertfail, __assertfail, true, NO_FLAG, P4, "Successful")
 
 ENTRY(cuGetExportTable, cuGetExportTable, true, NO_FLAG, P7, "Partial")
+
+// clang-format on

clang/lib/DPCT/APINamesMemory.inc

@@ -698,6 +698,107 @@ ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
                             ARG(0), ARG(1), ARG(2), ARG(3), ARG(4),
                             DEREF(makeCallArgCreatorWithCall(5))))))
 
+CONDITIONAL_FACTORY_ENTRY(
+    UseExpVirtualMemory,
+    ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        CALL_FACTORY_ENTRY("cuMemCreate",
+                           CALL(MapNames::getDpctNamespace() +
+                                    "experimental::mem_create",
+                                ARG(0), ARG(1), ARG(2), ARG(3))))),
+    UNSUPPORT_FACTORY_ENTRY("cuMemCreate",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cuMemCreate"),
+                            ARG("--use-experimental-features=virtual_memory")))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExpVirtualMemory,
+    ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        CALL_FACTORY_ENTRY("cuMemAddressReserve",
+                           CALL(MapNames::getDpctNamespace() +
+                                    "experimental::mem_address_reserve",
+                                ARG(0), ARG(1), ARG(2), ARG(3), ARG(4))))),
+    UNSUPPORT_FACTORY_ENTRY("cuMemAddressReserve",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cuMemAddressReserve"),
+                            ARG("--use-experimental-features=virtual_memory")))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExpVirtualMemory,
+    ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        CALL_FACTORY_ENTRY("cuMemAddressFree",
+                           CALL(MapNames::getDpctNamespace() +
+                                    "experimental::mem_address_free",
+                                ARG(0), ARG(1))))),
+    UNSUPPORT_FACTORY_ENTRY("cuMemAddressFree",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cuMemAddressFree"),
+                            ARG("--use-experimental-features=virtual_memory")))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExpVirtualMemory,
+    ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        CALL_FACTORY_ENTRY(
+            "cuMemGetAllocationGranularity",
+            CALL(MapNames::getDpctNamespace() +
+                     "experimental::mem_get_allocation_granularity",
+                 ARG(0), ARG(1), ARG(2))))),
+    UNSUPPORT_FACTORY_ENTRY("cuMemGetAllocationGranularity",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cuMemGetAllocationGranularity"),
+                            ARG("--use-experimental-features=virtual_memory")))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExpVirtualMemory,
+    ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        CALL_FACTORY_ENTRY("cuMemRelease", CALL(MapNames::getDpctNamespace() +
+                                                    "experimental::mem_release",
+                                                ARG(0))))),
+    UNSUPPORT_FACTORY_ENTRY("cuMemRelease",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cuMemRelease"),
+                            ARG("--use-experimental-features=virtual_memory")))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExpVirtualMemory,
+    ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        CALL_FACTORY_ENTRY("cuMemMap",
+                           CALL(MapNames::getDpctNamespace() +
+                                    "experimental::mem_map",
+                                ARG(0), ARG(1), ARG(2), ARG(3), ARG(4))))),
+    UNSUPPORT_FACTORY_ENTRY("cuMemMap", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cuMemMap"),
+                            ARG("--use-experimental-features=virtual_memory")))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExpVirtualMemory,
+    ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        CALL_FACTORY_ENTRY("cuMemUnmap", CALL(MapNames::getDpctNamespace() +
+                                                  "experimental::mem_unmap",
+                                              ARG(0), ARG(1))))),
+    UNSUPPORT_FACTORY_ENTRY("cuMemUnmap", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cuMemUnmap"),
+                            ARG("--use-experimental-features=virtual_memory")))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExpVirtualMemory,
+    ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        CALL_FACTORY_ENTRY("cuMemSetAccess",
+                           CALL(MapNames::getDpctNamespace() +
+                                    "experimental::mem_set_access",
+                                ARG(0), ARG(1), ARG(2), ARG(3))))),
+    UNSUPPORT_FACTORY_ENTRY("cuMemSetAccess",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cuMemSetAccess"),
+                            ARG("--use-experimental-features=virtual_memory")))
+
 CONDITIONAL_FACTORY_ENTRY(
     UsePeerAccess(),
     ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(

clang/lib/DPCT/ASTTraversal.cpp

@@ -10851,6 +10851,30 @@ void MemoryMigrationRule::registerMatcher(MatchFinder &MF) {
                            hasParent(callExpr(parentStmt()).bind("callExpr")))
           .bind("unresolvedCall"),
       this);
+  auto virtualmemoryAPI = [&]() {
+    return hasAnyName("cuMemCreate", "cuMemAddressReserve", "cuMemMap",
+                      "cuMemUnmap", "cuMemAddressFree", "cuMemRelease",
+                      "cuMemSetAccess", "cuMemGetAllocationGranularity");
+  };
+  auto virtualmemoryType = [&]() {
+    return hasAnyName("CUmemAllocationProp", "CUmemGenericAllocationHandle",
+                      "CUmemAccessDesc");
+  };
+  auto virtualmemoryEnum = [&]() {
+    return hasAnyName("CU_MEM_ALLOCATION_TYPE_PINNED",
+                      "CU_MEM_LOCATION_TYPE_DEVICE",
+                      "CU_MEM_ACCESS_FLAGS_PROT_READWRITE",
+                      "CU_MEM_ALLOC_GRANULARITY_RECOMMENDED");
+  };
+  MF.addMatcher(
+      callExpr(callee(functionDecl(virtualmemoryAPI()))).bind("vmCall"), this);
+  MF.addMatcher(
+      typeLoc(loc(qualType(hasDeclaration(namedDecl(virtualmemoryType())))))
+          .bind("vmType"),
+      this);
+  MF.addMatcher(
+      declRefExpr(to(enumConstantDecl(virtualmemoryEnum()))).bind("vmEnum"),
+      this);
 }
 
 void MemoryMigrationRule::runRule(const MatchFinder::MatchResult &Result) {
@@ -10953,6 +10977,58 @@ void MemoryMigrationRule::runRule(const MatchFinder::MatchResult &Result) {
       getAssistNodeAsType<CallExpr>(Result, "callExpr"),
       /* IsAssigned */ false,
       getAssistNodeAsType<UnresolvedLookupExpr>(Result, "unresolvedCall"));
+
+  auto &SM = DpctGlobalInfo::getSourceManager();
+  if (const CallExpr *CE = getNodeAsType<CallExpr>(Result, "vmCall")) {
+    ExprAnalysis EA(CE);
+    emplaceTransformation(EA.getReplacement());
+    EA.applyAllSubExprRepl();
+  }
+  if (auto TL = getNodeAsType<TypeLoc>(Result, "vmType")) {
+    auto TypeStr =
+        DpctGlobalInfo::getTypeName(TL->getType().getUnqualifiedType());
+    if (!DpctGlobalInfo::useExpVirtualMemory()) {
+      report(TL->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
+             TypeStr, "--use-experimental-features=virtual_memory");
+      return;
+    }
+    if (!DpctGlobalInfo::isInAnalysisScope(
+            SM.getSpellingLoc(TL->getBeginLoc()))) {
+      return;
+    }
+    auto Range = getDefinitionRange(TL->getBeginLoc(), TL->getEndLoc());
+    auto BeginLoc = Range.getBegin();
+    auto EndLoc = Range.getEnd();
+
+    if (SM.isWrittenInScratchSpace(SM.getSpellingLoc(TL->getBeginLoc()))) {
+      BeginLoc = SM.getExpansionRange(TL->getBeginLoc()).getBegin();
+      EndLoc = SM.getExpansionRange(TL->getBeginLoc()).getEnd();
+    }
+    std::string Str =
+        MapNames::findReplacedName(MapNames::TypeNamesMap, TypeStr);
+    if (!Str.empty()) {
+      auto Len = Lexer::MeasureTokenLength(
+          EndLoc, SM, DpctGlobalInfo::getContext().getLangOpts());
+      Len += SM.getDecomposedLoc(EndLoc).second -
+             SM.getDecomposedLoc(BeginLoc).second;
+      emplaceTransformation(new ReplaceText(BeginLoc, Len, std::move(Str)));
+      return;
+    }
+  }
+  if (auto *E = getNodeAsType<DeclRefExpr>(Result, "vmEnum")) {
+    std::string EnumName = E->getNameInfo().getName().getAsString();
+    if (!DpctGlobalInfo::useExpVirtualMemory()) {
+      report(E->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
+             EnumName, "--use-experimental-features=virtual_memory");
+      return;
+    }
+    auto Search = EnumConstantRule::EnumNamesMap.find(EnumName);
+    if (Search == EnumConstantRule::EnumNamesMap.end()) {
+      report(E->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, EnumName);
+      return;
+    }
+    emplaceTransformation(new ReplaceStmt(E, Search->second->NewName));
+  }
 }
 
 void MemoryMigrationRule::getSymbolAddressMigration(

clang/lib/DPCT/AnalysisInfo.h

@@ -1290,6 +1290,9 @@ class DpctGlobalInfo {
   static bool useExpDeviceGlobal() {
     return getUsingExperimental<ExperimentalFeatures::Exp_DeviceGlobal>();
   }
+  static bool useExpVirtualMemory() {
+    return getUsingExperimental<ExperimentalFeatures::Exp_VirtualMemory>();
+  }
   static bool useExpNonStandardSYCLBuiltins() {
     return getUsingExperimental<
         ExperimentalFeatures::Exp_NonStandardSYCLBuiltins>();

clang/lib/DPCT/CallExprRewriterCommon.h

@@ -1716,6 +1716,10 @@ inline auto UseSYCLCompat = [](const CallExpr *C) -> bool {
   return DpctGlobalInfo::useSYCLCompat();
 };
 
+inline auto UseExpVirtualMemory = [](const CallExpr *C) -> bool {
+  return DpctGlobalInfo::useExpVirtualMemory();
+};
+
 class CheckDerefedTypeBeforeCast {
   unsigned Idx;
   std::string TypeName;

clang/lib/DPCT/MapNames.cpp

@@ -449,14 +449,16 @@ void MapNames::setExplicitNamespaceMap(
                                           getLibraryHelperNamespace() +
                                           "sparse::optimize_info>",
                                       HelperFeatureEnum::device_ext)},
-      {"thrust::device_ptr",
-       std::make_shared<TypeNameRule>(getLibraryHelperNamespace() + "device_pointer",
-                                      HelperFeatureEnum::device_ext)},
+      {"thrust::device_ptr", std::make_shared<TypeNameRule>(
+                                 getLibraryHelperNamespace() + "device_pointer",
+                                 HelperFeatureEnum::device_ext)},
       {"thrust::device_reference",
-       std::make_shared<TypeNameRule>(getLibraryHelperNamespace() + "device_reference",
+       std::make_shared<TypeNameRule>(getLibraryHelperNamespace() +
+                                          "device_reference",
                                       HelperFeatureEnum::device_ext)},
       {"thrust::device_vector",
-       std::make_shared<TypeNameRule>(getLibraryHelperNamespace() + "device_vector",
+       std::make_shared<TypeNameRule>(getLibraryHelperNamespace() +
+                                          "device_vector",
                                       HelperFeatureEnum::device_ext)},
       {"thrust::device_malloc_allocator",
        std::make_shared<TypeNameRule>(getDpctNamespace() +
@@ -831,6 +833,15 @@ void MapNames::setExplicitNamespaceMap(
        std::make_shared<TypeNameRule>(
            getLibraryHelperNamespace() +
            "blas_gemm::experimental::transform_desc_ptr")},
+      {"CUmemAllocationProp",
+       std::make_shared<TypeNameRule>(getDpctNamespace() +
+                                      "experimental::mem_prop")},
+      {"CUmemGenericAllocationHandle",
+       std::make_shared<TypeNameRule>(getDpctNamespace() +
+                                      "experimental::mem_handle")},
+      {"CUmemAccessDesc",
+       std::make_shared<TypeNameRule>(getDpctNamespace() +
+                                      "experimental::mem_access_desc")},
       {"cudaGraphicsMapFlags", std::make_shared<TypeNameRule>("int")},
       {"cudaGraphicsRegisterFlags", std::make_shared<TypeNameRule>("int")},
       // ...
@@ -1438,6 +1449,54 @@ void MapNames::setExplicitNamespaceMap(
        std::make_shared<EnumNameRule>("0")},
       {"CU_MEM_ADVISE_SET_ACCESSED_BY", std::make_shared<EnumNameRule>("0")},
       {"CU_MEM_ADVISE_UNSET_ACCESSED_BY", std::make_shared<EnumNameRule>("0")},
+      {"CU_MEM_ALLOCATION_TYPE_PINNED",
+       std::make_shared<EnumNameRule>(
+           getDpctNamespace() +
+           "experimental::mem_allocation_type::MEM_ALLOCATION_TYPE_DEFAULT")},
+      {"CU_MEM_ALLOCATION_TYPE_INVALID",
+       std::make_shared<EnumNameRule>(
+           getDpctNamespace() +
+           "experimental::mem_allocation_type::MEM_ALLOCATION_TYPE_INVALID")},
+      {"CU_MEM_ALLOCATION_TYPE_MAX",
+       std::make_shared<EnumNameRule>(
+           getDpctNamespace() +
+           "experimental::mem_allocation_type::MEM_ALLOCATION_TYPE_MAX")},
+      {"CU_MEM_LOCATION_TYPE_DEVICE",
+       std::make_shared<EnumNameRule>(
+           getDpctNamespace() +
+           "experimental::mem_location_type::MEM_LOCATION_TYPE_DEVICE")},
+      {"CU_MEM_LOCATION_TYPE_INVALID",
+       std::make_shared<EnumNameRule>(
+           getDpctNamespace() +
+           "experimental::mem_location_type::MEM_LOCATION_TYPE_INVALID")},
+      {"CU_MEM_LOCATION_TYPE_MAX",
+       std::make_shared<EnumNameRule>(
+           getDpctNamespace() +
+           "experimental::mem_location_type::MEM_LOCATION_TYPE_MAX")},
+      {"CU_MEM_ACCESS_FLAGS_PROT_READWRITE",
+       std::make_shared<EnumNameRule>(getDpctNamespace() +
+                                      "experimental::address_access_flags::"
+                                      "ADDRESS_ACCESS_FLAGS_READ_WRITE")},
+      {"CU_MEM_ACCESS_FLAGS_PROT_NONE",
+       std::make_shared<EnumNameRule>(
+           getDpctNamespace() +
+           "experimental::address_access_flags::ADDRESS_ACCESS_FLAGS_NONE")},
+      {"CU_MEM_ACCESS_FLAGS_PROT_MAX",
+       std::make_shared<EnumNameRule>(
+           getDpctNamespace() +
+           "experimental::address_access_flags::ADDRESS_ACCESS_FLAGS_MAX")},
+      {"CU_MEM_ACCESS_FLAGS_PROT_READ",
+       std::make_shared<EnumNameRule>(
+           getDpctNamespace() +
+           "experimental::address_access_flags::ADDRESS_ACCESS_FLAGS_READ")},
+      {"CU_MEM_ALLOC_GRANULARITY_RECOMMENDED",
+       std::make_shared<EnumNameRule>(
+           getDpctNamespace() +
+           "experimental::granularity_flags::GRANULARITY_FLAGS_RECOMMENDED")},
+      {"CU_MEM_ALLOC_GRANULARITY_MINIMUM",
+       std::make_shared<EnumNameRule>(
+           getDpctNamespace() +
+           "experimental::granularity_flags::GRANULARITY_FLAGS_MINIMUM")},
       // enum Driver Device Attribute
       {"CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR",
        std::make_shared<EnumNameRule>("get_major_version",

clang/lib/DPCT/TypeNames.inc

@@ -137,11 +137,6 @@ ENTRY_TYPE(cooperative_groups::__v1::multi_grid_group, false, NO_FLAG, P4, "comm
 ENTRY_TYPE(cooperative_groups::__v1::thread_block_tile, true, NO_FLAG, P4, "Successful")
 ENTRY_TYPE(cooperative_groups::__v1::coalesced_group, true, NO_FLAG, P4, "Successful/DPCT1119")
 
-//CUmem
-ENTRY_TYPE(CUmemGenericAllocationHandle, false, NO_FLAG, P4, "comment")
-ENTRY_TYPE(CUmemAllocationProp, false, NO_FLAG, P4, "comment")
-ENTRY_TYPE(CUmemAccessDesc, false, NO_FLAG, P4, "comment")
-
 // cuRand
 ENTRY_TYPE(curandStateScrambledSobol64_t, false, NO_FLAG, P4, "comment")
 ENTRY_TYPE(curandStateSobol64_t, false, NO_FLAG, P4, "comment")

clang/lib/DPCT/ValidateArguments.h

@@ -93,6 +93,7 @@ enum class ExperimentalFeatures : unsigned int {
   Exp_Graph,
   Exp_NonUniformGroups,
   Exp_DeviceGlobal,
+  Exp_VirtualMemory,
   Exp_ExperimentalFeaturesEnumSize,
   Exp_NonStandardSYCLBuiltins,
   Exp_All