[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 ↗

(On Diff #59611)

Should we update this name and comment?

[tra]

Addressed Justin's comments.

[jlebar]

lib/Sema/SemaDeclAttr.cpp

4044 ↗

(On Diff #59624)

Nit, Oxford comma helps some here.

4046 ↗

(On Diff #59624)

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

4079 ↗

(On Diff #59624)

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

4046 ↗

(On Diff #59624)

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.

4079 ↗

(On Diff #59624)

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

4079 ↗

(On Diff #59631)

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

[tra]

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

cfe/trunk/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;
 
Diag(Attr.getLoc(), diag::err_attribute_regparm_invalid_number)
       << Context.getTargetInfo().getRegParmMax() << NumParamsExpr->getSourceRange();
     Attr.setInvalid();
     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,
-                                      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,
   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))
+  if (MinBlocks) {
+    MinBlocks = makeLaunchBoundsArgExpr(*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) ||

cfe/trunk/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 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

cfe/trunk/test/SemaCUDA/pr27778.cu

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