Fznamznon updated this revision to Diff 205831.
Fznamznon added a comment.
Fixed a couple coding style issues, renamed markDevice function with
markSYCLDevice.
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/Sema/Sema.h
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/Parse/ParseAST.cpp
clang/lib/Sema/CMakeLists.txt
clang/lib/Sema/Sema.cpp
clang/lib/Sema/SemaDeclAttr.cpp
clang/lib/Sema/SemaSYCL.cpp
clang/lib/Sema/SemaTemplateInstantiateDecl.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
clang/test/SemaSYCL/device-code-outlining.cpp
Index: clang/test/SemaSYCL/device-code-outlining.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/device-code-outlining.cpp
@@ -0,0 +1,37 @@
+// RUN: %clang_cc1 -std=c++11 -fsycl-is-device -ast-dump %s | FileCheck %s
+
+template <typename T>
+T bar(T arg);
+// CHECK: FunctionTemplateDecl {{.*}} bar
+// CHECK: SYCLDeviceAttr {{.*}} Implicit
+
+void foo() {
+ int a = 1 + 1 + bar(1);
+}
+// CHECK: FunctionDecl {{.*}} foo
+// CHECK: SYCLDeviceAttr {{.*}} Implicit
+
+template <typename T>
+T bar(T arg) {
+ return arg;
+}
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+ kernelFunc();
+}
+// CHECK: FunctionTemplateDecl {{.*}} kernel_single_task
+// CHECK: SYCLDeviceAttr {{.*}} Implicit
+
+void host_foo() {
+ int b = 0;
+}
+// CHECK: FunctionDecl {{.*}} host_foo
+// CHECK-NOT: SYCLDeviceAttr
+// CHECK: FunctionDecl {{.*}} main
+
+int main() {
+ kernel_single_task<class fake_kernel>([]() { foo(); });
+ host_foo();
+ return 0;
+}
Index: clang/test/SemaSYCL/device-attributes.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/device-attributes.cpp
@@ -0,0 +1,10 @@
+// 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)) void foo();
+[[clang::sycl_kernel]] void foo1();
+
+__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}}
Index: clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
@@ -0,0 +1,13 @@
+// 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@+6 {{'sycl_kernel' attribute ignored}}
+// expected-warning@+6 {{'sycl_kernel' attribute ignored}}
+#else
+// expected-no-diagnostics
+#endif
+
+__attribute__((sycl_kernel)) void foo();
+[[clang::sycl_kernel]] void foo2();
+
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
@@ -125,6 +125,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,29 @@
+// RUN: %clang_cc1 -triple spir64-unknown-unknown -std=c++11 -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();
+}
+
+int main() {
+ kernel_single_task<class fake_kernel>([]() { foo(); });
+ return 0;
+}
+// CHECK: define spir_func void @{{.*}}foo
+// CHECK: define linkonce_odr spir_func i32 @{{.*}}bar
+// CHECK: define internal spir_func void @{{.*}}kernel_single_task
+// FIXME: Next function is lambda () operator. spir_func calling convention
+// is missed for C++ methods.
+// CHECK: define internal void @"_ZZ4mainENK3$_0clEv"(%class.anon* %this)
Index: clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
===================================================================
--- clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -5532,14 +5532,30 @@
Function, [this, Inst, DefinitionRequired](FunctionDecl *CurFD) {
InstantiateFunctionDefinition(/*FIXME:*/ Inst.second, CurFD, true,
DefinitionRequired, true);
- if (CurFD->isDefined())
+ if (CurFD->isDefined()) {
+ // Because all SYCL kernel functions are template functions -
+ // they have deferred instantination. We need bodies of these
+ // functions so we are checking for the SYCL kernel attribute
+ // after instantination.
+ if (getLangOpts().SYCLIsDevice &&
+ CurFD->hasAttr<SYCLKernelAttr>())
+ constructOpenCLKernel(CurFD);
CurFD->setInstantiationIsPending(false);
+ }
});
} else {
InstantiateFunctionDefinition(/*FIXME:*/ Inst.second, Function, true,
DefinitionRequired, true);
- if (Function->isDefined())
+ if (Function->isDefined()) {
+ // Because all SYCL kernel functions are template functions - they
+ // have deferred instantination. We need bodies of these functions
+ // so we are checking for the SYCL kernel attribute after
+ // instantination.
+ if (getLangOpts().SYCLIsDevice &&
+ Function->hasAttr<SYCLKernelAttr>())
+ constructOpenCLKernel(Function);
Function->setInstantiationIsPending(false);
+ }
}
continue;
}
Index: clang/lib/Sema/SemaSYCL.cpp
===================================================================
--- /dev/null
+++ clang/lib/Sema/SemaSYCL.cpp
@@ -0,0 +1,74 @@
+//===- SemaSYCL.cpp - Semantic Analysis for SYCL constructs ---------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// This implements Semantic Analysis for SYCL constructs.
+//===----------------------------------------------------------------------===//
+
+#include "clang/AST/RecursiveASTVisitor.h"
+#include "clang/Sema/Sema.h"
+
+using namespace clang;
+
+class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
+public:
+ MarkDeviceFunction(Sema &S)
+ : RecursiveASTVisitor<MarkDeviceFunction>(), SemaRef(S) {}
+
+ bool VisitCallExpr(CallExpr *E) {
+ if (FunctionDecl *Callee = E->getDirectCallee()) {
+ Callee = Callee->getCanonicalDecl();
+ // Remember that all SYCL kernel functions have deferred
+ // instantiation as template functions. It means that
+ // all functions used by kernel have already been parsed and have
+ // definitions.
+ if (FunctionDecl *Def = Callee->getDefinition()) {
+ if (!Def->hasAttr<SYCLDeviceAttr>()) {
+ SemaRef.addSyclDeviceFunc(Def);
+ this->TraverseStmt(Def->getBody());
+ }
+ }
+ }
+ return true;
+ }
+
+ bool VisitCXXConstructExpr(CXXConstructExpr *E) {
+ CXXConstructorDecl *Ctor = E->getConstructor();
+
+ if (FunctionDecl *Def = Ctor->getDefinition())
+ SemaRef.addSyclDeviceFunc(Def);
+
+ const CXXRecordDecl *ConstructedType = Ctor->getParent();
+ if (ConstructedType->hasUserDeclaredDestructor()) {
+ CXXDestructorDecl *Dtor = ConstructedType->getDestructor();
+
+ if (FunctionDecl *Def = Dtor->getDefinition())
+ SemaRef.addSyclDeviceFunc(Def);
+ }
+ return true;
+ }
+
+private:
+ Sema &SemaRef;
+};
+
+void Sema::constructOpenCLKernel(FunctionDecl *KernelCallerFunc) {
+ addSyclDeviceFunc(KernelCallerFunc);
+}
+
+void Sema::markSYCLDevice(void) {
+ // Let's mark all called functions with the SYCL Device attribute.
+ MarkDeviceFunction Marker(*this);
+ for (const auto &Elt : syclDeviceFuncs()) {
+ if (auto *Func = dyn_cast<FunctionDecl>(Elt)) {
+ if (FunctionDecl *Def = Func->getDefinition()) {
+ if (!Def->hasAttr<SYCLDeviceAttr>())
+ addSyclDeviceFunc(Def);
+ Marker.TraverseStmt(Def->getBody());
+ }
+ }
+ }
+}
Index: clang/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- clang/lib/Sema/SemaDeclAttr.cpp
+++ clang/lib/Sema/SemaDeclAttr.cpp
@@ -6767,6 +6767,9 @@
case ParsedAttr::AT_Flatten:
handleSimpleAttribute<FlattenAttr>(S, D, AL);
break;
+ case ParsedAttr::AT_SYCLKernel:
+ handleSimpleAttribute<SYCLKernelAttr>(S, D, AL);
+ break;
case ParsedAttr::AT_Format:
handleFormatAttr(S, D, AL);
break;
Index: clang/lib/Sema/Sema.cpp
===================================================================
--- clang/lib/Sema/Sema.cpp
+++ clang/lib/Sema/Sema.cpp
@@ -905,6 +905,9 @@
PerformPendingInstantiations();
}
+ if (getLangOpts().SYCLIsDevice)
+ markSYCLDevice();
+
assert(LateParsedInstantiations.empty() &&
"end of TU template instantiation should not create more "
"late-parsed templates");
Index: clang/lib/Sema/CMakeLists.txt
===================================================================
--- clang/lib/Sema/CMakeLists.txt
+++ clang/lib/Sema/CMakeLists.txt
@@ -56,6 +56,7 @@
SemaStmt.cpp
SemaStmtAsm.cpp
SemaStmtAttr.cpp
+ SemaSYCL.cpp
SemaTemplate.cpp
SemaTemplateDeduction.cpp
SemaTemplateInstantiate.cpp
Index: clang/lib/Parse/ParseAST.cpp
===================================================================
--- clang/lib/Parse/ParseAST.cpp
+++ clang/lib/Parse/ParseAST.cpp
@@ -168,6 +168,12 @@
for (Decl *D : S.WeakTopLevelDecls())
Consumer->HandleTopLevelDecl(DeclGroupRef(D));
+ if (S.getLangOpts().SYCLIsDevice) {
+ for (Decl *D : S.syclDeviceFuncs()) {
+ Consumer->HandleTopLevelDecl(DeclGroupRef(D));
+ }
+ }
+
Consumer->HandleTranslationUnit(S.getASTContext());
// Finalize the template instantiation observer chain.
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -2405,6 +2405,11 @@
if (Global->hasAttr<IFuncAttr>())
return emitIFuncDefinition(GD);
+ // If this is SYCL device, only emit declarations marked with the SYCL device
+ // attribute.
+ if (LangOpts.SYCLIsDevice && !Global->hasAttr<SYCLDeviceAttr>())
+ return;
+
// If this is a cpu_dispatch multiversion function, emit the resolver.
if (Global->hasAttr<CPUDispatchAttr>())
return emitCPUDispatchDefinition(GD);
@@ -2519,6 +2524,10 @@
// The value must be emitted, but cannot be emitted eagerly.
assert(!MayBeEmittedEagerly(Global));
addDeferredDeclToEmit(GD);
+ } else if (LangOpts.SYCLIsDevice) {
+ // SYCL kernels can be templated and not called from anywhere in the
+ // module but should be emitted.
+ addDeferredDeclToEmit(GD);
} else {
// Otherwise, remember that we saw a deferred decl with this name. The
// first use of the mangled name will cause it to move into
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -11192,6 +11192,29 @@
ConstructorDestructor,
BuiltinFunction
};
+
+private:
+ /// Contains function declarations to be added to the SYCL device code.
+ /// In SYCL, when we generate device code, we don't know which functions we
+ /// will emit before we emit sycl kernels, so we add device functions to this
+ /// array and handle it in separate way.
+ SmallVector<Decl *, 4> SyclDeviceFunctions;
+
+public:
+ /// This function adds the function declaration to the SYCL device code.
+ void addSyclDeviceFunc(Decl *D) {
+ D->addAttr(SYCLDeviceAttr::CreateImplicit(Context));
+ SyclDeviceFunctions.push_back(D);
+ }
+ /// Access to SYCL device function decls.
+ SmallVectorImpl<Decl *> &syclDeviceFuncs() { return SyclDeviceFunctions; }
+
+ /// Constructs an OpenCL kernel using the KernelCaller function and adds it to
+ /// the SYCL device code.
+ void constructOpenCLKernel(FunctionDecl *KernelCallerFunc);
+ /// This function marks all functions accessible from SYCL kernels with the
+ /// SYCL device attribute and adds them to the SYCL device code.
+ void markSYCLDevice(void);
};
/// RAII object that enters a new expression evaluation context.
Index: clang/include/clang/Basic/AttrDocs.td
===================================================================
--- clang/include/clang/Basic/AttrDocs.td
+++ clang/include/clang/Basic/AttrDocs.td
@@ -253,6 +253,79 @@
}];
}
+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 OpenCL kernel.
+Here is a code example of the SYCL program, which demonstrates 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 entry point to the "device part"
+of the code. Compiler will traverse all symbols accessible from a
+"kernel function" and add them to the "device part" of the code. In this code
+example, compiler will add "foo" function to the "device part" of the code.
+More details about compilation of functions for device 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 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 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:
+
+- Function template with at least one parameter is expected. The compiler
+generates OpenCL kernel and uses first template parameter as unique name to the
+generated OpenCL kernel. Host application uses this unique name to invoke the
+OpenCL kernel generated for the ``sycl_kernel_function`` specialized by
+this name and ``KernelType`` (which might be a lambda type).
+- Function must have at least one parameter. First parameter expected to be a
+function object type (named or unnamed i.e. lambda). Compiler uses function
+object type fields to generate OpenCL kernel parameters.
+- Function must return void. Compiler re-uses body of marked function to
+generate OpenCL kernel body and 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
@@ -294,6 +294,7 @@
def MicrosoftExt : LangOpt<"MicrosoftExt">;
def Borland : LangOpt<"Borland">;
def CUDA : LangOpt<"CUDA">;
+def SYCL : LangOpt<"SYCLIsDevice">;
def COnly : LangOpt<"COnly", "!LangOpts.CPlusPlus">;
def CPlusPlus : LangOpt<"CPlusPlus">;
def OpenCL : LangOpt<"OpenCL">;
@@ -1007,6 +1008,20 @@
let Documentation = [Undocumented];
}
+def SYCLDevice : InheritableAttr {
+ let Spellings = [];
+ let Subjects = SubjectList<[Function, Var]>;
+ let LangOpts = [SYCL];
+ 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
