Visible to Intel only — GUID: GUID-96EF1537-DF86-4290-A7F6-2FDD8BC98F43
Execution Model Overview
Thread Mapping and GPU Occupancy
Kernels
Using Libraries for GPU Offload
Host/Device Memory, Buffer and USM
Host/Device Coordination
Using Multiple Heterogeneous Devices
Compilation
OpenMP Offloading Tuning Guide
Multi-GPU, Multi-Stack and Multi-C-Slice Architecture and Programming
Level Zero
Performance Profiling and Analysis
Configuring GPU Device
Sub-Groups and SIMD Vectorization
Removing Conditional Checks
Registerization and Avoiding Register Spills
Small Register Mode vs. Large Register Mode
Shared Local Memory
Pointer Aliasing and the Restrict Directive
Synchronization among Threads in a Kernel
Considerations for Selecting Work-Group Size
Reduction
Kernel Launch
Executing Multiple Kernels on the Device at the Same Time
Submitting Kernels to Multiple Queues
Avoiding Redundant Queue Constructions
Programming Intel® XMX Using SYCL Joint Matrix Extension
Doing I/O in the Kernel
Visible to Intel only — GUID: GUID-96EF1537-DF86-4290-A7F6-2FDD8BC98F43
Avoid moving data back and forth between host and device
The cost of moving data between host and device is quite high, especially in the case of discrete accelerators. So it is very important to avoid data transfers between host and device as much as possible. In some situations it may be required to bring the data that was computed by a kernel on the accelerator to the host and do some operation on it and send it back to the device for further processing. In such situation we will end up paying for the cost of device to host transfer and then again host to device transfer.
Consider the following example, where one kernel produces data through some operation (in this case vector add) into a new vector. This vector is then transformed into another vector by applying a function on each value and then fed as input into another kernel for some additional computation. This form of computation is quite common and occurs in many domains where algorithms are iterative and output from one computation needs to be fed as input into another computation. One classic example is machine learning models which are structured as layers of computation and output of one layer is input to the next layer.
double myFunc1(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b,
AlignedVector<int> &c, AlignedVector<int> &d,
AlignedVector<int> &res, int iter) {
sycl::range num_items{a.size()};
VectorAllocator<int> alloc;
AlignedVector<int> sum(a.size(), alloc);
const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
sycl::buffer a_buf(a, props);
sycl::buffer b_buf(b, props);
sycl::buffer c_buf(b, props);
sycl::buffer d_buf(b, props);
sycl::buffer res_buf(res, props);
sycl::buffer sum_buf(sum.data(), num_items, props);
Timer timer;
for (int i = 0; i < iter; i++) {
// kernel1
q.submit([&](auto &h) {
// Input accessors
sycl::accessor a_acc(a_buf, h, sycl::read_only);
sycl::accessor b_acc(b_buf, h, sycl::read_only);
// Output accessor
sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);
h.parallel_for(num_items,
[=](auto id) { sum_acc[id] = a_acc[id] + b_acc[id]; });
});
{
sycl::host_accessor h_acc(sum_buf);
for (int j = 0; j < a.size(); j++)
if (h_acc[j] > 10)
h_acc[j] = 1;
else
h_acc[j] = 0;
}
// kernel2
q.submit([&](auto &h) {
// Input accessors
sycl::accessor sum_acc(sum_buf, h, sycl::read_only);
sycl::accessor c_acc(c_buf, h, sycl::read_only);
sycl::accessor d_acc(d_buf, h, sycl::read_only);
// Output accessor
sycl::accessor res_acc(res_buf, h, sycl::write_only, sycl::no_init);
h.parallel_for(num_items, [=](auto id) {
res_acc[id] = sum_acc[id] * c_acc[id] + d_acc[id];
});
});
q.wait();
}
double elapsed = timer.Elapsed() / iter;
return (elapsed);
} // end myFunc1
Instead of bringing the data to the host and applying the function to the data and sending it back to the device in the second kernel, you can create a kernel to execute this function on the device itself. This has the advantage of avoiding the round trip of data from device to host. This technique is shown in the example below, which is functionally the same as the code before. We now introduce a third kernel kernel3 that operates on the intermediate data in accum_buf in between kernel1 and kernel2.
double myFunc2(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b,
AlignedVector<int> &c, AlignedVector<int> &d,
AlignedVector<int> &res, int iter) {
sycl::range num_items{a.size()};
VectorAllocator<int> alloc;
AlignedVector<int> sum(a.size(), alloc);
const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
sycl::buffer a_buf(a, props);
sycl::buffer b_buf(b, props);
sycl::buffer c_buf(b, props);
sycl::buffer d_buf(b, props);
sycl::buffer res_buf(res, props);
sycl::buffer sum_buf(sum.data(), num_items, props);
Timer timer;
for (int i = 0; i < iter; i++) {
// kernel1
q.submit([&](auto &h) {
// Input accessors
sycl::accessor a_acc(a_buf, h, sycl::read_only);
sycl::accessor b_acc(b_buf, h, sycl::read_only);
// Output accessor
sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);
h.parallel_for(num_items,
[=](auto i) { sum_acc[i] = a_acc[i] + b_acc[i]; });
});
// kernel3
q.submit([&](auto &h) {
sycl::accessor sum_acc(sum_buf, h, sycl::read_write);
h.parallel_for(num_items, [=](auto id) {
if (sum_acc[id] > 10)
sum_acc[id] = 1;
else
sum_acc[id] = 0;
});
});
// kernel2
q.submit([&](auto &h) {
// Input accessors
sycl::accessor sum_acc(sum_buf, h, sycl::read_only);
sycl::accessor c_acc(c_buf, h, sycl::read_only);
sycl::accessor d_acc(d_buf, h, sycl::read_only);
// Output accessor
sycl::accessor res_acc(res_buf, h, sycl::write_only, sycl::no_init);
h.parallel_for(num_items, [=](auto i) {
res_acc[i] = sum_acc[i] * c_acc[i] + d_acc[i];
});
});
q.wait();
}
double elapsed = timer.Elapsed() / iter;
return (elapsed);
} // end myFunc2
There are other ways to optimize this example. For instance, the clipping operation in kernel3 can be merged into the computation of kernel1 as shown below. This is kernel fusion and has the added advantage of not launching a third kernel. The SYCL compiler cannot do this kind of optimization. In some specific domains like machine learning, there are graph compilers that operate on the ML models and fuse the operations, which has the same impact.
double myFunc3(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b,
AlignedVector<int> &c, AlignedVector<int> &d,
AlignedVector<int> &res, int iter) {
sycl::range num_items{a.size()};
VectorAllocator<int> alloc;
AlignedVector<int> sum(a.size(), alloc);
const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
sycl::buffer a_buf(a, props);
sycl::buffer b_buf(b, props);
sycl::buffer c_buf(b, props);
sycl::buffer d_buf(b, props);
sycl::buffer res_buf(res, props);
sycl::buffer sum_buf(sum.data(), num_items, props);
Timer timer;
for (int i = 0; i < iter; i++) {
// kernel1
q.submit([&](auto &h) {
// Input accessors
sycl::accessor a_acc(a_buf, h, sycl::read_only);
sycl::accessor b_acc(b_buf, h, sycl::read_only);
// Output accessor
sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);
h.parallel_for(num_items, [=](auto i) {
int t = a_acc[i] + b_acc[i];
if (t > 10)
sum_acc[i] = 1;
else
sum_acc[i] = 0;
});
});
// kernel2
q.submit([&](auto &h) {
// Input accessors
sycl::accessor sum_acc(sum_buf, h, sycl::read_only);
sycl::accessor c_acc(c_buf, h, sycl::read_only);
sycl::accessor d_acc(d_buf, h, sycl::read_only);
// Output accessor
sycl::accessor res_acc(res_buf, h, sycl::write_only, sycl::no_init);
h.parallel_for(num_items, [=](auto i) {
res_acc[i] = sum_acc[i] * c_acc[i] + d_acc[i];
});
});
q.wait();
}
double elapsed = timer.Elapsed() / iter;
return (elapsed);
} // end myFunc3
We can take this kernel fusion one level further and fuse both kernel1 and kernel2 as shown in the code below. This gives very good performance since it avoids the intermediate accum_buf completely, saving memory in addition to launching an additional kernel. Most of the performance benefit in this case is due to improvement in locality of memory references.
double myFunc4(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b,
AlignedVector<int> &c, AlignedVector<int> &d,
AlignedVector<int> &res, int iter) {
sycl::range num_items{a.size()};
VectorAllocator<int> alloc;
const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
sycl::buffer a_buf(a, props);
sycl::buffer b_buf(b, props);
sycl::buffer c_buf(b, props);
sycl::buffer d_buf(b, props);
sycl::buffer res_buf(res, props);
Timer timer;
for (int i = 0; i < iter; i++) {
// kernel1
q.submit([&](auto &h) {
// Input accessors
sycl::accessor a_acc(a_buf, h, sycl::read_only);
sycl::accessor b_acc(b_buf, h, sycl::read_only);
sycl::accessor c_acc(c_buf, h, sycl::read_only);
sycl::accessor d_acc(d_buf, h, sycl::read_only);
// Output accessor
sycl::accessor res_acc(res_buf, h, sycl::write_only, sycl::no_init);
h.parallel_for(num_items, [=](auto i) {
int t = a_acc[i] + b_acc[i];
if (t > 10)
res_acc[i] = c_acc[i] + d_acc[i];
else
res_acc[i] = d_acc[i];
});
});
q.wait();
}
double elapsed = timer.Elapsed() / iter;
return (elapsed);
} // end myFunc4