如何在cuda中声明一个静态变量答案

【问题标题】：How to declare a static variable in cuda如何在cuda中声明一个静态变量
【发布时间】：2016-05-27 21:30:57
【问题描述】：

在java中我们可以声明一个静态变量。即，如果一个变量被声明为静态的，它在整个程序中都是一样的。如果我在 CUDA 内核中声明一个变量，每个内核都会创建该变量（多个单独的副本）。

例如，如果“int x=5”被初始化并且我们启动了两个内核。如果我们在一个内核中将 x 的值更改为 6 (x=6)。此更改在其他内核中不可见（x 的值仍为 5）。

我想在 CUDA 中声明一个静态变量，每个内核都应该能够访问该变量值，如果从一个内核对该变量进行了更改，它应该在另一个内核中可见（如果 x=6 在一个内核，其他内核应该更新到 x=6)。

我需要这个来查找矩阵中是否存在数字。例如，给出了一个矩阵和一个要查找的数字（比如 5）。

2 3 0 0 0

1 4 5 0 0

7 8 0 0 0

0 0 0 0 0

我应该得到肯定，row = 1 和 col = 2（假设行和列从 0 开始）。

【问题讨论】：

也许您应该在需要这种行为的地方提供更多详细信息，因为一般建议是：不要这样做，或者您使用atomic functions。后者，如果你不明白在哪里使用它们，不会给你满意的性能......
原子函数允许你以你想要的方式递增一个变量。但是：您的线程总是在一个扭曲中执行（一组 32 个线程，它们将同时执行相同的指令，即同时读取相同的值）。问题在这里暴露无遗：stackoverflow.com/a/21346015/4866974
@havogt，谢谢你的帮助，为什么我需要这个，我想搜索一个数字是否存在于矩阵中，如果该数字存在，则布尔值应设置为 true 和 false 它不't
也许你应该重新表述你的问题以允许一个合理的答案......
帮你做什么？我可以通过将您指向 minimal reproducible example 和 How do I ask a good question? 来帮助您学习如何使用为 SO 编写一个好问题

标签： cuda

【解决方案1】：

对于您描述的用例，可从设备代码访问的静态变量是使用 __device__ 限定符创建的。参考documentation。

此外，对于您描述的用例（线程间/块/内核通信），我还将使用 volatile 限定符标记该变量。参考documentation。

类似这样的：

__device__ volatile int found = 0;

或

__device__ volatile bool found = false;

Here 是编程指南中使用此构造进行线程间通信的示例。

然后您可以从任何线程将该变量设置为 1 或 true，然后再查询它的状态。

由于您的用例描述仅涉及将变量设置为单个值，因此无论哪个线程执行此操作，都无需担心来自多个线程的同时访问，只要唯一的操作您所做的是写入该变量：

found = 1;

（对于这种特定情况，volatile 也可能不是必需的，具体取决于您的具体用法。）

如果您知道只有一个线程会找到该项目，并且您还希望记录 x,y 坐标，那将是一个微不足道的扩展：

__device__ volatile int found = 0;
__device__ volatile int x = -1;
__device__ volatile int y = -1;

那么您的设备代码可能是：

if (item_found){
  found = 1;
  x = item.x;
  y = item.y;}

【讨论】：