Commit 6b20ea69 authored by Artem Belevich's avatar Artem Belevich
Browse files

[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
parent 22f01cd4

clang/include/clang/Sema/Sema.h

+2 −0
Original line number Diff line number Diff line
@@ -3891,6 +3891,8 @@ public: 
                                       SourceLocation LParenLoc,

                                       MultiExprArg Args,

                                       SourceLocation RParenLoc,

                                       Expr *ExecConfig = nullptr,

                                       bool IsExecConfig = false,

                                       bool AllowRecovery = false);

  ExprResult

  BuildCallToObjectOfClassType(Scope *S, Expr *Object, SourceLocation LParenLoc,

clang/lib/Sema/SemaExpr.cpp

+4 −2
Original line number Diff line number Diff line
@@ -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);















    }















  }















































clang/lib/Sema/SemaOverload.cpp



+3
−2






























Original line number
Diff line number
Diff line








@@ -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());

























clang/test/SemaCUDA/kernel-call.cu



+31
−0






























Original line number
Diff line number
Diff line








@@ -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); }













};