使用 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_ptrdevice_vector，因此这个向量将具有以下结果：

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

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

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

这个实现很糟糕，因为

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

这是每个块中有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 程序中非常常见的线程策略。

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

建议：

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