Fznamznon updated this revision to Diff 233013.
Fznamznon added a comment.
Updated tests using address space attributes added by D71005
<https://reviews.llvm.org/D71005>
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D71016/new/
https://reviews.llvm.org/D71016
Files:
clang/include/clang/Sema/Sema.h
clang/lib/AST/ASTContext.cpp
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/Parse/ParseAST.cpp
clang/lib/Sema/CMakeLists.txt
clang/lib/Sema/SemaSYCL.cpp
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
clang/test/CodeGenSYCL/Inputs/sycl.hpp
clang/test/CodeGenSYCL/basic-opencl-kernel.cpp
clang/test/CodeGenSYCL/device-functions.cpp
clang/test/SemaSYCL/Inputs/sycl.hpp
clang/test/SemaSYCL/accessors-targets.cpp
clang/test/SemaSYCL/basic-opencl-kernel.cpp
clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
clang/test/SemaSYCL/fake-accessors.cpp
clang/test/SemaSYCL/mangle-kernel.cpp
Index: clang/test/SemaSYCL/mangle-kernel.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/mangle-kernel.cpp
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 -triple spir64-unknown-unknown-unknown -I %S/Inputs -I %S/../Headers/Inputs/include/ -fsycl-is-device -ast-dump %s | FileCheck %s --check-prefix=CHECK-64
+// RUN: %clang_cc1 -triple spir-unknown-unknown-unknown -I %S/Inputs -I %S/../Headers/Inputs/include/ -fsycl-is-device -ast-dump %s | FileCheck %s --check-prefix=CHECK-32
+#include <sycl.hpp>
+#include <stdlib.h>
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+ kernelFunc();
+}
+
+template <typename T>
+class SimpleVadd;
+
+int main() {
+ kernel<class SimpleVadd<int>>(
+ [=](){});
+
+ kernel<class SimpleVadd<double>>(
+ [=](){});
+
+ kernel<class SimpleVadd<size_t>>(
+ [=](){});
+ return 0;
+}
+
+// CHECK: _ZTS10SimpleVaddIiE
+// CHECK: _ZTS10SimpleVaddIdE
+// CHECK-64: _ZTS10SimpleVaddImE
+// CHECK-32: _ZTS10SimpleVaddIjE
Index: clang/test/SemaSYCL/fake-accessors.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/fake-accessors.cpp
@@ -0,0 +1,56 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl-is-device -ast-dump %s | FileCheck %s
+
+#include <sycl.hpp>
+
+namespace foo {
+namespace cl {
+namespace sycl {
+class accessor {
+public:
+ int field;
+};
+} // namespace sycl
+} // namespace cl
+} // namespace foo
+
+class accessor {
+public:
+ int field;
+};
+
+typedef cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
+ cl::sycl::access::target::global_buffer>
+ MyAccessorTD;
+
+using MyAccessorA = cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
+ cl::sycl::access::target::global_buffer>;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+ kernelFunc();
+}
+
+int main() {
+ foo::cl::sycl::accessor acc = {1};
+ accessor acc1 = {1};
+
+ cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
+ cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorB;
+ cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorC;
+ kernel<class fake_accessors>(
+ [=]() {
+ accessorA.use((void*)(acc.field + acc1.field));
+ });
+ kernel<class accessor_typedef>(
+ [=]() {
+ accessorB.use((void*)(acc.field + acc1.field));
+ });
+ kernel<class accessor_alias>(
+ [=]() {
+ accessorC.use((void*)(acc.field + acc1.field));
+ });
+ return 0;
+}
+// CHECK: fake_accessors 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: accessor_typedef 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: accessor_alias 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
Index: clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
@@ -0,0 +1,70 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl-is-device -ast-dump %s | FileCheck %s
+
+// This test checks that compiler generates correct initialization for arguments
+// that have struct or built-in type inside the OpenCL kernel
+
+#include <sycl.hpp>
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+ kernelFunc();
+}
+
+struct test_struct {
+ int data;
+};
+
+void test(const int some_const) {
+ kernel<class kernel_const>(
+ [=]() {
+ int a = some_const;
+ });
+}
+
+int main() {
+ int data = 5;
+ test_struct s;
+ s.data = data;
+ kernel<class kernel_int>(
+ [=]() {
+ int kernel_data = data;
+ });
+ kernel<class kernel_struct>(
+ [=]() {
+ test_struct k_s;
+ k_s = s;
+ });
+ const int some_const = 10;
+ test(some_const);
+ return 0;
+}
+// Check kernel parameters
+// CHECK: FunctionDecl {{.*}}kernel_const{{.*}} 'void (const int)'
+// CHECK: ParmVarDecl {{.*}} used _arg_ 'const int'
+
+// Check that lambda field of const built-in type is initialized
+// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
+// CHECK-NEXT: InitListExpr
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'const int' lvalue ParmVar {{.*}} '_arg_' 'const int'
+
+// Check kernel parameters
+// CHECK: {{.*}}kernel_int{{.*}} 'void (int)'
+// CHECK: ParmVarDecl {{.*}} used _arg_ 'int'
+
+// Check that lambda field of built-in type is initialized
+// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
+// CHECK-NEXT: InitListExpr
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
+
+// Check kernel parameters
+// CHECK: {{.*}}kernel_struct{{.*}} 'void (test_struct)'
+// CHECK: ParmVarDecl {{.*}} used _arg_ 'test_struct'
+
+// Check that lambda field of struct type is initialized
+// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
+// CHECK-NEXT: InitListExpr
+// CHECK-NEXT: CXXConstructExpr {{.*}}'test_struct'{{.*}}void (const test_struct &)
+// CHECK-NEXT: ImplicitCastExpr {{.*}}'const test_struct' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'test_struct' lvalue ParmVar {{.*}} '_arg_' 'test_struct'
Index: clang/test/SemaSYCL/basic-opencl-kernel.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/basic-opencl-kernel.cpp
@@ -0,0 +1,74 @@
+// RUN: %clang_cc1 -std=c++11 -I %S/Inputs -fsycl-is-device -ast-dump %s | FileCheck %s
+
+// This test checks that compiler generates correct OpenCL kernel for basic
+// case.
+
+#include <sycl.hpp>
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+ kernelFunc();
+}
+
+int main() {
+ cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> acc;
+ kernel<class kernel>(
+ [=]() {
+ acc.use();
+ });
+}
+
+// Check declaration of the kernel
+
+// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+
+// Check parameters of the kernel
+
+// CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global int *'
+// CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
+// CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
+// CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>'
+
+// Check body of the kernel
+
+// Check lambda declaration inside the kernel
+
+// CHECK: DeclStmt
+// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}basic-opencl-kernel.cpp{{.*}})'
+
+// Check accessor initialization
+
+// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write>':'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-opencl-kernel.cpp{{.*}})' lvalue Var
+
+// CHECK-NEXT: ImplicitCastExpr {{.*}} <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__global int *'
+
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_AccessRange]]' 'cl::sycl::range<1>'
+
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_MemRange]]' 'cl::sycl::range<1>'
+
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'id<1>':'cl::sycl::id<1>'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::id<1>' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::id<1>' lvalue ParmVar {{.*}} '[[_arg_Offset]]' 'cl::sycl::id<1>'
+
+// Check that body of the kernel caller function is included into kernel
+
+// CHECK: CompoundStmt {{.*}}
+// CHECK-NEXT: CXXOperatorCallExpr {{.*}} 'void'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)() const' <FunctionToPointerDecay>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}basic-opencl-kernel.cpp{{.*}})' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-opencl-kernel.cpp{{.*}})' lvalue Var
+
+// Check kernel's attributes
+
+// CHECK: OpenCLKernelAttr {{.*}} Implicit
+// CHECK: AsmLabelAttr {{.*}} Implicit "{{.*}}kernel{{.*}}"
+// CHECK: ArtificialAttr {{.*}} Implicit
Index: clang/test/SemaSYCL/accessors-targets.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/accessors-targets.cpp
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl-is-device -ast-dump %s | FileCheck %s
+
+// This test checks that compiler generates correct OpenCL kernel arguments for
+// different accessors targets.
+
+#include <sycl.hpp>
+
+using namespace cl::sycl;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+ kernelFunc();
+}
+
+int main() {
+
+ accessor<int, 1, access::mode::read_write,
+ access::target::local>
+ local_acc;
+ accessor<int, 1, access::mode::read_write,
+ access::target::global_buffer>
+ global_acc;
+ accessor<int, 1, access::mode::read_write,
+ access::target::constant_buffer>
+ constant_acc;
+ kernel<class use_local>(
+ [=]() {
+ local_acc.use();
+ });
+ kernel<class use_global>(
+ [=]() {
+ global_acc.use();
+ });
+ kernel<class use_constant>(
+ [=]() {
+ constant_acc.use();
+ });
+}
+// CHECK: {{.*}}use_local 'void (__local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: {{.*}}use_global 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: {{.*}}use_constant 'void (__constant int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
Index: clang/test/SemaSYCL/Inputs/sycl.hpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -0,0 +1,87 @@
+#pragma once
+
+namespace cl {
+namespace sycl {
+namespace access {
+
+enum class target {
+ global_buffer = 2014,
+ constant_buffer,
+ local,
+ image,
+ host_buffer,
+ host_image,
+ image_array
+};
+
+enum class mode {
+ read = 1024,
+ write,
+ read_write,
+ discard_write,
+ discard_read_write,
+ atomic
+};
+
+enum class placeholder { false_t,
+ true_t };
+
+enum class address_space : int {
+ private_space = 0,
+ global_space,
+ constant_space,
+ local_space
+};
+} // namespace access
+
+template <int dim>
+struct range {
+};
+
+template <int dim>
+struct id {
+};
+
+template <int dim>
+struct _ImplT {
+ range<dim> AccessRange;
+ range<dim> MemRange;
+ id<dim> Offset;
+};
+
+template <typename dataT, access::target accessTarget>
+struct DeviceValueType;
+
+template <typename dataT>
+struct DeviceValueType<dataT, access::target::global_buffer> {
+ using type = __attribute__((opencl_global)) dataT;
+};
+
+template <typename dataT>
+struct DeviceValueType<dataT, access::target::constant_buffer> {
+ using type = __attribute__((opencl_constant)) dataT;
+};
+
+template <typename dataT>
+struct DeviceValueType<dataT, access::target::local> {
+ using type = __attribute__((opencl_local)) dataT;
+};
+
+template <typename dataT, int dimensions, access::mode accessmode,
+ access::target accessTarget = access::target::global_buffer,
+ access::placeholder isPlaceholder = access::placeholder::false_t>
+class accessor {
+
+public:
+ void use(void) const {}
+ void use(void *) const {}
+ _ImplT<dimensions> impl;
+
+private:
+ using PtrType = typename DeviceValueType<dataT, accessTarget>::type *;
+ void __init(PtrType Ptr, range<dimensions> AccessRange,
+ range<dimensions> MemRange, id<dimensions> Offset) {}
+};
+
+} // namespace sycl
+} // namespace cl
Index: clang/test/CodeGenSYCL/device-functions.cpp
===================================================================
--- /dev/null
+++ clang/test/CodeGenSYCL/device-functions.cpp
@@ -0,0 +1,41 @@
+// 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() {}
+
+// CHECK-LABEL: define spir_kernel void @{{.*}}test_kernel
+// 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/test/CodeGenSYCL/basic-opencl-kernel.cpp
===================================================================
--- /dev/null
+++ clang/test/CodeGenSYCL/basic-opencl-kernel.cpp
@@ -0,0 +1,52 @@
+// RUN: %clang_cc1 -I %S/Inputs -triple spir64-unknown-unknown -std=c++11 -fsycl-is-device -S -emit-llvm %s -o - | FileCheck %s
+
+// This test checks that compiler generates correct opencl kernel for basic
+// case.
+
+#include "sycl.hpp"
+
+template <typename Name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+ kernelFunc();
+}
+
+int main() {
+ cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
+ kernel<class kernel_function>(
+ [=]() {
+ accessorA.use();
+ });
+ return 0;
+}
+
+// CHECK: define spir_kernel void @{{.*}}kernel_function
+// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+_1]],
+// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+_2]],
+// CHECK-SAME: %"struct.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]])
+// Check alloca for pointer argument
+// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)*
+// Check lambda object alloca
+// CHECK: [[ANON:%[0-9]+]] = alloca %class.anon
+// Check allocas for ranges
+// CHECK: [[ARANGE:%agg.tmp.*]] = alloca %"struct.cl::sycl::range"
+// CHECK: [[MRANGE:%agg.tmp.*]] = alloca %"struct.cl::sycl::range"
+// CHECK: [[OID:%agg.tmp.*]] = alloca %"struct.cl::sycl::id"
+//
+// Check store of kernel pointer argument to alloca
+// CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)** [[MEM_ARG]].addr, align 8
+
+// Check for default constructor of accessor
+// CHECK: call spir_func {{.*}}accessor
+
+// Check accessor GEP
+// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon* [[ANON]], i32 0, i32 0
+
+// Check load from kernel pointer argument alloca
+// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG]].addr
+
+// Check accessor __init method call
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
+
+// Check lambda "()" operator call
+// CHECK-OLD: call spir_func void @{{.*}}(%class.anon* [[ANON]])
Index: clang/test/CodeGenSYCL/Inputs/sycl.hpp
===================================================================
--- /dev/null
+++ clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -0,0 +1,86 @@
+#pragma once
+
+namespace cl {
+namespace sycl {
+namespace access {
+
+enum class target {
+ global_buffer = 2014,
+ constant_buffer,
+ local,
+ image,
+ host_buffer,
+ host_image,
+ image_array
+};
+
+enum class mode {
+ read = 1024,
+ write,
+ read_write,
+ discard_write,
+ discard_read_write,
+ atomic
+};
+
+enum class placeholder {
+ false_t,
+ true_t
+};
+
+enum class address_space : int {
+ private_space = 0,
+ global_space,
+ constant_space,
+ local_space
+};
+} // namespace access
+
+template <int dim>
+struct id {
+ template <typename... T>
+ id(T... args) {} // fake constructor
+private:
+ // Some fake field added to see using of id arguments in the
+ // kernel wrapper
+ int Data;
+};
+
+template <int dim>
+struct range {
+ template <typename... T>
+ range(T... args) {} // fake constructor
+private:
+ // Some fake field added to see using of range arguments in the
+ // kernel wrapper
+ int Data;
+};
+
+template <int dim>
+struct _ImplT {
+ range<dim> AccessRange;
+ range<dim> MemRange;
+ id<dim> Offset;
+};
+
+template <typename dataT, int dimensions, access::mode accessmode,
+ access::target accessTarget = access::target::global_buffer,
+ access::placeholder isPlaceholder = access::placeholder::false_t>
+class accessor {
+
+public:
+ void use(void) const {}
+ template <typename... T>
+ void use(T... args) {}
+ template <typename... T>
+ void use(T... args) const {}
+ _ImplT<dimensions> impl;
+
+private:
+ void __init(__attribute__((opencl_global)) dataT *Ptr,
+ range<dimensions> AccessRange,
+ range<dimensions> MemRange, id<dimensions> Offset) {}
+};
+
+} // namespace sycl
+} // namespace cl
Index: clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
===================================================================
--- clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -17,6 +17,7 @@
#include "clang/AST/DependentDiagnostic.h"
#include "clang/AST/Expr.h"
#include "clang/AST/ExprCXX.h"
+#include "clang/AST/Mangle.h"
#include "clang/AST/PrettyDeclStackTrace.h"
#include "clang/AST/TypeLoc.h"
#include "clang/Sema/Initialization.h"
@@ -5612,6 +5613,8 @@
/// Performs template instantiation for all implicit template
/// instantiations we have seen until this point.
void Sema::PerformPendingInstantiations(bool LocalOnly) {
+ std::unique_ptr<MangleContext> MangleCtx(
+ getASTContext().createMangleContext());
while (!PendingLocalImplicitInstantiations.empty() ||
(!LocalOnly && !PendingInstantiations.empty())) {
PendingImplicitInstantiation Inst;
@@ -5630,17 +5633,25 @@
TSK_ExplicitInstantiationDefinition;
if (Function->isMultiVersion()) {
getASTContext().forEachMultiversionedFunctionVersion(
- Function, [this, Inst, DefinitionRequired](FunctionDecl *CurFD) {
+ Function, [this, Inst, DefinitionRequired,
+ MangleCtx = move(MangleCtx)](FunctionDecl *CurFD) {
InstantiateFunctionDefinition(/*FIXME:*/ Inst.second, CurFD, true,
DefinitionRequired, true);
- if (CurFD->isDefined())
+ if (CurFD->isDefined()) {
CurFD->setInstantiationIsPending(false);
+ if (getLangOpts().SYCLIsDevice &&
+ CurFD->hasAttr<SYCLKernelAttr>())
+ constructOpenCLKernel(CurFD, *MangleCtx);
+ }
});
} else {
InstantiateFunctionDefinition(/*FIXME:*/ Inst.second, Function, true,
DefinitionRequired, true);
- if (Function->isDefined())
+ if (Function->isDefined()) {
+ if (getLangOpts().SYCLIsDevice && Function->hasAttr<SYCLKernelAttr>())
+ constructOpenCLKernel(Function, *MangleCtx);
Function->setInstantiationIsPending(false);
+ }
}
continue;
}
Index: clang/lib/Sema/SemaSYCL.cpp
===================================================================
--- /dev/null
+++ clang/lib/Sema/SemaSYCL.cpp
@@ -0,0 +1,457 @@
+//===- 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 "TreeTransform.h"
+#include "clang/AST/AST.h"
+#include "clang/AST/Mangle.h"
+#include "clang/AST/QualTypeNames.h"
+#include "clang/Sema/Initialization.h"
+#include "clang/Sema/Sema.h"
+
+using namespace clang;
+
+using ParamDesc = std::tuple<QualType, IdentifierInfo *, TypeSourceInfo *>;
+
+/// Various utilities.
+class Util {
+public:
+ using DeclContextDesc = std::pair<clang::Decl::Kind, StringRef>;
+
+ /// Checks whether given clang type is a full specialization of the SYCL
+ /// accessor class.
+ static bool isSyclAccessorType(const QualType &Ty);
+
+ /// Checks whether given clang type is declared in the given hierarchy of
+ /// declaration contexts.
+ /// \param Ty the clang type being checked
+ /// \param Scopes the declaration scopes leading from the type to the
+ /// translation unit (excluding the latter)
+ static bool matchQualifiedTypeName(const QualType &Ty,
+ ArrayRef<Util::DeclContextDesc> Scopes);
+};
+
+static CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) {
+ return (*Caller->param_begin())->getType()->getAsCXXRecordDecl();
+}
+
+class KernelBodyTransform : public TreeTransform<KernelBodyTransform> {
+public:
+ KernelBodyTransform(std::pair<DeclaratorDecl *, DeclaratorDecl *> &MPair,
+ Sema &S)
+ : TreeTransform<KernelBodyTransform>(S), MappingPair(MPair), SemaRef(S) {}
+ bool AlwaysRebuild() { return true; }
+
+ ExprResult TransformDeclRefExpr(DeclRefExpr *DRE) {
+ auto Ref = dyn_cast<DeclaratorDecl>(DRE->getDecl());
+ if (Ref && Ref == MappingPair.first) {
+ auto NewDecl = MappingPair.second;
+ return DeclRefExpr::Create(
+ SemaRef.getASTContext(), DRE->getQualifierLoc(),
+ DRE->getTemplateKeywordLoc(), NewDecl, false, DRE->getNameInfo(),
+ NewDecl->getType(), DRE->getValueKind());
+ }
+ return DRE;
+ }
+
+private:
+ std::pair<DeclaratorDecl *, DeclaratorDecl *> MappingPair;
+ Sema &SemaRef;
+};
+
+static FunctionDecl *
+CreateOpenCLKernelDeclaration(ASTContext &Context, StringRef Name,
+ ArrayRef<ParamDesc> ParamDescs) {
+
+ DeclContext *DC = Context.getTranslationUnitDecl();
+ QualType RetTy = Context.VoidTy;
+ SmallVector<QualType, 8> ArgTys;
+
+ // Extract argument types from the descriptor array:
+ std::transform(
+ ParamDescs.begin(), ParamDescs.end(), std::back_inserter(ArgTys),
+ [](const ParamDesc &PD) -> QualType { return std::get<0>(PD); });
+ FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel);
+ QualType FuncTy = Context.getFunctionType(RetTy, ArgTys, Info);
+ DeclarationName DN = DeclarationName(&Context.Idents.get(Name));
+
+ FunctionDecl *OpenCLKernel = FunctionDecl::Create(
+ Context, DC, SourceLocation(), SourceLocation(), DN, FuncTy,
+ Context.getTrivialTypeSourceInfo(RetTy), SC_None);
+
+ llvm::SmallVector<ParmVarDecl *, 16> Params;
+ int i = 0;
+ for (const auto &PD : ParamDescs) {
+ auto P = ParmVarDecl::Create(Context, OpenCLKernel, SourceLocation(),
+ SourceLocation(), std::get<1>(PD),
+ std::get<0>(PD), std::get<2>(PD), SC_None, 0);
+ P->setScopeInfo(0, i++);
+ P->setIsUsed();
+ Params.push_back(P);
+ }
+ OpenCLKernel->setParams(Params);
+
+ OpenCLKernel->addAttr(OpenCLKernelAttr::CreateImplicit(Context));
+ OpenCLKernel->addAttr(AsmLabelAttr::CreateImplicit(Context, Name));
+ OpenCLKernel->addAttr(ArtificialAttr::CreateImplicit(Context));
+
+ // Add kernel to translation unit to see it in AST-dump
+ DC->addDecl(OpenCLKernel);
+ return OpenCLKernel;
+}
+
+/// Return __init method
+static CXXMethodDecl *getInitMethod(const CXXRecordDecl *CRD) {
+ CXXMethodDecl *InitMethod;
+ auto It = std::find_if(CRD->methods().begin(), CRD->methods().end(),
+ [](const CXXMethodDecl *Method) {
+ return Method->getNameAsString() == "__init";
+ });
+ InitMethod = (It != CRD->methods().end()) ? *It : nullptr;
+ return InitMethod;
+}
+
+// Creates body for new OpenCL kernel. This body contains initialization of SYCL
+// kernel object fields with kernel parameters and a little bit transformed body
+// of the kernel caller function.
+static CompoundStmt *CreateOpenCLKernelBody(Sema &S,
+ FunctionDecl *KernelCallerFunc,
+ DeclContext *KernelDecl) {
+ llvm::SmallVector<Stmt *, 16> BodyStmts;
+ CXXRecordDecl *LC = getKernelObjectType(KernelCallerFunc);
+ assert(LC && "Kernel object must be available");
+ TypeSourceInfo *TSInfo = LC->isLambda() ? LC->getLambdaTypeInfo() : nullptr;
+
+ // Create a local kernel object (lambda or functor) assembled from the
+ // incoming formal parameters.
+ auto KernelObjClone = VarDecl::Create(
+ S.Context, KernelDecl, SourceLocation(), SourceLocation(),
+ LC->getIdentifier(), QualType(LC->getTypeForDecl(), 0), TSInfo, SC_None);
+ Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone),
+ SourceLocation(), SourceLocation());
+ BodyStmts.push_back(DS);
+ auto KernelObjCloneRef =
+ DeclRefExpr::Create(S.Context, NestedNameSpecifierLoc(), SourceLocation(),
+ KernelObjClone, false, DeclarationNameInfo(),
+ QualType(LC->getTypeForDecl(), 0), VK_LValue);
+
+ auto KernelFuncDecl = cast<FunctionDecl>(KernelDecl);
+ auto KernelFuncParam =
+ KernelFuncDecl->param_begin(); // Iterator to ParamVarDecl (VarDecl)
+ if (KernelFuncParam) {
+ llvm::SmallVector<Expr *, 16> InitExprs;
+ InitializedEntity VarEntity =
+ InitializedEntity::InitializeVariable(KernelObjClone);
+ for (auto Field : LC->fields()) {
+ // Creates Expression for special SYCL object accessor.
+ // All special SYCL objects must have __init method, here we use it to
+ // initialize them. We create call of __init method and pass built kernel
+ // arguments as parameters to the __init method.
+ auto getExprForSpecialSYCLObj = [&](const QualType ¶mTy,
+ FieldDecl *Field,
+ const CXXRecordDecl *CRD,
+ Expr *Base) {
+ // All special SYCL objects must have __init method.
+ CXXMethodDecl *InitMethod = getInitMethod(CRD);
+ assert(InitMethod &&
+ "__init method is expected.");
+ unsigned NumParams = InitMethod->getNumParams();
+ llvm::SmallVector<Expr *, 4> ParamDREs(NumParams);
+ auto KFP = KernelFuncParam;
+ for (size_t I = 0; I < NumParams; ++KFP, ++I) {
+ QualType ParamType = (*KFP)->getOriginalType();
+ ParamDREs[I] = DeclRefExpr::Create(
+ S.Context, NestedNameSpecifierLoc(), SourceLocation(), *KFP,
+ false, DeclarationNameInfo(), ParamType, VK_LValue);
+ }
+
+ if (NumParams)
+ std::advance(KernelFuncParam, NumParams - 1);
+
+ DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none);
+ // [kernel_obj].special_obj
+ auto SpecialObjME = MemberExpr::Create(
+ S.Context, Base, false, SourceLocation(), NestedNameSpecifierLoc(),
+ SourceLocation(), Field, FieldDAP,
+ DeclarationNameInfo(Field->getDeclName(), SourceLocation()),
+ nullptr, Field->getType(), VK_LValue, OK_Ordinary, NOUR_None);
+
+ // [kernel_obj].special_obj.__init
+ DeclAccessPair MethodDAP = DeclAccessPair::make(InitMethod, AS_none);
+ auto ME = MemberExpr::Create(
+ S.Context, SpecialObjME, false, SourceLocation(),
+ NestedNameSpecifierLoc(), SourceLocation(), InitMethod, MethodDAP,
+ DeclarationNameInfo(InitMethod->getDeclName(), SourceLocation()),
+ nullptr, InitMethod->getType(), VK_LValue, OK_Ordinary, NOUR_None);
+
+ // Not referenced -> not emitted
+ S.MarkFunctionReferenced(SourceLocation(), InitMethod, true);
+
+ QualType ResultTy = InitMethod->getReturnType();
+ ExprValueKind VK = Expr::getValueKindForType(ResultTy);
+ ResultTy = ResultTy.getNonLValueExprType(S.Context);
+
+ llvm::SmallVector<Expr *, 4> ParamStmts;
+ const auto *Proto = cast<FunctionProtoType>(InitMethod->getType());
+ S.GatherArgumentsForCall(SourceLocation(), InitMethod, Proto, 0,
+ ParamDREs, ParamStmts);
+ // [kernel_obj].special_obj.__init(_ValueType*,
+ // range<int>, range<int>, id<int>)
+ CXXMemberCallExpr *Call = CXXMemberCallExpr::Create(
+ S.Context, ME, ParamStmts, ResultTy, VK, SourceLocation());
+ BodyStmts.push_back(Call);
+ };
+
+ // Run through kernel object fields and add initialization for them using
+ // built kernel parameters. There are a several possible cases:
+ // - Kernel object field is a SYCL special object (SYCL accessor).
+ // These objects has a special initialization scheme - using
+ // __init method.
+ // - Kernel object field has a scalar type. In this case we should add
+ // simple initialization.
+ // - Kernel object field has a structure or class type. Same handling as
+ // a scalar.
+ QualType FieldType = Field->getType();
+ CXXRecordDecl *CRD = FieldType->getAsCXXRecordDecl();
+ InitializedEntity Entity =
+ InitializedEntity::InitializeMember(Field, &VarEntity);
+ if (Util::isSyclAccessorType(FieldType)) {
+ // Initialize kernel object field with the default constructor and
+ // construct a call of __init method.
+ InitializationKind InitKind =
+ InitializationKind::CreateDefault(SourceLocation());
+ InitializationSequence InitSeq(S, Entity, InitKind, None);
+ ExprResult MemberInit = InitSeq.Perform(S, Entity, InitKind, None);
+ InitExprs.push_back(MemberInit.get());
+ getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef);
+ } else if (CRD || FieldType->isScalarType()) {
+ // If field has built-in or a structure/class type just initialize
+ // this field with corresponding kernel argument using copy
+ // initialization.
+ QualType ParamType = (*KernelFuncParam)->getOriginalType();
+ Expr *DRE =
+ DeclRefExpr::Create(S.Context, NestedNameSpecifierLoc(),
+ SourceLocation(), *KernelFuncParam, false,
+ DeclarationNameInfo(), ParamType, VK_LValue);
+
+ InitializationKind InitKind =
+ InitializationKind::CreateCopy(SourceLocation(), SourceLocation());
+ InitializationSequence InitSeq(S, Entity, InitKind, DRE);
+
+ ExprResult MemberInit = InitSeq.Perform(S, Entity, InitKind, DRE);
+ InitExprs.push_back(MemberInit.get());
+
+ } else
+ llvm_unreachable("Unsupported field type");
+ KernelFuncParam++;
+ }
+ Expr *ILE = new (S.Context)
+ InitListExpr(S.Context, SourceLocation(), InitExprs, SourceLocation());
+ ILE->setType(QualType(LC->getTypeForDecl(), 0));
+ KernelObjClone->setInit(ILE);
+ }
+
+ // In the kernel caller function kernel object is a function parameter, so we
+ // need to replace all refs to this kernel oject with refs to our clone
+ // declared inside the kernel body.
+ Stmt *FunctionBody = KernelCallerFunc->getBody();
+ ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin());
+
+ // DeclRefExpr with a valid source location but with decl which is not marked
+ // as used becomes invalid.
+ KernelObjClone->setIsUsed();
+ std::pair<DeclaratorDecl *, DeclaratorDecl *> MappingPair;
+ MappingPair.first = KernelObjParam;
+ MappingPair.second = KernelObjClone;
+
+ // Function scope might be empty, so we do push
+ S.PushFunctionScope();
+ KernelBodyTransform KBT(MappingPair, S);
+ Stmt *NewBody = KBT.TransformStmt(FunctionBody).get();
+ BodyStmts.push_back(NewBody);
+ return CompoundStmt::Create(S.Context, BodyStmts, SourceLocation(),
+ SourceLocation());
+}
+
+/// Creates a kernel parameter descriptor
+/// \param Src field declaration to construct name from
+/// \param Ty the desired parameter type
+/// \return the constructed descriptor
+static ParamDesc makeParamDesc(const FieldDecl *Src, QualType Ty) {
+ ASTContext &Ctx = Src->getASTContext();
+ std::string Name = (Twine("_arg_") + Src->getName()).str();
+ return std::make_tuple(Ty, &Ctx.Idents.get(Name),
+ Ctx.getTrivialTypeSourceInfo(Ty));
+}
+
+// Creates list of kernel parameters descriptors using KernelObj (kernel
+// object). Fields of kernel object must be initialized with SYCL kernel
+// arguments so in the following function we extract types of kernel object
+// fields and add it to the array with kernel parameters descriptors.
+static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
+ SmallVectorImpl<ParamDesc> &ParamDescs) {
+ auto CreateAndAddPrmDsc = [&](const FieldDecl *Fld, const QualType &ArgType) {
+ // Create a parameter descriptor and append it to the result
+ ParamDescs.push_back(makeParamDesc(Fld, ArgType));
+ };
+
+ // Creates a parameter descriptor for SYCL special object - SYCL accessor.
+ // All special SYCL objects must have __init method. We extract types for
+ // kernel parameters from __init method parameters. We will use __init method
+ // and kernel parameters which we build here to initialize special objects in
+ // the kernel body.
+ auto createSpecialSYCLObjParamDesc = [&](const FieldDecl *Fld,
+ const QualType &ArgTy) {
+ const auto *RecordDecl = ArgTy->getAsCXXRecordDecl();
+ assert(RecordDecl && "Special SYCL object must be of a record type");
+
+ CXXMethodDecl *InitMethod = getInitMethod(RecordDecl);
+ assert(InitMethod && "__init method is expected.");
+ unsigned NumParams = InitMethod->getNumParams();
+ for (size_t I = 0; I < NumParams; ++I) {
+ ParmVarDecl *PD = InitMethod->getParamDecl(I);
+ CreateAndAddPrmDsc(Fld, PD->getType().getCanonicalType());
+ }
+ };
+
+ // Run through kernel object fields and create corresponding kernel
+ // parameters descriptors. There are a several possible cases:
+ // - Kernel object field is a SYCL special object (SYCL accessor).
+ // These objects has a special initialization scheme - using
+ // __init method.
+ // - Kernel object field has a scalar type. In this case we should add
+ // kernel parameter with the same type.
+ // - Kernel object field has a structure or class type. Same handling as a
+ // scalar but we should check if this structure/class contains accessors
+ // and add parameter decriptor for them properly.
+ for (const auto *Fld : KernelObj->fields()) {
+ QualType ArgTy = Fld->getType();
+ if (Util::isSyclAccessorType(ArgTy))
+ createSpecialSYCLObjParamDesc(Fld, ArgTy);
+ else if (ArgTy->isStructureOrClassType())
+ CreateAndAddPrmDsc(Fld, ArgTy);
+ else if (ArgTy->isScalarType())
+ CreateAndAddPrmDsc(Fld, ArgTy);
+ else
+ llvm_unreachable("Unsupported kernel parameter type");
+ }
+}
+
+// Creates a mangled kernel name for given kernel name type
+static std::string constructKernelName(QualType KernelNameType,
+ MangleContext &MC) {
+ SmallString<256> Result;
+ llvm::raw_svector_ostream Out(Result);
+
+ MC.mangleTypeName(KernelNameType, Out);
+ return Out.str();
+}
+
+// Generates the OpenCL kernel using KernelCallerFunc (kernel caller
+// function) defined is SYCL headers.
+// Generated OpenCL kernel contains the body of the kernel caller function,
+// receives OpenCL like parameters and additionally does some manipulation to
+// initialize captured lambda/functor fields with these parameters.
+// SYCL runtime marks kernel caller function with sycl_kernel attribute.
+// To be able to generate OpenCL kernel from KernelCallerFunc we put
+// the following requirements to the function which SYCL runtime can mark with
+// sycl_kernel attribute:
+// - Must be template function with at least two template parameters.
+// First parameter must represent "unique kernel name"
+// Second parameter must be the function object type
+// - Must have only one function parameter - function object.
+//
+// Example of kernel caller function:
+// template <typename KernelName, typename KernelType/*, ...*/>
+// __attribute__((sycl_kernel)) void kernel_caller_function(KernelType
+// KernelFuncObj) {
+// KernelFuncObj();
+// }
+//
+//
+void Sema::constructOpenCLKernel(FunctionDecl *KernelCallerFunc,
+ MangleContext &MC) {
+ CXXRecordDecl *LE = getKernelObjectType(KernelCallerFunc);
+ assert(LE && "invalid kernel caller");
+
+ // Build list of kernel arguments.
+ llvm::SmallVector<ParamDesc, 16> ParamDescs;
+ buildArgTys(getASTContext(), LE, ParamDescs);
+
+ // Extract name from kernel caller parameters and mangle it.
+ const TemplateArgumentList *TemplateArgs =
+ KernelCallerFunc->getTemplateSpecializationArgs();
+ assert(TemplateArgs && "No template argument info");
+ QualType KernelNameType = TypeName::getFullyQualifiedType(
+ TemplateArgs->get(0).getAsType(), getASTContext(), true);
+ std::string Name = constructKernelName(KernelNameType, MC);
+
+ FunctionDecl *OpenCLKernel =
+ CreateOpenCLKernelDeclaration(getASTContext(), Name, ParamDescs);
+
+ // Let's copy source location of a functor/lambda to emit nicer diagnostics.
+ OpenCLKernel->setLocation(LE->getLocation());
+
+ CompoundStmt *OpenCLKernelBody =
+ CreateOpenCLKernelBody(*this, KernelCallerFunc, OpenCLKernel);
+ OpenCLKernel->setBody(OpenCLKernelBody);
+
+ addSYCLKernel(OpenCLKernel);
+}
+
+// -----------------------------------------------------------------------------
+// Utility class methods
+// -----------------------------------------------------------------------------
+
+bool Util::isSyclAccessorType(const QualType &Ty) {
+ static std::array<DeclContextDesc, 3> Scopes = {
+ Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
+ Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
+ Util::DeclContextDesc{clang::Decl::Kind::ClassTemplateSpecialization,
+ "accessor"}};
+ return matchQualifiedTypeName(Ty, Scopes);
+}
+
+bool Util::matchQualifiedTypeName(const QualType &Ty,
+ ArrayRef<Util::DeclContextDesc> Scopes) {
+ // The idea: check the declaration context chain starting from the type
+ // itself. At each step check the context is of expected kind
+ // (namespace) and name.
+ const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
+
+ if (!RecTy)
+ return false; // only classes/structs supported
+ const auto *Ctx = dyn_cast<DeclContext>(RecTy);
+ StringRef Name = "";
+
+ for (const auto &Scope : llvm::reverse(Scopes)) {
+ clang::Decl::Kind DK = Ctx->getDeclKind();
+
+ if (DK != Scope.first)
+ return false;
+
+ switch (DK) {
+ case clang::Decl::Kind::ClassTemplateSpecialization:
+ // ClassTemplateSpecializationDecl inherits from CXXRecordDecl
+ case clang::Decl::Kind::CXXRecord:
+ Name = cast<CXXRecordDecl>(Ctx)->getName();
+ break;
+ case clang::Decl::Kind::Namespace:
+ Name = cast<NamespaceDecl>(Ctx)->getName();
+ break;
+ default:
+ llvm_unreachable("matchQualifiedTypeName: decl kind not supported");
+ }
+ if (Name != Scope.second)
+ return false;
+ Ctx = Ctx->getParent();
+ }
+ return Ctx->isTranslationUnit();
+}
+
Index: clang/lib/Sema/CMakeLists.txt
===================================================================
--- clang/lib/Sema/CMakeLists.txt
+++ clang/lib/Sema/CMakeLists.txt
@@ -57,6 +57,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,10 @@
for (Decl *D : S.WeakTopLevelDecls())
Consumer->HandleTopLevelDecl(DeclGroupRef(D));
+ if (S.getLangOpts().SYCLIsDevice)
+ for (Decl *D : S.getSYCLKernels())
+ 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
@@ -2474,6 +2474,12 @@
}
}
+ if (LangOpts.SYCLIsDevice && Global->hasAttr<OpenCLKernelAttr>() &&
+ MustBeEmitted(Global)) {
+ 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
@@ -10047,6 +10047,10 @@
if (D->hasAttr<AliasAttr>() || D->hasAttr<UsedAttr>())
return true;
+ // If SYCL, only kernels are required.
+ if (LangOpts.SYCLIsDevice && !(D->hasAttr<OpenCLKernelAttr>()))
+ return false;
+
if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
// Forward declarations aren't required.
if (!FD->doesThisDeclarationHaveABody())
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -11728,6 +11728,19 @@
ConstructorDestructor,
BuiltinFunction
};
+
+private:
+ /// Contains generated OpenCL kernel functions for SYCL.
+ SmallVector<Decl *, 4> SYCLKernels;
+
+public:
+ void addSYCLKernel(Decl *D) { SYCLKernels.push_back(D); }
+ /// Access to SYCL kernels.
+ SmallVectorImpl<Decl *> &getSYCLKernels() { return SYCLKernels; }
+
+ /// Constructs an OpenCL kernel using the KernelCaller function and adds it to
+ /// the SYCL device code.
+ void constructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC);
};
/// RAII object that enters a new expression evaluation context.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
