[clang][cuda/hip] Allow `__noinline__` lambdas

[Pierre-vh]

Summary

D124866 seem to have had an unintended side effect: noinline on lambdas was no longer accepted.

This fixes the regression and adds a test case for it.

[yaxunl]

need a CodeGenCUDA test

[tra]

LGTM in principle. Please wait for @rsmith's OK.

clang/lib/Parse/ParseExprCXX.cpp

1297–1308

Nit. I'd rearrange it a bit.

while (true) {
  if (Tok.is(tok::kw___noinline__)) {
    IdentifierInfo *AttrName = Tok.getIdentifierInfo();
    SourceLocation AttrNameLoc = ConsumeToken();
    Attr.addNew(AttrName, AttrNameLoc, nullptr, AttrNameLoc, nullptr, 0,
                ParsedAttr::AS_Keyword);
  } else if (Tok.is(tok::kw___attribute))
    ParseGNUAttributes(Attr, nullptr, &D);
  else 
    break;
}

[aaron.ballman]

Can you also add a release note for the fix?

clang/lib/Parse/ParseExprCXX.cpp

1300

Any other keyword attributes that are missing?

alignas/_Alignas
__forceinline
__cdecl/__stdcall/etc

[rsmith]

In D137251#3902835, @tra wrote:

LGTM in principle. Please wait for @rsmith's OK.

I'm happy to defer to @aaron.ballman on this :)

[Pierre-vh]

Comments

Not sure if the release note is in the right place though.
As for the test, I did something quite targeted/minimal, hope it's fine?

clang/lib/Parse/ParseExprCXX.cpp

1300

I'm not too familiar with how those attributes work yet, so maybe there's more to handle but I don't have any concrete example of it being an issue and would rather not touch those unless needed

[yaxunl]

clang/lib/Parse/ParseExprCXX.cpp

1300

CUDA only allows __attribute__ for lambda. __noinline__ used to be an attribute. That's why people use it with lambda. I don't think people use other keywords with lambda.

[yaxunl]

In D137251#3904618, @Pierre-vh wrote:

Comments

Not sure if the release note is in the right place though.
As for the test, I did something quite targeted/minimal, hope it's fine?

LGTM. Thanks.

[aaron.ballman]

LGTM aside from some very minor nits.

clang/docs/ReleaseNotes.rst

641

785–786

Unrelated whitespace changes that can be rolled back.

clang/lib/Parse/ParseExprCXX.cpp

1300

CUDA only allows attribute for lambda. noinline used to be an attribute. That's why people use it with lambda. I don't think people use other keywords with lambda.

Ahhh, thank you, that helps!

clang/docs/ReleaseNotes.rst

 - Support label at end of compound statement (`P2324 <https://wg21.link/p2324r2>`_).
 - Implemented `P1169R4: static operator() <https://wg21.link/P1169R4>`_.
 - Implemented "char8_t Compatibility and Portability Fix" (`P2513R3 <https://wg21.link/P2513R3>`_).
   This Change was applied to C++20 as a Defect Report.
 
 CUDA/HIP Language Changes in Clang
 ----------------------------------
 
+ - Allow the use of ``__noinline__`` as a keyword (instead of ``__attribute__((noinline))``)
+   in lambda declarations.
+
 Objective-C Language Changes in Clang
 -------------------------------------
 
 OpenCL C Language Changes in Clang
 ----------------------------------
 
 ...
 

 - Introduced the new function ``clang_getNonReferenceType``, which mimics
   the behavior of ``QualType::getNonReferenceType`` for ``CXType``.
 - Introduced the new function ``clang_CXXMethod_isDeleted``, which queries
   whether the method is declared ``= delete``.
 - Introduced the new function ``clang_CXXMethod_isCopyAssignmentOperator``,
   which identifies whether a method cursor is a copy-assignment
   operator.
 - ``clang_Cursor_getNumTemplateArguments``, ``clang_Cursor_getTemplateArgumentKind``, 
   ``clang_Cursor_getTemplateArgumentType``, ``clang_Cursor_getTemplateArgumentValue`` and 
   ``clang_Cursor_getTemplateArgumentUnsignedValue`` now work on struct, class,
   and partial template specialization cursors in addition to function cursors.
 
 Static Analyzer
 ---------------
 - Removed the deprecated ``-analyzer-store`` and
   ``-analyzer-opt-analyze-nested-blocks`` analyzer flags.
   ``scanbuild`` was also updated accordingly.
   Passing these flags will result in a hard error.

clang/lib/Parse/ParseExprCXX.cpp

ExprResult Parser::ParseLambdaExpressionAfterIntroducer(
   Declarator D(DS, ParsedAttributesView::none(), DeclaratorContext::LambdaExpr);
   TemplateParameterDepthRAII CurTemplateDepthTracker(TemplateParameterDepth);
   Actions.PushLambdaScope();
 
   ParsedAttributes Attr(AttrFactory);
   if (getLangOpts().CUDA) {
     // In CUDA code, GNU attributes are allowed to appear immediately after the
     // "[...]", even if there is no "(...)" before the lambda body.
-    MaybeParseGNUAttributes(D);
+    //
+    // Note that we support __noinline__ as a keyword in this mode and thus
+    // it has to be separately handled.
+    while (true) {
+      if (Tok.is(tok::kw___noinline__)) {
+        IdentifierInfo *AttrName = Tok.getIdentifierInfo();
+        SourceLocation AttrNameLoc = ConsumeToken();
+        Attr.addNew(AttrName, AttrNameLoc, nullptr, AttrNameLoc, nullptr, 0,
+                    ParsedAttr::AS_Keyword);
+      } else if (Tok.is(tok::kw___attribute))
+        ParseGNUAttributes(Attr, nullptr, &D);
+      else
+        break;
+    }
+
+    D.takeAttributes(Attr);
   }
 
   // Helper to emit a warning if we see a CUDA host/device/global attribute
   // after '(...)'. nvcc doesn't accept this.
   auto WarnIfHasCUDATargetAttr = [&] {
     if (getLangOpts().CUDA)
       for (const ParsedAttr &A : Attr)
         if (A.getKind() == ParsedAttr::AT_CUDADevice ||

clang/test/CodeGenCUDA/lambda-noinline.cu

+// RUN: %clang_cc1 -no-opaque-pointers -x hip -emit-llvm -std=c++11 %s -o - \
+// RUN:   -triple x86_64-linux-gnu \
+// RUN:   | FileCheck -check-prefix=HOST %s
+// RUN: %clang_cc1 -no-opaque-pointers -x hip -emit-llvm -std=c++11 %s -o - \
+// RUN:   -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:   | FileCheck -check-prefix=DEV %s
+
+#include "Inputs/cuda.h"
+
+// Checks noinline is correctly added to the lambda function.
+
+// HOST: define{{.*}}@_ZZ4HostvENKUlvE_clEv({{.*}}) #[[ATTR:[0-9]+]]
+// HOST: attributes #[[ATTR]]{{.*}}noinline
+
+// DEV: define{{.*}}@_ZZ6DevicevENKUlvE_clEv({{.*}}) #[[ATTR:[0-9]+]]
+// DEV: attributes #[[ATTR]]{{.*}}noinline
+
+__device__ int a;
+int b;
+
+__device__ int Device() { return ([&] __device__ __noinline__ (){ return a; })(); }
+
+__host__ int Host() { return ([&] __host__ __noinline__ (){ return b; })(); }

clang/test/Parser/lambda-attr.cu

__attribute__((device)) void device_attr() {
   ([&]() __attribute__((device)){ device_fn(); })();
   // expected-warning@-1 {{nvcc does not allow '__device__' to appear after the parameter list in lambdas}}
   ([&] __attribute__((device)) () { device_fn(); })();
   ([&] __attribute__((device)) { device_fn(); })();
 
   ([&](int) __attribute__((device)){ device_fn(); })(0);
   // expected-warning@-1 {{nvcc does not allow '__device__' to appear after the parameter list in lambdas}}
   ([&] __attribute__((device)) (int) { device_fn(); })(0);
+
+  // test that noinline can appear anywhere.
+  ([&] __attribute__((device)) __noinline__ () { device_fn(); })();
+  ([&] __noinline__ __attribute__((device)) () { device_fn(); })();
 }
 
 __attribute__((host)) __attribute__((device)) void host_device_attrs() {
   ([]() __attribute__((host)) __attribute__((device)){ hd_fn(); })();
   // expected-warning@-1 {{nvcc does not allow '__host__' to appear after the parameter list in lambdas}}
   // expected-warning@-2 {{nvcc does not allow '__device__' to appear after the parameter list in lambdas}}
   ([] __attribute__((host)) __attribute__((device)) () { hd_fn(); })();
   ([] __attribute__((host)) __attribute__((device)) { hd_fn(); })();
 
   ([&]() __attribute__((host)) __attribute__((device)){ hd_fn(); })();
   // expected-warning@-1 {{nvcc does not allow '__host__' to appear after the parameter list in lambdas}}
   // expected-warning@-2 {{nvcc does not allow '__device__' to appear after the parameter list in lambdas}}
   ([&] __attribute__((host)) __attribute__((device)) () { hd_fn(); })();
   ([&] __attribute__((host)) __attribute__((device)) { hd_fn(); })();
 
   ([&](int) __attribute__((host)) __attribute__((device)){ hd_fn(); })(0);
   // expected-warning@-1 {{nvcc does not allow '__host__' to appear after the parameter list in lambdas}}
   // expected-warning@-2 {{nvcc does not allow '__device__' to appear after the parameter list in lambdas}}
   ([&] __attribute__((host)) __attribute__((device)) (int) { hd_fn(); })(0);
+
+  // test that noinline can also appear anywhere.
+  ([] __attribute__((host)) __attribute__((device)) () { hd_fn(); })();
+  ([] __attribute__((host)) __noinline__ __attribute__((device)) () { hd_fn(); })();
+  ([] __attribute__((host)) __attribute__((device)) __noinline__ () { hd_fn(); })();
 }
 
 // TODO: Add tests for __attribute__((global)) once we support global lambdas.