Visible to Intel only — GUID: GUID-B2E04ED2-8B89-4CF9-8870-23AACE592D01
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
Prefetch
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-B2E04ED2-8B89-4CF9-8870-23AACE592D01
Optimizing Memory Movement Between Host and Accelerator
Buffers can be created using properties to control how they are allocated. One such property is use_host_ptr which informs the runtime that if possible, the host memory should be directly used by the buffer instead of a copy. This avoids the need to copy the content of the buffer back and forth between the host memory and the buffer memory, potentially saving time during buffer creation and destruction.
To take another case, when the GPU and CPU have shared memory, it is possible to avoid copies of memory through sharing of pages. But for page sharing to be possible, the allocated memory needs to have some properties like being aligned on page boundary. In case of discrete devices, the benefit may not be realized because any memory operation by the accelerator will have to go across PCIe or some other slower interface than the memory of the accelerator.
The following code shows how to print the memory addresses on the host, inside the buffer, and on the accelerator device inside the kernel.
int VectorAdd0(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b,
AlignedVector<int> &sum, int iter) {
sycl::range num_items{a.size()};
const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
for (int i = 0; i < iter; i++) {
sycl::buffer a_buf(a, props);
sycl::buffer b_buf(b, props);
sycl::buffer sum_buf(sum.data(), num_items, props);
{
sycl::host_accessor a_host_acc(a_buf);
std::cout << "add0: buff memory address =" << a_host_acc.get_pointer()
<< "\n";
std::cout << "add0: address of vector a = " << a.data() << "\n";
}
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);
sycl::stream out(1024 * 1024, 1 * 128, h);
h.parallel_for(num_items, [=](auto i) {
if (i[0] == 0)
out << "add0: dev addr = " << a_acc.get_pointer() << "\n";
sum_acc[i] = a_acc[i] + b_acc[i];
});
});
}
q.wait();
return (0);
}
When this program is run, it can be seen that the addresses for all three (on the host, in the buffer, and on the accelerator) are the same when the property use_host_ptr is set for integrated GPU devices. But for discrete GPU devices, the buffer and device addresses will be different. Also note that in line 1, none of the incoming arguments are declared to be const. If these are declared const then during buffer creation they are copied and new memory is allocated instead of reusing the memory in the host vectors. The code snippet below demonstrates this. When this code is executed, we see that the addresses associated with the incoming vectors are different from the memory present in the buffer and also the memory present in the accelerator device.
int VectorAdd1(sycl::queue &q, const AlignedVector<int> &a,
const AlignedVector<int> &b, AlignedVector<int> &sum, int iter) {
sycl::range num_items{a.size()};
const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
for (int i = 0; i < iter; i++) {
sycl::buffer a_buf(a, props);
sycl::buffer b_buf(b, props);
sycl::buffer sum_buf(sum.data(), num_items, props);
{
sycl::host_accessor a_host_acc(a_buf);
std::cout << "add1: buff memory address =" << a_host_acc.get_pointer()
<< "\n";
std::cout << "add1: address of vector aa = " << a.data() << "\n";
}
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);
sycl::stream out(16 * 1024, 16 * 1024, h);
h.parallel_for(num_items, [=](auto i) {
if (i[0] == 0)
out << "add1: dev addr = " << a_acc.get_pointer() << "\n";
sum_acc[i] = a_acc[i] + b_acc[i];
});
});
}
q.wait();
return (0);
}
The kernel vectorAdd3 will not incur the cost of copying the memory contents from the buffer to the accelerator device because the use_host_ptr property is set while creating the buffers, and the buffers are aligned on a page boundary for an integrated GPU device. If memory pointed to by a buffer is not aligned on a page boundary, then new memory is allocated that aligns on a page boundary and the contents of the buffer are copied into that memory. This new memory from the buffer is then shared with the accelerator either by copying the contents from the buffer on host to the device (for accelerators that do not share any memory) or by using the page tables to avoid a physical copy of memory available on the device (for accelerators that share memory).
int VectorAdd2(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b,
AlignedVector<int> &sum, int iter) {
sycl::range num_items{a.size()};
const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
auto start = std::chrono::steady_clock::now();
for (int i = 0; i < iter; i++) {
sycl::buffer a_buf(a, props);
sycl::buffer b_buf(b, props);
sycl::buffer sum_buf(sum.data(), num_items, props);
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]; });
});
}
q.wait();
auto end = std::chrono::steady_clock::now();
std::cout << "Vector add2 completed on device - took "
<< (end - start).count() << " u-secs\n";
return ((end - start).count());
}
The kernel below will incur the cost of copying memory contents between the host and buffer, and also from the buffer to the accelerator.
int VectorAdd3(sycl::queue &q, const AlignedVector<int> &a,
const AlignedVector<int> &b, AlignedVector<int> &sum, int iter) {
sycl::range num_items{a.size()};
auto start = std::chrono::steady_clock::now();
for (int i = 0; i < iter; i++) {
sycl::buffer a_buf(a);
sycl::buffer b_buf(b);
sycl::buffer sum_buf(sum.data(), num_items);
auto e = 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]; });
});
}
q.wait();
auto end = std::chrono::steady_clock::now();
std::cout << "Vector add3 completed on device - took "
<< (end - start).count() << " u-secs\n";
return ((end - start).count());
}
Care must be taken to ensure that unnecessary copies are avoided during the creation of buffers and passing the memory from the buffers to the kernels. Even when the accelerator shares memory with the host, a few additional conditions must be satisfied to avoid these extra copies.