关于Intel oneAPI指南中两个SYCL示例内存传输逻辑的确认问询

我最近在研究Intel oneAPI指南里的两个SYCL代码示例，对其中的内存传输触发逻辑有些困惑，想和大家确认我的理解是否正确，同时也想搞清楚指南里的描述是不是存在表述偏差。

第一个示例：缓冲区位于main全局作用域 #

先贴出指南里的这段代码：

#include <stdio.h>
#include <sycl/sycl.hpp>
constexpr int N = 64;
int main() {
  int AData[N];
  int CData[N];
  std::cout << "AData:\n";
  for(int i=0; i<N; i++) std::cout << AData[i] << "\n";
  
  sycl::queue Q;
  sycl::buffer<int> ABuf(&AData[0], N);
  sycl::buffer<int> CBuf(&CData[0], N);

  // Kernel1 initializes the data in the device buffers
  Q.submit([&](auto &h) {
    sycl::accessor aA(ABuf, h, sycl::write_only, sycl::no_init);
    sycl::accessor aC(CBuf, h, sycl::write_only, sycl::no_init);
    h.parallel_for(N, [=](auto i) {
      aA[i] = 11;
      aC[i] = i;
    });
  });

  // Kernel2 does work
  Q.submit([&](auto &h) {
    sycl::accessor aA(ABuf, h, sycl::read_only);
    sycl::accessor aC(CBuf, h);
    h.parallel_for(N, [=](auto i) {
      aC[i] += aA[i];
    });
  });

  // host_accessor gets the data from the device
  sycl::host_accessor h_acc(CBuf);
  std::cout << "CBuf:\n";
  for (int i = 0; i < N; i++) {
    printf("%d\n", h_acc[i]);
  }

  std::cout << "AData:\n";
  for(int i=0; i<N; i++) std::cout << AData[i] << "\n";
  return 0;
}

根据我对SYCL buffer/accessor模型的理解，这个示例应该只触发1次主机-设备内存传输：也就是在构造host_accessor h_acc的时候，把设备上CBuf的计算结果同步回主机的CData数组。理由如下：

• Kernel1使用了write_only + no_init属性的accessor，这告诉SYCL runtime不需要从主机拷贝AData/CData的初始值到设备，直接在设备端分配内存并初始化即可，避免了不必要的主机到设备的传输。
• Kernel1和Kernel2都操作同一个设备上的buffer，SYCL runtime会自动跟踪数据依赖，在设备内部完成数据流转，完全不需要把数据同步回主机再重新拷贝到设备——毕竟两个kernel都提交到了同一个队列，runtime能感知到数据一直在设备侧。
• 实际在Intel Tiber云的Xeon+Max GPU上运行的结果也验证了我的理解：程序开头和结尾打印的AData内容完全一致（大部分是0或随机值），说明AData从未被设备同步回主机；而CData通过host_accessor拿到了正确的计算结果（11+i）。

指南里的矛盾描述与第二个示例 #

但指南里提到，第一个示例在Kernel1和Kernel2之间存在内存传输，而第二个示例（缓冲区放在临时作用域）避免了这些传输。这和我的理解完全相反，让我有点困惑。

第二个示例的唯一区别是把ABuf和CBuf放在了临时代码块中：

int AData[N];
int CData[N];
sycl::queue Q;
{
  sycl::buffer<int> ABuf(&AData[0], N);
  sycl::buffer<int> CBuf(&CData[0], N);
  // Kernel1 initializes the data
  Q.submit([&](auto &h) {...});
  // Kernel2 does work
  Q.submit([&](auto &h) {...});
} // Buffers are destroyed, the data is synced
// so CData is up to date
for (int i = 0; i < N; i++) {
  printf("%d\n", CData[i]);
}

根据SYCL规则，缓冲区销毁时会自动将设备上的数据同步回绑定的主机数组，所以实际运行时AData会被同步为全11（Kernel1的赋值结果），CData也能拿到正确的计算结果，这点我是理解的。

我的疑问 #

1. 指南里说第一个示例在Kernel1和Kernel2之间存在内存传输，这是不是描述错误？还是我对SYCL accessor的内存触发逻辑有误解？
2. 另外，当N<64时，程序开头打印的AData全是同一个值，我猜测这是Linux系统对栈内存的初始化机制导致的，这个理解应该没问题吧？

内容来源于stack exchange