2

假设我们有这个排序数组

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

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

    0 1 5 10 11

我知道有一些库（推力）可以实现这一点，但是我想为教育目的创建自己的有效实现。

你可以在这里找到整个代码：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. Ndevice_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因此该向量将具有以下结果：

  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. 内核内部创建了很多分支，因此会出现低效的换行处理
  2. 我假设每个线程仅使用 2 个值，这是非常低效的，因为创建了太多线程
  3. 我创建并传输一个device_vector与输入一样大并且驻留在全局内存中的 a。每个线程访问这个向量以便写入结果，这是非常低效的。

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

     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 实现可能会变得更快，但是我相信更有效的方法可能会使实现更快，即使对于较小的输入也是如此。

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

先感谢您

4

1 回答 1

4

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

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

建议：

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