[SYCLomatic] Fix the migration bug of __attribute__(__noinline__). (#2387)

Signed-off-by: Tang, Jiajun [email protected]

Author: tangjj11

clang/include/clang/Lex/Preprocessor.h

@@ -3020,6 +3020,10 @@ class Preprocessor {
   bool setDeserializedSafeBufferOptOutMap(
       const SmallVectorImpl<SourceLocation> &SrcLocSeqs);
 
+#ifdef SYCLomatic_CUSTOMIZATION
+  bool IsInAttr = false;
+#endif // SYCLomatic_CUSTOMIZATION
+
 private:
   /// Helper functions to forward lexing to the actual lexer. They all share the
   /// same signature.

clang/lib/DPCT/ASTTraversal.cpp

@@ -279,6 +279,12 @@ void IncludesCallbacks::MacroDefined(const Token &MacroNameTok,
     if (!II)
       continue;
 
+    // The "__noinline__" macro is re-defined and it is used in
+    // "__attribute__()", do not migrate it.
+    if (II->hasMacroDefinition() && (II->getName() == "__noinline__")) {
+      continue;
+    }
+
     auto ItRule = MapNames::MacroRuleMap.find(II->getName().str());
     if (ItRule != MapNames::MacroRuleMap.end()) {
       TransformSet.emplace_back(

clang/lib/Lex/Preprocessor.cpp

@@ -815,6 +815,15 @@ bool Preprocessor::HandleIdentifier(Token &Identifier) {
     HandlePoisonedIdentifier(Identifier);
   }
 
+#ifdef SYCLomatic_CUSTOMIZATION
+  // The "__noinline__" macro is re-defined and it is used in "__attribute__()",
+  // do not handle it.
+  if (LangOpts.CUDA && Identifier.is(tok::TokenKind::kw___noinline__) &&
+      IsInAnalysisScopeFunc(Identifier.getLocation()) && IsInAttr) {
+    return true;
+  }
+#endif // SYCLomatic_CUSTOMIZATION
+
   // If this is a macro to be expanded, do it.
   if (const MacroDefinition MD = getMacroDefinition(&II)) {
     const auto *MI = MD.getMacroInfo();

clang/lib/Parse/ParseDecl.cpp

@@ -193,6 +193,11 @@ void Parser::ParseGNUAttributes(ParsedAttributes &Attrs,
 
   SourceLocation StartLoc = Tok.getLocation();
   SourceLocation EndLoc = StartLoc;
+#ifdef SYCLomatic_CUSTOMIZATION
+  // Set the 'IsInAttr' to true before parse the arguments of __attribute__(()).
+  // Then we can handle the macro in __attribute__ arguments specially.
+  llvm::SaveAndRestore<bool> SART(PP.IsInAttr, true);
+#endif // SYCLomatic_CUSTOMIZATION
 
   while (Tok.is(tok::kw___attribute)) {
     SourceLocation AttrTokLoc = ConsumeToken();

clang/test/dpct/noinline.cu

@@ -49,3 +49,10 @@ __noinline__ __host__ __device__ scalar_t calc_igammac(scalar_t a, scalar_t b) {
   scalar_t c = a + b;
   return c;
 }
+
+// CHECK: __attribute__((__noinline__)) void macro_with_attr() {}
+__attribute__((__noinline__)) void macro_with_attr() {}
+// CHECK: #define NOINLINE __attribute__((__noinline__))
+#define NOINLINE __attribute__((__noinline__))
+// CHECK: NOINLINE void macro_in_macro_with_attr() {}
+NOINLINE void macro_in_macro_with_attr() {}