[clang] afc9d67 - [CUDA][HIP] support __noinline__ as keyword
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Tue May 10 11:34:31 PDT 2022
Author: Yaxun (Sam) Liu
Date: 2022-05-10T14:32:27-04:00
New Revision: afc9d674fe5a14b95c50a38d8605a159c2460427
URL: https://github.com/llvm/llvm-project/commit/afc9d674fe5a14b95c50a38d8605a159c2460427
DIFF: https://github.com/llvm/llvm-project/commit/afc9d674fe5a14b95c50a38d8605a159c2460427.diff
LOG: [CUDA][HIP] support __noinline__ as keyword
CUDA/HIP programs use __noinline__ like a keyword e.g.
__noinline__ void foo() {} since __noinline__ is defined
as a macro __attribute__((noinline)) in CUDA/HIP runtime
header files.
However, gcc and clang supports __attribute__((__noinline__))
the same as __attribute__((noinline)). Some C++ libraries
use __attribute__((__noinline__)) in their header files.
When CUDA/HIP programs include such header files,
clang will emit error about invalid attributes.
This patch fixes this issue by supporting __noinline__ as
a keyword, so that CUDA/HIP runtime could remove
the macro definition.
Reviewed by: Aaron Ballman, Artem Belevich
Differential Revision: https://reviews.llvm.org/D124866
Added:
clang/test/CodeGenCUDA/noinline.cu
clang/test/Lexer/has_feature.cu
clang/test/SemaCUDA/noinline.cu
Modified:
clang/docs/ReleaseNotes.rst
clang/include/clang/Basic/Attr.td
clang/include/clang/Basic/AttrDocs.td
clang/include/clang/Basic/Features.def
clang/include/clang/Basic/TokenKinds.def
clang/include/clang/Parse/Parser.h
clang/lib/Basic/IdentifierTable.cpp
clang/lib/Parse/ParseDecl.cpp
Removed:
################################################################################
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index c3605ab9bda3..ff7b4285b7d7 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -351,8 +351,11 @@ C++2b Feature Support
- Implemented `P0849R8: auto(x): decay-copy in the language <https://wg21.link/P0849R8>`_.
- Implemented `P2242R3: Non-literal variables (and labels and gotos) in constexpr functions <https://wg21.link/P2242R3>`_.
-CUDA Language Changes in Clang
-------------------------------
+CUDA/HIP Language Changes in Clang
+----------------------------------
+
+- Added `__noinline__` as a keyword to avoid diagnostics due to usage of
+ `__attribute__((__noinline__))` in CUDA/HIP programs.
Objective-C Language Changes in Clang
-------------------------------------
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index eb896fe210f9..39359f414ae7 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1775,8 +1775,9 @@ def Convergent : InheritableAttr {
}
def NoInline : DeclOrStmtAttr {
- let Spellings = [GCC<"noinline">, CXX11<"clang", "noinline">,
- C2x<"clang", "noinline">, Declspec<"noinline">];
+ let Spellings = [Keyword<"__noinline__">, GCC<"noinline">,
+ CXX11<"clang", "noinline">, C2x<"clang", "noinline">,
+ Declspec<"noinline">];
let Accessors = [Accessor<"isClangNoInline", [CXX11<"clang", "noinline">,
C2x<"clang", "noinline">]>];
let Documentation = [NoInlineDocs];
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 02b6031fd5aa..c7ef52f67afa 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -538,6 +538,10 @@ spellings of the attribute are not supported on statements. If a statement is
marked ``[[clang::noinline]]`` and contains calls, those calls inside the
statement will not be inlined by the compiler.
+``__noinline__`` can be used as a keyword in CUDA/HIP languages. This is to
+avoid diagnostics due to usage of ``__attribute__((__noinline__))``
+with ``__noinline__`` defined as a macro as ``__attribute__((noinline))``.
+
.. code-block:: c
int example(void) {
diff --git a/clang/include/clang/Basic/Features.def b/clang/include/clang/Basic/Features.def
index fbaa6174f9bf..c3f3fe79770d 100644
--- a/clang/include/clang/Basic/Features.def
+++ b/clang/include/clang/Basic/Features.def
@@ -270,5 +270,8 @@ EXTENSION(cxx_attributes_on_using_declarations, LangOpts.CPlusPlus11)
FEATURE(cxx_abi_relative_vtable, LangOpts.CPlusPlus && LangOpts.RelativeCXXABIVTables)
+// CUDA/HIP Features
+FEATURE(cuda_noinline_keyword, LangOpts.CUDA)
+
#undef EXTENSION
#undef FEATURE
diff --git a/clang/include/clang/Basic/TokenKinds.def b/clang/include/clang/Basic/TokenKinds.def
index 093389615c26..7b65a1537805 100644
--- a/clang/include/clang/Basic/TokenKinds.def
+++ b/clang/include/clang/Basic/TokenKinds.def
@@ -599,6 +599,9 @@ KEYWORD(pipe , KEYOPENCLC | KEYOPENCLCXX)
// C++ for OpenCL s2.3.1: addrspace_cast operator
KEYWORD(addrspace_cast , KEYOPENCLCXX)
+// CUDA/HIP function attributes
+KEYWORD(__noinline__ , KEYCUDA)
+
// OpenMP Type Traits
UNARY_EXPR_OR_TYPE_TRAIT(__builtin_omp_required_simd_align, OpenMPRequiredSimdAlign, KEYALL)
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index 99fe375f9145..caa58d926062 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -2827,6 +2827,7 @@ class Parser : public CodeCompletionHandler {
void ParseOpenCLKernelAttributes(ParsedAttributes &attrs);
void ParseOpenCLQualifiers(ParsedAttributes &Attrs);
void ParseNullabilityTypeSpecifiers(ParsedAttributes &attrs);
+ void ParseCUDAFunctionAttributes(ParsedAttributes &attrs);
VersionTuple ParseVersionTuple(SourceRange &Range);
void ParseAvailabilityAttribute(IdentifierInfo &Availability,
diff --git a/clang/lib/Basic/IdentifierTable.cpp b/clang/lib/Basic/IdentifierTable.cpp
index b86cb7af69bd..af19de44bed8 100644
--- a/clang/lib/Basic/IdentifierTable.cpp
+++ b/clang/lib/Basic/IdentifierTable.cpp
@@ -108,6 +108,7 @@ namespace {
KEYOPENCLCXX = 0x400000,
KEYMSCOMPAT = 0x800000,
KEYSYCL = 0x1000000,
+ KEYCUDA = 0x2000000,
KEYALLCXX = KEYCXX | KEYCXX11 | KEYCXX20,
KEYALL = (0x1ffffff & ~KEYNOMS18 &
~KEYNOOPENCL) // KEYNOMS18 and KEYNOOPENCL are used to exclude.
@@ -158,6 +159,8 @@ static KeywordStatus getKeywordStatus(const LangOptions &LangOpts,
return KS_Future;
if (LangOpts.isSYCL() && (Flags & KEYSYCL))
return KS_Enabled;
+ if (LangOpts.CUDA && (Flags & KEYCUDA))
+ return KS_Enabled;
return KS_Disabled;
}
diff --git a/clang/lib/Parse/ParseDecl.cpp b/clang/lib/Parse/ParseDecl.cpp
index 8d6e84b924c3..89e13cf5c90f 100644
--- a/clang/lib/Parse/ParseDecl.cpp
+++ b/clang/lib/Parse/ParseDecl.cpp
@@ -897,6 +897,15 @@ void Parser::ParseOpenCLKernelAttributes(ParsedAttributes &attrs) {
}
}
+void Parser::ParseCUDAFunctionAttributes(ParsedAttributes &attrs) {
+ while (Tok.is(tok::kw___noinline__)) {
+ IdentifierInfo *AttrName = Tok.getIdentifierInfo();
+ SourceLocation AttrNameLoc = ConsumeToken();
+ attrs.addNew(AttrName, AttrNameLoc, nullptr, AttrNameLoc, nullptr, 0,
+ ParsedAttr::AS_Keyword);
+ }
+}
+
void Parser::ParseOpenCLQualifiers(ParsedAttributes &Attrs) {
IdentifierInfo *AttrName = Tok.getIdentifierInfo();
SourceLocation AttrNameLoc = Tok.getLocation();
@@ -3690,6 +3699,11 @@ void Parser::ParseDeclarationSpecifiers(DeclSpec &DS,
ParseOpenCLKernelAttributes(DS.getAttributes());
continue;
+ // CUDA/HIP single token adornments.
+ case tok::kw___noinline__:
+ ParseCUDAFunctionAttributes(DS.getAttributes());
+ continue;
+
// Nullability type specifiers.
case tok::kw__Nonnull:
case tok::kw__Nullable:
diff --git a/clang/test/CodeGenCUDA/noinline.cu b/clang/test/CodeGenCUDA/noinline.cu
new file mode 100644
index 000000000000..41e8231246b2
--- /dev/null
+++ b/clang/test/CodeGenCUDA/noinline.cu
@@ -0,0 +1,34 @@
+// Uses -O2 since the defalt -O0 option adds noinline to all functions.
+
+// RUN: %clang_cc1 -triple nvptx-nvidia-cuda -fcuda-is-device \
+// RUN: -O2 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN: -O2 -disable-llvm-passes -emit-llvm -o - -x hip %s | FileCheck %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-gnu-linux \
+// RUN: -O2 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+__noinline__ __device__ __host__ void fun1() {}
+
+__attribute__((noinline)) __device__ __host__ void fun2() {}
+
+__attribute__((__noinline__)) __device__ __host__ void fun3() {}
+
+[[gnu::__noinline__]] __device__ __host__ void fun4() {}
+
+#define __noinline__ __attribute__((__noinline__))
+__noinline__ __device__ __host__ void fun5() {}
+
+__device__ __host__ void fun6() {}
+
+// CHECK: define{{.*}}@_Z4fun1v{{.*}}#[[ATTR1:[0-9]*]]
+// CHECK: define{{.*}}@_Z4fun2v{{.*}}#[[ATTR1:[0-9]*]]
+// CHECK: define{{.*}}@_Z4fun3v{{.*}}#[[ATTR1:[0-9]*]]
+// CHECK: define{{.*}}@_Z4fun4v{{.*}}#[[ATTR1:[0-9]*]]
+// CHECK: define{{.*}}@_Z4fun5v{{.*}}#[[ATTR1:[0-9]*]]
+// CHECK: define{{.*}}@_Z4fun6v{{.*}}#[[ATTR2:[0-9]*]]
+// CHECK: attributes #[[ATTR1]] = {{.*}}noinline
+// CHECK-NOT: attributes #[[ATTR2]] = {{.*}}noinline
diff --git a/clang/test/Lexer/has_feature.cu b/clang/test/Lexer/has_feature.cu
new file mode 100644
index 000000000000..5cb653593ecd
--- /dev/null
+++ b/clang/test/Lexer/has_feature.cu
@@ -0,0 +1,8 @@
+// RUN: %clang_cc1 -E -triple x86_64-linux-gnu %s -o - | FileCheck %s
+
+// CHECK: has_noinline_keyword
+#if __has_feature(cuda_noinline_keyword)
+int has_noinline_keyword();
+#else
+int no_noinine_keyword();
+#endif
diff --git a/clang/test/SemaCUDA/noinline.cu b/clang/test/SemaCUDA/noinline.cu
new file mode 100644
index 000000000000..bd963436c213
--- /dev/null
+++ b/clang/test/SemaCUDA/noinline.cu
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -fsyntax-only -verify=cuda %s
+// RUN: %clang_cc1 -fsyntax-only -verify=cuda -pedantic %s
+// RUN: %clang_cc1 -fsyntax-only -verify=cpp -x c++ %s
+
+// cuda-no-diagnostics
+
+__noinline__ void fun1() { } // cpp-error {{unknown type name '__noinline__'}}
+
+__attribute__((noinline)) void fun2() { }
+__attribute__((__noinline__)) void fun3() { }
+[[gnu::__noinline__]] void fun4() { }
+
+#define __noinline__ __attribute__((__noinline__))
+__noinline__ void fun5() {}
+
+#undef __noinline__
+#10 "cuda.h" 3
+#define __noinline__ __attribute__((__noinline__))
+__noinline__ void fun6() {}
More information about the cfe-commits
mailing list