[clang] 6b20ea6 - [CUDA] Pass ExecConfig through BuildCallToMemberFunction

[Artem Belevich via cfe-commits]

Author: Artem Belevich
Date: 2021-09-16T11:18:12-07:00
New Revision: 6b20ea6963561f2c91490c0993390b7f2ff8f71c

URL: https://github.com/llvm/llvm-project/commit/6b20ea6963561f2c91490c0993390b7f2ff8f71c
DIFF: https://github.com/llvm/llvm-project/commit/6b20ea6963561f2c91490c0993390b7f2ff8f71c.diff

LOG: [CUDA] Pass ExecConfig through BuildCallToMemberFunction

Otherwise, we fail to compile calls to CUDA kernels that are static members.

Differential Revision: https://reviews.llvm.org/D108787

Added: 
    

Modified: 
    clang/include/clang/Sema/Sema.h
    clang/lib/Sema/SemaExpr.cpp
    clang/lib/Sema/SemaOverload.cpp
    clang/test/SemaCUDA/kernel-call.cu

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 710abeb1ea514..6214ef6503848 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -3891,6 +3891,8 @@ class Sema final {
                                        SourceLocation LParenLoc,
                                        MultiExprArg Args,
                                        SourceLocation RParenLoc,
+                                       Expr *ExecConfig = nullptr,
+                                       bool IsExecConfig = false,
                                        bool AllowRecovery = false);
   ExprResult
   BuildCallToObjectOfClassType(Scope *S, Expr *Object, SourceLocation LParenLoc,

diff  --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 7c8d753730a0a..7991e4c5e8b0a 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -6454,7 +6454,8 @@ ExprResult Sema::BuildCallExpr(Scope *Scope, Expr *Fn, SourceLocation LParenLoc,
 
     if (Fn->getType() == Context.BoundMemberTy) {
       return BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
-                                       RParenLoc, AllowRecovery);
+                                       RParenLoc, ExecConfig, IsExecConfig,
+                                       AllowRecovery);
     }
   }
 
@@ -6473,7 +6474,8 @@ ExprResult Sema::BuildCallExpr(Scope *Scope, Expr *Fn, SourceLocation LParenLoc,
             Scope, Fn, ULE, LParenLoc, ArgExprs, RParenLoc, ExecConfig,
             /*AllowTypoCorrection=*/true, find.IsAddressOfOperand);
       return BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
-                                       RParenLoc, AllowRecovery);
+                                       RParenLoc, ExecConfig, IsExecConfig,
+                                       AllowRecovery);
     }
   }
 

diff  --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index 4b1b61a390276..e12413ca994f4 100644
--- a/clang/lib/Sema/SemaOverload.cpp
+++ b/clang/lib/Sema/SemaOverload.cpp
@@ -14233,6 +14233,7 @@ ExprResult Sema::BuildCallToMemberFunction(Scope *S, Expr *MemExprE,
                                            SourceLocation LParenLoc,
                                            MultiExprArg Args,
                                            SourceLocation RParenLoc,
+                                           Expr *ExecConfig, bool IsExecConfig,
                                            bool AllowRecovery) {
   assert(MemExprE->getType() == Context.BoundMemberTy ||
          MemExprE->getType() == Context.OverloadTy);
@@ -14428,8 +14429,8 @@ ExprResult Sema::BuildCallToMemberFunction(Scope *S, Expr *MemExprE,
     // If overload resolution picked a static member, build a
     // non-member call based on that function.
     if (Method->isStatic()) {
-      return BuildResolvedCallExpr(MemExprE, Method, LParenLoc, Args,
-                                   RParenLoc);
+      return BuildResolvedCallExpr(MemExprE, Method, LParenLoc, Args, RParenLoc,
+                                   ExecConfig, IsExecConfig);
     }
 
     MemExpr = cast<MemberExpr>(MemExprE->IgnoreParens());

diff  --git a/clang/test/SemaCUDA/kernel-call.cu b/clang/test/SemaCUDA/kernel-call.cu
index b2433c956e251..136844c7c8d65 100644
--- a/clang/test/SemaCUDA/kernel-call.cu
+++ b/clang/test/SemaCUDA/kernel-call.cu
@@ -26,3 +26,34 @@ int main(void) {
 
   g1<<<undeclared, 1>>>(42); // expected-error {{use of undeclared identifier 'undeclared'}}
 }
+
+// Make sure we can call static member kernels.
+template <typename > struct a0 {
+  template <typename T> static __global__ void Call(T);
+};
+struct a1 {
+  template <typename T> static __global__ void Call(T);
+};
+template <typename T> struct a2 {
+  static __global__ void Call(T);
+};
+struct a3 {
+  static __global__ void Call(int);
+  static __global__ void Call(void*);
+};
+
+struct b {
+  template <typename c> void d0(c arg) {
+    a0<c>::Call<<<0, 0>>>(arg);
+    a1::Call<<<0,0>>>(arg);
+    a2<c>::Call<<<0,0>>>(arg);
+    a3::Call<<<0, 0>>>(arg);
+  }
+  void d1(void* arg) {
+    a0<void*>::Call<<<0, 0>>>(arg);
+    a1::Call<<<0,0>>>(arg);
+    a2<void*>::Call<<<0,0>>>(arg);
+    a3::Call<<<0, 0>>>(arg);
+  }
+  void e() { d0(1); }
+};