本章节翻译by weavingtime@formail.com 原文 Avoiding Moving Data Back and Forth between Host and Device (intel.com)
在主机和设备之间移动数据的成本相当高,特别是在独立加速器(比如独显)的情况下。 因此, 尽可能避免主机和设备之间的数据传输非常重要。 在某些情况下,可能需要将加速器上的kernel计算的数据带到主机并对其进行一些操作, 然后将其发送回设备以进行进一步处理。 在这种情况下,我们最终将产生设备到主机传输的开销,然后再次产生主机到设备传输的开销。
考虑以下示例,其中一个kernel通过某种操作(在本例中为vector相加)将数据生成到新vector中。 然后,通过对每个值应用函数,将这个新vector转换为第三个vector, 最后将这个第三个vector作为输入传到另一个kernel中以进行一些额外的计算。 在许多领域中这种形式的计算非常常见,在这些领域中,算法是迭代的,并且一个计算的输出需要作为另一个计算的输入。 例如,在机器学习中,模型被构造为计算层,一层的输出被输入到下一层。
1double myFunc1(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b, 2 AlignedVector<int> &c, AlignedVector<int> &d, 3 AlignedVector<int> &res, int iter) { 4 sycl::range num_items{a.size()}; 5 VectorAllocator<int> alloc; 6 AlignedVector<int> sum(a.size(), alloc); 7 8 const sycl::property_list props = {sycl::property::buffer::use_host_ptr()}; 9 sycl::buffer a_buf(a, props); 10 sycl::buffer b_buf(b, props); 11 sycl::buffer c_buf(b, props); 12 sycl::buffer d_buf(b, props); 13 sycl::buffer res_buf(res, props); 14 sycl::buffer sum_buf(sum.data(), num_items, props); 15 16 Timer timer; 17 for (int i = 0; i < iter; i++) { 18 // kernel1 19 q.submit([&](auto &h) { 20 // Input accessors 21 sycl::accessor a_acc(a_buf, h, sycl::read_only); 22 sycl::accessor b_acc(b_buf, h, sycl::read_only); 23 // Output accessor 24 sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init); 25 26 h.parallel_for(num_items, 27 [=](auto id) { sum_acc[id] = a_acc[id] + b_acc[id]; }); 28 }); 29 30 { 31 sycl::host_accessor h_acc(sum_buf); 32 for (int j = 0; j < a.size(); j++) 33 if (h_acc[j] > 10) 34 h_acc[j] = 1; 35 else 36 h_acc[j] = 0; 37 } 38 39 // kernel2 40 q.submit([&](auto &h) { 41 // Input accessors 42 sycl::accessor sum_acc(sum_buf, h, sycl::read_only); 43 sycl::accessor c_acc(c_buf, h, sycl::read_only); 44 sycl::accessor d_acc(d_buf, h, sycl::read_only); 45 // Output accessor 46 sycl::accessor res_acc(res_buf, h, sycl::write_only, sycl::no_init); 47 48 h.parallel_for(num_items, [=](auto id) { 49 res_acc[id] = sum_acc[id] * c_acc[id] + d_acc[id]; 50 }); 51 }); 52 q.wait(); 53 } 54 double elapsed = timer.Elapsed() / iter; 55 return (elapsed); 56} // end myFunc1
你可以创建一个 kernel3
来在设备上执行此函数,而不是将数据传送到主机并将函数应用于数据 然后将其发送回设备,如以下示例所示。 内核 kernel3
对 kernel1
和 kernel2
之间 的 accum_buf
中的中间数据进行操作,避免了设备和主机之间数据传输的往返。
1double myFunc2(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b, 2 AlignedVector<int> &c, AlignedVector<int> &d, 3 AlignedVector<int> &res, int iter) { 4 sycl::range num_items{a.size()}; 5 VectorAllocator<int> alloc; 6 AlignedVector<int> sum(a.size(), alloc); 7 8 const sycl::property_list props = {sycl::property::buffer::use_host_ptr()}; 9 sycl::buffer a_buf(a, props); 10 sycl::buffer b_buf(b, props); 11 sycl::buffer c_buf(b, props); 12 sycl::buffer d_buf(b, props); 13 sycl::buffer res_buf(res, props); 14 sycl::buffer sum_buf(sum.data(), num_items, props); 15 16 Timer timer; 17 for (int i = 0; i < iter; i++) { 18 // kernel1 19 q.submit([&](auto &h) { 20 // Input accessors 21 sycl::accessor a_acc(a_buf, h, sycl::read_only); 22 sycl::accessor b_acc(b_buf, h, sycl::read_only); 23 // Output accessor 24 sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init); 25 26 h.parallel_for(num_items, 27 [=](auto i) { sum_acc[i] = a_acc[i] + b_acc[i]; }); 28 }); 29 30 // kernel3 31 q.submit([&](auto &h) { 32 sycl::accessor sum_acc(sum_buf, h, sycl::read_write); 33 h.parallel_for(num_items, [=](auto id) { 34 if (sum_acc[id] > 10) 35 sum_acc[id] = 1; 36 else 37 sum_acc[id] = 0; 38 }); 39 }); 40 41 // kernel2 42 q.submit([&](auto &h) { 43 // Input accessors 44 sycl::accessor sum_acc(sum_buf, h, sycl::read_only); 45 sycl::accessor c_acc(c_buf, h, sycl::read_only); 46 sycl::accessor d_acc(d_buf, h, sycl::read_only); 47 // Output accessor 48 sycl::accessor res_acc(res_buf, h, sycl::write_only, sycl::no_init); 49 50 h.parallel_for(num_items, [=](auto i) { 51 res_acc[i] = sum_acc[i] * c_acc[i] + d_acc[i]; 52 }); 53 }); 54 q.wait(); 55 } 56 double elapsed = timer.Elapsed() / iter; 57 return (elapsed); 58} // end myFunc2
还有其他方法可以优化此示例。 例如, kernel3
中的裁剪操作可以合并到 kernel1
的计算中,如下所示。 这就是kernel融合,并且还有一个好处就是不用启动第三个内核。 SYCL编译器无法进行这种优化。 在机器学习等某些特定领域,图编译器可以对机器学习模型进行操作并融合操作,这和kernel融合有一样的作用。
1double myFunc3(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b, 2 AlignedVector<int> &c, AlignedVector<int> &d, 3 AlignedVector<int> &res, int iter) { 4 sycl::range num_items{a.size()}; 5 VectorAllocator<int> alloc; 6 AlignedVector<int> sum(a.size(), alloc); 7 8 const sycl::property_list props = {sycl::property::buffer::use_host_ptr()}; 9 sycl::buffer a_buf(a, props); 10 sycl::buffer b_buf(b, props); 11 sycl::buffer c_buf(b, props); 12 sycl::buffer d_buf(b, props); 13 sycl::buffer res_buf(res, props); 14 sycl::buffer sum_buf(sum.data(), num_items, props); 15 16 Timer timer; 17 for (int i = 0; i < iter; i++) { 18 // kernel1 19 q.submit([&](auto &h) { 20 // Input accessors 21 sycl::accessor a_acc(a_buf, h, sycl::read_only); 22 sycl::accessor b_acc(b_buf, h, sycl::read_only); 23 // Output accessor 24 sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init); 25 26 h.parallel_for(num_items, [=](auto i) { 27 int t = a_acc[i] + b_acc[i]; 28 if (t > 10) 29 sum_acc[i] = 1; 30 else 31 sum_acc[i] = 0; 32 }); 33 }); 34 35 // kernel2 36 q.submit([&](auto &h) { 37 // Input accessors 38 sycl::accessor sum_acc(sum_buf, h, sycl::read_only); 39 sycl::accessor c_acc(c_buf, h, sycl::read_only); 40 sycl::accessor d_acc(d_buf, h, sycl::read_only); 41 // Output accessor 42 sycl::accessor res_acc(res_buf, h, sycl::write_only, sycl::no_init); 43 44 h.parallel_for(num_items, [=](auto i) { 45 res_acc[i] = sum_acc[i] * c_acc[i] + d_acc[i]; 46 }); 47 }); 48 q.wait(); 49 } 50 double elapsed = timer.Elapsed() / iter; 51 return (elapsed); 52} // 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.
我们可以将此kernel融合进一步提升一层,融合 kernel1
和 kernel2
,如下面的代码所示。 这个改动提供了非常好的性能, 因为它完全避免了中间的 accum_buf
,除了启动额外的kernel之外还节省了内存。 在这种情况下,大部分性能优势是来自内存引用局部性的改进。
1double myFunc4(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b, 2 AlignedVector<int> &c, AlignedVector<int> &d, 3 AlignedVector<int> &res, int iter) { 4 sycl::range num_items{a.size()}; 5 VectorAllocator<int> alloc; 6 7 const sycl::property_list props = {sycl::property::buffer::use_host_ptr()}; 8 sycl::buffer a_buf(a, props); 9 sycl::buffer b_buf(b, props); 10 sycl::buffer c_buf(b, props); 11 sycl::buffer d_buf(b, props); 12 sycl::buffer res_buf(res, props); 13 14 Timer timer; 15 for (int i = 0; i < iter; i++) { 16 // kernel1 17 q.submit([&](auto &h) { 18 // Input accessors 19 sycl::accessor a_acc(a_buf, h, sycl::read_only); 20 sycl::accessor b_acc(b_buf, h, sycl::read_only); 21 sycl::accessor c_acc(c_buf, h, sycl::read_only); 22 sycl::accessor d_acc(d_buf, h, sycl::read_only); 23 // Output accessor 24 sycl::accessor res_acc(res_buf, h, sycl::write_only, sycl::no_init); 25 26 h.parallel_for(num_items, [=](auto i) { 27 int t = a_acc[i] + b_acc[i]; 28 if (t > 10) 29 res_acc[i] = c_acc[i] + d_acc[i]; 30 else 31 res_acc[i] = d_acc[i]; 32 }); 33 }); 34 q.wait(); 35 } 36 double elapsed = timer.Elapsed() / iter; 37 return (elapsed); 38} // end myFunc4