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

[ESIMD] Make profiling support optional in ESIMD tests #1604

Merged
merged 3 commits into from
Feb 20, 2023
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
30 changes: 18 additions & 12 deletions SYCL/ESIMD/BitonicSortK.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -606,6 +606,9 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
num_iters = 2;
}

const bool profiling =
pQueue_->has_property<property::queue::enable_profiling>();

// num_iters + 1, iteration#0 is for warmup
for (int iter = 0; iter <= num_iters; ++iter) {
try {
Expand All @@ -622,9 +625,11 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
});
});
e.wait();
double etime = esimd_test::report_time("kernel1 time", e, e);
if (iter > 0)
kernel_times += etime;
if (profiling) {
double etime = esimd_test::report_time("kernel1 time", e, e);
if (iter > 0)
kernel_times += etime;
}
} catch (sycl::exception const &e) {
std::cout << "SYCL exception caught: " << e.what() << '\n';
return 0;
Expand Down Expand Up @@ -673,19 +678,21 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
}

mergeEvent[k - 1].wait();
double etime = esimd_test::report_time("kernel2 time", mergeEvent[0],
mergeEvent[k - 1]);
if (iter > 0)
kernel_times += etime;
else
if (profiling) {
double etime = esimd_test::report_time("kernel2 time", mergeEvent[0],
mergeEvent[k - 1]);
if (iter > 0)
kernel_times += etime;
}
if (iter == 0)
start = timer.Elapsed();
}

// End timer.
double end = timer.Elapsed();

esimd_test::display_timing_stats(kernel_times, num_iters,
(end - start) * 1000);
esimd_test::display_timing_stats(profiling ? &kernel_times : nullptr,
num_iters, (end - start) * 1000);
return 1;
}

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

queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler(),
property::queue::enable_profiling{});
queue q = esimd_test::createQueue();

BitonicSort bitonicSort;

Expand Down
32 changes: 17 additions & 15 deletions SYCL/ESIMD/BitonicSortKv2.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -516,6 +516,8 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
// Start Timer
esimd_test::Timer timer;
double start;
const bool profiling =
pQueue_->has_property<sycl::property::queue::enable_profiling>();

// Launches the task on the GPU.
double kernel_times = 0;
Expand All @@ -539,9 +541,11 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
});
});
e.wait();
double etime = esimd_test::report_time("kernel1 time", e, e);
if (iter > 0)
kernel_times += etime;
if (profiling) {
double etime = esimd_test::report_time("kernel1 time", e, e);
if (iter > 0)
kernel_times += etime;
}
} catch (sycl::exception const &e) {
std::cout << "SYCL exception caught: " << e.what() << '\n';
return 0;
Expand Down Expand Up @@ -589,19 +593,21 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
}

mergeEvent[k - 1].wait();
double etime = esimd_test::report_time("kernel2 time", mergeEvent[0],
mergeEvent[k - 1]);
if (iter > 0)
kernel_times += etime;
else
if (profiling) {
double etime = esimd_test::report_time("kernel2 time", mergeEvent[0],
mergeEvent[k - 1]);
if (iter > 0)
kernel_times += etime;
}
if (iter == 0)
start = timer.Elapsed();
}

// End timer.
double end = timer.Elapsed();

esimd_test::display_timing_stats(kernel_times, num_iters,
(end - start) * 1000);
esimd_test::display_timing_stats(profiling ? &kernel_times : nullptr,
num_iters, (end - start) * 1000);
return 1;
}

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

sycl::property_list props{sycl::property::queue::enable_profiling{},
sycl::property::queue::in_order()};

queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler(),
props);
queue q = esimd_test::createQueue(/*inOrder*/ true);

BitonicSort bitonicSort;

Expand Down
19 changes: 11 additions & 8 deletions SYCL/ESIMD/Prefix_Local_sum1.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -125,8 +125,7 @@ int main(int argc, char *argv[]) {

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

queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler(),
property::queue::enable_profiling{});
queue q = esimd_test::createQueue();

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

double kernel_times = 0;
unsigned num_iters = 10;
const bool profiling =
q.has_property<sycl::property::queue::enable_profiling>();

try {
for (int iter = 0; iter <= num_iters; ++iter) {
Expand All @@ -168,10 +169,12 @@ int main(int argc, char *argv[]) {
});
});
e0.wait();
double etime = esimd_test::report_time("kernel time", e0, e0);
if (iter > 0)
kernel_times += etime;
else
if (profiling) {
double etime = esimd_test::report_time("kernel time", e0, e0);
if (iter > 0)
kernel_times += etime;
}
if (iter == 0)
start = timer.Elapsed();
}
} catch (sycl::exception const &e) {
Expand All @@ -185,8 +188,8 @@ int main(int argc, char *argv[]) {
// End timer.
double end = timer.Elapsed();

esimd_test::display_timing_stats(kernel_times, num_iters,
(end - start) * 1000);
esimd_test::display_timing_stats(profiling ? &kernel_times : nullptr,
num_iters, (end - start) * 1000);

bool pass = memcmp(pDeviceOutputs, pExpectOutputs,
size * TUPLE_SZ * sizeof(unsigned int)) == 0;
Expand Down
19 changes: 11 additions & 8 deletions SYCL/ESIMD/Prefix_Local_sum2.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -120,8 +120,7 @@ int main(int argc, char *argv[]) {

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

queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler(),
property::queue::enable_profiling{});
queue q = esimd_test::createQueue();

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

double kernel_times = 0;
unsigned num_iters = 10;
const bool profiling =
q.has_property<sycl::property::queue::enable_profiling>();

try {
for (int iter = 0; iter <= num_iters; ++iter) {
Expand All @@ -162,10 +163,12 @@ int main(int argc, char *argv[]) {
});
});
e1.wait();
double etime = esimd_test::report_time("kernel time", e1, e1);
if (iter > 0)
kernel_times += etime;
else
if (profiling) {
double etime = esimd_test::report_time("kernel time", e1, e1);
if (iter > 0)
kernel_times += etime;
}
if (iter == 0)
start = timer.Elapsed();
}
} catch (sycl::exception const &e) {
Expand All @@ -179,8 +182,8 @@ int main(int argc, char *argv[]) {
// End timer.
double end = timer.Elapsed();

esimd_test::display_timing_stats(kernel_times, num_iters,
(end - start) * 1000);
esimd_test::display_timing_stats(profiling ? &kernel_times : nullptr,
num_iters, (end - start) * 1000);

bool pass = memcmp(pDeviceOutputs, pExpectOutputs,
size * TUPLE_SZ * sizeof(unsigned int)) == 0;
Expand Down
24 changes: 15 additions & 9 deletions SYCL/ESIMD/Prefix_Local_sum3.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -236,6 +236,8 @@ double hierarchical_prefix(queue &q, unsigned *buf, unsigned elem_stride,
unsigned thrd_stride, unsigned n_entries,
unsigned entry_per_th) {
double kernel_times = 0;
const bool profiling =
q.has_property<sycl::property::queue::enable_profiling>();
try {
if (n_entries <= REMAINING_ENTRIES) {
#ifdef DEBUG_DUMPS
Expand All @@ -252,7 +254,8 @@ double hierarchical_prefix(queue &q, unsigned *buf, unsigned elem_stride,
});
});
e.wait();
kernel_times += esimd_test::report_time("kernel1 time", e, e);
if (profiling)
kernel_times += esimd_test::report_time("kernel1 time", e, e);
return kernel_times;
}

Expand All @@ -272,7 +275,8 @@ double hierarchical_prefix(queue &q, unsigned *buf, unsigned elem_stride,
});
});
e.wait();
kernel_times += esimd_test::report_time("kernel2 time", e, e);
if (profiling)
kernel_times += esimd_test::report_time("kernel2 time", e, e);
} else {
auto e = q.submit([&](handler &cgh) {
cgh.parallel_for<class Accum_iterative2>(
Expand All @@ -283,7 +287,8 @@ double hierarchical_prefix(queue &q, unsigned *buf, unsigned elem_stride,
});
});
e.wait();
kernel_times += esimd_test::report_time("kernel3 time", e, e);
if (profiling)
kernel_times += esimd_test::report_time("kernel3 time", e, e);
}
} catch (sycl::exception const &e) {
std::cout << "SYCL exception caught: " << e.what() << '\n';
Expand Down Expand Up @@ -324,8 +329,7 @@ int main(int argc, char *argv[]) {

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

queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler(),
property::queue::enable_profiling{});
queue q = esimd_test::createQueue();

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

double kernel_times = 0;
unsigned num_iters = 10;
const bool profiling =
q.has_property<sycl::property::queue::enable_profiling>();

for (int iter = 0; iter <= num_iters; ++iter) {
memcpy(pDeviceOutputs, pInputs, size * TUPLE_SZ * sizeof(unsigned));
double etime = hierarchical_prefix(q, pDeviceOutputs, 1, PREFIX_ENTRIES,
size, PREFIX_ENTRIES);
if (iter > 0)
if (profiling && iter > 0)
kernel_times += etime;
else
if (iter == 0)
start = timer.Elapsed();
}

// End timer.
double end = timer.Elapsed();

esimd_test::display_timing_stats(kernel_times, num_iters,
(end - start) * 1000);
esimd_test::display_timing_stats(profiling ? &kernel_times : nullptr,
num_iters, (end - start) * 1000);

compute_local_prefixsum(pExpectOutputs, size, 1, PREFIX_ENTRIES,
PREFIX_ENTRIES);
Expand Down
19 changes: 11 additions & 8 deletions SYCL/ESIMD/Stencil.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -88,8 +88,7 @@ int main(int argc, char *argv[]) {
<< std::endl;
sycl::range<2> LocalRange{1, 1};

queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler(),
property::queue::enable_profiling{});
queue q = esimd_test::createQueue();

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

double kernel_times = 0;
unsigned num_iters = 10;
const bool profiling =
q.has_property<sycl::property::queue::enable_profiling>();

try {
for (int iter = 0; iter <= num_iters; ++iter) {
Expand Down Expand Up @@ -185,10 +186,12 @@ int main(int argc, char *argv[]) {
});
});
e.wait();
double etime = esimd_test::report_time("kernel time", e, e);
if (iter > 0)
kernel_times += etime;
else
if (profiling) {
double etime = esimd_test::report_time("kernel time", e, e);
if (iter > 0)
kernel_times += etime;
}
if (iter == 0)
start = timer.Elapsed();
}
} catch (sycl::exception const &e) {
Expand All @@ -201,8 +204,8 @@ int main(int argc, char *argv[]) {
// End timer.
double end = timer.Elapsed();

esimd_test::display_timing_stats(kernel_times, num_iters,
(end - start) * 1000);
esimd_test::display_timing_stats(profiling ? &kernel_times : nullptr,
num_iters, (end - start) * 1000);

// check result
bool passed = CheckResults(outputMatrix, inputMatrix, DIM_SIZE);
Expand Down
Loading