[SYCL] Define SYCL_EXTERNAL macro as prescribed by the spec

Signed-off-by: Sergey Dmitriev <[email protected]>

Author: sndmitriev

clang/lib/Frontend/InitPreprocessor.cpp

@@ -1074,6 +1074,7 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
   // SYCL device compiler which doesn't produce host binary.
   if (LangOpts.SYCLIsDevice) {
     Builder.defineMacro("__SYCL_DEVICE_ONLY__", "1");
+    Builder.defineMacro("SYCL_EXTERNAL", "__attribute__((sycl_device))");
     if (!getenv("DISABLE_INFER_AS"))
       Builder.defineMacro("__SYCL_ENABLE_INFER_AS__", "1");
   }

clang/test/Preprocessor/sycl-macro.cpp

@@ -3,9 +3,11 @@
 // RUN: %clang_cc1 %s -fsycl -E -dM | FileCheck --check-prefix=CHECK-ANY-SYCL %s
 // RUN: %clang_cc1 %s -fsycl-is-device -E -dM -fms-compatibility | FileCheck --check-prefix=CHECK-MSVC %s
 // CHECK-NOT:#define __SYCL_DEVICE_ONLY__ 1
+// CHECK-NOT:#define SYCL_EXTERNAL
 // CHECK-NOT:#define CL_SYCL_LANGUAGE_VERSION 121
 // CHECK-ANY-SYCL-NOT:#define __SYCL_DEVICE_ONLY__ 1
 // CHECK-ANY-SYCL:#define CL_SYCL_LANGUAGE_VERSION 121
 // CHECK-SYCL:#define CL_SYCL_LANGUAGE_VERSION 121
+// CHECK-SYCL:#define SYCL_EXTERNAL __attribute__((sycl_device))
 // CHECK-MSVC-NOT: __GNUC__
 // CHECK-MSVC-NOT: __STDC__

sycl/include/CL/sycl/detail/common.hpp

@@ -88,6 +88,10 @@ static inline std::string codeToString(cl_int code){
 #define ALWAYS_INLINE
 #endif
 
+#ifndef SYCL_EXTERNAL
+#define SYCL_EXTERNAL
+#endif
+
 namespace cl {
 namespace sycl {
 namespace detail {

sycl/test/separate-compile/sycl-external.cpp

@@ -0,0 +1,73 @@
+// Test1 - check that kernel can call a SYCL_EXTERNAL function defined in a
+// different object file.
+// RUN: %clangxx -fsycl -DSOURCE1 -c %s -o %t1.o
+// RUN: %clangxx -fsycl -DSOURCE2 -c %s -o %t2.o
+// RUN: %clangxx -fsycl %t1.o %t2.o -o %t.exe
+// RUN: %CPU_RUN_PLACEHOLDER %t.exe
+// RUN: %GPU_RUN_PLACEHOLDER %t.exe
+// RUN: %ACC_RUN_PLACEHOLDER %t.exe
+//
+// Test2 - check that kernel can call a SYCL_EXTERNAL function defined in a
+// static library.
+// RUN: rm -f %t.a
+// RUN: llvm-ar crv %t.a %t1.o
+// RUN: %clangxx -fsycl %t2.o -foffload-static-lib=%t.a -o %t.exe
+// RUN: %CPU_RUN_PLACEHOLDER %t.exe
+// RUN: %GPU_RUN_PLACEHOLDER %t.exe
+// RUN: %ACC_RUN_PLACEHOLDER %t.exe
+
+#include <CL/sycl.hpp>
+#include <iostream>
+
+#ifdef SOURCE1
+int bar(int b);
+
+SYCL_EXTERNAL
+int foo(int a, int b) {
+  return a + bar(b);
+}
+
+int bar(int b) {
+  return b + 5;
+}
+#endif // SOURCE1
+
+#ifdef SOURCE2
+SYCL_EXTERNAL
+int foo(int A, int B);
+
+int main(void) {
+  constexpr unsigned Size = 4;
+  int A[Size] = {1, 2, 3, 4};
+  int B[Size] = {1, 2, 3, 4};
+  int C[Size];
+
+  {
+    cl::sycl::range<1> range{Size};
+    cl::sycl::buffer<int, 1> bufA(A, range);
+    cl::sycl::buffer<int, 1> bufB(B, range);
+    cl::sycl::buffer<int, 1> bufC(C, range);
+
+    cl::sycl::queue().submit([&](cl::sycl::handler &cgh) {
+      auto accA = bufA.get_access<cl::sycl::access::mode::read>(cgh);
+      auto accB = bufB.get_access<cl::sycl::access::mode::read>(cgh);
+      auto accC = bufC.get_access<cl::sycl::access::mode::write>(cgh);
+
+      cgh.parallel_for<class Test>(range, [=](cl::sycl::id<1> ID) {
+        accC[ID] = foo(accA[ID], accB[ID]);
+      });
+    });
+  }
+
+  for (unsigned I = 0; I < Size; ++I) {
+    int Ref = foo(A[I], B[I]);
+    if (C[I] != Ref) {
+      std::cout << "fail: [" << I << "] == " << C[I] << ", expected " << Ref
+                << "\n";
+      return 1;
+    }
+  }
+  std::cout << "pass\n";
+  return 0;
+}
+#endif // SOURCE2