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

Posted

技术标签:

【中文标题】sycl/dpc++ 使用全局指针访问全局变量【英文标题】:sycl/dpc++ accessing global variable using global pointer 【发布时间】:2021-06-03 18:29:42 【问题描述】:

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

#include<CL/sycl.hpp>
using namespace sycl;
int g_var = 22;
int * const g_ptr = &g_var;
int main() 
    queue qhost_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++ 编译器是在设备内部创建全局变量的副本还是在编译期间传播常量值以便在运行时不访问全局内存?

【问题讨论】:

【参考方案1】:

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

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

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

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

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

  auto defaultQueue = sycl::queue;

  
    auto bufA = sycl::buffer&a, sycl::range1;
    auto bufB = sycl::buffer&b, sycl::range1;
    auto bufR = sycl::buffer&r, sycl::range1;

    defaultQueue
        .submit([&](sycl::handler &cgh) 
          auto accA = sycl::accessorbufA, cgh, sycl::read_only;
          auto accB = sycl::accessorbufB, cgh, sycl::read_only;
          auto accR = sycl::accessorbufR, cgh, sycl::write_only;

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

使用 USM:

 auto usmQueue = sycl::queueusm_selector, asyncHandler;

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

 usmQueue
     .parallel_for<vector_add>(sycl::rangedataSize,
                                  [=](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();

【讨论】:

感谢您的详细解释和代码示例。我将来会使用 Buffers 或 USM。只是试图解释我得到的输出。

以上是关于sycl/dpc++ 使用全局指针访问全局变量的主要内容,如果未能解决你的问题,请参考以下文章

如何访问 C 中的阴影全局变量?

SYCL/DPC++ cpu 版本给出了正确的结果,但 gpu 给出了不正确的数据

全局变量比C中的局部变量快吗? [关闭]

不同风格的递归,引用/全局/指针变量的使用

C语言中用const声明全局变量赋初值和不赋初值有何区别?变量存放位置有啥不一样?

全局外部指针变量在 DLL 中不可见