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,
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,
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
