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

[Artem Belevich via Phabricator via cfe-commits]

tra updated this revision to Diff 369140.
tra edited the summary of this revision.
tra added a comment.

Added more tests


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D108787/new/

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,34 @@
 
   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); }
+};
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
@@ -6498,7 +6498,8 @@
 
     if (Fn->getType() == Context.BoundMemberTy) {
       return BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
-                                       RParenLoc, AllowRecovery);
+                                       RParenLoc, ExecConfig, IsExecConfig,
+                                       AllowRecovery);
     }
   }
 
@@ -6517,7 +6518,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
@@ -3887,6 +3887,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.369140.patch
Type: text/x-patch
Size: 3899 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210827/33cee8b3/attachment.bin>