[PATCH] D123976: [CUDA][HIP] Fix delete operator for -fopenmp

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added a reviewer: tra.
Herald added subscribers: mattd, guansong.
Herald added a project: All.
yaxunl requested review of this revision.
Herald added a reviewer: jdoerfert.
Herald added a subscriber: sstefan1.

When new operator is called in OpenMP parallel region,
delete operator is resolved and checked. Due to similar
issue fixed by https://reviews.llvm.org/D121765,
when resolving delete operator, the caller was not
determined correctly, which results in error as
shown in https://godbolt.org/z/jKhd8qKos.

This patch fixes the issue in a similar way as
https://reviews.llvm.org/D121765


https://reviews.llvm.org/D123976

Files:
  clang/lib/Sema/SemaExprCXX.cpp
  clang/test/SemaCUDA/openmp-parallel.cu


Index: clang/test/SemaCUDA/openmp-parallel.cu
===================================================================
--- clang/test/SemaCUDA/openmp-parallel.cu
+++ clang/test/SemaCUDA/openmp-parallel.cu
@@ -1,4 +1,5 @@
 // RUN: %clang_cc1 -fopenmp -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fopenmp -fexceptions -fsyntax-only -verify %s
 
 #include "Inputs/cuda.h"
 
@@ -7,13 +8,17 @@
 
 int main() {
   #pragma omp parallel
-  for (int i = 0; i < 100; i++)
+  for (int i = 0; i < 100; i++) {
     foo(1); // expected-error {{no matching function for call to 'foo'}}
-  
+    new int;
+  }
+
   auto Lambda = []() {
     #pragma omp parallel
-    for (int i = 0; i < 100; i++)
+    for (int i = 0; i < 100; i++) {
       foo(1); // expected-error {{reference to __device__ function 'foo' in __host__ __device__ function}}
-    };
+      new int;
+    }
+  };
   Lambda(); // expected-note {{called by 'main'}}
 }
Index: clang/lib/Sema/SemaExprCXX.cpp
===================================================================
--- clang/lib/Sema/SemaExprCXX.cpp
+++ clang/lib/Sema/SemaExprCXX.cpp
@@ -1597,7 +1597,8 @@
 
 bool Sema::isUsualDeallocationFunction(const CXXMethodDecl *Method) {
   // [CUDA] Ignore this function, if we can't call it.
-  const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
+  const FunctionDecl *Caller =
+      dyn_cast<FunctionDecl>(getCurFunctionDecl(/*AllowLambda=*/true));
   if (getLangOpts().CUDA) {
     auto CallPreference = IdentifyCUDAPreference(Caller, Method);
     // If it's not callable at all, it's not the right function.
@@ -1691,7 +1692,8 @@
 
       // In CUDA, determine how much we'd like / dislike to call this.
       if (S.getLangOpts().CUDA)
-        if (auto *Caller = dyn_cast<FunctionDecl>(S.CurContext))
+        if (auto *Caller = dyn_cast<FunctionDecl>(
+                S.getCurFunctionDecl(/*AllowLambda=*/true)))
           CUDAPref = S.IdentifyCUDAPreference(Caller, FD);
     }
 
@@ -2830,7 +2832,9 @@
     }
 
     if (getLangOpts().CUDA)
-      EraseUnwantedCUDAMatches(dyn_cast<FunctionDecl>(CurContext), Matches);
+      EraseUnwantedCUDAMatches(
+          dyn_cast<FunctionDecl>(getCurFunctionDecl(/*AllowLambda=*/true)),
+          Matches);
   } else {
     // C++1y [expr.new]p22:
     //   For a non-placement allocation function, the normal deallocation


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D123976.423513.patch
Type: text/x-patch
Size: 2338 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20220419/6a7e2947/attachment-0001.bin>