c++ - 使用 CUDA 以有效计算元素更改的已排序数组的位置

Question

假设我们有这个排序数组

     0 1 1 1 1 2 2 2 2 2 3 10 10 10

我想高效地找到元素变化的位置。例如在我们的数组中，位置如下:

    0 1 5 10 11

我知道有一些库 (Thrust) 可以实现这一点，但是我想出于教育目的创建自己的高效实现。

您可以在这里找到完整的代码:http://pastebin.com/Wu34F4M2

它也包括验证。

内核是以下函数:

__global__ void findPositions(int *device_data, 
         int totalAmountOfValuesPerThread, int* pos_ptr, int N){

   int res1 = 9999999;
   int res2 = 9999999;
   int index = totalAmountOfValuesPerThread*(threadIdx.x + 
                  blockIdx.x*blockDim.x);
   int start = index; //from this index each thread will begin searching
   if(start < N){ //if the index is out of bounds do nothing
      if(start!=0){ //if start is not in the beginning, check the previous value
        if(device_data[start-1] != device_data[start]){
        res1 = start;
        }
      }
      else res1 = start; //since it's the 
          //beginning we update the first output buffer for the thread
      pos_ptr[index] = res1;

      start++; //move to the next place and see if the 
      //second output buffer needs updating or not

      if(start < N && device_data[start] != device_data[start-1]){
         res2 = start;
      }

      if((index+1) < N)
        pos_ptr[index+ 1] = res2;
      }
}

我创建了很多线程，因此每个线程都必须处理数组的两个值。

1. device_data 数组中存储了所有的数字
2. totalAmountOfValuesPerThread 在这种情况下是每个线程必须处理的值的总量
3. pos_ptr 与device_data 的长度相同，每个线程将缓冲区的结果写入此device_vector
4. N是device_data数组中的数字总数

在名为 res1 和 res2 的输出缓冲区中，每个线程要么保存一个之前未找到的位置，要么保持原样。

例子:

  0   <---- thread 1
  1
  1   <---- thread 2
  1
  2   <---- thread 3
  2
  3   <---- thread 4

每个线程的输出缓冲区，假设大数字 9999999 是 inf 将是:

  thread1 => {res1=0, res2=1}
  thread2 => {res1=inf, res2=inf}
  thread3 => {res1=4, res2=inf}
  thread4 => {res1=6, res2=inf}

每个线程都会更新 pos_ptr device_vector 因此该 vector 将具有以下结果:

  pos_ptr =>{0, 1, inf, inf, 4, inf, 6, inf}

完成内核后，我使用库 Thrust 对 vector 进行排序，并将结果保存在名为 host_pos 的主机 vector 中。所以 host_pos vector 将具有以下内容:

  host_pos => {0, 1, 4, 6, inf, inf, inf, inf}

这个实现很糟糕，因为

1. 在内核中创建了很多分支，因此会出现低效的换行处理
2. 我假设每个线程只使用 2 个值，这是非常低效的，因为创建了太多线程
3. 我创建并传输了一个 device_vector，它与输入一样大并且也驻留在全局内存中。每个线程访问这个 vector 以写入结果，这是非常低效的。

这是当每个 block 中有 512 线程时输入大小 1 000 000 的测试用例。

     CPU time: 0.000875688 seconds
     GPU time: 1.35816 seconds

输入大小 10 000 000 的另一个测试用例

     CPU time: 0.0979209
     GPU time: 1.41298 seconds

请注意，CPU 版本几乎慢了 100 倍，而 GPU 的速度几乎相同。

不幸的是我的 GPU 没有足够的内存，所以让我们试试 50 000 000

     CPU time: 0.459832 seconds
     GPU time: 1.59248 seconds

据我了解，对于大量输入，我的 GPU 实现可能会变得更快，但我相信更高效的方法可能会使实现更快，即使对于较小的输入也是如此。

为了让我的算法运行得更快，你会建议什么设计？不幸的是，我想不出更好的办法。

提前致谢

Answer 1

我真的不明白您认为这很糟糕的任何原因。线程太多？线程过多的定义是什么？每个输入元素一个线程是 CUDA 程序中非常常见的线程策略。

因为您似乎愿意考虑在大部分工作中使用 thrust(例如，您愿意在完成数据标记后调用 thrust::sort)并考虑到 Benc 的观察(您正在花费很多时间试图优化总运行时间的 3%)也许你可以通过更好地利用推力来产生更大的影响。

建议:

1. 让你的内核尽可能简单。只是让每个线程看起来在一个元素上，并根据与前一个元素。我不确定是否取得了重大进展通过让每个线程处理 2 个元素。或者，有一个内核创建的 block 数量多，但让它们循环遍历整个 device_data 数组，并在它们移动时标记边界。这可能会显着改善您的内核。但同样，优化 3% 不一定是您要花费大量精力的地方。
2. 您的内核将受内存带宽限制。因此，我不会担心分支之类的事情，而是专注于内存的有效使用，即最小化对全局内存的读写，并寻找机会确保您的读写合并。独立于程序的其余部分测试您的内核，并使用可视化分析器告诉您您是否在内存操作方面做得很好。
3. 考虑使用共享内存。通过让每个线程将其各自的元素加载到共享内存中，您可以轻松合并所有全局读取(并确保您只读取每个全局元素一次，或几乎每个元素一次)然后在共享内存之外运行，即让每个线程将它的元素与共享内存中的前一个元素进行比较。
4. 一旦你创建了你的 pos_ptr 数组，让我们注意除了inf 值它已经排序。所以也许有更聪明的选项而不是“推力::排序”，然后修剪数组，以产生结果。看一下推力函数，例如remove_if 和 copy_if。我没有对它进行基准测试，但我的猜测是它们会比 sort 便宜得多，其次是修剪数组(删除 inf 值)。