Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit 2b46e42

Browse files
sergey-semenovsergey-semenov
authored
[ESIMD] Make profiling support optional in ESIMD tests (#1604)
Change ESIMD tests to only use profiling if the used device supports it. Fixes failures on ESIMD emulator, which does not support profiling.
1 parent 432cbbc commit 2b46e42

24 files changed

+280
-177
lines changed

‎SYCL/ESIMD/BitonicSortK.cpp

Lines changed: 18 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -606,6 +606,9 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
606606
num_iters = 2;
607607
}
608608

609+
const bool profiling =
610+
pQueue_->has_property<property::queue::enable_profiling>();
611+
609612
// num_iters + 1, iteration#0 is for warmup
610613
for (int iter = 0; iter <= num_iters; ++iter) {
611614
try {
@@ -622,9 +625,11 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
622625
});
623626
});
624627
e.wait();
625-
double etime = esimd_test::report_time("kernel1 time", e, e);
626-
if (iter > 0)
627-
kernel_times += etime;
628+
if (profiling) {
629+
double etime = esimd_test::report_time("kernel1 time", e, e);
630+
if (iter > 0)
631+
kernel_times += etime;
632+
}
628633
} catch (sycl::exception const &e) {
629634
std::cout << "SYCL exception caught: " << e.what() << '\n';
630635
return 0;
@@ -673,19 +678,21 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
673678
}
674679

675680
mergeEvent[k - 1].wait();
676-
double etime = esimd_test::report_time("kernel2 time", mergeEvent[0],
677-
mergeEvent[k - 1]);
678-
if (iter > 0)
679-
kernel_times += etime;
680-
else
681+
if (profiling) {
682+
double etime = esimd_test::report_time("kernel2 time", mergeEvent[0],
683+
mergeEvent[k - 1]);
684+
if (iter > 0)
685+
kernel_times += etime;
686+
}
687+
if (iter == 0)
681688
start = timer.Elapsed();
682689
}
683690

684691
// End timer.
685692
double end = timer.Elapsed();
686693

687-
esimd_test::display_timing_stats(kernel_times, num_iters,
688-
(end - start) * 1000);
694+
esimd_test::display_timing_stats(profiling ? &kernel_times : nullptr,
695+
num_iters, (end - start) * 1000);
689696
return 1;
690697
}
691698

@@ -700,8 +707,7 @@ int main(int argc, char *argv[]) {
700707
int size = 1 << LOG2_ELEMENTS;
701708
cout << "BitonicSort (" << size << ") Start..." << std::endl;
702709

703-
queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler(),
704-
property::queue::enable_profiling{});
710+
queue q = esimd_test::createQueue();
705711

706712
BitonicSort bitonicSort;
707713

‎SYCL/ESIMD/BitonicSortKv2.cpp

Lines changed: 17 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -516,6 +516,8 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
516516
// Start Timer
517517
esimd_test::Timer timer;
518518
double start;
519+
const bool profiling =
520+
pQueue_->has_property<sycl::property::queue::enable_profiling>();
519521

520522
// Launches the task on the GPU.
521523
double kernel_times = 0;
@@ -539,9 +541,11 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
539541
});
540542
});
541543
e.wait();
542-
double etime = esimd_test::report_time("kernel1 time", e, e);
543-
if (iter > 0)
544-
kernel_times += etime;
544+
if (profiling) {
545+
double etime = esimd_test::report_time("kernel1 time", e, e);
546+
if (iter > 0)
547+
kernel_times += etime;
548+
}
545549
} catch (sycl::exception const &e) {
546550
std::cout << "SYCL exception caught: " << e.what() << '\n';
547551
return 0;
@@ -589,19 +593,21 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
589593
}
590594

591595
mergeEvent[k - 1].wait();
592-
double etime = esimd_test::report_time("kernel2 time", mergeEvent[0],
593-
mergeEvent[k - 1]);
594-
if (iter > 0)
595-
kernel_times += etime;
596-
else
596+
if (profiling) {
597+
double etime = esimd_test::report_time("kernel2 time", mergeEvent[0],
598+
mergeEvent[k - 1]);
599+
if (iter > 0)
600+
kernel_times += etime;
601+
}
602+
if (iter == 0)
597603
start = timer.Elapsed();
598604
}
599605

600606
// End timer.
601607
double end = timer.Elapsed();
602608

603-
esimd_test::display_timing_stats(kernel_times, num_iters,
604-
(end - start) * 1000);
609+
esimd_test::display_timing_stats(profiling ? &kernel_times : nullptr,
610+
num_iters, (end - start) * 1000);
605611
return 1;
606612
}
607613

@@ -617,11 +623,7 @@ int main(int argc, char *argv[]) {
617623
int size = 1 << LOG2_ELEMENTS;
618624
cout << "BitonicSort (" << size << ") Start..." << std::endl;
619625

620-
sycl::property_list props{sycl::property::queue::enable_profiling{},
621-
sycl::property::queue::in_order()};
622-
623-
queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler(),
624-
props);
626+
queue q = esimd_test::createQueue(/*inOrder*/ true);
625627

626628
BitonicSort bitonicSort;
627629

‎SYCL/ESIMD/Prefix_Local_sum1.cpp

Lines changed: 11 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -125,8 +125,7 @@ int main(int argc, char *argv[]) {
125125

126126
sycl::range<2> LocalRange{1, 1};
127127

128-
queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler(),
129-
property::queue::enable_profiling{});
128+
queue q = esimd_test::createQueue();
130129

131130
auto dev = q.get_device();
132131
std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
@@ -157,6 +156,8 @@ int main(int argc, char *argv[]) {
157156

158157
double kernel_times = 0;
159158
unsigned num_iters = 10;
159+
const bool profiling =
160+
q.has_property<sycl::property::queue::enable_profiling>();
160161

161162
try {
162163
for (int iter = 0; iter <= num_iters; ++iter) {
@@ -168,10 +169,12 @@ int main(int argc, char *argv[]) {
168169
});
169170
});
170171
e0.wait();
171-
double etime = esimd_test::report_time("kernel time", e0, e0);
172-
if (iter > 0)
173-
kernel_times += etime;
174-
else
172+
if (profiling) {
173+
double etime = esimd_test::report_time("kernel time", e0, e0);
174+
if (iter > 0)
175+
kernel_times += etime;
176+
}
177+
if (iter == 0)
175178
start = timer.Elapsed();
176179
}
177180
} catch (sycl::exception const &e) {
@@ -185,8 +188,8 @@ int main(int argc, char *argv[]) {
185188
// End timer.
186189
double end = timer.Elapsed();
187190

188-
esimd_test::display_timing_stats(kernel_times, num_iters,
189-
(end - start) * 1000);
191+
esimd_test::display_timing_stats(profiling ? &kernel_times : nullptr,
192+
num_iters, (end - start) * 1000);
190193

191194
bool pass = memcmp(pDeviceOutputs, pExpectOutputs,
192195
size * TUPLE_SZ * sizeof(unsigned int)) == 0;

‎SYCL/ESIMD/Prefix_Local_sum2.cpp

Lines changed: 11 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -120,8 +120,7 @@ int main(int argc, char *argv[]) {
120120

121121
sycl::range<2> LocalRange{1, 1};
122122

123-
queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler(),
124-
property::queue::enable_profiling{});
123+
queue q = esimd_test::createQueue();
125124

126125
auto dev = q.get_device();
127126
std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
@@ -149,6 +148,8 @@ int main(int argc, char *argv[]) {
149148

150149
double kernel_times = 0;
151150
unsigned num_iters = 10;
151+
const bool profiling =
152+
q.has_property<sycl::property::queue::enable_profiling>();
152153

153154
try {
154155
for (int iter = 0; iter <= num_iters; ++iter) {
@@ -162,10 +163,12 @@ int main(int argc, char *argv[]) {
162163
});
163164
});
164165
e1.wait();
165-
double etime = esimd_test::report_time("kernel time", e1, e1);
166-
if (iter > 0)
167-
kernel_times += etime;
168-
else
166+
if (profiling) {
167+
double etime = esimd_test::report_time("kernel time", e1, e1);
168+
if (iter > 0)
169+
kernel_times += etime;
170+
}
171+
if (iter == 0)
169172
start = timer.Elapsed();
170173
}
171174
} catch (sycl::exception const &e) {
@@ -179,8 +182,8 @@ int main(int argc, char *argv[]) {
179182
// End timer.
180183
double end = timer.Elapsed();
181184

182-
esimd_test::display_timing_stats(kernel_times, num_iters,
183-
(end - start) * 1000);
185+
esimd_test::display_timing_stats(profiling ? &kernel_times : nullptr,
186+
num_iters, (end - start) * 1000);
184187

185188
bool pass = memcmp(pDeviceOutputs, pExpectOutputs,
186189
size * TUPLE_SZ * sizeof(unsigned int)) == 0;

‎SYCL/ESIMD/Prefix_Local_sum3.cpp

Lines changed: 15 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -236,6 +236,8 @@ double hierarchical_prefix(queue &q, unsigned *buf, unsigned elem_stride,
236236
unsigned thrd_stride, unsigned n_entries,
237237
unsigned entry_per_th) {
238238
double kernel_times = 0;
239+
const bool profiling =
240+
q.has_property<sycl::property::queue::enable_profiling>();
239241
try {
240242
if (n_entries <= REMAINING_ENTRIES) {
241243
#ifdef DEBUG_DUMPS
@@ -252,7 +254,8 @@ double hierarchical_prefix(queue &q, unsigned *buf, unsigned elem_stride,
252254
});
253255
});
254256
e.wait();
255-
kernel_times += esimd_test::report_time("kernel1 time", e, e);
257+
if (profiling)
258+
kernel_times += esimd_test::report_time("kernel1 time", e, e);
256259
return kernel_times;
257260
}
258261

@@ -272,7 +275,8 @@ double hierarchical_prefix(queue &q, unsigned *buf, unsigned elem_stride,
272275
});
273276
});
274277
e.wait();
275-
kernel_times += esimd_test::report_time("kernel2 time", e, e);
278+
if (profiling)
279+
kernel_times += esimd_test::report_time("kernel2 time", e, e);
276280
} else {
277281
auto e = q.submit([&](handler &cgh) {
278282
cgh.parallel_for<class Accum_iterative2>(
@@ -283,7 +287,8 @@ double hierarchical_prefix(queue &q, unsigned *buf, unsigned elem_stride,
283287
});
284288
});
285289
e.wait();
286-
kernel_times += esimd_test::report_time("kernel3 time", e, e);
290+
if (profiling)
291+
kernel_times += esimd_test::report_time("kernel3 time", e, e);
287292
}
288293
} catch (sycl::exception const &e) {
289294
std::cout << "SYCL exception caught: " << e.what() << '\n';
@@ -324,8 +329,7 @@ int main(int argc, char *argv[]) {
324329

325330
sycl::range<2> LocalRange{1, 1};
326331

327-
queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler(),
328-
property::queue::enable_profiling{});
332+
queue q = esimd_test::createQueue();
329333

330334
auto dev = q.get_device();
331335
std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
@@ -353,22 +357,24 @@ int main(int argc, char *argv[]) {
353357

354358
double kernel_times = 0;
355359
unsigned num_iters = 10;
360+
const bool profiling =
361+
q.has_property<sycl::property::queue::enable_profiling>();
356362

357363
for (int iter = 0; iter <= num_iters; ++iter) {
358364
memcpy(pDeviceOutputs, pInputs, size * TUPLE_SZ * sizeof(unsigned));
359365
double etime = hierarchical_prefix(q, pDeviceOutputs, 1, PREFIX_ENTRIES,
360366
size, PREFIX_ENTRIES);
361-
if (iter > 0)
367+
if (profiling && iter > 0)
362368
kernel_times += etime;
363-
else
369+
if (iter == 0)
364370
start = timer.Elapsed();
365371
}
366372

367373
// End timer.
368374
double end = timer.Elapsed();
369375

370-
esimd_test::display_timing_stats(kernel_times, num_iters,
371-
(end - start) * 1000);
376+
esimd_test::display_timing_stats(profiling ? &kernel_times : nullptr,
377+
num_iters, (end - start) * 1000);
372378

373379
compute_local_prefixsum(pExpectOutputs, size, 1, PREFIX_ENTRIES,
374380
PREFIX_ENTRIES);

‎SYCL/ESIMD/Stencil.cpp

Lines changed: 11 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -88,8 +88,7 @@ int main(int argc, char *argv[]) {
8888
<< std::endl;
8989
sycl::range<2> LocalRange{1, 1};
9090

91-
queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler(),
92-
property::queue::enable_profiling{});
91+
queue q = esimd_test::createQueue();
9392

9493
auto dev = q.get_device();
9594
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
@@ -109,6 +108,8 @@ int main(int argc, char *argv[]) {
109108

110109
double kernel_times = 0;
111110
unsigned num_iters = 10;
111+
const bool profiling =
112+
q.has_property<sycl::property::queue::enable_profiling>();
112113

113114
try {
114115
for (int iter = 0; iter <= num_iters; ++iter) {
@@ -185,10 +186,12 @@ int main(int argc, char *argv[]) {
185186
});
186187
});
187188
e.wait();
188-
double etime = esimd_test::report_time("kernel time", e, e);
189-
if (iter > 0)
190-
kernel_times += etime;
191-
else
189+
if (profiling) {
190+
double etime = esimd_test::report_time("kernel time", e, e);
191+
if (iter > 0)
192+
kernel_times += etime;
193+
}
194+
if (iter == 0)
192195
start = timer.Elapsed();
193196
}
194197
} catch (sycl::exception const &e) {
@@ -201,8 +204,8 @@ int main(int argc, char *argv[]) {
201204
// End timer.
202205
double end = timer.Elapsed();
203206

204-
esimd_test::display_timing_stats(kernel_times, num_iters,
205-
(end - start) * 1000);
207+
esimd_test::display_timing_stats(profiling ? &kernel_times : nullptr,
208+
num_iters, (end - start) * 1000);
206209

207210
// check result
208211
bool passed = CheckResults(outputMatrix, inputMatrix, DIM_SIZE);

0 commit comments

Comments
 (0)