[PATCH] D108787: [CUDA] Pass ExecConfig through BuildCallToMemberFunction

[Artem Belevich via Phabricator via cfe-commits]

tra created this revision.
tra added reviewers: jlebar, yaxunl.
Herald added subscribers: bixia, inglorion.
tra requested review of this revision.
Herald added a project: clang.

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


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D108787

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


Index: clang/test/SemaCUDA/kernel-call.cu
===================================================================
--- clang/test/SemaCUDA/kernel-call.cu
+++ clang/test/SemaCUDA/kernel-call.cu
@@ -26,3 +26,14 @@
 
   g1<<<undeclared, 1>>>(42); // expected-error {{use of undeclared identifier 'undeclared'}}
 }
+
+// Make sure we can call static member kernels.
+template <typename > struct a {
+  template <typename T> static __global__ void Call(T);
+};
+struct b {
+  template <typename c> void d(c arg) {
+    a<c>::Call<<<0, 0>>>(arg);
+  }
+  void e() { d(1); }
+};
Index: clang/lib/Sema/SemaOverload.cpp
===================================================================
--- clang/lib/Sema/SemaOverload.cpp
+++ clang/lib/Sema/SemaOverload.cpp
@@ -14166,6 +14166,7 @@
                                            SourceLocation LParenLoc,
                                            MultiExprArg Args,
                                            SourceLocation RParenLoc,
+                                           Expr *ExecConfig, bool IsExecConfig,
                                            bool AllowRecovery) {
   assert(MemExprE->getType() == Context.BoundMemberTy ||
          MemExprE->getType() == Context.OverloadTy);
@@ -14361,8 +14362,8 @@
     // 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());
Index: clang/lib/Sema/SemaExpr.cpp
===================================================================
--- clang/lib/Sema/SemaExpr.cpp
+++ clang/lib/Sema/SemaExpr.cpp
@@ -6505,7 +6505,8 @@
 
     if (Fn->getType() == Context.BoundMemberTy) {
       return BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
-                                       RParenLoc, AllowRecovery);
+                                       RParenLoc, ExecConfig, IsExecConfig,
+                                       AllowRecovery);
     }
   }
 
@@ -6524,7 +6525,8 @@
             Scope, Fn, ULE, LParenLoc, ArgExprs, RParenLoc, ExecConfig,
             /*AllowTypoCorrection=*/true, find.IsAddressOfOperand);
       return BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
-                                       RParenLoc, AllowRecovery);
+                                       RParenLoc, ExecConfig, IsExecConfig,
+                                       AllowRecovery);
     }
   }
 
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -3885,6 +3885,8 @@
                                        SourceLocation LParenLoc,
                                        MultiExprArg Args,
                                        SourceLocation RParenLoc,
+                                       Expr *ExecConfig = nullptr,
+                                       bool IsExecConfig = false,
                                        bool AllowRecovery = false);
   ExprResult
   BuildCallToObjectOfClassType(Scope *S, Expr *Object, SourceLocation LParenLoc,


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D108787.368979.patch
Type: text/x-patch
Size: 3400 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210826/9dc62113/attachment.bin>