4

这个内核正在做正确的事情,给了我正确的结果。如果我想提高性能,我的问题更多在于 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
}
4

3 回答 3

4

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

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

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

于 2012-04-12T12:56:02.840 回答
2

一些东西:

  1. 您遗漏了用于在设备上声明全局数组的代码。拥有这些信息会很有帮助。
  2. 当使用多个块时,您的算法不是线程安全的。换句话说,如果您正在运行多个块,它们不仅会做冗余工作(因此不会给您带来任何收益),而且它们还可能在某些时候尝试写入相同的全局内存位置,从而产生错误。
  3. 因此,当仅使用一个块时,您的代码是正确的,但这使得它变得毫无意义......您正在并行设备上运行串行或轻线程操作。您不能在所有可用资源上运行(多个 SMP 上的多个块而没有内存冲突(见下文)...

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

  1. int i = (threadIdx.x)+2;...产生2单个线程的起始索引;2以及单个3块中的两个线程,依此类推。我怀疑这是您想要的,因为前两个位置(, )永远不会得到解决。(请记住,数组从C 中的索引开始。) 010

  2. 此外,如果您包含多个块(例如 2 个块,每个块有一个线程),那么您将有多个重复的索引(例如,对于 2 bx 1 t --> 索引 b1t1: 2, b1t2: 2),当您使用索引写入时全局内存会产生冲突和错误。做类似的事情int i = threadIdx.x + blockDim.x * blockIdx.x;将是正确计算索引以避免此问题的典型方法。

  3. 您的最终表达式i += blockDim.x * gridDim.x;没问题,因为它向 i 添加了一个等于线程总数的数字,因此不会产生额外的冲突或重叠。

  4. 为什么要使用 GPU 洗牌内存并进行微不足道的计算?当您考虑将阵列进出设备的时间时,与快速 CPU 相比,您可能看不到多少加速。

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

于 2012-04-12T17:54:43.833 回答
1

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

一个简单的例子:

序列号:

__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);
}
于 2012-04-12T18:13:29.130 回答