Author: tra Date: Mon Jun 6 17:54:57 2016 New Revision: 271951 URL: http://llvm.org/viewvc/llvm-project?rev=271951&view=rev Log: [CUDA] Add implicit conversion of __launch_bounds__ arguments to rvalue.
Fixes clang crash reported in PR27778. Differential Revision: http://reviews.llvm.org/D20985 Added: cfe/trunk/test/SemaCUDA/pr27778.cu Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp cfe/trunk/test/CodeGenCUDA/launch-bounds.cu Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=271951&r1=271950&r2=271951&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original) +++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Mon Jun 6 17:54:57 2016 @@ -28,6 +28,7 @@ #include "clang/Lex/Preprocessor.h" #include "clang/Sema/DeclSpec.h" #include "clang/Sema/DelayedDiagnostic.h" +#include "clang/Sema/Initialization.h" #include "clang/Sema/Lookup.h" #include "clang/Sema/Scope.h" #include "llvm/ADT/StringExtras.h" @@ -4039,48 +4040,60 @@ bool Sema::CheckRegparmAttr(const Attrib return false; } -// Checks whether an argument of launch_bounds attribute is acceptable -// May output an error. -static bool checkLaunchBoundsArgument(Sema &S, Expr *E, - const CUDALaunchBoundsAttr &Attr, - const unsigned Idx) { +// Checks whether an argument of launch_bounds attribute is +// acceptable, performs implicit conversion to Rvalue, and returns +// non-nullptr Expr result on success. Otherwise, it returns nullptr +// and may output an error. +static Expr *makeLaunchBoundsArgExpr(Sema &S, Expr *E, + const CUDALaunchBoundsAttr &Attr, + const unsigned Idx) { if (S.DiagnoseUnexpandedParameterPack(E)) - return false; + return nullptr; // Accept template arguments for now as they depend on something else. // We'll get to check them when they eventually get instantiated. if (E->isValueDependent()) - return true; + return E; llvm::APSInt I(64); if (!E->isIntegerConstantExpr(I, S.Context)) { S.Diag(E->getExprLoc(), diag::err_attribute_argument_n_type) << &Attr << Idx << AANT_ArgumentIntegerConstant << E->getSourceRange(); - return false; + return nullptr; } // Make sure we can fit it in 32 bits. if (!I.isIntN(32)) { S.Diag(E->getExprLoc(), diag::err_ice_too_large) << I.toString(10, false) << 32 << /* Unsigned */ 1; - return false; + return nullptr; } if (I < 0) S.Diag(E->getExprLoc(), diag::warn_attribute_argument_n_negative) << &Attr << Idx << E->getSourceRange(); - return true; + // We may need to perform implicit conversion of the argument. + InitializedEntity Entity = InitializedEntity::InitializeParameter( + S.Context, S.Context.getConstType(S.Context.IntTy), /*consume*/ false); + ExprResult ValArg = S.PerformCopyInitialization(Entity, SourceLocation(), E); + assert(!ValArg.isInvalid() && + "Unexpected PerformCopyInitialization() failure."); + + return ValArg.getAs<Expr>(); } void Sema::AddLaunchBoundsAttr(SourceRange AttrRange, Decl *D, Expr *MaxThreads, Expr *MinBlocks, unsigned SpellingListIndex) { CUDALaunchBoundsAttr TmpAttr(AttrRange, Context, MaxThreads, MinBlocks, SpellingListIndex); - - if (!checkLaunchBoundsArgument(*this, MaxThreads, TmpAttr, 0)) + MaxThreads = makeLaunchBoundsArgExpr(*this, MaxThreads, TmpAttr, 0); + if (MaxThreads == nullptr) return; - if (MinBlocks && !checkLaunchBoundsArgument(*this, MinBlocks, TmpAttr, 1)) - return; + if (MinBlocks) { + MinBlocks = makeLaunchBoundsArgExpr(*this, MinBlocks, TmpAttr, 1); + if (MinBlocks == nullptr) + return; + } D->addAttr(::new (Context) CUDALaunchBoundsAttr( AttrRange, Context, MaxThreads, MinBlocks, SpellingListIndex)); Modified: cfe/trunk/test/CodeGenCUDA/launch-bounds.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/launch-bounds.cu?rev=271951&r1=271950&r2=271951&view=diff ============================================================================== --- cfe/trunk/test/CodeGenCUDA/launch-bounds.cu (original) +++ cfe/trunk/test/CodeGenCUDA/launch-bounds.cu Mon Jun 6 17:54:57 2016 @@ -79,3 +79,8 @@ Kernel7() } // CHECK: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel7{{.*}}, !"maxntidx", // CHECK-NOT: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel7{{.*}}, !"minctasm", + +const char constchar = 12; +__global__ void __launch_bounds__(constint, constchar) Kernel8() {} +// CHECK: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel8{{.*}}, !"maxntidx", i32 100 +// CHECK: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel8{{.*}}, !"minctasm", i32 12 Added: cfe/trunk/test/SemaCUDA/pr27778.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/pr27778.cu?rev=271951&view=auto ============================================================================== --- cfe/trunk/test/SemaCUDA/pr27778.cu (added) +++ cfe/trunk/test/SemaCUDA/pr27778.cu Mon Jun 6 17:54:57 2016 @@ -0,0 +1,6 @@ +// RUN: %clang_cc1 -fsyntax-only %s + +#include "Inputs/cuda.h" + +const int constint = 512; +__launch_bounds__(constint) void TestConstInt(void) {} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
