From 800f26386cd9054ceed68e481612908f635b9718 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu@amd.com>
Date: Mon, 18 Apr 2022 22:21:47 -0400
Subject: [PATCH] [CUDA][HIP] Fix delete operator for -fopenmp

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

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D123976
---
 clang/lib/Sema/SemaExprCXX.cpp         |  7 ++++---
 clang/test/SemaCUDA/openmp-parallel.cu | 13 +++++++++----
 2 files changed, 13 insertions(+), 7 deletions(-)

diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp
index e765d1ff9760d..77353d8307bc6 100644
--- a/clang/lib/Sema/SemaExprCXX.cpp
+++ b/clang/lib/Sema/SemaExprCXX.cpp
@@ -1597,7 +1597,7 @@ Sema::BuildCXXTypeConstructExpr(TypeSourceInfo *TInfo,
 
 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 = 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 +1691,7 @@ namespace {
 
       // 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 = S.getCurFunctionDecl(/*AllowLambda=*/true))
           CUDAPref = S.IdentifyCUDAPreference(Caller, FD);
     }
 
@@ -2830,7 +2830,8 @@ bool Sema::FindAllocationFunctions(SourceLocation StartLoc, SourceRange Range,
     }
 
     if (getLangOpts().CUDA)
-      EraseUnwantedCUDAMatches(dyn_cast<FunctionDecl>(CurContext), Matches);
+      EraseUnwantedCUDAMatches(getCurFunctionDecl(/*AllowLambda=*/true),
+                               Matches);
   } else {
     // C++1y [expr.new]p22:
     //   For a non-placement allocation function, the normal deallocation
diff --git a/clang/test/SemaCUDA/openmp-parallel.cu b/clang/test/SemaCUDA/openmp-parallel.cu
index 2e519e903b25c..79dc1fee4bb1b 100644
--- a/clang/test/SemaCUDA/openmp-parallel.cu
+++ b/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 @@ __device__ void foo(int) {} // expected-note {{candidate function not viable: ca
 
 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'}}
 }