[PATCH] D137251: [clang][cuda/hip] Allow `__noinline__` lambdas

[Pierre van Houtryve via Phabricator via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Pierre-vh marked 3 inline comments as done.
Closed by commit rGc05f1639f7f4: [clang][cuda/hip] Allow `__noinline__`  lambdas (authored by Pierre-vh).

Changed prior to commit:
  https://reviews.llvm.org/D137251?vs=472875&id=473145#toc

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D137251/new/

https://reviews.llvm.org/D137251

Files:
  clang/docs/ReleaseNotes.rst
  clang/lib/Parse/ParseExprCXX.cpp
  clang/test/CodeGenCUDA/lambda-noinline.cu
  clang/test/Parser/lambda-attr.cu


Index: clang/test/Parser/lambda-attr.cu
===================================================================
--- clang/test/Parser/lambda-attr.cu
+++ clang/test/Parser/lambda-attr.cu
@@ -18,6 +18,10 @@
   ([&](int) __attribute__((device)){ device_fn(); })(0);
   // expected-warning at -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() {
@@ -37,6 +41,11 @@
   // expected-warning at -1 {{nvcc does not allow '__host__' to appear after the parameter list in lambdas}}
   // expected-warning at -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.
Index: clang/test/CodeGenCUDA/lambda-noinline.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/lambda-noinline.cu
@@ -0,0 +1,23 @@
+// 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; })(); }
Index: clang/lib/Parse/ParseExprCXX.cpp
===================================================================
--- clang/lib/Parse/ParseExprCXX.cpp
+++ clang/lib/Parse/ParseExprCXX.cpp
@@ -1291,7 +1291,22 @@
   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
Index: clang/docs/ReleaseNotes.rst
===================================================================
--- clang/docs/ReleaseNotes.rst
+++ clang/docs/ReleaseNotes.rst
@@ -638,6 +638,9 @@
 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
 -------------------------------------
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D137251.473145.patch
Type: text/x-patch
Size: 4003 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20221104/1fb4c850/attachment-0001.bin>