[SYCLomatic] CodePin fix malloc instrumentation location issue in mac…

…ro context (#1984) Signed-off-by: Huang, Andy <[email protected]>
oneapi-src · May 15, 2024 · 8ea3832 · 8ea3832
1 parent faf16d7
commit 8ea3832
Show file tree

Hide file tree

Showing 2 changed files with 13 additions and 3 deletions.
diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp
@@ -9872,10 +9872,19 @@ bool MemoryMigrationRule::canUseTemplateStyleMigration(
 
 void MemoryMigrationRule::instrumentAddressToSizeRecordForCodePin(
     const CallExpr *C, int PtrArgLoc, int AllocMemSizeLoc) {
+  auto &SM = dpct::DpctGlobalInfo::getSourceManager(); 
   if (DpctGlobalInfo::isCodePinEnabled()) {
+    SourceLocation CallEnd = C->getEndLoc();
+    if(SM.isMacroArgExpansion(C->getEndLoc())){
+      CallEnd = SM.getExpansionRange(C->getEndLoc()).getEnd();
+    } else {
+      CallEnd = getDefinitionRange(C->getBeginLoc(), C->getEndLoc()).getEnd();
+    }
+
     auto PtrSizeLoc = Lexer::findLocationAfterToken(
-        C->getEndLoc(), tok::semi, DpctGlobalInfo::getSourceManager(),
-        DpctGlobalInfo::getContext().getLangOpts(), false);
+        CallEnd, tok::semi, SM, DpctGlobalInfo::getContext().getLangOpts(),
+        false);
+
     emplaceTransformation(new InsertText(
         PtrSizeLoc,
         std::string(getNL()) + "dpctexp::codepin::get_ptr_size_map()[" +

diff --git a/clang/test/dpct/debug_test/vector_add/vectorAdd_build_in_type.cu b/clang/test/dpct/debug_test/vector_add/vectorAdd_build_in_type.cu
@@ -19,13 +19,14 @@ __global__ void VectorAddKernel(float *A, float *B, float *C) {
   B[threadIdx.x] = threadIdx.x + 1.0f;
   C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x];
 }
+#define CALL(x) if(0!=x){int a=4;}
 
 int main() {
   float *d_A, *d_B, *d_C;
   cudaError_t status;
 
   //CHECK: dpctexp::codepin::get_ptr_size_map()[d_A] = VECTOR_SIZE * sizeof(float);
-  cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float));
+  CALL(cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)));
   //CHECK: dpctexp::codepin::get_ptr_size_map()[d_B] = VECTOR_SIZE * sizeof(float);
   cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float));
   //CHECK: dpctexp::codepin::get_ptr_size_map()[d_C] = VECTOR_SIZE * sizeof(float);