Skip to content

Commit d442364

Browse files
s-kanaevbader
s-kanaev
authored and
bader
committed
[SYCL] Fix program build results caching (#865)
This patch fixes regression introduced by #847: programs and kernels are cached when program is built with sequence of compile_with_kernel_type()/compile_with_source() and then linked with link(). By design, caching of kernels and programs is permitted only when program is built with build_with_kernel_type() using the default options. Signed-off-by: Sergey Kanaev <[email protected]>
1 parent 0c717f8 commit d442364

File tree

2 files changed

+256
-18
lines changed

2 files changed

+256
-18
lines changed

‎sycl/include/CL/sycl/detail/program_impl.hpp

Lines changed: 10 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -44,13 +44,13 @@ class program_impl {
4444
program_impl(const context &Context, vector_class<device> DeviceList)
4545
: Context(Context), Devices(DeviceList) {}
4646

47-
// Don't allow kernels caching for linked programs due to only compiled
47+
// Kernels caching for linked programs won't be allowed due to only compiled
4848
// state of each and every program in the list and thus unknown state of
4949
// caching resolution
5050
program_impl(vector_class<std::shared_ptr<program_impl>> ProgramList,
5151
string_class LinkOptions = "")
5252
: State(program_state::linked), LinkOptions(LinkOptions),
53-
BuildOptions(LinkOptions), AllowKernelsCaching(false) {
53+
BuildOptions(LinkOptions) {
5454
// Verify arguments
5555
if (ProgramList.empty()) {
5656
throw runtime_error("Non-empty vector of programs expected");
@@ -96,10 +96,9 @@ class program_impl {
9696
}
9797
}
9898

99-
// Disallow kernels caching for programs created by interoperability c-tor
99+
// Kernel caching for programs created by interoperability c-tor isn't allowed
100100
program_impl(const context &Context, RT::PiProgram Program)
101-
: Program(Program), Context(Context), IsLinkable(true),
102-
AllowKernelsCaching(false) {
101+
: Program(Program), Context(Context), IsLinkable(true) {
103102

104103
// TODO handle the case when cl_program build is in progress
105104
cl_uint NumDevices;
@@ -210,13 +209,12 @@ class program_impl {
210209
if (!is_host()) {
211210
OSModuleHandle M = OSUtil::getOSModuleHandle(AddressInThisModule);
212211
// If there are no build options, program can be safely cached
213-
if (is_cacheable_with_build_options(BuildOptions)) {
212+
if (is_cacheable_with_options(BuildOptions)) {
213+
IsProgramAndKernelCachingAllowed = true;
214214
Program =
215215
ProgramManager::getInstance().getBuiltOpenCLProgram(M, Context);
216216
PI_CALL(piProgramRetain)(Program);
217217
} else {
218-
AllowKernelsCaching = false;
219-
220218
create_cl_program_with_il(M);
221219
build(BuildOptions);
222220
}
@@ -227,9 +225,6 @@ class program_impl {
227225
void build_with_source(string_class KernelSource,
228226
string_class BuildOptions = "") {
229227
throw_if_state_is_not(program_state::none);
230-
231-
AllowKernelsCaching = false;
232-
233228
// TODO should it throw if it's host?
234229
if (!is_host()) {
235230
create_cl_program_with_source(KernelSource);
@@ -425,12 +420,12 @@ class program_impl {
425420
}
426421

427422
bool is_cacheable() const {
428-
return is_cacheable_with_build_options(BuildOptions) && AllowKernelsCaching;
423+
return IsProgramAndKernelCachingAllowed;
429424
}
430425

431426
static bool
432-
is_cacheable_with_build_options(const string_class &BuildOptions) {
433-
return BuildOptions.empty();
427+
is_cacheable_with_options(const string_class &Options) {
428+
return Options.empty();
434429
}
435430

436431
RT::PiKernel get_pi_kernel(const string_class &KernelName) const {
@@ -488,7 +483,7 @@ class program_impl {
488483
// Only allow kernel caching for programs constructed with context only (or
489484
// device list and context) and built with build_with_kernel_type with
490485
// default build options
491-
bool AllowKernelsCaching = true;
486+
bool IsProgramAndKernelCachingAllowed = false;
492487
};
493488

494489
template <>

‎sycl/test/kernel-and-program/cache.cpp

Lines changed: 246 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,13 @@
1010

1111
#include <CL/sycl.hpp>
1212

13+
namespace pi = cl::sycl::detail::pi;
14+
namespace RT = cl::sycl::RT;
15+
16+
#define KERNEL_NAME_SRC "kernel_source"
17+
#define TEST_SOURCE "kernel void " KERNEL_NAME_SRC "(global int* a) " \
18+
"{ a[get_global_id(0)] += 1; }\n"
19+
1320
class Functor {
1421
public:
1522
void operator()(cl::sycl::item<1> Item) { (void)Item; }
@@ -37,6 +44,21 @@ struct TestContext {
3744
return std::move(Prog);
3845
}
3946

47+
cl::sycl::program
48+
getProgramWSource(const cl::sycl::string_class &BuildOptions = "") {
49+
cl::sycl::program Prog(Queue.get_context());
50+
51+
Prog.build_with_source(TEST_SOURCE, BuildOptions);
52+
53+
assert(Prog.get_state() == cl::sycl::program_state::linked &&
54+
"Linked state was expected");
55+
56+
assert(Prog.has_kernel<class SingleTask>() &&
57+
"Expecting SingleTask kernel exists");
58+
59+
return std::move(Prog);
60+
}
61+
4062
cl::sycl::program getCompiledProgram() {
4163
cl::sycl::program Prog(Queue.get_context());
4264

@@ -48,6 +70,43 @@ struct TestContext {
4870
return std::move(Prog);
4971
}
5072

73+
cl::sycl::program
74+
getCompiledAndLinkedProgram(const cl::sycl::string_class &CompileOptions = "",
75+
const cl::sycl::string_class &LinkOptions = "") {
76+
cl::sycl::program Prog(Queue.get_context());
77+
78+
Prog.compile_with_kernel_type<class SingleTask>(CompileOptions);
79+
80+
assert(Prog.get_state() == cl::sycl::program_state::compiled &&
81+
"Compiled state was expected");
82+
83+
Prog.link(LinkOptions);
84+
85+
assert(Prog.get_state() == cl::sycl::program_state::linked &&
86+
"Linked state was expected");
87+
88+
return std::move(Prog);
89+
}
90+
91+
cl::sycl::program
92+
getCompiledAndLinkedProgramWSource(
93+
const cl::sycl::string_class &CompileOptions = "",
94+
const cl::sycl::string_class &LinkOptions = "") {
95+
cl::sycl::program Prog(Queue.get_context());
96+
97+
Prog.compile_with_source(TEST_SOURCE, CompileOptions);
98+
99+
assert(Prog.get_state() == cl::sycl::program_state::compiled &&
100+
"Compiled state was expected");
101+
102+
Prog.link(LinkOptions);
103+
104+
assert(Prog.get_state() == cl::sycl::program_state::linked &&
105+
"Linked state was expected");
106+
107+
return std::move(Prog);
108+
}
109+
51110
cl::sycl::kernel getKernel(cl::sycl::program &Prog) {
52111
auto Kernel = Prog.get_kernel<class SingleTask>();
53112

@@ -58,10 +117,13 @@ struct TestContext {
58117

59118
return std::move(Kernel);
60119
}
61-
};
62120

63-
namespace pi = cl::sycl::detail::pi;
64-
namespace RT = cl::sycl::RT;
121+
cl::sycl::kernel getKernelWSource(cl::sycl::program &Prog) {
122+
auto Kernel = Prog.get_kernel(KERNEL_NAME_SRC);
123+
124+
return std::move(Kernel);
125+
}
126+
};
65127

66128
static void testProgramCachePositive() {
67129
TestContext TestCtx;
@@ -90,6 +152,86 @@ static void testProgramCacheNegativeCustomBuildOptions() {
90152
"Expecting empty program cache");
91153
}
92154

155+
static void testProgramCacheNegativeCompileLinkCustomOpts() {
156+
TestContext TestCtx;
157+
158+
{
159+
auto Prog = TestCtx.getCompiledAndLinkedProgram();
160+
161+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
162+
163+
assert(Ctx->getCachedPrograms().size() == 0 &&
164+
"Expecting empty program cache");
165+
}
166+
167+
{
168+
auto Prog = TestCtx.getCompiledAndLinkedProgram("-g", "-cl-no-signed-zeroes");
169+
170+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
171+
172+
assert(Ctx->getCachedPrograms().size() == 0 &&
173+
"Expecting empty program cache");
174+
}
175+
176+
{
177+
auto Prog = TestCtx.getCompiledAndLinkedProgram("", "-cl-no-signed-zeroes");
178+
179+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
180+
181+
assert(Ctx->getCachedPrograms().size() == 0 &&
182+
"Expecting empty program cache");
183+
}
184+
185+
{
186+
auto Prog = TestCtx.getCompiledAndLinkedProgram("-g", "");
187+
188+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
189+
190+
assert(Ctx->getCachedPrograms().size() == 0 &&
191+
"Expecting empty program cache");
192+
}
193+
}
194+
195+
static void testProgramCacheNegativeCompileLinkSource() {
196+
TestContext TestCtx;
197+
198+
{
199+
auto Prog = TestCtx.getCompiledAndLinkedProgramWSource();
200+
201+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
202+
203+
assert(Ctx->getCachedPrograms().size() == 0 &&
204+
"Expecting empty program cache");
205+
}
206+
207+
{
208+
auto Prog = TestCtx.getCompiledAndLinkedProgramWSource("-g", "-cl-no-signed-zeroes");
209+
210+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
211+
212+
assert(Ctx->getCachedPrograms().size() == 0 &&
213+
"Expecting empty program cache");
214+
}
215+
216+
{
217+
auto Prog = TestCtx.getCompiledAndLinkedProgramWSource("", "-cl-no-signed-zeroes");
218+
219+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
220+
221+
assert(Ctx->getCachedPrograms().size() == 0 &&
222+
"Expecting empty program cache");
223+
}
224+
225+
{
226+
auto Prog = TestCtx.getCompiledAndLinkedProgramWSource("-g", "");
227+
228+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
229+
230+
assert(Ctx->getCachedPrograms().size() == 0 &&
231+
"Expecting empty program cache");
232+
}
233+
}
234+
93235
static void testKernelCachePositive() {
94236
TestContext TestCtx;
95237

@@ -163,14 +305,115 @@ void testKernelCacheNegativeCustomBuildOptions() {
163305
}
164306
}
165307

308+
void testKernelCacheNegativeCompileLink() {
309+
TestContext TestCtx;
310+
311+
{
312+
auto Prog = TestCtx.getCompiledAndLinkedProgram();
313+
auto Kernel = TestCtx.getKernel(Prog);
314+
315+
if (!TestCtx.Queue.is_host()) {
316+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
317+
assert(Ctx->getCachedKernels().size() == 0 &&
318+
"Unexpected data in kernels cache");
319+
}
320+
}
321+
322+
{
323+
TestContext TestCtx1;
324+
auto Prog = TestCtx1.getCompiledAndLinkedProgram("-g", "-cl-no-signed-zeroes");
325+
auto Kernel = TestCtx1.getKernel(Prog);
326+
327+
if (!TestCtx1.Queue.is_host()) {
328+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
329+
assert(Ctx->getCachedKernels().size() == 0 &&
330+
"Unexpected data in kernels cache");
331+
}
332+
}
333+
334+
{
335+
auto Prog = TestCtx.getCompiledAndLinkedProgram("-g", "");
336+
auto Kernel = TestCtx.getKernel(Prog);
337+
338+
if (!TestCtx.Queue.is_host()) {
339+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
340+
assert(Ctx->getCachedKernels().size() == 0 &&
341+
"Unexpected data in kernels cache");
342+
}
343+
}
344+
345+
{
346+
auto Prog = TestCtx.getCompiledAndLinkedProgram("", "-cl-no-signed-zeroes");
347+
auto Kernel = TestCtx.getKernel(Prog);
348+
349+
if (!TestCtx.Queue.is_host()) {
350+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
351+
assert(Ctx->getCachedKernels().size() == 0 &&
352+
"Unexpected data in kernels cache");
353+
}
354+
}
355+
}
356+
357+
void testKernelCacheNegativeCompileLinkSource() {
358+
TestContext TestCtx;
359+
360+
{
361+
auto Prog = TestCtx.getCompiledAndLinkedProgramWSource();
362+
auto Kernel = TestCtx.getKernelWSource(Prog);
363+
364+
if (!TestCtx.Queue.is_host()) {
365+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
366+
assert(Ctx->getCachedKernels().size() == 0 &&
367+
"Unexpected data in kernels cache");
368+
}
369+
}
370+
371+
{
372+
auto Prog = TestCtx.getCompiledAndLinkedProgramWSource("-g", "-cl-no-signed-zeroes");
373+
auto Kernel = TestCtx.getKernelWSource(Prog);
374+
375+
if (!TestCtx.Queue.is_host()) {
376+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
377+
assert(Ctx->getCachedKernels().size() == 0 &&
378+
"Unexpected data in kernels cache");
379+
}
380+
}
381+
382+
{
383+
auto Prog = TestCtx.getCompiledAndLinkedProgramWSource("-g", "");
384+
auto Kernel = TestCtx.getKernelWSource(Prog);
385+
386+
if (!TestCtx.Queue.is_host()) {
387+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
388+
assert(Ctx->getCachedKernels().size() == 0 &&
389+
"Unexpected data in kernels cache");
390+
}
391+
}
392+
393+
{
394+
auto Prog = TestCtx.getCompiledAndLinkedProgramWSource("", "-cl-no-signed-zeroes");
395+
auto Kernel = TestCtx.getKernelWSource(Prog);
396+
397+
if (!TestCtx.Queue.is_host()) {
398+
auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
399+
assert(Ctx->getCachedKernels().size() == 0 &&
400+
"Unexpected data in kernels cache");
401+
}
402+
}
403+
}
404+
166405
int main() {
167406
testProgramCachePositive();
168407
testProgramCacheNegativeCustomBuildOptions();
408+
testProgramCacheNegativeCompileLinkCustomOpts();
409+
testProgramCacheNegativeCompileLinkSource();
169410

170411
testKernelCachePositive();
171412
testKernelCacheNegativeLinkedProgram();
172413
testKernelCacheNegativeOCLProgram();
173414
testKernelCacheNegativeCustomBuildOptions();
415+
testKernelCacheNegativeCompileLink();
416+
testKernelCacheNegativeCompileLinkSource();
174417

175418
return 0;
176419
}

0 commit comments

Comments
 (0)