[SYCL] Set preprocessor macros in integration header (#4823)

Predefines which need to be set for custom host compilation
must be defined in integration header.

Signed-off-by: Elizabeth Andrews <[email protected]>

Author: elizabethandrews

clang/include/clang/Basic/Version.h

@@ -15,7 +15,9 @@
 #ifndef LLVM_CLANG_BASIC_VERSION_H
 #define LLVM_CLANG_BASIC_VERSION_H
 
+#include "clang/Basic/LangOptions.h"
 #include "clang/Basic/Version.inc"
+#include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringRef.h"
 
 namespace clang {
@@ -56,6 +58,11 @@ namespace clang {
   /// for use in the CPP __VERSION__ macro, which includes the clang version
   /// number, the repository version, and the vendor tag.
   std::string getClangFullCPPVersion();
+
+  /// Retrieves a string representing the SYCL standard version for use in
+  /// the CL_SYCL_LANGUAGE_VERSION and SYCL_LANGUAGE_VERSION macros.
+  llvm::SmallVector<std::pair<llvm::StringRef, llvm::StringRef>, 2>
+  getSYCLVersionMacros(const LangOptions &LangOpts);
 }
 
 #endif // LLVM_CLANG_BASIC_VERSION_H

clang/lib/Basic/Version.cpp

@@ -123,4 +123,13 @@ std::string getClangFullCPPVersion() {
   return buf;
 }
 
+llvm::SmallVector<std::pair<llvm::StringRef, llvm::StringRef>, 2>
+getSYCLVersionMacros(const LangOptions &LangOpts) {
+  if (LangOpts.getSYCLVersion() == LangOptions::SYCL_2017)
+    return {{"CL_SYCL_LANGUAGE_VERSION", "121"},
+            {"SYCL_LANGUAGE_VERSION", "201707"}};
+  if (LangOpts.getSYCLVersion() == LangOptions::SYCL_2020)
+    return {{"SYCL_LANGUAGE_VERSION", "202001"}};
+  llvm_unreachable("SYCL standard should be set");
+}
 } // end namespace clang

clang/lib/Frontend/InitPreprocessor.cpp

@@ -491,11 +491,9 @@ static void InitializeStandardPredefinedMacros(const TargetInfo &TI,
 
   if (LangOpts.SYCLIsDevice || LangOpts.SYCLIsHost) {
     // SYCL Version is set to a value when building SYCL applications
-    if (LangOpts.getSYCLVersion() == LangOptions::SYCL_2017) {
-      Builder.defineMacro("CL_SYCL_LANGUAGE_VERSION", "121");
-      Builder.defineMacro("SYCL_LANGUAGE_VERSION", "201707");
-    } else if (LangOpts.getSYCLVersion() == LangOptions::SYCL_2020)
-      Builder.defineMacro("SYCL_LANGUAGE_VERSION", "202001");
+    for (const std::pair<StringRef, StringRef> &Macro :
+         getSYCLVersionMacros(LangOpts))
+      Builder.defineMacro(Macro.first, Macro.second);
 
     if (LangOpts.SYCLValueFitInMaxInt)
       Builder.defineMacro("__SYCL_ID_QUERIES_FIT_IN_INT__");

clang/lib/Sema/SemaSYCL.cpp

@@ -20,6 +20,7 @@
 #include "clang/Basic/Attributes.h"
 #include "clang/Basic/Builtins.h"
 #include "clang/Basic/Diagnostic.h"
+#include "clang/Basic/Version.h"
 #include "clang/Sema/Initialization.h"
 #include "clang/Sema/Sema.h"
 #include "llvm/ADT/APSInt.h"
@@ -4594,6 +4595,21 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   Policy.SuppressUnwrittenScope = true;
   SYCLFwdDeclEmitter FwdDeclEmitter(O, S.getLangOpts());
 
+  // Predefines which need to be set for custom host compilation
+  // must be defined in integration header.
+  for (const std::pair<StringRef, StringRef> &Macro :
+       getSYCLVersionMacros(S.getLangOpts())) {
+    O << "#ifndef " << Macro.first << '\n';
+    O << "#define " << Macro.first << " " << Macro.second << '\n';
+    O << "#endif //" << Macro.first << "\n\n";
+  }
+
+  if (S.getLangOpts().SYCLDisableRangeRounding) {
+    O << "#ifndef __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ \n";
+    O << "#define __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ 1\n";
+    O << "#endif //__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__\n\n";
+  }
+
   if (SpecConsts.size() > 0) {
     O << "// Forward declarations of templated spec constant types:\n";
     for (const auto &SC : SpecConsts)

clang/test/CodeGenSYCL/integration_header_ppmacros.cpp

@@ -0,0 +1,36 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -sycl-std=2020 -fsycl-int-header=%t.h %s
+// RUN: FileCheck -input-file=%t.h %s --check-prefix=CHECK-SYCL2020
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -sycl-std=2017 -fsycl-int-header=%t.h %s
+// RUN: FileCheck -input-file=%t.h %s --check-prefix=CHECK-SYCL2017
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-disable-range-rounding -fsycl-int-header=%t.h %s
+// RUN: FileCheck -input-file=%t.h %s --check-prefix=CHECK-RANGE
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-int-header=%t.h %s
+// RUN: FileCheck -input-file=%t.h %s --check-prefix=CHECK-NO-RANGE
+
+// Test verifying predefines which need to be set for host and device compilation.
+// Preprocessor macros which are required when using custom host compiler must be
+// defined in integration header.
+
+#include "Inputs/sycl.hpp"
+
+int main() {
+  cl::sycl::kernel_single_task<class first_kernel>([]() {});
+}
+// CHECK-SYCL2020: #ifndef SYCL_LANGUAGE_VERSION
+// CHECK-SYCL2020-NEXT: #define SYCL_LANGUAGE_VERSION 202001
+// CHECK-SYCL2020-NEXT: #endif //SYCL_LANGUAGE_VERSION
+// CHECK-SYCL2020-NOT: #define CL_SYCL_LANGUAGE_VERSION 121
+// CHECK-SYCL2020-NOT: #define SYCL_LANGUAGE_VERSION 201707
+
+// CHECK-SYCL2017: #ifndef CL_SYCL_LANGUAGE_VERSION
+// CHECK-SYCL2017-NEXT: #define CL_SYCL_LANGUAGE_VERSION 121
+// CHECK-SYCL2017-NEXT: #endif //CL_SYCL_LANGUAGE_VERSION
+// CHECK-SYCL2017: #ifndef SYCL_LANGUAGE_VERSION
+// CHECK-SYCL2017-NEXT: #define SYCL_LANGUAGE_VERSION 201707
+// CHECK-SYCL2017-NEXT: #endif //SYCL_LANGUAGE_VERSION
+// CHECK-SYCL2017-NOT: #define SYCL_LANGUAGE_VERSION 202001
+
+// CHECK-RANGE: #ifndef __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__
+// CHECK-RANGE-NEXT: #define __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ 1
+// CHECK-RANGE-NEXT: #endif //__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__
+// CHECK-NO-RANGE-NOT: #define __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ 1

sycl/include/CL/sycl/handler.hpp

@@ -1062,7 +1062,7 @@ class __SYCL_EXPORT handler {
   }
 
 #ifdef SYCL_LANGUAGE_VERSION
-#define __SYCL_KERNEL_ATTR__ __attribute__((sycl_kernel))
+#define __SYCL_KERNEL_ATTR__ [[clang::sycl_kernel]]
 #else
 #define __SYCL_KERNEL_ATTR__
 #endif