sycl/dpc++ 使用全局指针访问全局变量

Question

我正在尝试从 sycl 内核内部访问一个全局变量。使用此模式的代码及其输出如下。

#include<CL/sycl.hpp>
using namespace sycl;
int g_var = 22;
int * const g_ptr = &g_var;
int main() {
    queue q{host_selector{}};
    g_var = 27;
    std::cout<<"global var changed : " <<g_var<<" "<<*g_ptr<<"\n";
    q.submit( [=] (handler &h) {
        stream os(1024, 128, h);
        h.single_task([=] () {
            os<<"global var : "<<*g_ptr<<"\n";
        });
    }).wait();
    return 0;
}

它的输出如下。

$$ dpcpp test.cpp; ./a.out
global var changed : 27 27
global var : 22

即使 g_var 被更改为 27，它在内核中还是以 22 的初始值打印出来的。这是预期的行为吗？

通常 lambdas 不会创建全局变量的副本。 dpc++ 编译器是在设备内部创建全局变量的副本还是在编译期间传播常量值以便在运行时不访问全局内存？

Answer 1

SYCL 通过“主机”和“设备”来划分执行和数据。 “主机”上的数据通常位于 CPU 上，需要传输到“设备”（通常是 GPU）才能被访问以在 SYCL 内核中使用。

SYCL 使用缓冲区和访问器或“统一共享内存”为“设备”端代码（即在 GPU 上运行的代码）传输和提供对数据的访问。所以你的指针永远不会被发送到设备上，因此永远不会被修改。

因此，您应该使用缓冲区或 USM，而不是使用 int * const g_ptr = &g_var;。 SYCL Academy 有一些关于缓冲区和 USM 的课程。

例如使用缓冲区和访问器：

int a = 18, b = 24, r = 0;

  auto defaultQueue = sycl::queue{};

  {
    auto bufA = sycl::buffer{&a, sycl::range{1}};
    auto bufB = sycl::buffer{&b, sycl::range{1}};
    auto bufR = sycl::buffer{&r, sycl::range{1}};

    defaultQueue
        .submit([&](sycl::handler &cgh) {
          auto accA = sycl::accessor{bufA, cgh, sycl::read_only};
          auto accB = sycl::accessor{bufB, cgh, sycl::read_only};
          auto accR = sycl::accessor{bufR, cgh, sycl::write_only};

          cgh.single_task<scalar_add>([=] { accR[0] = accA[0] + accB[0]; });
        })
        .wait();
  }

使用 USM：

 auto usmQueue = sycl::queue{usm_selector{}, asyncHandler};

 usmQueue.memcpy(devicePtrA, a, sizeof(float) * dataSize).wait();
 usmQueue.memcpy(devicePtrB, b, sizeof(float) * dataSize).wait();

 usmQueue
     .parallel_for<vector_add>(sycl::range{dataSize},
                                  [=](sycl::id<1> idx) {
                                    auto globalId = idx[0];
                                    devicePtrR[globalId] =
                                        devicePtrA[globalId] +
                                        devicePtrB[globalId];
                                  })
     .wait();

 usmQueue.memcpy(r, devicePtrR, sizeof(float) * dataSize).wait();

 sycl::free(devicePtrA, usmQueue);
 sycl::free(devicePtrB, usmQueue);
 sycl::free(devicePtrR, usmQueue);

 usmQueue.throw_asynchronous();