全局设备变量 OpenCL

Posted 2023-02-16

【中文标题】全局设备变量 OpenCL

【英文标题】：Global device variable OpenCL

【发布时间】：2013-05-14 15:04:40

【问题描述】：

如何声明一个对 OpenCL 中的所有线程都是全局的设备变量？ 我正在将一些代码从 CUDA 移植到 OpenCL。在我的 CUDA 实现中，我有类似的东西

...
...

__device__ int d_var;

...
...
void foo() 
   int h_var = 0
   cudaMemcpyToSymbol(d_var, h_var, sizeof(int));
   do
      //launch kernel, inside kernel d_var is modified
      cudaMemcpyFromSymbol(h_var, d_var, sizeof(int));
   while(h_var != 0);

我一直在阅读 OpenCL 示例代码，但不知道如何执行此操作。 任何建议都会很棒！

【参考方案1】：

不幸的是，这并不是直接移植。虽然我不完全确定您是否至少可以在 OpenCL 程序中定义和使用（读/写）这样一个全局变量（但我不这么认为），但绝对没有办法访问（读/写）来自主机的这个变量。

你要做的就是把它作为一个额外的内核参数放入内核。如果它仅由主机写入并由内核读取，则普通的int 变量就足够了（因此每个内核都有一个副本）。但在您的情况下，因为它也是写入内核并从主机读取的，所以它必须是一个 __global 指针，并绑定了一个适当的缓冲区：

__kernel void func(..., __global int *d_var)  ... 

然后您可以使用通常的缓冲区功能在主机上读取和写入（当然，在内核中通过经典指针取消引用*d_var，但请记住，如果并发写入此变量会产生未指定的结果不是原子的）。

【讨论】：

抱歉，我的代码出错了。有一个主机 h_var 和一个设备 d_var，然后我从 device2host 来回复制。我想我将不得不按照您的建议创建另一个缓冲区。我尝试使用__kernel void foo(..., __constant int d_var)，但出现错误，是因为 d_var 不是指针吗？

@BRabbit27 “是因为 d_var 不是指针吗？” - 是的。我理解你的代码示例，答案仍然适用。您必须为您的 d_var 创建一个缓冲区（只是单个 int 大缓冲区，或较大缓冲区的子缓冲区）并从那里复制到 h_var，反之亦然，使用通常的缓冲区读/写操作。然后像往常一样将此缓冲区作为*d_var 参数的内核参数传递。顺便说一句，如果你想在内核中写信给d_var，就像你在问题中写的那样，你需要__global而不是__constant。

我必须把它作为__global 放在内核中？如果我不输入预选赛怎么办？它转换为局部变量？两种方法都试过了，但都没有奏效。好吧，我的循环只运行一次，但我确信它必须运行更多次。

我又错了。我忘了取消引用内核内的指针。 *d_change = 1 我有它d_change = 1。工作完美:)

【参考方案2】：

对于适合变量的小项目，将它们作为参数传递。

对于较大的常量项，使用 OpenCL“常量”地址空间。它针对广播使用进行了优化。

在您的 CUDA 示例中，您将数据读回。在 OpenCL 中，您需要为此使用常规缓冲区对象，因为不会将参数读回主机，并且常量内存是只读的。