[CUDA][HIP] support __noinline__ as keyword

[yaxunl]

Summary

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.

[yaxunl]

add feature cuda_noinline_keyword to facilitate CUDA/HIP headers removing noinline macro

[aaron.ballman]

I don't know how language extensions come about in CUDA or HIP -- is there an appropriate standards body (or something similar) that's aware of this extension and supports it?

The changes should likely come with a release note entry about the new functionality, and some documentation changes as well.

clang/include/clang/Basic/Attr.td

1777–1780

clang/include/clang/Basic/Features.def

274

Do the CUDA or HIP specs define __noinline__ as a keyword specifically? If not, this isn't a FEATURE, it's an EXTENSION because it's specific to Clang, not the language standard.

clang/lib/Parse/ParseDecl.cpp

902

I think we should we be issuing a pedantic "this is a clang extension" warning here, WDYT?

clang/test/SemaCUDA/noinline.cu

9

I think there should also be a test like:

[[gnu::__noinline__]] void fun4() {}

to verify that the double square bracket syntax also correctly handles this being a keyword now (I expect the test to pass).

[tra]

I don't know how language extensions come about in CUDA or HIP -- is there an appropriate standards body (or something similar) that's aware of this extension and supports it?

Summoning @rsmith for his language lawyer expertise.

[yaxunl]

clang/include/clang/Basic/Attr.td

1777–1780

will do

clang/include/clang/Basic/Features.def

274

CUDA/HIP do not have language spec. In their programming guide, they do not define __noinline__ as a keyword.

Will make it an extension.

clang/lib/Parse/ParseDecl.cpp

902

will do

clang/test/SemaCUDA/noinline.cu

9

will do

[yaxunl]

revised by Aaron's comments

[aaron.ballman]

clang/include/clang/Basic/Features.def

274

CUDA/HIP do not have language spec.

Then what body of people governs changes to the language? Basically, I'm trying to understand whether this patch meets the community requirements for adding an extension: https://clang.llvm.org/get_involved.html#criteria, specifically #4 (though the rest of the points are worth keeping in mind). I don't want to Clang ending up stepping on toes by defining this extension only to accidentally frustrate the CUDA community.

[yaxunl]

clang/include/clang/Basic/Features.def

274

specific to __noinline__, it is largely determined by the existing behaviour of CUDA SDK.

The CUDA SDK defines __noinline__ as a macro __attribute__((noinline)). However, it is not compatible with some C++ headers which use __attribute__((__noinline__)).

This patch will not change the usage pattern of __noinline__. It is equivalent to the original behaviour with the benefit of being compatible with C++ headers.

[yaxunl]

added release note and documentation

[tra]

CUDA/HIP do not have language spec.

Well. It's not completely true. CUDA programming guide does serve as the de-facto spec for CUDA. It's far from perfect, but it does mention __noinline__ and __forceinline__ as function qualifiers: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#noinline-and-forceinline

[yaxunl]

In D124866#3497439, @tra wrote:

CUDA/HIP do not have language spec.

Well. It's not completely true. CUDA programming guide does serve as the de-facto spec for CUDA. It's far from perfect, but it does mention __noinline__ and __forceinline__ as function qualifiers: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#noinline-and-forceinline

Thanks for the pointer. I missed that part.

CUDA SDK implements __noinline__ as attribute __attribute__((noinline)) though. Some requirements may not have diagnostics.

[aaron.ballman]

In D124866#3497439, @tra wrote:

CUDA/HIP do not have language spec.

Well. It's not completely true. CUDA programming guide does serve as the de-facto spec for CUDA. It's far from perfect, but it does mention __noinline__ and __forceinline__ as function qualifiers: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#noinline-and-forceinline

Thank you, that's the magic words I was hoping for -- because they're described as function qualifiers, I think it's justifiable to add them as a keyword implementation in Clang and not worry about stepping on the toes of the CUDA spec (it's adhering to what the current spec requires).

Should we do __forceinline__ at the same time so that there's consistency?

clang/lib/Parse/ParseDecl.cpp

902

I'm questioning whether my advice here was good or not -- now that I see the CUDA spec already calls these function qualifiers... it's debatable whether this is a Clang extension or just the way in which Clang implements the CUDA function qualifiers. @tra -- do you have opinions?

I'm sort of leaning towards dropping the extension warning, but the only reason I can think of for keeping it is if Clang is the only CUDA compiler that doesn't require you to include a header before using the function qualifiers. If that's the case, there is a portability concern.

[yaxunl]

In D124866#3500761, @aaron.ballman wrote:

Should we do __forceinline__ at the same time so that there's consistency?

__forceinline__ does not have the issue as __noinline__ has since it is not a GCC attribute. The current CUDA/HIP implementation of __forceinline__ in header files is sufficient. I do not see the benefit of implementing __forceinline__ as a keyword.

[aaron.ballman]

In D124866#3501181, @yaxunl wrote:

In D124866#3500761, @aaron.ballman wrote:

Should we do __forceinline__ at the same time so that there's consistency?

__forceinline__ does not have the issue as __noinline__ has since it is not a GCC attribute. The current CUDA/HIP implementation of __forceinline__ in header files is sufficient. I do not see the benefit of implementing __forceinline__ as a keyword.

Primarily to reduce user confusion. It's kind of weird for __noinline__ to be a keyword and __forceinline__ to not be a keyword when they're both defined the same way by the CUDA spec. This means you can #undef one of them but not the other, that sort of thing.

clang/test/CodeGenCUDA/noinline.cu

2

I've asked @erichkeane to weigh in on whether there's a better approach here than specifying an optimization level.

clang/test/SemaCUDA/noinline.cu

9

Ah, I just noticed we also have no tests for the behavior of the keyword in the presence of the macro being defined. e.g.,

#define __noinline__ __attribute__((__noinline__))
__noinline__ void fun5() {}

[erichkeane]

clang/test/CodeGenCUDA/noinline.cu

2

You don't need to do this, it looks like all you're trying to do is keep 'clang' out of O0 mode. However, what you do NOT want is the optimizations to run. The common way to do that is to combine O1/O2/etc like: -O2 -disable-llvm-passes

This will keep clang in O2 mode, but will keep the optimizer from running anything, which might mess with the test later on.

[tra]

In D124866#3501203, @aaron.ballman wrote:

Should we do __forceinline__ at the same time so that there's consistency?

Primarily to reduce user confusion. It's kind of weird for __noinline__ to be a keyword and __forceinline__ to not be a keyword when they're both defined the same way by the CUDA spec. This means you can #undef one of them but not the other, that sort of thing.

I'm slightly biased towards making them both a keyword. That said, I may be convinced otherwise if we discover that it may break some assumptions in existing C++ code. I just don't know enough.

clang/lib/Parse/ParseDecl.cpp

902

I'm not sure if such a warning would be useful.

the only reason I can think of for keeping it is if Clang is the only CUDA compiler that doesn't require you to include a header before using the function qualifiers. If that's the case, there is a portability concern.

I don't think it's an issue.

We already have similar divergence between nvcc/clang. E.g. built-in variables like threadIdx. Clang implements them in a header, but NVCC provides them by compiler itself.
With both compilers the variables are available by the time we get to compile user code. Virtually all CUDA compilations are done with tons of CUDA headers pre-included by compiler. Those that do not do that are already on their own and have to provide many other 'standard' CUDA things like target attributes. I don't think we need to worry about that.

[yaxunl]

In D124866#3501203, @aaron.ballman wrote:

__forceinline__ does not have the issue as __noinline__ has since it is not a GCC attribute. The current CUDA/HIP implementation of __forceinline__ in header files is sufficient. I do not see the benefit of implementing __forceinline__ as a keyword.

Primarily to reduce user confusion. It's kind of weird for __noinline__ to be a keyword and __forceinline__ to not be a keyword when they're both defined the same way by the CUDA spec. This means you can #undef one of them but not the other, that sort of thing.

If we are to add __forceinline__ as a keyword, I feel it better be a separate patch to be cleaner.

clang/lib/Parse/ParseDecl.cpp

902

I can remove the diagnostics since it seems unnecessary.

I tend to treat it as an extension since nvcc is the de facto standard implementation, which does not implement it as a keyword. Compared to that, this is like an extension.

clang/test/CodeGenCUDA/noinline.cu

2

will use -O2 -disable-llvm-passes

clang/test/SemaCUDA/noinline.cu

9

will do

[yaxunl]

removed diagnostics and added more tests

[tra]

In D124866#3501641, @yaxunl wrote:

If we are to add __forceinline__ as a keyword, I feel it better be a separate patch to be cleaner.

Fine with me.

clang/lib/Parse/ParseDecl.cpp

902

I'd argue that NVCC does implement it (as in "documents and makes it available"). Providing the documented functionality using a different implementation does not reach the point of being an extension, IMO. While there are observable differences between implementations, depending on them would be a portability error for the user.

[aaron.ballman]

In D124866#3501641, @yaxunl wrote:

If we are to add __forceinline__ as a keyword, I feel it better be a separate patch to be cleaner.

I'm fine with that.

A few nits and a question about the test recently added.

clang/docs/ReleaseNotes.rst

353–358

clang/include/clang/Basic/AttrDocs.td

543

clang/test/SemaCUDA/noinline.cu

9

I missed an important detail -- I think this is now going to generate a warning in -pedantic mode (through -Wkeyword-macro) when compiling for CUDA; is that going to be a problem for CUDA headers, or are those always included as a system header (and so the diagnostics will be suppressed)?

[yaxunl]

clang/docs/ReleaseNotes.rst

353–358

will fix

clang/include/clang/Basic/AttrDocs.td

543

will fix.

clang/lib/Parse/ParseDecl.cpp

902

that makes sense. will change the extension to feature

clang/test/SemaCUDA/noinline.cu

9

I could not find how clang driver adds CUDA include path

https://github.com/llvm/llvm-project/blob/main/clang/lib/Driver/ToolChains/Cuda.cpp#L284

@tra do you know how CUDA include path is added? is it done by CMake?

For HIP the HIP include path is added as a system include path by clang driver.

[aaron.ballman]

clang/test/SemaCUDA/noinline.cu

9

Whatever we find out, we can emulate its behavior here in the test file to see what the diagnostic behavior will be (you can use GNU linemarkers to convince the compiler parts of the source are in a system header).

[yaxunl]

clang/test/SemaCUDA/noinline.cu

9

will add tests for that.

It seems no matter it is system header or normal header, no warnings are emitted even with -pedantic.

[aaron.ballman]

clang/test/SemaCUDA/noinline.cu

9

Excellent, thank you!

[yaxunl]

make it a feature, add tests for pedantic, fix release notes and doecumentation

[tra]

clang/test/SemaCUDA/noinline.cu

9

CUDA includes are added via -internal-isystem here: https://github.com/llvm/llvm-project/blob/main/clang/lib/Driver/ToolChains/Cuda.cpp#L892

[delcypher]

clang/lib/Basic/IdentifierTable.cpp

111

@yaxunl Is it intentional that you didn't update KEYALL here? That means KEYALL doesn't include the bit for KEYCUDA.

If that was your intention then this will break if someone adds a new key. E.g.

KEYCUDA = 0x2000000,
KEYSOMENEWTHING = 0x4000000,
// ...
// KEYALL now includes `KEYCUDA`, whereas it didn't before.
// KEYALL includes KEYSOMENEWTHING 
KEYALL = (0x7ffffff & ~KEYNOMS18 &
              ~KEYNOOPENCL) // KEYNOMS18 and KEYNOOPENCL are used to exclude.
...

1. Updating the 0x1ffffff constant to 0x3ffffff so that KEYALL includes KEYCUDA
2. If your intention is to not have KEYCUDA set in KEYALL then amend KEYALL to be.

KEYALL = (0x7ffffff & ~KEYNOMS18 &
              ~KEYNOOPENCL & ~KEYCUDA ) // KEYNOMS18 and KEYNOOPENCL are used to exclude.
// KEYCUDA is not included in KEYALL

[yaxunl]

clang/lib/Basic/IdentifierTable.cpp

111

My intention is not to include KEYCUDA in KEYALL.

Should I change KEYALL to

KEYALL = (0x3ffffff & ~KEYNOMS18 &
              ~KEYNOOPENCL & ~KEYCUDA ) // KEYNOMS18 and KEYNOOPENCL are used to exclude.
// KEYCUDA is not included in KEYALL

instead of

KEYALL = (0x7ffffff & ~KEYNOMS18 &
              ~KEYNOOPENCL & ~KEYCUDA ) // KEYNOMS18 and KEYNOOPENCL are used to exclude.
// KEYCUDA is not included in KEYALL

since the current maximum mask is 0x3ffffff instead of 0x7ffffff

[delcypher]

clang/lib/Basic/IdentifierTable.cpp

111

Oops, you're right it would be 0x3ffffff. I wonder though if we should clean this up so we don't need to manually update the bit mask every time... what if it was written like this?

 enum {
    KEYC99        = 0x1,
    KEYCXX        = 0x2,
    KEYCXX11      = 0x4,
    ....
    KEYSYCL       = 0x1000000,
    KEYCUDA       = 0x2000000,
    KEYMAX = KEYCUDA, // Must be set to the largest KEY enum value
    KEYALLCXX = KEYCXX | KEYCXX11 | KEYCXX20,

    // KEYNOMS18 and KEYNOOPENCL are used to exclude.
    // KEYCUDA is not included in KEYALL because <FIXME add reason here>
    KEYALL = (((KEYMAX & (KEYMAX-1)) & ~KEYNOMS18 & ~KEYNOOPENCL & ~KEYCUDA)
};

[yaxunl]

clang/lib/Basic/IdentifierTable.cpp

111

On second thought, KEYALL does not need to exclude KEYCUDA.

However, it would be good to set KEYALL in a generic approach. I will open a separate review.

[yaxunl]

clang/lib/Basic/IdentifierTable.cpp

111

opened https://reviews.llvm.org/D125396 to fix KEYALL

[delcypher]

clang/lib/Basic/IdentifierTable.cpp

111

Oops that should say

KEYALL = (((KEYMAX | (KEYMAX-1)) & ~KEYNOMS18 & ~KEYNOOPENCL & ~KEYCUDA)

clang/docs/ReleaseNotes.rst

- The mangling scheme for C++20 modules has incompatibly changed. The
   symbols with named module attachment.
 
 C++2b Feature Support
 ^^^^^^^^^^^^^^^^^^^^^
 
 - Implemented `P2128R6: Multidimensional subscript operator <https://wg21.link/P2128R6>`_.
 - 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
 -------------------------------------
 
 OpenCL C Language Changes in Clang
 ----------------------------------
 
 ...

clang/include/clang/Basic/Attr.td

 
 def Convergent : InheritableAttr {
   let Spellings = [Clang<"convergent">];
   let Subjects = SubjectList<[Function]>;
   let Documentation = [ConvergentDocs];
   let SimpleHandler = 1;
 }
 
 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];
   let Subjects = SubjectList<[Function, Stmt], WarnDiag,
                              "functions and statements">;
   let SimpleHandler = 1;
 }
 

clang/include/clang/Basic/AttrDocs.td

 This function attribute suppresses the inlining of a function at the call sites
 of the function.
 
 ``[[clang::noinline]]`` spelling can be used as a statement attribute; other
 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) {
     int r;
     [[clang::noinline]] foo();
     [[clang::noinline]] r = bar();
     return r;
   }

clang/include/clang/Basic/Features.def

 EXTENSION(gnu_asm, LangOpts.GNUAsm)
 EXTENSION(gnu_asm_goto_with_outputs, LangOpts.GNUAsm)
 EXTENSION(matrix_types, LangOpts.MatrixTypes)
 EXTENSION(matrix_types_scalar_division, true)
 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

clang/include/clang/Basic/TokenKinds.def

 KEYWORD(__builtin_astype            , KEYOPENCLC | KEYOPENCLCXX)
 UNARY_EXPR_OR_TYPE_TRAIT(vec_step, VecStep, KEYOPENCLC | KEYOPENCLCXX | KEYALTIVEC | KEYZVECTOR)
 #define GENERIC_IMAGE_TYPE(ImgType, Id) KEYWORD(ImgType##_t, KEYOPENCLC | KEYOPENCLCXX)
 #include "clang/Basic/OpenCLImageTypes.def"
 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)
 
 // Borland Extensions.
 KEYWORD(__pascal                    , KEYALL)
 
 // Altivec Extension.
 KEYWORD(__vector                    , KEYALTIVEC|KEYZVECTOR)

clang/include/clang/Parse/Parser.h

private:
   void ParseMicrosoftTypeAttributes(ParsedAttributes &attrs);
   void DiagnoseAndSkipExtendedMicrosoftTypeAttributes();
   SourceLocation SkipExtendedMicrosoftTypeAttributes();
   void ParseMicrosoftInheritanceClassAttributes(ParsedAttributes &attrs);
   void ParseBorlandTypeAttributes(ParsedAttributes &attrs);
   void ParseOpenCLKernelAttributes(ParsedAttributes &attrs);
   void ParseOpenCLQualifiers(ParsedAttributes &Attrs);
   void ParseNullabilityTypeSpecifiers(ParsedAttributes &attrs);
+  void ParseCUDAFunctionAttributes(ParsedAttributes &attrs);
 
   VersionTuple ParseVersionTuple(SourceRange &Range);
   void ParseAvailabilityAttribute(IdentifierInfo &Availability,
                                   SourceLocation AvailabilityLoc,
                                   ParsedAttributes &attrs,
                                   SourceLocation *endLoc,
                                   IdentifierInfo *ScopeName,
                                   SourceLocation ScopeLoc,

clang/lib/Basic/IdentifierTable.cpp

enum {
     KEYOBJC       = 0x20000,
     KEYZVECTOR    = 0x40000,
     KEYCOROUTINES = 0x80000,
     KEYMODULES    = 0x100000,
     KEYCXX20      = 0x200000,
     KEYOPENCLCXX  = 0x400000,
     KEYMSCOMPAT   = 0x800000,
     KEYSYCL       = 0x1000000,
+    KEYCUDA       = 0x2000000,
     KEYALLCXX = KEYCXX | KEYCXX11 | KEYCXX20,
     KEYALL = (0x1ffffff & ~KEYNOMS18 &
               ~KEYNOOPENCL) // KEYNOMS18 and KEYNOOPENCL are used to exclude.
   };
 
   /// How a keyword is treated in the selected standard.
   enum KeywordStatus {
     KS_Disabled,    // Disabled
static KeywordStatus getKeywordStatus(const LangOptions &LangOpts,
   if (LangOpts.CPlusPlus20 && (Flags & KEYCONCEPTS)) return KS_Enabled;
   if (LangOpts.Coroutines && (Flags & KEYCOROUTINES)) return KS_Enabled;
   if (LangOpts.ModulesTS && (Flags & KEYMODULES)) return KS_Enabled;
   if (LangOpts.CPlusPlus && (Flags & KEYALLCXX)) return KS_Future;
   if (LangOpts.CPlusPlus && !LangOpts.CPlusPlus20 && (Flags & CHAR8SUPPORT))
     return KS_Future;
   if (LangOpts.isSYCL() && (Flags & KEYSYCL))
     return KS_Enabled;
+  if (LangOpts.CUDA && (Flags & KEYCUDA))
+    return KS_Enabled;
   return KS_Disabled;
 }
 
 /// AddKeyword - This method is used to associate a token ID with specific
 /// identifiers because they are language keywords.  This causes the lexer to
 /// automatically map matching identifiers to specialized token codes.
 static void AddKeyword(StringRef Keyword,
                        tok::TokenKind TokenCode, unsigned Flags,

clang/lib/Parse/ParseDecl.cpp

void Parser::ParseOpenCLKernelAttributes(ParsedAttributes &attrs) {
   while (Tok.is(tok::kw___kernel)) {
     IdentifierInfo *AttrName = Tok.getIdentifierInfo();
     SourceLocation AttrNameLoc = ConsumeToken();
     attrs.addNew(AttrName, AttrNameLoc, nullptr, AttrNameLoc, nullptr, 0,
                  ParsedAttr::AS_Keyword);
   }
 }
 
+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();
   Attrs.addNew(AttrName, AttrNameLoc, nullptr, AttrNameLoc, nullptr, 0,
                ParsedAttr::AS_Keyword);
 }
 
 void Parser::ParseNullabilityTypeSpecifiers(ParsedAttributes &attrs) {
case tok::kw___pascal:
       ParseBorlandTypeAttributes(DS.getAttributes());
       continue;
 
     // OpenCL single token adornments.
     case tok::kw___kernel:
       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:
     case tok::kw__Nullable_result:
     case tok::kw__Null_unspecified:
       ParseNullabilityTypeSpecifiers(DS.getAttributes());
       continue;
 

clang/test/CodeGenCUDA/noinline.cu

+// 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

clang/test/Lexer/has_feature.cu

+// 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

clang/test/SemaCUDA/noinline.cu

+// 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() {}