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上的代码)的数据的访问。因此，您的指针永远不会发送到设备，因此也不会被修改。

因此，您应该使用缓冲区或int * const g_ptr = &g_var;，而不是使用USM。在SYCL Academy中有一些关于buffers和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();