yaxunl updated this revision to Diff 426794.
yaxunl added a comment.
add feature cuda_noinline_keyword to facilitate CUDA/HIP headers removing
__noinline__ macro
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D124866/new/
https://reviews.llvm.org/D124866
Files:
clang/include/clang/Basic/Attr.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
clang/test/CodeGenCUDA/noinline.cu
clang/test/Lexer/has_feature.cu
clang/test/SemaCUDA/noinline.cu
Index: clang/test/SemaCUDA/noinline.cu
===================================================================
--- /dev/null
+++ clang/test/SemaCUDA/noinline.cu
@@ -0,0 +1,8 @@
+// RUN: %clang_cc1 -fsyntax-only -verify=cuda %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() { }
Index: clang/test/Lexer/has_feature.cu
===================================================================
--- /dev/null
+++ 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
Index: clang/test/CodeGenCUDA/noinline.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/noinline.cu
@@ -0,0 +1,27 @@
+// optimization is needed, otherwise by default all functions have noinline.
+
+// RUN: %clang_cc1 -triple nvptx-nvidia-cuda -fcuda-is-device \
+// RUN: -O2 -emit-llvm -o - %s | FileCheck %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN: -O2 -emit-llvm -o - -x hip %s | FileCheck %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-gnu-linux \
+// RUN: -O2 -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() {}
+
+__device__ __host__ void fun4() {}
+
+// CHECK: define{{.*}}@_Z4fun1v{{.*}}#[[ATTR1:[0-9]*]]
+// CHECK: define{{.*}}@_Z4fun2v{{.*}}#[[ATTR1:[0-9]*]]
+// CHECK: define{{.*}}@_Z4fun3v{{.*}}#[[ATTR1:[0-9]*]]
+// CHECK: define{{.*}}@_Z4fun4v{{.*}}#[[ATTR2:[0-9]*]]
+// CHECK: attributes #[[ATTR1]] = {{.*}}noinline
+// CHECK-NOT: attributes #[[ATTR2]] = {{.*}}noinline
Index: clang/lib/Parse/ParseDecl.cpp
===================================================================
--- clang/lib/Parse/ParseDecl.cpp
+++ clang/lib/Parse/ParseDecl.cpp
@@ -897,6 +897,15 @@
}
}
+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 @@
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:
Index: clang/lib/Basic/IdentifierTable.cpp
===================================================================
--- clang/lib/Basic/IdentifierTable.cpp
+++ clang/lib/Basic/IdentifierTable.cpp
@@ -108,6 +108,7 @@
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 @@
return KS_Future;
if (LangOpts.isSYCL() && (Flags & KEYSYCL))
return KS_Enabled;
+ if (LangOpts.CUDA && (Flags & KEYCUDA))
+ return KS_Enabled;
return KS_Disabled;
}
Index: clang/include/clang/Parse/Parser.h
===================================================================
--- clang/include/clang/Parse/Parser.h
+++ clang/include/clang/Parse/Parser.h
@@ -2824,6 +2824,7 @@
void ParseOpenCLKernelAttributes(ParsedAttributes &attrs);
void ParseOpenCLQualifiers(ParsedAttributes &Attrs);
void ParseNullabilityTypeSpecifiers(ParsedAttributes &attrs);
+ void ParseCUDAFunctionAttributes(ParsedAttributes &attrs);
VersionTuple ParseVersionTuple(SourceRange &Range);
void ParseAvailabilityAttribute(IdentifierInfo &Availability,
Index: clang/include/clang/Basic/TokenKinds.def
===================================================================
--- clang/include/clang/Basic/TokenKinds.def
+++ clang/include/clang/Basic/TokenKinds.def
@@ -599,6 +599,9 @@
// 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)
Index: clang/include/clang/Basic/Features.def
===================================================================
--- clang/include/clang/Basic/Features.def
+++ clang/include/clang/Basic/Features.def
@@ -270,5 +270,8 @@
FEATURE(cxx_abi_relative_vtable, LangOpts.CPlusPlus && LangOpts.RelativeCXXABIVTables)
+// CUDA/HIP Features
+FEATURE(cuda_noinline_keyword, true)
+
#undef EXTENSION
#undef FEATURE
Index: clang/include/clang/Basic/Attr.td
===================================================================
--- clang/include/clang/Basic/Attr.td
+++ clang/include/clang/Basic/Attr.td
@@ -1775,7 +1775,7 @@
}
def NoInline : DeclOrStmtAttr {
- let Spellings = [GCC<"noinline">, CXX11<"clang", "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">]>];
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
