[SYCL] Add support for SYCL_EXTERNAL

clang/include/clang/Basic/Attr.td

@@ -1025,9 +1025,9 @@ def CUDAShared : InheritableAttr {
 
 def SYCLDevice : InheritableAttr {
   let Spellings = [GNU<"sycl_device">];
-  let Subjects = SubjectList<[Function, Var]>;
+  let Subjects = SubjectList<[Function]>;
   let LangOpts = [SYCLIsDevice];
-  let Documentation = [Undocumented];
+  let Documentation = [SYCLDeviceDocs];
 }
 
 def SYCLKernel : InheritableAttr {

clang/include/clang/Basic/AttrDocs.td

@@ -1845,6 +1845,17 @@ function pointer for the specified function.
   }];
 }
 
+def SYCLDeviceDocs : Documentation {
+  let Category = DocCatFunction;
+  let Heading = "sycl_device";
+  let Content = [{
+This attribute can only be applied to functions and indicates that the
+function must be treated as a device function and must be emitted even if it has
+no direct uses from other device functions. All ``sycl_device`` function callees
+implicitly inherit this attribute.
+  }];
+}
+
 def RISCVInterruptDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "interrupt (RISCV)";

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -9803,10 +9803,12 @@ def err_sycl_non_std_layout_type : Error<
   "kernel parameter has non-standard layout class/struct type">;
 def err_conflicting_sycl_kernel_attributes : Error<
   "conflicting attributes applied to a SYCL kernel">;
-def err_sycl_device_indirectly_callable_cannot_be_applied_here
-    : Error<"device_indirectly_callable attribute cannot be applied to a "
+def err_sycl_attibute_cannot_be_applied_here
+    : Error<"%0 attribute cannot be applied to a "
             "%select{static function or function in an anonymous namespace"
-            "|class member function}0">;
+            "|class member function"
+            "|function with a raw pointer return type"
+            "|function with a raw pointer parameter type}1">;
 
 def err_bit_cast_non_trivially_copyable : Error<
   "__builtin_bit_cast %select{source|destination}0 type must be trivially copyable">;

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/lib/Sema/SemaDeclAttr.cpp

@@ -4417,19 +4417,45 @@ static void handleOptimizeNoneAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
     D->addAttr(Optnone);
 }
 
+static void handleSYCLDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+  auto *FD = cast<FunctionDecl>(D);
+  if (!FD->isExternallyVisible()) {
+    S.Diag(AL.getLoc(), diag::err_sycl_attibute_cannot_be_applied_here)
+        << AL << 0 /* static function or anonymous namespace */;
+    return;
+  }
+  if (isa<CXXMethodDecl>(FD)) {
+    S.Diag(AL.getLoc(), diag::err_sycl_attibute_cannot_be_applied_here)
+        << AL << 1 /* class member function */;
+    return;
+  }
+  if (FD->getReturnType()->isPointerType()) {
+    S.Diag(AL.getLoc(), diag::err_sycl_attibute_cannot_be_applied_here)
+        << AL << 2 /* function with a raw pointer return type */;
+    return;
+  }
+  for (const ParmVarDecl *Param : FD->parameters())
+    if (Param->getType()->isPointerType()) {
+      S.Diag(AL.getLoc(), diag::err_sycl_attibute_cannot_be_applied_here)
+          << AL << 3 /* function with a raw pointer parameter type */;
+      return;
+    }
+
+  S.addSyclDeviceDecl(D);
+  handleSimpleAttribute<SYCLDeviceAttr>(S, D, AL);
+}
+
 static void handleSYCLDeviceIndirectlyCallableAttr(Sema &S, Decl *D,
                                                    const ParsedAttr &AL) {
   auto *FD = cast<FunctionDecl>(D);
   if (!FD->isExternallyVisible()) {
-    S.Diag(AL.getLoc(),
-           diag::err_sycl_device_indirectly_callable_cannot_be_applied_here)
-        << 0 /* static function or anonymous namespace */;
+    S.Diag(AL.getLoc(), diag::err_sycl_attibute_cannot_be_applied_here)
+        << AL << 0 /* static function or anonymous namespace */;
     return;
   }
   if (isa<CXXMethodDecl>(FD)) {
-    S.Diag(AL.getLoc(),
-           diag::err_sycl_device_indirectly_callable_cannot_be_applied_here)
-        << 1 /* class member function */;
+    S.Diag(AL.getLoc(), diag::err_sycl_attibute_cannot_be_applied_here)
+        << AL << 1 /* class member function */;
     return;
   }
 
@@ -7116,6 +7142,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_SYCLKernel:
     handleSimpleAttribute<SYCLKernelAttr>(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLDevice:
+    handleSYCLDeviceAttr(S, D, AL);
+    break;
   case ParsedAttr::AT_SYCLDeviceIndirectlyCallable:
     handleSYCLDeviceIndirectlyCallableAttr(S, D, AL);
     break;

clang/test/CodeGenSYCL/sycl-device.cpp

@@ -0,0 +1,52 @@
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s
+// Test code generation for sycl_device attribute.
+
+int bar(int b);
+
+// CHECK-DAG: define spir_func i32 @_Z3fooii
+__attribute__((sycl_device))
+int foo(int a, int b) { return a + bar(b); }
+
+// CHECK-DAG: define spir_func i32 @_Z3bari
+int bar(int b) { return b; }
+
+// CHECK-DAG: define spir_func i32 @_Z3fari
+int far(int b) { return b; }
+
+// CHECK-DAG: define spir_func i32 @_Z3booii
+__attribute__((sycl_device))
+int boo(int a, int b) { return a + far(b); }
+
+// CHECK-DAG: define spir_func i32 @_Z3cari
+__attribute__((sycl_device))
+int car(int b);
+int car(int b) { return b; }
+
+// CHECK-DAG: define spir_func i32 @_Z3cazi
+int caz(int b);
+__attribute__((sycl_device))
+int caz(int b) { return b; }
+
+template<typename T>
+__attribute__((sycl_device))
+void taf(T t) {}
+
+// CHECK-DAG: define weak_odr spir_func void @_Z3tafIiEvT_
+template void taf<int>(int t);
+
+// CHECK-DAG: define spir_func void @_Z3tafIcEvT_
+template<> void taf<char>(char t) {}
+
+template<typename T>
+void tar(T t) {}
+
+// CHECK-DAG: define spir_func void @_Z3tarIcEvT_
+template<>
+__attribute__((sycl_device))
+void tar<char>(char t) {}
+
+// CHECK-NOT: @_Z3tarIiEvT_
+template void tar<int>(int t);
+
+// CHECK-NOT: @_Z3gooi
+int goo(int b) { return b; }

clang/test/Misc/pragma-attribute-supported-attributes-list.test

@@ -126,7 +126,7 @@
 // CHECK-NEXT: ReturnTypestate (SubjectMatchRule_function, SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: ReturnsNonNull (SubjectMatchRule_objc_method, SubjectMatchRule_function)
 // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
-// CHECK-NEXT: SYCLDevice (SubjectMatchRule_function, SubjectMatchRule_variable)
+// CHECK-NEXT: SYCLDevice (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLKernel (SubjectMatchRule_function)
 // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record)

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__

clang/test/SemaSYCL/device-indirectly-callable-attr.cpp

@@ -10,19 +10,19 @@ int N;
 [[intel::device_indirectly_callable(3)]] // expected-error {{'device_indirectly_callable' attribute takes no arguments}}
 void bar() {}
 
-[[intel::device_indirectly_callable]] // expected-error {{device_indirectly_callable attribute cannot be applied to a static function or function in an anonymous namespace}}
+[[intel::device_indirectly_callable]] // expected-error {{'device_indirectly_callable' attribute cannot be applied to a static function or function in an anonymous namespace}}
 static void func1() {}
 
 namespace {
-  [[intel::device_indirectly_callable]] // expected-error {{device_indirectly_callable attribute cannot be applied to a static function or function in an anonymous namespace}}
+  [[intel::device_indirectly_callable]] // expected-error {{'device_indirectly_callable' attribute cannot be applied to a static function or function in an anonymous namespace}}
   void func2() {}
 }
 
 class A {
-  [[intel::device_indirectly_callable]] // expected-error {{device_indirectly_callable attribute cannot be applied to a class member function}}
+  [[intel::device_indirectly_callable]] // expected-error {{'device_indirectly_callable' attribute cannot be applied to a class member function}}
   A() {}
 
-  [[intel::device_indirectly_callable]] // expected-error {{device_indirectly_callable attribute cannot be applied to a class member function}}
+  [[intel::device_indirectly_callable]] // expected-error {{'device_indirectly_callable' attribute cannot be applied to a class member function}}
   int func3() {}
 };
 

clang/test/SemaSYCL/sycl-device.cpp

@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -verify -DNO_SYCL %s
+
+#ifndef NO_SYCL
+
+__attribute__((sycl_device)) // expected-warning {{'sycl_device' attribute only applies to functions}}
+int N;
+
+__attribute__((sycl_device(3))) // expected-error {{'sycl_device' attribute takes no arguments}}
+void bar() {}
+
+__attribute__((sycl_device)) // expected-error {{'sycl_device' attribute cannot be applied to a static function or function in an anonymous namespace}}
+static void func1() {}
+
+namespace {
+  __attribute__((sycl_device)) // expected-error {{'sycl_device' attribute cannot be applied to a static function or function in an anonymous namespace}}
+  void func2() {}
+}
+
+class A {
+  __attribute__((sycl_device)) // expected-error {{'sycl_device' attribute cannot be applied to a class member function}}
+  A() {}
+
+  __attribute__((sycl_device)) // expected-error {{'sycl_device' attribute cannot be applied to a class member function}}
+  int func3() {}
+};
+
+__attribute__((sycl_device)) // expected-error {{'sycl_device' attribute cannot be applied to a function with a raw pointer return type}}
+int* func3() { return nullptr; }
+
+__attribute__((sycl_device)) // expected-error {{'sycl_device' attribute cannot be applied to a function with a raw pointer parameter type}}
+void func3(int *) {}
+
+#else
+
+__attribute__((sycl_device)) // expected-warning {{'sycl_device' attribute ignored}}
+void baz() {}
+
+#endif // NO_SYCL

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