[SYCLomatic] Migrate cudaMallocPitch with rewriter (#753)

Signed-off-by: Huang, Andy <[email protected]>

Author: AndyCHHuang

clang/lib/DPCT/APINamesMemory.inc

@@ -362,3 +362,20 @@ ASSIGNABLE_FACTORY(
         HelperFeatureEnum::Memory_pointer_attributes,
         MEMBER_CALL_FACTORY_ENTRY(
             "cudaPointerGetAttributes", DEREF(0), false, "init", ARG(1))))
+
+ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+    HelperFeatureEnum::Memory_dpct_malloc_2d,
+    CONDITIONAL_FACTORY_ENTRY(
+        CheckDerefedTypeBeforeCast(0, "void *"),
+        ASSIGN_FACTORY_ENTRY(
+            "cudaMallocPitch", DEREF(makeDerefArgCreatorWithCall(0)),
+            CALL(MapNames::getDpctNamespace() + "dpct_malloc",
+                 DEREF(makeCallArgCreatorWithCall(1)),
+                 makeCallArgCreatorWithCall(2), makeDerefArgCreatorWithCall(3))),
+        ASSIGN_FACTORY_ENTRY(
+            "cudaMallocPitch", DEREF(makeDerefArgCreatorWithCall(0)),
+            CAST(getDerefedType(0),
+                 CALL(MapNames::getDpctNamespace() + "dpct_malloc",
+                      DEREF(makeCallArgCreatorWithCall(1)),
+                      makeCallArgCreatorWithCall(2),
+                      makeCallArgCreatorWithCall(3)))))))

clang/lib/DPCT/ASTTraversal.cpp

@@ -10332,7 +10332,7 @@ void MemoryMigrationRule::mallocMigration(
     }
   } else if (Name == "cudaHostAlloc" || Name == "cudaMallocHost" ||
              Name == "cuMemHostAlloc" || Name == "cuMemAllocHost_v2" ||
-             Name == "cuMemAllocPitch_v2") {
+             Name == "cuMemAllocPitch_v2" || Name == "cudaMallocPitch") {
     ExprAnalysis EA(C);
     emplaceTransformation(EA.getReplacement());
     EA.applyAllSubExprRepl();
@@ -10401,7 +10401,7 @@ void MemoryMigrationRule::mallocMigration(
         report(C->getBeginLoc(), Diagnostics::NOERROR_RETURN_COMMA_OP, false);
       }
     }
-  } else if (Name == "cudaMallocPitch" || Name == "cudaMalloc3D") {
+  } else if (Name == "cudaMalloc3D") {
     std::ostringstream OS;
     std::string Type;
     if (IsAssigned)
@@ -10420,9 +10420,6 @@ void MemoryMigrationRule::mallocMigration(
     emplaceTransformation(removeArg(C, 0, *Result.SourceManager));
     std::ostringstream OS2;
     printDerefOp(OS2, C->getArg(1));
-    if (Name == "cudaMallocPitch") {
-      emplaceTransformation(new ReplaceStmt(C->getArg(1), OS2.str()));
-    }
     if (IsAssigned) {
       emplaceTransformation(new InsertAfterStmt(C, ", 0)"));
       report(C->getBeginLoc(), Diagnostics::NOERROR_RETURN_COMMA_OP, false);

clang/lib/DPCT/CallExprRewriter.h

@@ -636,6 +636,7 @@ class DerefExpr {
     if (!AddrOfRemoved && !IgnoreDerefOp)
       Stream << "*";
 
+    AA.setCallSpelling(C);
     printWithParens(Stream, AA, P);
   }
 
@@ -644,7 +645,7 @@ class DerefExpr {
 public:
   template <class StreamT>
   void printArg(StreamT &Stream, ArgumentAnalysis &A) const {
-    print(Stream, A, false);
+    print(Stream);
   }
   template <class StreamT> void printMemberBase(StreamT &Stream) const {
     ExprAnalysis EA;
@@ -658,6 +659,7 @@ class DerefExpr {
       print(Stream, EA, false);
     } else {
       ArgumentAnalysis AA;
+
       std::pair<const CallExpr*, const Expr*> ExprPair(C, E);
       print(Stream, AA, false, ExprPair);
     }

clang/test/dpct/memory_management.cu

@@ -195,6 +195,8 @@ void checkError(cudaError_t err) {
 
 void cuCheckError(CUresult err) {
 }
+//CHECK: #define PITCH(a,b,c,d) a = (float *)dpct::dpct_malloc(b, c, d);
+#define PITCH(a,b,c,d) cudaMallocPitch(a, b, c, d);
 
 void testCommas() {
   size_t size = 1234567 * sizeof(float);
@@ -219,10 +221,13 @@ void testCommas() {
 
   // CHECK: d_A = (float *)dpct::dpct_malloc(size, size, size);
   cudaMallocPitch((void **)&d_A, &size, size, size);
+
+  // CHECK: PITCH((void **)&d_A, &size, size, size);
+  PITCH((void **)&d_A, &size, size, size);
   int sz;
-  // CHECK: d_A = (float *)dpct::dpct_malloc(*((size_t *)&size), size, size);
+  // CHECK: d_A = (float *)dpct::dpct_malloc(*(size_t *)&size, size, size);
   cudaMallocPitch((void **)&d_A, (size_t *)&size, size, size);
-  // CHECK: d_A = (float *)dpct::dpct_malloc(*((size_t *)&sz), size, size);
+  // CHECK: d_A = (float *)dpct::dpct_malloc(*(size_t *)&sz, size, size);
   cudaMallocPitch((void **)&d_A, (size_t *)&sz, size, size);
   // CHECK:/*
   // CHECK-NEXT:DPCT1003:{{[0-9]+}}: Migrated API does not return error code. (*, 0) is inserted. You may need to rewrite this code.