Fznamznon created this revision.
Herald added subscribers: cfe-commits, Anastasia, ebevhan, jfb, yaxunl, mgorny.
Herald added a project: clang.
Fznamznon added reviewers: bader, Naghasan, ABataev.
All SYCL memory objects shared between host and device (buffers/images, these
objects map to OpenCL buffers and images) must be accessed through special
accessor classes. The "device" side implementation of these classes contain
pointers to the device memory. As there is no way in OpenCL to pass
structures with pointers inside as kernel arguments, all memory objects
shared between host and device must be passed to the kernel as raw
pointers. SYCL also has a special mechanism for passing kernel arguments
from host to the device. In OpenCL kernel arguments are set by calling
`clSetKernelArg` function for each kernel argument, meanwhile in SYCL all the
kernel arguments are fields of "SYCL kernel function" which can be defined
as a lambda function or a named function object and passed as an argument
to SYCL function for invoking kernels (such as `parallel_for` or `single_task`).
To facilitate the mapping of SYCL kernel data members to OpenCL kernel
arguments and overcome OpenCL limitations we added the generation of an
OpenCL kernel function inside the compiler. An OpenCL kernel function
contains the body of the SYCL kernel function, receives OpenCL-like
parameters and additionally does some manipulation to initialize SYCL
kernel data members with these parameters. In some pseudo code the OpenCL
kernel function can look like this:
// SYCL kernel is defined in SYCL headers:
template <typename KernelName, typename KernelType/*, ...*/>
__attribute__((sycl_kernel)) void sycl_kernel_function(KernelType
KernelFuncObj) {
// ...
KernelFuncObj();
}
// Generated OpenCL kernel function
__kernel KernelName(global int* a) {
KernelType KernelFuncObj; // Actually kernel function object declaration
// doesn't have a name in AST.
// Let the kernel function object have one captured field - accessor A.
// We need to init it with global pointer from arguments:
KernelFuncObj.A.__init(a);
// Body of the SYCL kernel from SYCL headers:
{
KernelFuncObj();
}
}
OpenCL kernel function is generated by the compiler inside the Sema
using AST nodes.
Repository:
rG LLVM Github Monorepo
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 (__attribute__((address_space(1))) int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: accessor_typedef 'void (__attribute__((address_space(1))) int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: accessor_alias 'void (__attribute__((address_space(1))) 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 (__attribute__((address_space(1))) 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_]+]] '__attribute__((address_space(1))) 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 {{.*}} '__attribute__((address_space(1))) int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__attribute__((address_space(1))) 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 (__attribute__((address_space(3))) int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: {{.*}}use_global 'void (__attribute__((address_space(1))) int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: {{.*}}use_constant 'void (__attribute__((address_space(2))) 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__((address_space(1))) dataT;
+};
+
+template <typename dataT>
+struct DeviceValueType<dataT, access::target::constant_buffer> {
+ using type = __attribute__((address_space(2))) dataT;
+};
+
+template <typename dataT>
+struct DeviceValueType<dataT, access::target::local> {
+ using type = __attribute__((address_space(3))) 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__((address_space(1))) 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"
@@ -5610,6 +5611,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;
@@ -5628,17 +5631,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
@@ -10023,6 +10023,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
@@ -11632,6 +11632,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
