diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 8416edf3899cb..d3f2501284eac 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/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 { diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 2a5e7493614bf..ae9b64234bcad 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/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)"; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 9c244567bc09f..d48c1d159e067 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/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">; diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index f369d2a087ed2..d95b91c940591 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/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"); } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index d81be80223d48..4fdf91fb5f3c1 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/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(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(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(S, D, AL); +} + static void handleSYCLDeviceIndirectlyCallableAttr(Sema &S, Decl *D, const ParsedAttr &AL) { auto *FD = cast(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(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(S, D, AL); break; + case ParsedAttr::AT_SYCLDevice: + handleSYCLDeviceAttr(S, D, AL); + break; case ParsedAttr::AT_SYCLDeviceIndirectlyCallable: handleSYCLDeviceIndirectlyCallableAttr(S, D, AL); break; diff --git a/clang/test/CodeGenSYCL/sycl-device.cpp b/clang/test/CodeGenSYCL/sycl-device.cpp new file mode 100644 index 0000000000000..424cacff09fcb --- /dev/null +++ b/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 +__attribute__((sycl_device)) +void taf(T t) {} + +// CHECK-DAG: define weak_odr spir_func void @_Z3tafIiEvT_ +template void taf(int t); + +// CHECK-DAG: define spir_func void @_Z3tafIcEvT_ +template<> void taf(char t) {} + +template +void tar(T t) {} + +// CHECK-DAG: define spir_func void @_Z3tarIcEvT_ +template<> +__attribute__((sycl_device)) +void tar(char t) {} + +// CHECK-NOT: @_Z3tarIiEvT_ +template void tar(int t); + +// CHECK-NOT: @_Z3gooi +int goo(int b) { return b; } diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index 89b8639783f96..eba50b7de3116 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ 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) diff --git a/clang/test/Preprocessor/sycl-macro.cpp b/clang/test/Preprocessor/sycl-macro.cpp index 9e2ebf25bb864..378eb1dc49bfd 100644 --- a/clang/test/Preprocessor/sycl-macro.cpp +++ b/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__ diff --git a/clang/test/SemaSYCL/device-indirectly-callable-attr.cpp b/clang/test/SemaSYCL/device-indirectly-callable-attr.cpp index ed1913e21bb05..772015fb9b694 100644 --- a/clang/test/SemaSYCL/device-indirectly-callable-attr.cpp +++ b/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() {} }; diff --git a/clang/test/SemaSYCL/sycl-device.cpp b/clang/test/SemaSYCL/sycl-device.cpp new file mode 100644 index 0000000000000..2eb2ef364c61e --- /dev/null +++ b/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 diff --git a/sycl/include/CL/sycl/detail/common.hpp b/sycl/include/CL/sycl/detail/common.hpp index f90f6dd6e1ebb..e60e7c01b4571 100644 --- a/sycl/include/CL/sycl/detail/common.hpp +++ b/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 { diff --git a/sycl/test/separate-compile/sycl-external.cpp b/sycl/test/separate-compile/sycl-external.cpp new file mode 100644 index 0000000000000..bb46ffdae4c4e --- /dev/null +++ b/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 +#include + +#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 bufA(A, range); + cl::sycl::buffer bufB(B, range); + cl::sycl::buffer bufC(C, range); + + cl::sycl::queue().submit([&](cl::sycl::handler &cgh) { + auto accA = bufA.get_access(cgh); + auto accB = bufB.get_access(cgh); + auto accC = bufC.get_access(cgh); + + cgh.parallel_for(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