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