Fznamznon updated this revision to Diff 203785.
Fznamznon added a comment.
Applied comments from @Anastasia
- Added link to documentation for `sycl_device` attribute
- Removed redundant comment from test
@Anastasia, do you have additional comments?
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
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
+
+#if not defined(__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
@@ -5517,14 +5517,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 SYCL kernel attribute after instantination.
+ if (getLangOpts().SYCLIsDevice &&
+ CurFD->hasAttr<SYCLKernelAttr>()) {
+ ConstructSYCLKernel(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 SYCL kernel attribute after instantination.
+ if (getLangOpts().SYCLIsDevice &&
+ Function->hasAttr<SYCLKernelAttr>()) {
+ ConstructSYCLKernel(Function);
+ }
Function->setInstantiationIsPending(false);
+ }
}
continue;
}
Index: clang/lib/Sema/SemaSYCL.cpp
===================================================================
--- /dev/null
+++ clang/lib/Sema/SemaSYCL.cpp
@@ -0,0 +1,80 @@
+//===- 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"
+
+#include <array>
+
+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 auto *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::ConstructSYCLKernel(FunctionDecl *KernelCallerFunc) {
+ AddSyclDeviceFunc(KernelCallerFunc);
+}
+
+void Sema::MarkDevice(void) {
+ // Let's mark all called functions with SYCL Device attribute.
+ MarkDeviceFunction Marker(*this);
+ for (const auto &elt : SyclDeviceFuncs()) {
+ if (FunctionDecl *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
@@ -904,6 +904,9 @@
PerformPendingInstantiations();
}
+ if (getLangOpts().SYCLIsDevice)
+ MarkDevice();
+
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
@@ -51,6 +51,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
@@ -2407,6 +2407,13 @@
if (Global->hasAttr<IFuncAttr>())
return emitIFuncDefinition(GD);
+ // If this is SYCL device, only emit declarations marked with SYCL device
+ // attribute.
+ if (LangOpts.SYCLIsDevice) {
+ if (!Global->hasAttr<SYCLDeviceAttr>())
+ return;
+ }
+
// If this is a cpu_dispatch multiversion function, emit the resolver.
if (Global->hasAttr<CPUDispatchAttr>())
return emitCPUDispatchDefinition(GD);
@@ -2521,6 +2528,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
@@ -11177,6 +11177,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 function declaration to the SYCL device code.
+ void AddSyclDeviceFunc(Decl *D) {
+ D->addAttr(SYCLDeviceAttr::CreateImplicit(Context));
+ SyclDeviceFunctions.push_back(D);
+ }
+ /// SyclDeviceFuncs - access to SYCL device function decls.
+ SmallVector<Decl *, 4> &SyclDeviceFuncs() { return SyclDeviceFunctions; }
+
+ /// Constructs SYCL kernel which is compatible with OpenCL from SYCL "kernel
+ /// function" and adds it to the SYCL device code.
+ void ConstructSYCLKernel(FunctionDecl *KernelCallerFunc);
+ /// This function marks all functions accessible from SYCL kernels with SYCL
+ /// device attribute and adds them to the SYCL device code.
+ void MarkDevice(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,41 @@
}];
}
+def SYCLKernelDocs : Documentation {
+ let Category = DocCatFunction;
+ let Content = [{
+The ``sycl_kernel`` attribute specifies that a function is SYCL "kernel
+function". SYCL "kernel function" defines an entry point to a kernel.
+Kernel is a function which will be compiled for the device and defines some
+entry point to device code i.e. will be called by host in run time.
+Here is a code example of the SYCL program, which demonstrates the need for
+this attribute:
+.. 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 SYCL
+"kernel function".
+The SYCL Runtime implementation will use ``sycl_kernel`` attribute to mark this
+lambda as SYCL "kernel function". Compiler is supposed to construct a kernel
+from "kernel function", add it to the "device part" of code and traverse all
+symbols accessible from "kernel function" and add them to the "device part" of
+the code. In this code example compiler is supposed to add "foo" function to the
+"device part" of the code.
+More details can be found in the SYCL 1.2.1 specification, Sections 4.8.9 and
+6.4.
+ }];
+}
+
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
