Skip to content

Commit effac35

Browse files
vladimirlazvladimirlaz
committed
[SYCL] Update integration header format to match SYCL static
library compilation flow. - Use definitions for kernel_param_kind_t and kernel_param_desc_t types from SYCL header. - Add compiler LIT test. Signed-off-by: Vladimir Lazarev <vladimir.lazarev@intel.com>
1 parent 5a3040d commit effac35

File tree

3 files changed

+147
-37
lines changed

3 files changed

+147
-37
lines changed

‎clang/include/clang/Sema/Sema.h‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -298,8 +298,7 @@ class SYCLIntegrationHeader {
298298
// kernel lambda or function object
299299
enum kernel_param_kind_t {
300300
kind_first,
301-
kind_none = kind_first,
302-
kind_accessor,
301+
kind_accessor = kind_first,
303302
kind_scalar,
304303
kind_struct,
305304
kind_sampler,
@@ -333,7 +332,7 @@ class SYCLIntegrationHeader {
333332
// Kernel actual parameter descriptor.
334333
struct KernelParamDesc {
335334
// Represents a parameter kind.
336-
kernel_param_kind_t Kind = kind_none;
335+
kernel_param_kind_t Kind = kind_last;
337336
// If Kind is kind_scalar or kind_struct, then
338337
// denotes parameter size in bytes (includes padding for structs)
339338
// If Kind is kind_accessor

‎clang/lib/Sema/SemaSYCL.cpp‎

Lines changed: 10 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -412,16 +412,17 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name,
412412
QualType NameType, CXXRecordDecl *Lambda) {
413413
ASTContext &Ctx = Lambda->getASTContext();
414414
const ASTRecordLayout &Layout = Ctx.getASTRecordLayout(Lambda);
415-
KernelParamKind Knd = SYCLIntegrationHeader::kind_none;
415+
KernelParamKind Knd = SYCLIntegrationHeader::kind_last;
416416
H.startKernel(Name, NameType);
417417
unsigned Offset = 0;
418418
int Info = 0;
419419

420420
auto Vis = std::make_tuple(
421421
// pre_visit
422422
[&](int CaptureN, VarDecl *CapturedVar, FieldDecl *CapturedVal) {
423+
// Set offset in bytes
423424
Offset = static_cast<unsigned>(
424-
Layout.getFieldOffset(CapturedVal->getFieldIndex()));
425+
Layout.getFieldOffset(CapturedVal->getFieldIndex()))/8;
425426
},
426427
// visit_accessor
427428
[&](int CaptureN, target AccTrg, QualType PointeeType,
@@ -516,7 +517,6 @@ static const char *paramKind2Str(KernelParamKind K) {
516517
case SYCLIntegrationHeader::kind_##x: \
517518
return "kind_" #x
518519
switch (K) {
519-
CASE(none);
520520
CASE(accessor);
521521
CASE(scalar);
522522
CASE(struct);
@@ -664,37 +664,15 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
664664
}
665665
O << "\n";
666666

667+
O << "#include <CL/sycl/detail/kernel_desc.hpp>\n";
668+
669+
O << "\n";
670+
667671
O << "namespace cl {\n";
668672
O << "namespace sycl {\n";
669673
O << "namespace detail {\n";
670674

671-
O << "// kernel parameter kinds\n";
672-
O << "enum kernel_param_kind_t {\n";
673-
674-
for (int I = SYCLIntegrationHeader::kind_first;
675-
I <= SYCLIntegrationHeader::kind_last; I++) {
676-
KernelParamKind It = static_cast<KernelParamKind>(I);
677-
O << " " << std::string(paramKind2Str(It));
678-
if (I < SYCLIntegrationHeader::kind_last)
679-
O << ",";
680-
O << "\n";
681-
}
682-
O << "};\n";
683675
O << "\n";
684-
O << "// describes a kernel parameter\n";
685-
O << "struct kernel_param_desc_t {\n";
686-
O << " // parameter kind\n";
687-
O << " kernel_param_kind_t kind;\n";
688-
O << " // kind == kind_scalar, kind_struct\n";
689-
O << " // parameter size in bytes (includes padding for structs)\n";
690-
O << " // kind == kind_accessor\n";
691-
O << " // access target; possible access targets are defined in "
692-
"access/access.hpp\n";
693-
O << " int info;\n";
694-
O << " // offset of the captured value of the parameter in the lambda or "
695-
"function object\n";
696-
O << " int offs;\n";
697-
O << "};\n\n";
698676

699677
O << "// names of all kernels defined in the corresponding source\n";
700678
O << "static constexpr\n";
@@ -720,11 +698,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
720698

721699
for (const auto &P : K.Params) {
722700
std::string TyStr = paramKind2Str(P.Kind);
723-
O << " { " << TyStr << ", " << P.Info << ", " << P.Offset << " },\n";
701+
O << " { kernel_param_kind_t::" << TyStr << ", ";
702+
O << P.Info << ", " << P.Offset << " },\n";
724703
}
725-
O << " { kind_none, 0, 0 }";
726-
if (I < KernelDescs.size() - 1)
727-
O << ",";
728704
O << "\n";
729705
}
730706
O << "};\n\n";
@@ -772,7 +748,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
772748
O << " return kernel_signatures[i+" << CurStart << "];\n";
773749
O << " }\n";
774750
O << "};\n";
775-
CurStart += N + 1;
751+
CurStart += N;
776752
}
777753
O << "\n";
778754
O << "} // namespace detail\n";

‎clang/test/CodeGenSYCL/integration_header.cpp‎

Lines changed: 135 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,135 @@
1+
// RUN: %clang --sycl -Xclang -fsycl-int-header=%t.h %s -c -o %T/kernel.spv
2+
// RUN: FileCheck -input-file=%t.h %s
3+
//
4+
// CHECK: class first_kernel;
5+
// CHECK-NEXT: template <typename T> class second_kernel;
6+
// CHECK-NEXT: struct X;
7+
// CHECK-NEXT: template <typename T> struct point ;
8+
// CHECK-NEXT: template <int a, typename T1, typename T2> class third_kernel;
9+
//
10+
// CHECK: #include <CL/sycl/detail/kernel_desc.hpp>
11+
//
12+
// CHECK: static constexpr
13+
// CHECK-NEXT: const char* const kernel_names[] = {
14+
// CHECK-NEXT: "first_kernel",
15+
// CHECK-NEXT: "second_namespace::second_kernel<char>",
16+
// CHECK-NEXT: "third_kernel<1, int, point< X> >"
17+
// CHECK-NEXT: };
18+
//
19+
// CHECK: const kernel_param_desc_t kernel_signatures[] = {
20+
// CHECK-NEXT: //--- first_kernel
21+
// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 4, 0 },
22+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2014, 4 },
23+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 5 },
24+
// CHECK-EMPTY:
25+
// CHECK-NEXT: //--- second_namespace::second_kernel<char>
26+
// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 4, 0 },
27+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 4 },
28+
// CHECK-EMPTY:
29+
// CHECK-NEXT: //--- third_kernel<1, int, point< X> >
30+
// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 4, 0 },
31+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 4 },
32+
// CHECK-EMPTY:
33+
// CHECK-NEXT: };
34+
//
35+
// CHECK: template <class KernelNameType> struct KernelInfo;
36+
// CHECK: template <> struct KernelInfo<class first_kernel> {
37+
// CHECK: template <> struct KernelInfo<class second_namespace::second_kernel<char>> {
38+
// CHECK: template <> struct KernelInfo<class third_kernel<1, int, struct point<struct X> >> {
39+
40+
namespace cl {
41+
namespace sycl {
42+
namespace access {
43+
44+
enum class target {
45+
global_buffer = 2014,
46+
constant_buffer,
47+
local,
48+
image,
49+
host_buffer,
50+
host_image,
51+
image_array
52+
};
53+
54+
enum class mode {
55+
read = 1024,
56+
write,
57+
read_write,
58+
discard_write,
59+
discard_read_write,
60+
atomic
61+
};
62+
63+
enum class placeholder { false_t,
64+
true_t };
65+
66+
enum class address_space : int {
67+
private_space = 0,
68+
global_space,
69+
constant_space,
70+
local_space
71+
};
72+
} // namespace access
73+
template <typename dataT, int dimensions, access::mode accessmode,
74+
access::target accessTarget = access::target::global_buffer,
75+
access::placeholder isPlaceholder = access::placeholder::false_t>
76+
class accessor {
77+
78+
public:
79+
void use(void) const {}
80+
};
81+
} // namespace sycl
82+
} // namespace cl
83+
84+
template <typename KernelName, typename KernelType>
85+
__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) {
86+
kernelFunc();
87+
}
88+
struct x {};
89+
template <typename T>
90+
struct point {};
91+
namespace second_namespace {
92+
template <typename T>
93+
class second_kernel;
94+
}
95+
96+
template <int a, typename T1, typename T2>
97+
class third_kernel;
98+
99+
int main() {
100+
101+
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> acc1;
102+
cl::sycl::accessor<float, 2, cl::sycl::access::mode::write,
103+
cl::sycl::access::target::local,
104+
cl::sycl::access::placeholder::true_t>
105+
acc2;
106+
int i = 13;
107+
// TODO: Uncomemnt when structures in kernel arguments are correctly processed
108+
// by SYCL compiler
109+
/* struct {
110+
char c;
111+
int i;
112+
} test_s;
113+
test_s.c = 14;*/
114+
kernel_single_task<class first_kernel>([=]() {
115+
if (i == 13 /*&& test_s.c == 14*/) {
116+
117+
acc1.use();
118+
acc2.use();
119+
}
120+
});
121+
122+
kernel_single_task<class second_namespace::second_kernel<char>>([=]() {
123+
if (i == 13) {
124+
acc2.use();
125+
}
126+
});
127+
kernel_single_task<class third_kernel<1, int,point<struct X>>>([=]() {
128+
if (i == 13) {
129+
acc2.use();
130+
}
131+
});
132+
133+
return 0;
134+
}
135+

0 commit comments

Comments
 (0)