[CUDA] Add implicit conversion of __launch_bounds__ arguments to rvalue.

[tra]

Summary

Fixes clang crash reported in PR27778.

[jlebar]

How is this different from test/SemaCUDA/launch_bounds.cu:27-28? It does

const int constint = 512;
__launch_bounds__(constint) void TestConstInt(void);

which looks verbatim the same as this testcase.

[tra]

In D20985#448822, @jlebar wrote:

How is this different from test/SemaCUDA/launch_bounds.cu:27-28? It does

const int constint = 512;
__launch_bounds__(constint) void TestConstInt(void);

which looks verbatim the same as this testcase.

Existing test is a declaration of the function which did not trigger the crash.
Second issue is that -verify interferes with reproduction case -- the crash does not happen if any //expect-* are seen before it.
Plus, the outcome of the failing test is a crash which would prevent reports of other failures.
Separate test file makes the crash isolated and reliably reproducible.

[jlebar]

In D20985#448836, @tra wrote:

In D20985#448822, @jlebar wrote:

How is this different from test/SemaCUDA/launch_bounds.cu:27-28? It does

const int constint = 512;
__launch_bounds__(constint) void TestConstInt(void);

which looks verbatim the same as this testcase.

Existing test is a declaration of the function which did not trigger the crash.
Second issue is that -verify interferes with reproduction case -- the crash does not happen if any //expect-* are seen before it.
Plus, the outcome of the failing test is a crash which would prevent reports of other failures.
Separate test file makes the crash isolated and reliably reproducible.

Got it, thanks.

Should we have a test that passes a char or a short and ensures that we do the correct implicit conversion there?

lib/Sema/SemaDeclAttr.cpp

4045

Should we update this name and comment?

[tra]

Addressed Justin's comments.

[jlebar]

lib/Sema/SemaDeclAttr.cpp

4046

Nit, Oxford comma helps some here.

4048

Presumably it "returns nullptr and outputs an error" otherwise? Like, we get nullptr iff it outputs an error?

4077

Do we need to output an error here, or is does PerformCopyInitialization do so for us? In any case, is it covered by a test?

[tra]

lib/Sema/SemaDeclAttr.cpp

4048

It returns nullptr without error message in case of unexpanded parameter pack.
Arguments for that case will be checked after template instantiation is done. See instantiateDependentCUDALaunchBoundsAttr() in SemaTemplateInstantiateDecl.cpp.

4077

Actually, by this point we've verified that Expr is an ICE. I assume that PerformCopyInitialization() should always succeed for such an expression.

[tra]

Rephrased comments

[jlebar]

lib/Sema/SemaDeclAttr.cpp

4077

OK, so then we want an assert, not an if?

[tra]

Replaced if() with assert() to catch unexpected PerformCopyInitialization() failures.

lib/Sema/SemaDeclAttr.cpp

 #include "clang/AST/Mangle.h"
 #include "clang/AST/ASTMutationListener.h"
 #include "clang/Basic/CharInfo.h"
 #include "clang/Basic/SourceManager.h"
 #include "clang/Basic/TargetInfo.h"
 #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"
 #include "llvm/Support/MathExtras.h"
 
 using namespace clang;
 using namespace sema;
 
if (numParams > Context.getTargetInfo().getRegParmMax()) {
     return true;
   }
 
   return false;
 }
 
 // Checks whether an argument of launch_bounds attribute is acceptable
 // May output an error.
-static bool checkLaunchBoundsArgument(Sema &S, Expr *E,
+static Expr *checkLaunchBoundsArgument(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);
+  if (ValArg.isInvalid())
+    return nullptr;
+
+  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 = checkLaunchBoundsArgument(*this, MaxThreads, TmpAttr, 0);
+  if (MaxThreads == nullptr)
     return;
 
-  if (MinBlocks && !checkLaunchBoundsArgument(*this, MinBlocks, TmpAttr, 1))
+  if (MinBlocks) {
+    MinBlocks = checkLaunchBoundsArgument(*this, MinBlocks, TmpAttr, 1);
+    if (MinBlocks == nullptr)
       return;
+  }
 
   D->addAttr(::new (Context) CUDALaunchBoundsAttr(
       AttrRange, Context, MaxThreads, MinBlocks, SpellingListIndex));
 }
 
 static void handleLaunchBoundsAttr(Sema &S, Decl *D,
                                    const AttributeList &Attr) {
   if (!checkAttributeAtLeastNumArgs(S, Attr, 1) ||

test/CodeGenCUDA/launch-bounds.cu

 
 __global__ void
 __launch_bounds__( MAX_THREADS_PER_BLOCK, -MIN_BLOCKS_PER_MP )
 Kernel7()
 {
 }
 // CHECK:     !{{[0-9]+}} = !{void ()* @{{.*}}Kernel7{{.*}}, !"maxntidx",
 // CHECK-NOT: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel7{{.*}}, !"minctasm",
+
+const int max_threads_per_block = 11;
+const int min_blocks_per_mp = 12;
+__global__ void __launch_bounds__(max_threads_per_block, min_blocks_per_mp)
+    Kernel8() {}
+// CHECK:     !{{[0-9]+}} = !{void ()* @{{.*}}Kernel8{{.*}}, !"maxntidx", i32 11
+// CHECK:     !{{[0-9]+}} = !{void ()* @{{.*}}Kernel8{{.*}}, !"minctasm", i32 12

test/SemaCUDA/pr27778.cu

+// RUN: %clang_cc1 -fsyntax-only %s
+
+#include "Inputs/cuda.h"
+
+const int constint = 512;
+__launch_bounds__(constint) void TestConstInt(void) {}