<div dir="ltr">Hi,<br><div class="gmail_extra"><br><div class="gmail_quote">On Mon, Apr 13, 2015 at 9:47 AM, Aaron Ballman <span dir="ltr"><<a href="mailto:aaron@aaronballman.com" target="_blank">aaron@aaronballman.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex"><div><div><span style="color:rgb(34,34,34)">> Index: include/clang/Basic/Attr.td</span><br></div></div>
> ===================================================================<br>
> --- include/clang/Basic/Attr.td<br>
> +++ include/clang/Basic/Attr.td<br>
> @@ -581,7 +581,7 @@<br>
><br>
> def CUDALaunchBounds : InheritableAttr {<br>
> let Spellings = [GNU<"launch_bounds">];<br>
> - let Args = [IntArgument<"MaxThreads">, DefaultIntArgument<"MinBlocks", 0>];<br>
> + let Args = [ExprArgument<"MaxThreads">, ExprArgument<"MinBlocks", 1>];<br>
<br>
On the wish list for someday: a ConstIntExprArgument that does<br>
boilerplate for you. I don't expect that for your patch, but boy, it<br>
sure would be nice. ;-)<br></blockquote><div><br></div><div>Indeed. Alas, I've got to jump through the hoops for now.</div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex">
> Index: lib/CodeGen/TargetInfo.cpp<br>
> ===================================================================<br>
> --- lib/CodeGen/TargetInfo.cpp<br>
> +++ lib/CodeGen/TargetInfo.cpp<br>
> @@ -5079,16 +5079,24 @@<br>
> }<br>
> if (FD->hasAttr<CUDALaunchBoundsAttr>()) {<br>
<span>> // Create !{<func-ref>, metadata !"maxntidx", i32 <val>} node<br>
> - addNVVMMetadata(F, "maxntidx",<br>
> - FD->getAttr<CUDALaunchBoundsAttr>()->getMaxThreads());<br>
> - // min blocks is a default argument for CUDALaunchBoundsAttr, so getting a<br>
> - // zero value from getMinBlocks either means it was not specified in<br>
> - // __launch_bounds__ or the user specified a 0 value. In both cases, we<br>
> + CUDALaunchBoundsAttr *Attr = FD->getAttr<CUDALaunchBoundsAttr>();<br>
<br>
</span>Push this up into the if block instead of checking hasAttr and then<br>
calling getAttr.<br>
<br></blockquote><div> </div><div>Done.</div><div><br></div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex">
> + llvm::APSInt MaxThreads;<br>
> + if (Attr->getMaxThreads()->EvaluateAsInt(MaxThreads, M.getContext()))<br>
> + addNVVMMetadata(F, "maxntidx", MaxThreads.getExtValue());<br>
<br>
Are a negative number of max threads allowed? I noticed that -1 seems<br>
to get passed through unhindered.<br></blockquote><div><br></div><div>checkUInt32Argument in the original code casts it to uint32_t so -1 ends up being 0xffffffff.</div><div>Nvidia's compiler silently ignores non-positive arguments. Neither behavior is ideal, IMO.</div><div><br></div><div>I guess it would be more appropriate to issue a warning if an argument is negative but ignore it otherwise.</div><div> <br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex">
<br>
> + else<br>
> + llvm_unreachable("launch_bounds arg 1 evaluation failed.");<br>
<span>> +<br>
> + // min blocks is a default argument for CUDALaunchBoundsAttr. If it was<br>
> + // not specified in __launch_bounds__ or the user specified a 0 value, we<br>
</span>> // don't have to add a PTX directive.<br>
> - int MinCTASM = FD->getAttr<CUDALaunchBoundsAttr>()->getMinBlocks();<br>
> - if (MinCTASM > 0) {<br>
> - // Create !{<func-ref>, metadata !"minctasm", i32 <val>} node<br>
> - addNVVMMetadata(F, "minctasm", MinCTASM);<br>
> + if (Attr->getMinBlocks()) {<br>
> + llvm::APSInt MinBlocks;<br>
> + if (Attr->getMinBlocks()->EvaluateAsInt(MinBlocks, M.getContext())) {<br>
> + if (MinBlocks > 0)<br>
> + // Create !{<func-ref>, metadata !"minctasm", i32 <val>} node<br>
> + addNVVMMetadata(F, "minctasm", MinBlocks.getExtValue());<br>
<br>
Silently disallowing negative or zero values is unkind; that should<br>
really be a warning with a test case.<br>
<br></blockquote><div> </div><div><div style="color:rgb(34,34,34);font-family:arial,sans-serif;font-size:small;font-style:normal;font-variant:normal;font-weight:normal;letter-spacing:normal;line-height:normal;text-align:start;text-indent:0px;text-transform:none;white-space:normal;word-spacing:0px;background-color:rgb(255,255,255)">I've updated the patch to warn on negative values and disable PTX directive emission for non-positive values.</div><br></div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex">
> + } else<br>
> + llvm_unreachable("launch_bounds arg 2 evaluation failed.");<br>
> }<br>
> }<br>
> }<br>
> Index: lib/Sema/SemaDeclAttr.cpp<br>
> ===================================================================<br>
> --- lib/Sema/SemaDeclAttr.cpp<br>
> +++ lib/Sema/SemaDeclAttr.cpp<br>
> @@ -3457,20 +3457,38 @@<br>
> return false;<br>
> }<br>
><br>
> +static bool isAcceptableLaunchBoundsArgument(Sema &S, Expr *E) {<br>
> + return E->getType()->isIntegerType() &&<br>
> + !E->containsUnexpandedParameterPack() &&<br>
> + (E->isInstantiationDependent() || E->isEvaluatable(S.Context));<br>
> +}<br>
> +<br>
> static void handleLaunchBoundsAttr(Sema &S, Decl *D,<br>
> const AttributeList &Attr) {<br>
> - uint32_t MaxThreads, MinBlocks = 0;<br>
> - if (!checkUInt32Argument(S, Attr, Attr.getArgAsExpr(0), MaxThreads, 1))<br>
> + if (!checkAttributeAtLeastNumArgs(S, Attr, 1) ||<br>
> + !checkAttributeAtMostNumArgs(S, Attr, 2))<br>
> return;<br>
> - if (Attr.getNumArgs() > 1 && !checkUInt32Argument(S, Attr,<br>
> - Attr.getArgAsExpr(1),<br>
> - MinBlocks, 2))<br>
<br>
One thing the old code did that the new code does not is test the<br>
validity of the value (for instance, that it's 32-bit instead of<br>
64-bit) and diagnose.<br></blockquote><div><br></div><div>Hmm. Nvidia docs don't seem to specify any particular limit on the values. I guess c<span style="color:rgb(34,34,34);font-family:arial,sans-serif;font-size:small;font-style:normal;font-variant:normal;font-weight:normal;letter-spacing:normal;line-height:normal;text-align:start;text-indent:0px;text-transform:none;white-space:normal;word-spacing:0px;float:none;display:inline!important;background-color:rgb(255,255,255)">heckUInt32Argument() was used for convenience. </span>I've changed the code so that it only checks whether it's an integer constant expression. If it has to be limited to any particular value, it can be easily added.</div><div> <br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex">
<br>
> +<br>
> + Expr *MaxThreads = Attr.getArgAsExpr(0);<br>
> + if (!isAcceptableLaunchBoundsArgument(S, MaxThreads)) {<br>
> + S.Diag(Attr.getLoc(), diag::err_attribute_argument_n_type)<br>
> + << Attr.getName() << 0 << AANT_ArgumentIntegerConstant<br>
> + << MaxThreads->getSourceRange();<br>
> + return;<br>
> + }<br>
> +<br>
> + Expr *MinBlocks = Attr.getNumArgs() > 1 ? Attr.getArgAsExpr(1) : nullptr;<br>
> + if (MinBlocks && !isAcceptableLaunchBoundsArgument(S, MinBlocks)) {<br>
> + S.Diag(Attr.getLoc(), diag::err_attribute_argument_n_type)<br>
> + << Attr.getName() << 1 << AANT_ArgumentIntegerConstant<br>
> + << MinBlocks->getSourceRange();<br>
> return;<br>
> + }<br>
><br>
> - D->addAttr(::new (S.Context)<br>
> - CUDALaunchBoundsAttr(Attr.getRange(), S.Context,<br>
> - MaxThreads, MinBlocks,<br>
> - Attr.getAttributeSpellingListIndex()));<br>
> + D->addAttr(::new (S.Context) CUDALaunchBoundsAttr(<br>
> + Attr.getRange(), S.Context, MaxThreads, MinBlocks,<br>
> + Attr.getAttributeSpellingListIndex()));<br>
> + return;<br>
> }<br>
><br>
> static void handleArgumentWithTypeTagAttr(Sema &S, Decl *D,<br>
> Index: lib/Sema/SemaTemplateInstantiateDecl.cpp<br>
> ===================================================================<br>
> --- lib/Sema/SemaTemplateInstantiateDecl.cpp<br>
> +++ lib/Sema/SemaTemplateInstantiateDecl.cpp<br>
> @@ -202,6 +202,37 @@<br>
> New->addAttr(EIA);<br>
<span>> }<br>
><br>
> +static void instantiateDependentCUDALaunchBoundsAttr(<br>
> + Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,<br>
</span>> + const CUDALaunchBoundsAttr *A, Decl *New) {<br>
> +<br>
> + // LaunchBounds arguments are constant expressions<br>
> + EnterExpressionEvaluationContext Unevaluated(S, Sema::ConstantEvaluated);<br>
> + ExprResult MaxThreadsResult = S.SubstExpr(A->getMaxThreads(), TemplateArgs);<br>
> + if (MaxThreadsResult.isInvalid()) {<br>
> + S.Diag(A->getLocation(), diag::err_attribute_argument_n_type)<br>
> + << A->getSpelling() << 0 << AANT_ArgumentIntegerConstant<br>
> + << A->getMaxThreads()->getSourceRange();<br>
<br>
You should be able to just pass A instead of callign getSpelling on it<br>
(this will also properly quote things).<br>
<br></blockquote><div><br></div><div>Done.</div><div><br></div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex">
> + return;<br>
> + }<br>
> +<br>
> + Expr *MinBlocksExpr = A->getMinBlocks();<br>
> + if (MinBlocksExpr) {<br>
> + ExprResult MinBlocksResult = S.SubstExpr(A->getMinBlocks(), TemplateArgs);<br>
> + if (MinBlocksResult.isInvalid()) {<br>
> + S.Diag(A->getLocation(), diag::err_attribute_argument_n_type)<br>
> + << A->getSpelling() << 1 << AANT_ArgumentIntegerConstant<br>
> + << A->getMinBlocks()->getSourceRange();<br>
<br>
Same here as above.<br>
<br></blockquote><div>Done.</div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex">
> + return;<br>
> + }<br>
> + MinBlocksExpr = MinBlocksResult.getAs<Expr>();<br>
> + }<br>
> +<br>
> + New->addAttr(::new (S.Context) CUDALaunchBoundsAttr(<br>
> + A->getRange(), S.Context, MaxThreadsResult.getAs<Expr>(), MinBlocksExpr,<br>
> + A->getSpellingListIndex()));<br>
> +}<br>
> +<br>
> void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,<br>
> const Decl *Tmpl, Decl *New,<br>
> LateInstantiatedAttrVec *LateAttrs,<br>
> @@ -233,6 +264,13 @@<br>
> continue;<br>
> }<br>
><br>
> + const CUDALaunchBoundsAttr *CUDALaunchBounds =<br>
<span>> + dyn_cast<CUDALaunchBoundsAttr>(TmplAttr);<br>
> + if (CUDALaunchBounds) {<br>
<br>
</span>Fold the dyn_cast into the if statement, use auto.<br>
<br></blockquote><div>Done.</div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex">
> + instantiateDependentCUDALaunchBoundsAttr(*this, TemplateArgs,<br>
> + CUDALaunchBounds, New);<br>
> + continue;<br>
> + }<br>
> // Existing DLL attribute on the instantiation takes precedence.<br>
> if (TmplAttr->getKind() == attr::DLLExport ||<br>
> TmplAttr->getKind() == attr::DLLImport) {<br>
> Index: test/CodeGenCUDA/<a href="http://launch-bounds.cu" target="_blank">launch-bounds.cu</a><br>
> ===================================================================<br>
> --- test/CodeGenCUDA/<a href="http://launch-bounds.cu" target="_blank">launch-bounds.cu</a><br>
> +++ test/CodeGenCUDA/<a href="http://launch-bounds.cu" target="_blank">launch-bounds.cu</a><br>
> @@ -28,3 +28,23 @@<br>
> }<br>
><br>
> // CHECK: !{{[0-9]+}} = !{void ()* @Kernel2, !"maxntidx", i32 256}<br>
> +<br>
> +template <int max_threads_per_block><br>
> +__global__ void<br>
> +__launch_bounds__(max_threads_per_block)<br>
> +Kernel3()<br>
> +{<br>
> +}<br>
> +<br>
> +template void Kernel3<MAX_THREADS_PER_BLOCK>();<br>
> +// CHECK: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel3{{.*}}, !"maxntidx", i32 256}<br>
> +<br>
> +template <int max_threads_per_block, int min_blocks_per_mp><br>
> +__global__ void<br>
> +__launch_bounds__(max_threads_per_block, min_blocks_per_mp)<br>
> +Kernel4()<br>
> +{<br>
> +}<br>
> +template void Kernel4<MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP>();<br>
> +// CHECK: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel4{{.*}}, !"maxntidx", i32 256}<br>
> +// CHECK: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel4{{.*}}, !"minctasm", i32 2}<br>
> Index: test/SemaCUDA/<a href="http://launch_bounds.cu" target="_blank">launch_bounds.cu</a><br>
> ===================================================================<br>
> --- test/SemaCUDA/<a href="http://launch_bounds.cu" target="_blank">launch_bounds.cu</a><br>
> +++ test/SemaCUDA/<a href="http://launch_bounds.cu" target="_blank">launch_bounds.cu</a><br>
> @@ -9,3 +9,9 @@<br>
> __launch_bounds__() void Test4(void); // expected-error {{'launch_bounds' attribute takes at least 1 argument}}<br>
><br>
> int Test5 __launch_bounds__(128, 7); // expected-warning {{'launch_bounds' attribute only applies to functions and methods}}<br>
<span>> +<br>
> +template <int a, int b> __launch_bounds__(a, b) void Test6(void) {}<br>
> +template void Test6<128,7>(void);<br>
</span>> +<br>
> +template <int a> __launch_bounds__(a) void Test7(void) {}<br>
> +template void Test7<128>(void);<br>
><br>
<br>
You're missing tests that use actual expressions.It would also be</blockquote><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex">
good to capture some tests with expressions you would expect to fail.<br>
<span><font color="#888888"><br></font></span></blockquote><div> </div><div>Done and done. </div><div><br></div></div>-- <br><div><div dir="ltr">--Artem Belevich</div></div>
</div></div>