[SYCLomatic] Add migration for 2 virtual memory related enums by removing them(#2764)

Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>

Author: zhiweij1

clang/lib/DPCT/RuleInfra/MapNames.cpp

@@ -1145,6 +1145,8 @@ void MapNames::setExplicitNamespaceMap(
       {"CU_MEM_ADVISE_UNSET_ACCESSED_BY", std::make_shared<EnumNameRule>("0")},
       {"CU_MEM_ALLOCATION_TYPE_PINNED", std::make_shared<EnumNameRule>("0")},
       {"CU_MEM_ALLOCATION_TYPE_INVALID", std::make_shared<EnumNameRule>("1")},
+      {"CU_MEM_ALLOCATION_COMP_NONE", std::make_shared<EnumNameRule>("0")},
+      {"CU_MEM_ALLOCATION_COMP_GENERIC", std::make_shared<EnumNameRule>("1")},
       {"CU_MEM_ALLOCATION_TYPE_MAX",
        std::make_shared<EnumNameRule>("0xFFFFFFFF")},
       {"CU_MEM_LOCATION_TYPE_DEVICE", std::make_shared<EnumNameRule>("1")},

clang/lib/DPCT/RulesLang/RulesLang.cpp

@@ -8167,7 +8167,8 @@ void VirtualMemRule::registerMatcher(ast_matchers::MatchFinder &MF) {
         "CU_MEM_ACCESS_FLAGS_PROT_NONE", "CU_MEM_ACCESS_FLAGS_PROT_READ",
         "CU_MEM_ACCESS_FLAGS_PROT_READWRITE",
         "CU_MEM_ALLOC_GRANULARITY_RECOMMENDED",
-        "CU_MEM_ALLOC_GRANULARITY_MINIMUM");
+        "CU_MEM_ALLOC_GRANULARITY_MINIMUM", "CU_MEM_ALLOCATION_COMP_NONE",
+        "CU_MEM_ALLOCATION_COMP_GENERIC");
   };
   MF.addMatcher(
       callExpr(callee(functionDecl(virtualmemoryAPI()))).bind("vmCall"), this);
@@ -8178,11 +8179,47 @@ void VirtualMemRule::registerMatcher(ast_matchers::MatchFinder &MF) {
   MF.addMatcher(
       declRefExpr(to(enumConstantDecl(virtualmemoryEnum()))).bind("vmEnum"),
       this);
+  MF.addMatcher(
+      memberExpr(
+          hasObjectExpression(hasType(qualType(
+              hasCanonicalType(asString("struct CUmemAllocationProp_st"))))),
+          member(hasName("allocFlags")),
+          hasParent(memberExpr(anyOf(
+              hasParent(implicitCastExpr(
+                            hasCastKind(CK_LValueToRValue),
+                            hasParent(binaryOperator(isAssignmentOperator())))
+                            .bind("replaceWithZero")),
+              hasParent(
+                  binaryOperator(isAssignmentOperator()).bind("removeBO")))))),
+      this);
 }
 
 void VirtualMemRule::runRule(
     const ast_matchers::MatchFinder::MatchResult &Result) {
   auto &SM = DpctGlobalInfo::getSourceManager();
+
+  const Expr *ICE = getNodeAsType<ImplicitCastExpr>(Result, "replaceWithZero");
+  const Expr *BO = getNodeAsType<BinaryOperator>(Result, "removeBO");
+  const Expr *E = ICE ? ICE : BO;
+  if (E) {
+    // Process pattern like:
+    // prop.allocFlags.compressionType = CU_MEM_ALLOCATION_COMP_NONE;
+    // uc = prop.allocFlags.compressionType;
+    auto Range = getDefinitionRange(E->getBeginLoc(), E->getEndLoc());
+    auto Begin = Range.getBegin();
+    auto End = Range.getEnd();
+    auto Length = Lexer::MeasureTokenLength(
+        End, SM, dpct::DpctGlobalInfo::getContext().getLangOpts());
+    Length +=
+        SM.getDecomposedLoc(End).second - SM.getDecomposedLoc(Begin).second;
+    if (ICE) {
+      emplaceTransformation(new ReplaceText(Begin, Length, std::move("0")));
+    } else {
+      emplaceTransformation(
+          new ReplaceText(Begin, Length, std::move("(void)0")));
+    }
+    return;
+  }
   if (const CallExpr *CE = getNodeAsType<CallExpr>(Result, "vmCall")) {
     ExprAnalysis EA(CE);
     emplaceTransformation(EA.getReplacement());

clang/test/dpct/virtual_memory.cu

@@ -13,15 +13,20 @@ int main() {
     cuDeviceGet(&device, 0);
     CUcontext context;
     cuCtxCreate(&context, 0, device);
+    unsigned char uc;
 
 // CHECK:  dpct::experimental::mem_prop prop = {};
 // CHECK:  prop.type = 0;
 // CHECK:  prop.location.type = 1;
 // CHECK:  prop.location.id = device;
+// CHECK:  (void)0;
+// CHECK:  uc = 0;
     CUmemAllocationProp prop = {};
     prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
     prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
     prop.location.id = device;
+    prop.allocFlags.compressionType = CU_MEM_ALLOCATION_COMP_NONE;
+    uc = prop.allocFlags.compressionType;
     size_t granularity;
 // CHECK:    granularity = sycl::ext::oneapi::experimental::get_mem_granularity(dpct::get_device(prop.location.id), dpct::get_device(prop.location.id).get_context(), sycl::ext::oneapi::experimental::granularity_mode::minimum);
     cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM);