CUDA：while循环索引正确性答案

【问题标题】：CUDA: while loop index correctnessCUDA：while循环索引正确性
【发布时间】：2012-04-18 20:11:54
【问题描述】：

这个内核正在做正确的事情，给了我正确的结果。如果我想提高性能，我的问题更多在于 while 循环的正确性。我尝试了几种块和线程的配置，但如果我要更改它们，while 循环不会给我正确的结果。我在更改内核配置时获得的结果是 firstArray 和 secondArray 不会被完全填充（它们在单元格内将有 0）。两个数组都必须填充从 if 循环中获得的 curValue。

欢迎任何建议:)

提前谢谢你

#define N 65536

__global__ void whileLoop(int* firstArray_device, int* secondArray_device)
{   
    int curValue = 0;
    int curIndex = 1;

    int i = (threadIdx.x)+2;

    while(i < N) {
        if (i % curIndex == 0) {
            curValue = curValue + curIndex;
            curIndex *= 2;
        }
        firstArray_device[i] = curValue;
        secondArray_device[i] = curValue;
        i += blockDim.x * gridDim.x;
    }
}

int main(){

  firstArray_host[0] = 0;
  firstArray_host[1] = 1;

  secondArray_host[0] = 0;
  secondArray_host[1] = 1;


  // memory allocation + copy on GPU

  // definition number of blocks and threads
  dim3 dimBlock(1, 1);
  dim3 dimGrid(1, 1);

  whileLoop<<<dimGrid, dimBlock>>>(firstArray_device, secondArray_device);

  // copy back to CPU + free memory
}

【问题讨论】：

这不是一个真正的优化问题——它是一个正确性问题，不是吗？
我猜你是对的。我将更改帖子的标题。
我不太明白这个问题...您到底想要什么建议？你说你的内核工作正常但是你的“正确性”有问题？
我的意思是，如果我更改内核的设置以提高性能，我将不会在数组中拥有正确的索引。所以我想知道是否有另一种方法可以在具有不同内核设置的数组中获得相同的结果。
@RogerDahl：基本上他的代码只能在串行运行（即一个线程）时才能工作。尝试运行更多，它会以多种方式中断，这都是因为代码错误和算法不正确。

标签： optimization cuda while-loop

【解决方案1】：

您在这里遇到了数据依赖性问题，这会阻碍您进行一些有意义的优化。变量 curValue 和 curIndex 在 while 循环中被更改并前馈到下一次运行。一旦您尝试优化循环，您就会发现您处于此变量具有不同状态并且结果发生更改的情况。

我真的不知道你试图实现什么，但尝试使 while 循环独立于循环之前运行的值以避免依赖关系。尝试将数据分成线程和数据块，以在 threadIdx、blockDim、gridDim 等环境状态上计算 indizes 和 value...

还要尽量避免条件循环。最好使用具有恒定运行次数的 for 循环。这也更容易优化。

【讨论】：

我只是一个好奇的读者，在没有任何 CUDA 知识的情况下偶然发现了这一点。为什么说“尽量避免条件循环”？这在CUDA中客观上正确吗？还是用其他任何语言？这是某种最佳实践吗？为什么？
It 技术相关和最佳实践： 1) 如果您在 for 循环中确实有恒定的运行次数，则编译有机会进行安全优化，如循环注册、缓存优化、...对于 CUDA 和 OpenCL 编译器也是如此。 2）如果有条件循环，条件本身就很容易成为数据依赖。这对于 c 风格的 for 循环也是如此，但是当您将其用作计数循环时，情况并非如此。 3) CUDA 和 OpenCL 最好的方法是有一组恒定的输入数据和一个固定的函数来计算输出。
好点 Rick-Rainer！此外，条件循环是创建线程分歧的好方法，它会通过造成调度困难而扼杀并行性能。即使代码是正确的，你也希望你的所有线程都做同样的工作。所有线程评估的短条件并不是什么大问题，但是一个循环在不同的时间完成不同的线程可能会导致很多分歧和性能下降。一般来说，避免发散的条件逻辑是一种很好的 CUDA 范例。

【解决方案2】：

一些事情：

您遗漏了用于在设备。拥有这些信息会很有帮助。
您的算法是使用多个块时不是线程安全的。换句话说，如果您正在运行多个块，他们不仅会做多余的工作（因此给你没有收获），但他们也可能在某个时候尝试写到相同的全局内存位置，从而产生错误。
您的代码是这样的当只使用一个块时是正确的，但这使得它变得毫无意义......您正在并行设备上运行串行或轻线程操作。您不能在所有可用资源上运行（多个 SMP 上的多个块而不会发生内存冲突（见下文）...

目前，从并行的角度来看，此代码存在两个主要问题：

int i = (threadIdx.x)+2; ...产生2 的起始索引单线程； 2 和 3 用于 single 块中的两个线程，依此类推。我怀疑这是你想要的前两个职位（0，1）永远不会得到解决。（请记住，数组从 C 中的索引 0 开始。）
此外，如果您包含多个块（例如 2 个块每个都有一个线程），那么您将有多个重复的索引（例如，对于 2 b x 1 t --> 索引 b1t1：2，b1t2：2），当您使用索引时写入全局内存会产生冲突和错误。执行int i = threadIdx.x + blockDim.x * blockIdx.x; 之类的操作将是正确计算索引以避免此问题的典型方法。
你的最终表达式i += blockDim.x * gridDim.x; 可以，因为它将与线程总数相等的数字添加到 i ，因此不会产生额外的冲突或重叠。
为什么要使用 GPU 洗牌内存并进行微不足道的计算？当您考虑将阵列装入和取出设备的时间时，与快速 CPU 相比，您可能看不到多少加速。

如果您愿意，可以解决问题 1 和 2，但除此之外，请考虑您的总体目标以及您正在尝试优化的具体算法类型，并提出一个对并行更友好的解决方案 - 或者考虑 GPU 计算是否真的可以感觉你的问题。

【讨论】：

很好的分析！我没有时间做这个。在我看来，你是绝对正确的。
你总是可以，然后让我回答。 ;)
@JasonR.Mick 感谢您的分析！非常有帮助！但我有话要说：关于第 1 点，我需要从 2 开始获取 threadIdx.x，因为数组的前两个单元格是 0 和 1（我的错误是我没有在帖子中指定 - 我会更改）。最后你说了一些我可能会遵循的东西。在过去的三天里，我真的在想是否值得将 GPU 用于此类计算。无论如何，再次感谢您

【解决方案3】：

要并行化这个算法，你需要想出一个公式，可以直接计算数组中给定索引的值。因此，在数组范围内选择一个随机索引，然后考虑确定该位置的值的因素是什么。找到公式后，通过将随机索引的输出值与串行算法的计算值进行比较来测试它。如果正确，则创建一个内核，该内核首先根据其线程和块索引选择唯一索引。然后计算该索引的值并将其存储在数组中的相应索引中。

一个简单的例子：

序列号：

__global__ void serial(int* array)
{
  int j(0);
  for (int i(0); i < 1024; ++i) {
    array[i] = j;
    j += 5;
}

int main() {
  dim3 dimBlock(1);
  dim3 dimGrid(1);
  serial<<<dimGrid, dimBlock>>>(array);
}

平行：

__global__ void parallel(int* array)
{
  int i(threadIdx.x + blockDim.x * blockIdx.x);
  int j(i * 5);
  array[i] = j;
}

int main(){
  dim3 dimBlock(256);
  dim3 dimGrid(1024 / 256);
  parallel<<<dimGrid, dimBlock>>>(array);
}

【讨论】：