bader updated this revision to Diff 230310.
bader added a comment.
Applied code review comments from Aaron.
Allow template template parameters for function templates marked with
`sycl_kernel` attribute.
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/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 function templates}}
+__attribute__((sycl_kernel)) int gv3 = 0; // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
+
+__attribute__((sycl_kernel(1))) void foo(); // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
+[[clang::sycl_kernel(1)]] void foo2(); // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
+
+// Only function templates
+__attribute__((sycl_kernel)) void foo(); // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
+[[clang::sycl_kernel]] void foo1(); // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
+
+// At least two template parameters
+template <typename T>
+__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}}
+template <typename T>
+[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}}
+
+// First two template parameters can't be non-type template parameters
+template <typename T, int A>
+__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute can't be a non-type template parameter}}
+template <int A, typename T>
+[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute can't be a non-type template parameter}}
+
+// Must return void
+template <typename T, typename A>
+__attribute__((sycl_kernel)) int foo(T P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}}
+template <typename T, typename A>
+[[clang::sycl_kernel]] int foo1(T P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}}
+
+// Must take at least one argument
+template <typename T, typename A>
+__attribute__((sycl_kernel)) void foo(); // expected-warning {{function template with 'sycl_kernel' attribute must have a single parameter}}
+template <typename T, typename A>
+[[clang::sycl_kernel]] void foo1(T t, A a); // expected-warning {{function template with 'sycl_kernel' attribute must have a single parameter}}
+
+// No diagnostics
+template <typename T, typename A>
+__attribute__((sycl_kernel)) void foo(T P);
+template <typename T, typename A, int I>
+[[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/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,45 @@
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();
+ assert(FT && "Function template is expected");
+
+ // Function template must have at least two template parameters.
+ const TemplateParameterList *TL = FT->getTemplateParameters();
+ if (TL->size() < 2) {
+ S.Diag(FT->getLocation(), diag::warn_sycl_kernel_num_of_template_params);
+ return;
+ }
+
+ // Template parameters must be typenames.
+ for (unsigned I = 0; I < 2; ++I) {
+ const NamedDecl *TParam = TL->getParam(I);
+ if (isa<NonTypeTemplateParmDecl>(TParam)) {
+ S.Diag(FT->getLocation(),
+ diag::warn_sycl_kernel_invalid_template_param_type);
+ return;
+ }
+ }
+
+ // Function must have at least one argument.
+ if (getFunctionOrMethodNumParams(D) != 1) {
+ S.Diag(FT->getLocation(), diag::warn_sycl_kernel_num_of_function_params);
+ return;
+ }
+
+ // Function must return void.
+ QualType RetTy = getFunctionOrMethodResultType(D);
+ if (!RetTy->isVoidType()) {
+ S.Diag(FT->getLocation(), diag::warn_sycl_kernel_return_type);
+ 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 +6721,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,19 @@
"__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_num_of_template_params : Warning<
+ "'sycl_kernel' attribute only applies to a function template with at least"
+ " two template parameters">, InGroup<IgnoredAttributes>;
+def warn_sycl_kernel_invalid_template_param_type : Warning<
+ "template parameter of a function template with the 'sycl_kernel' attribute"
+ " can't be a non-type template parameter">, InGroup<IgnoredAttributes>;
+def warn_sycl_kernel_num_of_function_params : Warning<
+ "function template with 'sycl_kernel' attribute must have a single parameter">,
+ InGroup<IgnoredAttributes>;
+def warn_sycl_kernel_return_type : Warning<
+ "function template with 'sycl_kernel' attribute must have a 'void' return type">,
+ 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
@@ -121,6 +121,11 @@
def InlineFunction : SubsetSubject<Function,
[{S->isInlineSpecified()}], "inline functions">;
+def FunctionTmpl
+ : SubsetSubject<Function, [{S->getTemplatedKind() ==
+ FunctionDecl::TK_FunctionTemplate}],
+ "function templates">;
+
// FIXME: this hack is needed because DeclNodes.td defines the base Decl node
// type to be a class, not a definition. This makes it impossible to create an
// attribute subject which accepts a Decl. Normally, this is not a problem,
@@ -296,6 +301,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 +1061,13 @@
let Documentation = [Undocumented];
}
+def SYCLKernel : InheritableAttr {
+ let Spellings = [Clang<"sycl_kernel">];
+ let Subjects = SubjectList<[FunctionTmpl]>;
+ let LangOpts = [SYCL];
+ let Documentation = [SYCLKernelDocs];
+}
+
def C11NoReturn : InheritableAttr {
let Spellings = [Keyword<"_Noreturn">];
let Subjects = SubjectList<[Function], ErrorDiag>;
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
