[CUDA][HIP] Fix delete operator for -fopenmp

[yaxunl]

Summary

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

[yaxunl]

remove redundant casts.

clang/lib/Sema/SemaExprCXX.cpp

Result = CXXFunctionalCastExpr::Create(
         Locs.getBegin(), Locs.getEnd());
   }
 
   return Result;
 }
 
 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.
     if (CallPreference < CFP_WrongSide)
       return false;
     if (CallPreference == CFP_WrongSide) {
       // Maybe. We have to check if there are better alternatives.
       DeclContext::lookup_result R =
UsualDeallocFnInfo(Sema &S, DeclAccessPair Found)
       if (NumBaseParams < FD->getNumParams() &&
           FD->getParamDecl(NumBaseParams)->getType()->isAlignValT()) {
         ++NumBaseParams;
         HasAlignValT = true;
       }
 
       // 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);
     }
 
     explicit operator bool() const { return FD; }
 
     bool isBetterThan(const UsualDeallocFnInfo &Other, bool WantSize,
                       bool WantAlign) const {
       // C++ P0722:
for (LookupResult::iterator D = FoundDelete.begin(),
       if (Context.hasSameType(adjustCCAndNoReturn(Fn->getType(),
                                                   ExpectedFunctionType,
                                                   /*AdjustExcpetionSpec*/true),
                               ExpectedFunctionType))
         Matches.push_back(std::make_pair(D.getPair(), Fn));
     }
 
     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
     //   function lookup is used
     //
     // Per [expr.delete]p10, this lookup prefers a member operator delete
     // without a size_t argument, but prefers a non-member operator delete
     // with a size_t where possible (which it always is in this case).

clang/test/SemaCUDA/openmp-parallel.cu

 // RUN: %clang_cc1 -fopenmp -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fopenmp -fexceptions -fsyntax-only -verify %s
 
 #include "Inputs/cuda.h"
 
 __device__ void foo(int) {} // expected-note {{candidate function not viable: call to __device__ function from __host__ function}}
 // expected-note@-1 {{'foo' declared here}}
 
 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'}}
 }