本章节翻译by chenshusmail@163.com 原文：Submitting Kernels to Multiple Queues (intel.com)
queue 提供了一个用于提交多个 kernel 到加速器进行执行的方式。 queue 还有一个描述设备状态的 context。这个状态包括 buffer 的内容和执行 kernel 所需的任何内存。 runtime 跟踪当前的设备 context，并避免主机和设备之间不必要的内存传输。因此， 最好一起从一个 context 提交和启动 kernel，而不是在不同的 context 中交错提交 kernel。
以下示例提交了 30 个独立的 kernel，它们使用相同的 buffer 作为输入， 将结果计算到不同的输出 buffer 中。所有这些 kernel 都是完全独立的，有可能并发执行和无序执行。 这些 kernel 被提交到三个 queue 中，每个 kernel 的执行将根据 queue 的创建方式产生不同的开销。
int VectorAdd(sycl::queue &q1, sycl::queue &q2, sycl::queue &q3,
const IntArray &a, const IntArray &b) {
sycl::buffer a_buf(a);
sycl::buffer b_buf(b);
sycl::buffer<int> *sum_buf[3 * iter];
for (size_t i = 0; i < (3 * iter); i++)
sum_buf[i] = new sycl::buffer<int>(256);
size_t num_groups = 1;
size_t wg_size = 256;
auto start = std::chrono::steady_clock::now();
for (int i = 0; i < iter; i++) {
q1.submit([&](auto &h) {
sycl::accessor a_acc(a_buf, h, sycl::read_only);
sycl::accessor b_acc(b_buf, h, sycl::read_only);
auto sum_acc = sum_buf[3 * i]->get_access<sycl::access::mode::write>(h);
h.parallel_for(sycl::nd_range<1>(num_groups * wg_size, wg_size),
[=](sycl::nd_item<1> index) {
size_t loc_id = index.get_local_id();
sum_acc[loc_id] = 0;
for (size_t i = loc_id; i < array_size; i += wg_size) {
sum_acc[loc_id] += a_acc[i] + b_acc[i];
}
});
});
q2.submit([&](auto &h) {
sycl::accessor a_acc(a_buf, h, sycl::read_only);
sycl::accessor b_acc(b_buf, h, sycl::read_only);
auto sum_acc =
sum_buf[3 * i + 1]->get_access<sycl::access::mode::write>(h);
h.parallel_for(sycl::nd_range<1>(num_groups * wg_size, wg_size),
[=](sycl::nd_item<1> index) {
size_t loc_id = index.get_local_id();
sum_acc[loc_id] = 0;
for (size_t i = loc_id; i < array_size; i += wg_size) {
sum_acc[loc_id] += a_acc[i] + b_acc[i];
}
});
});
q3.submit([&](auto &h) {
sycl::accessor a_acc(a_buf, h, sycl::read_only);
sycl::accessor b_acc(b_buf, h, sycl::read_only);
auto sum_acc =
sum_buf[3 * i + 2]->get_access<sycl::access::mode::write>(h);
h.parallel_for(sycl::nd_range<1>(num_groups * wg_size, wg_size),
[=](sycl::nd_item<1> index) {
size_t loc_id = index.get_local_id();
sum_acc[loc_id] = 0;
for (size_t i = loc_id; i < array_size; i += wg_size) {
sum_acc[loc_id] += a_acc[i] + b_acc[i];
}
});
});
}
q1.wait();
q2.wait();
q3.wait();
auto end = std::chrono::steady_clock::now();
std::cout << "Vector add completed on device - took " << (end - start).count()
<< " u-secs\n";
// check results
for (size_t i = 0; i < (3 * iter); i++)
delete sum_buf[i];
return ((end - start).count());
} // end VectorAdd将 kernel 提交到同一个 queue 可以获得最佳性能，因为所有的 kernel 都能在开始时一次性传输所需的输入， 并完成所有的计算。
VectorAdd(q, q, q, a, b);如果将 kernel 提交到共享相同 context 的不同 queue，其性能与提交到一个 queue 类似。 需要注意的问题是，当一个 kernel 被提交到具有不同 context 的新 queue 时， JIT 过程会将 kernel 编译到与 context 关联的新设备。如果不计算这个 JIT 编译时间， kernel 的实际执行是类似的。
sycl::queue q1(sycl::default_selector_v);
sycl::queue q2(q1.get_context(), sycl::default_selector_v);
sycl::queue q3(q1.get_context(), sycl::default_selector_v);
VectorAdd(q1, q2, q3, a, b);如果将 kernel 提交到具有三个不同 context 的三个不同 queue，性能会下降， 因为在调用 kernel 时，runtime 需要每次都将所有输入 buffer 传输到加速器。 此外，每个 context 都会对 kernel 进行 JIT。
sycl::queue q4(sycl::default_selector_v);
sycl::queue q5(sycl::default_selector_v);
sycl::queue q6(sycl::default_selector_v);
VectorAdd(q4, q5, q6, a, b);如果由于某种原因你需要使用不同的 queue，可以通过创建具有共享 context 的 queue 来缓解这个问题。 这将防止需要传输输入 buffer，但是 kernel 的内存占用将会增加， 因为所有的输出 buffer 必须同时驻留在 context 中，而在之前使用相同 queue 的方法中， 设备上的相同内存可以用于输出 buffer。另一件需要记住的事情是 kernel 中的内存与计算比例问题。 在上面的示例中，kernel 的计算需求较低，所以整体执行时间主要由内存传输主导。 当计算量大时，这些传输对整体执行时间的贡献不大。
以下的示例中说明了这一点，其中 kernel 的计算量增加了千倍，因此 runtime 将会有所不同。
int VectorAdd(sycl::queue &q1, sycl::queue &q2, sycl::queue &q3,
const IntArray &a, const IntArray &b) {
sycl::buffer a_buf(a);
sycl::buffer b_buf(b);
sycl::buffer<int> *sum_buf[3 * iter];
for (size_t i = 0; i < (3 * iter); i++)
sum_buf[i] = new sycl::buffer<int>(256);
size_t num_groups = 1;
size_t wg_size = 256;
auto start = std::chrono::steady_clock::now();
for (int i = 0; i < iter; i++) {
q1.submit([&](auto &h) {
sycl::accessor a_acc(a_buf, h, sycl::read_only);
sycl::accessor b_acc(b_buf, h, sycl::read_only);
auto sum_acc = sum_buf[3 * i]->get_access<sycl::access::mode::write>(h);
h.parallel_for(sycl::nd_range<1>(num_groups * wg_size, wg_size),
[=](sycl::nd_item<1> index) {
size_t loc_id = index.get_local_id();
sum_acc[loc_id] = 0;
for (int j = 0; j < 1000; j++)
for (size_t i = loc_id; i < array_size; i += wg_size) {
sum_acc[loc_id] += a_acc[i] + b_acc[i];
}
});
});
q2.submit([&](auto &h) {
sycl::accessor a_acc(a_buf, h, sycl::read_only);
sycl::accessor b_acc(b_buf, h, sycl::read_only);
auto sum_acc =
sum_buf[3 * i + 1]->get_access<sycl::access::mode::write>(h);
h.parallel_for(sycl::nd_range<1>(num_groups * wg_size, wg_size),
[=](sycl::nd_item<1> index) {
size_t loc_id = index.get_local_id();
sum_acc[loc_id] = 0;
for (int j = 0; j < 1000; j++)
for (size_t i = loc_id; i < array_size; i += wg_size) {
sum_acc[loc_id] += a_acc[i] + b_acc[i];
}
});
});
q3.submit([&](auto &h) {
sycl::accessor a_acc(a_buf, h, sycl::read_only);
sycl::accessor b_acc(b_buf, h, sycl::read_only);
auto sum_acc =
sum_buf[3 * i + 2]->get_access<sycl::access::mode::write>(h);
h.parallel_for(sycl::nd_range<1>(num_groups * wg_size, wg_size),
[=](sycl::nd_item<1> index) {
size_t loc_id = index.get_local_id();
sum_acc[loc_id] = 0;
for (int j = 0; j < 1000; j++)
for (size_t i = loc_id; i < array_size; i += wg_size) {
sum_acc[loc_id] += a_acc[i] + b_acc[i];
}
});
});
}
q1.wait();
q2.wait();
q3.wait();
auto end = std::chrono::steady_clock::now();
std::cout << "Vector add completed on device - took " << (end - start).count()
<< " u-secs\n";
// check results
for (size_t i = 0; i < (3 * iter); i++)
delete sum_buf[i];
return ((end - start).count());
} // end VectorAdd
41
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
