bader updated this revision to Diff 228286.
bader added a comment.
Applied comments from Aaron.
Two comments left unresolved:
- Split diagnostic message for sycl_kernel attribute into multiple messages.
Will do tomorrow.
- Change attribute "subject" in TableGen file from "Function" to
"FunctionTemplate". I need guidance here as I'm not sure how to make it work.
Refactored patch to re-use CodeGen infrastructure for emitting SYCL device code.
It turned out to be quite simple change - just two one-liner changes in
ASTContext to say that only SYCL kernels must be emitted when we compile for
SYCL device + similar change in the CodeGen to mark symbols which must be
emitted.
Removed `sycl_device` attribute, which was required by previous implementation
for device code outlining. I think we still might need this attribute to mark
"non-kernel" symbols as "device code", so the compiler will emit even though
they are not used by SYCL kernels. Anyway it's not required for device code
outlining and shouldn't be part of this patch.
Enhanced CodeGen test to check that host part of the code is not emitted.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D60455/new/
https://reviews.llvm.org/D60455
Files:
clang/include/clang/Basic/Attr.td
clang/include/clang/Basic/AttrDocs.td
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/AST/ASTContext.cpp
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/Sema/SemaDeclAttr.cpp
clang/test/CodeGenSYCL/device-functions.cpp
clang/test/Misc/pragma-attribute-supported-attributes-list.test
clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
clang/test/SemaSYCL/device-attributes.cpp
Index: clang/test/SemaSYCL/device-attributes.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/device-attributes.cpp
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fsycl-is-device -verify %s
+
+[[clang::sycl_kernel]] int gv2 = 0; // expected-warning {{'sycl_kernel' attribute only applies to functions}}
+__attribute__((sycl_kernel)) int gv3 = 0; // expected-warning {{'sycl_kernel' attribute only applies to functions}}
+
+__attribute__((sycl_kernel(1))) void foo(); // expected-error {{'sycl_kernel' attribute takes no arguments}}
+[[clang::sycl_kernel(1)]] void foo2(); // expected-error {{'sycl_kernel' attribute takes no arguments}}
+
+// Only template functions
+__attribute__((sycl_kernel)) void foo(); // expected-warning {{'sycl_kernel' attribute only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+[[clang::sycl_kernel]] void foo1(); // expected-warning {{'sycl_kernel' attribute only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+
+// At least two template parameters
+template <typename T>
+__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{'sycl_kernel' attribute only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+template <typename T>
+[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{'sycl_kernel' attribute only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+
+// Both first two template parameters must be a typenames
+template <typename T, int A>
+__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{'sycl_kernel' attribute only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+template <typename T, int A>
+[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{'sycl_kernel' attribute only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+
+// Must return void
+template <typename T, typename A>
+__attribute__((sycl_kernel)) int foo(T P); // expected-warning {{'sycl_kernel' attribute only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+template <typename T, typename A>
+[[clang::sycl_kernel]] int foo1(T P); // expected-warning {{'sycl_kernel' attribute only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+
+// Must take at least one argument
+template <typename T, typename A>
+__attribute__((sycl_kernel)) void foo(); // expected-warning {{'sycl_kernel' attribute only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+template <typename T, typename A>
+[[clang::sycl_kernel]] void foo1(); // expected-warning {{'sycl_kernel' attribute only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+
+// No diagnosticts
+template <typename T, typename A, int B>
+__attribute__((sycl_kernel)) void foo(T P);
+template <typename T, typename A, int B>
+[[clang::sycl_kernel]] void foo1(T P);
Index: clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fsycl-is-device -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -x c++ %s
+
+#ifndef __SYCL_DEVICE_ONLY__
+// expected-warning@+7 {{'sycl_kernel' attribute ignored}}
+// expected-warning@+8 {{'sycl_kernel' attribute ignored}}
+#else
+// expected-no-diagnostics
+#endif
+
+template <typename T, typename A, int B>
+__attribute__((sycl_kernel)) void foo(T P);
+template <typename T, typename A, int B>
+[[clang::sycl_kernel]] void foo1(T P);
Index: clang/test/Misc/pragma-attribute-supported-attributes-list.test
===================================================================
--- clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -131,6 +131,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: SYCLKernel (SubjectMatchRule_function)
// CHECK-NEXT: ScopedLockable (SubjectMatchRule_record)
// CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property)
// CHECK-NEXT: SetTypestate (SubjectMatchRule_function_is_member)
Index: clang/test/CodeGenSYCL/device-functions.cpp
===================================================================
--- /dev/null
+++ clang/test/CodeGenSYCL/device-functions.cpp
@@ -0,0 +1,42 @@
+// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -S -emit-llvm %s -o - | FileCheck %s
+
+template <typename T>
+T bar(T arg);
+
+void foo() {
+ int a = 1 + 1 + bar(1);
+}
+
+template <typename T>
+T bar(T arg) {
+ return arg;
+}
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+ kernelFunc();
+}
+
+// Make sure that definitions for the types not used in SYCL kernels are not
+// emitted
+// CHECK-NOT: %struct.A
+// CHECK-NOT: @a = {{.*}} %struct.A
+struct A {
+ int x = 10;
+} a;
+
+int main() {
+ a.x = 8;
+ kernel_single_task<class test_kernel>([]() { foo(); });
+ return 0;
+}
+
+// baz is not called from the SYCL kernel, so it must not be emitted
+// CHECK-NOT: define {{.*}} @{{.*}}baz
+void baz() {}
+
+// FIXME: calling convention for kernel_single_task must be set to spir_kernel
+// CHECK-LABEL: define internal spir_func void @{{.*}}kernel_single_task
+// CHECK-LABEL: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%class.anon* %this)
+// CHECK-LABEL: define spir_func void @{{.*}}foo
+// CHECK-LABEL: define linkonce_odr spir_func i32 @{{.*}}bar
Index: clang/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- clang/lib/Sema/SemaDeclAttr.cpp
+++ clang/lib/Sema/SemaDeclAttr.cpp
@@ -6358,6 +6358,49 @@
D->addAttr(::new (S.Context) OpenCLAccessAttr(S.Context, AL));
}
+static void handleSYCLKernelAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+ // The 'sycl_kernel' attribute applies only to functions.
+ const auto *FD = cast<FunctionDecl>(D);
+ const FunctionTemplateDecl *FT = FD->getDescribedFunctionTemplate();
+
+ // Function template is expected
+ if (!FT) {
+ S.Diag(AL.getLoc(), diag::warn_sycl_kernel_attribute_invalid);
+ return;
+ }
+
+ // Function template must have at least two template parameters.
+ const TemplateParameterList *TL = FT->getTemplateParameters();
+ if (TL->size() < 2) {
+ S.Diag(AL.getLoc(), diag::warn_sycl_kernel_attribute_invalid);
+ return;
+ }
+
+ // Template parameters must be typenames.
+ for (unsigned I = 0; I < 2; ++I) {
+ const NamedDecl *TParam = TL->getParam(I);
+ if (!isa<TemplateTypeParmDecl>(TParam)) {
+ S.Diag(AL.getLoc(), diag::warn_sycl_kernel_attribute_invalid);
+ return;
+ }
+ }
+
+ // Function must have at least one argument.
+ if (getFunctionOrMethodNumParams(D) < 1) {
+ S.Diag(AL.getLoc(), diag::warn_sycl_kernel_attribute_invalid);
+ return;
+ }
+
+ // Function must return void.
+ QualType RetTy = getFunctionOrMethodResultType(D);
+ if (!RetTy->isVoidType()) {
+ S.Diag(AL.getLoc(), diag::warn_sycl_kernel_attribute_invalid);
+ return;
+ }
+
+ handleSimpleAttribute<SYCLKernelAttr>(S, D, AL);
+}
+
static void handleDestroyAttr(Sema &S, Decl *D, const ParsedAttr &A) {
if (!cast<VarDecl>(D)->hasGlobalStorage()) {
S.Diag(D->getLocation(), diag::err_destroy_attr_on_non_static_var)
@@ -6682,6 +6725,9 @@
case ParsedAttr::AT_Flatten:
handleSimpleAttribute<FlattenAttr>(S, D, AL);
break;
+ case ParsedAttr::AT_SYCLKernel:
+ handleSYCLKernelAttr(S, D, AL);
+ break;
case ParsedAttr::AT_Format:
handleFormatAttr(S, D, AL);
break;
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -2474,6 +2474,13 @@
}
}
+ if (LangOpts.SYCLIsDevice && Global->hasAttr<SYCLKernelAttr>()) {
+ // SYCL kernels can be templated and not called from anywhere in the
+ // module but should be emitted.
+ addDeferredDeclToEmit(GD);
+ return;
+ }
+
// Ignore declarations, they will be emitted on their first use.
if (const auto *FD = dyn_cast<FunctionDecl>(Global)) {
// Forward declarations are emitted lazily on first use.
Index: clang/lib/AST/ASTContext.cpp
===================================================================
--- clang/lib/AST/ASTContext.cpp
+++ clang/lib/AST/ASTContext.cpp
@@ -10023,6 +10023,10 @@
if (D->hasAttr<AliasAttr>() || D->hasAttr<UsedAttr>())
return true;
+ // If SYCL, only kernels are required.
+ if (LangOpts.SYCLIsDevice && !D->hasAttr<SYCLKernelAttr>())
+ return false;
+
if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
// Forward declarations aren't required.
if (!FD->doesThisDeclarationHaveABody())
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -10070,4 +10070,10 @@
"__builtin_bit_cast %select{source|destination}0 type must be trivially copyable">;
def err_bit_cast_type_size_mismatch : Error<
"__builtin_bit_cast source size does not equal destination size (%0 vs %1)">;
+
+// SYCL-specific diagnostics
+def warn_sycl_kernel_attribute_invalid : Warning<
+ "'sycl_kernel' attribute only applies to template funtions with special prototype, "
+ "please refer 'sycl_kernel' attribute documentation">, InGroup<IgnoredAttributes>;
+
} // end of sema component.
Index: clang/include/clang/Basic/AttrDocs.td
===================================================================
--- clang/include/clang/Basic/AttrDocs.td
+++ clang/include/clang/Basic/AttrDocs.td
@@ -253,6 +253,80 @@
}];
}
+def SYCLKernelDocs : Documentation {
+ let Category = DocCatFunction;
+ let Content = [{
+The ``sycl_kernel`` attribute specifies that a function will be used by the
+compiler to outline device code and to generate an OpenCL kernel.
+Here is a code example of the SYCL program, which demonstrates the compiler's
+outlining job:
+.. code-block:: c++
+
+ int foo(int x) { return ++x; }
+
+ using namespace cl::sycl;
+ queue Q;
+ buffer<int, 1> a(range<1>{1024});
+ Q.submit([&](handler& cgh) {
+ auto A = a.get_access<access::mode::write>(cgh);
+ cgh.parallel_for<init_a>(range<1>{1024}, [=](id<1> index) {
+ A[index] = index[0] * 2 + index[1] + foo(42);
+ });
+ }
+
+The lambda that is passed to the ``parallel_for`` is called a SYCL "kernel
+function". A SYCL "kernel function" defines the entry point to the "device
+part" of the code. The compiler will emit all symbols accessible from a "kernel
+function". In this code example, the compiler will emit "foo" function.
+More details about the compilation of functions for the device part can be
+found in the SYCL 1.2.1 specification Section 6.4.
+To show to the compiler entry point to the "device part" of the code, the SYCL
+runtime can use the ``sycl_kernel`` attribute in the following way:
+.. code-block:: c++
+namespace cl {
+namespace sycl {
+class handler {
+ template <typename KernelName, typename KernelType/*, ...*/>
+ __attribute__((sycl_kernel)) void sycl_kernel_function(KernelType KernelFuncObj) {
+ // ...
+ KernelFuncObj();
+ }
+
+ template <typename KernelName, typename KernelType, int Dims>
+ void parallel_for(range<Dims> NumWorkItems, KernelType KernelFunc) {
+#ifdef __SYCL_DEVICE_ONLY__
+ sycl_kernel_function<KernelName, KernelType, Dims>(KernelFunc);
+#else
+ // Host implementation
+#endif
+ }
+};
+} // namespace sycl
+} // namespace cl
+
+The compiler will also generate an OpenCL kernel using the function marked with
+the ``sycl_kernel`` attribute.
+Here is the list of SYCL device compiler expectations with regard to the
+function marked with the ``sycl_kernel`` attribute:
+
+- The function must be a template with at least two template parameters is
+ expected. The compiler generates an OpenCL kernel and uses the first template
+ parameter as a unique name to the generated OpenCL kernel. The host
+ application uses this unique name to invoke the OpenCL kernel generated for
+ the ``sycl_kernel_function`` specialized by this name and second template
+ parameter ``KernelType`` (which might be a lambda type).
+- The function must have at least one parameter. The first parameter is
+ required to be a function object type (named or unnamed i.e. lambda).
+ The compiler uses function object type fields to generate OpenCL kernel
+ parameters.
+- The function must return void. The compiler reuses the body of marked functions to
+ generate the OpenCL kernel body, and the OpenCL kernel must return `void`.
+
+The ``sycl_kernel_function`` in the previous code sample meets these
+expectations.
+ }];
+}
+
def C11NoReturnDocs : Documentation {
let Category = DocCatFunction;
let Content = [{
Index: clang/include/clang/Basic/Attr.td
===================================================================
--- clang/include/clang/Basic/Attr.td
+++ clang/include/clang/Basic/Attr.td
@@ -296,6 +296,7 @@
def Borland : LangOpt<"Borland">;
def CUDA : LangOpt<"CUDA">;
def HIP : LangOpt<"HIP">;
+def SYCL : LangOpt<"SYCLIsDevice">;
def COnly : LangOpt<"COnly", "!LangOpts.CPlusPlus">;
def CPlusPlus : LangOpt<"CPlusPlus">;
def OpenCL : LangOpt<"OpenCL">;
@@ -1055,6 +1056,13 @@
let Documentation = [Undocumented];
}
+def SYCLKernel : InheritableAttr {
+ let Spellings = [Clang<"sycl_kernel">];
+ let Subjects = SubjectList<[Function]>;
+ let LangOpts = [SYCL];
+ let Documentation = [SYCLKernelDocs];
+}
+
def C11NoReturn : InheritableAttr {
let Spellings = [Keyword<"_Noreturn">];
let Subjects = SubjectList<[Function], ErrorDiag>;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
