Pierre-vh updated this revision to Diff 472875.
Pierre-vh marked 2 inline comments as done.
Pierre-vh added a comment.
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?
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@-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@-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.
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
@@ -615,6 +615,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
-------------------------------------
@@ -751,8 +754,8 @@
- 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_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.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
