【问题标题】:Calculating differences between consecutive indices fast快速计算连续索引之间的差异
【发布时间】:2012-02-07 10:28:11
【问题描述】:

鉴于我有数组

Let Sum be 16
dintptr = { 0 , 2, 8,11,13,15} 

我想使用 GPU 计算连续索引之间的差异。所以最终的数组应该如下:

count = { 2, 6,3,2,2,1}

下面是我的内核:

//for this function n is 6

__global__ void kernel(int *dintptr, int * count, int n){

   int id = blockDim.x * blockIdx.x + threadIdx.x;
   __shared__ int indexes[256];
   int need = (n % 256 ==0)?0:1;
   int allow = 256 * ( n/256 + need);
   while(id < allow){
     if(id < n ){
       indexes[threadIdx.x] = dintptr[id];

     }
     __syncthreads();
     if(id < n - 1 ){
       if(threadIdx.x % 255 == 0 ){
            count[id] = indexes[threadIdx.x + 1] - indexes[threadIdx.x];
       }else{
            count[id] = dintptr[id+1] - dintptr[id];
       }


    }//end if id<n-1
      __syncthreads();
     id+=(gridDim.x * blockDim.x);
    }//end while
}//end kernel
// For last element explicitly set count[n-1] = SUm - dintptr[n-1]

2 个问题:

  1. 这个内核速度快吗?您能否提出一个更快的实施方案?
  2. 此内核是否处理任意大小的数组(我认为可以)

【问题讨论】:

  • 您是否测试/计时/分析了您的代码?你学到了什么?
  • @programmer:我认为您错过了第三个问题“代码是否正确”,答案是否定的,因为编写时可能会出现死锁。
  • 它还假设块大小为 256(我认为)。
  • @talonmies:很好的建议!我通过将第二个 __syncthreads 移到 if(id
  • 程序员@talonmies 和我在其他问题中为您提供了相关代码。真的没有必要提出不同的问题。

标签: cuda gpu nvidia


【解决方案1】:

我会咬人的。

__global__ void kernel(int *dintptr, int * count, int n)
{
    for (int id = blockDim.x * blockIdx.x + threadIdx.x; 
         id < n-1; 
         id += gridDim.x * blockDim.x)
        count[id] = dintptr[id+1] - dintptr[i];
}

(既然你说你“明确”设置了最后一个元素的值,而你没有在你的内核中,我也懒得在这里设置它。)

我没有看到像你那样在这个内核中使用共享内存有很多优势:Fermi 上的 L1 缓存应该给你几乎相同的优势,因为你的局部性高而重用率低。

您的内核和我的内核似乎都可以处理任意大小的数组。但是,您的似乎假设 blockDim.x == 256。

【讨论】:

  • 马克,你比我更勇敢!我认为提问者计划在较旧的硬件上运行它,因此共享内存版本可能比 Fermi 之前的版本具有优势,但我同意,在 Fermi L1 缓存上对于差分和类似的效率惊人。
  • @harrism:我的第一个解决方案和你的一样。我只是好奇如何对共享内存进行编码:)
  • @harrism:为什么有人会在这么有趣的问题上给我-1?
  • 不确定。也许是因为你在问只有你能/应该回答的问题。快吗?那么,相比什么?真正的问题是“速度够快吗?”如果它不是您应用中的瓶颈,那么它就足够快了。
  • 程序员:你运行代码了吗? Talonmies 是发现您的死锁问题的人。@Pavan:由于您是推力用户,如果您有功能请求,请在推力谷歌代码网站上提出问题。
猜你喜欢
  • 1970-01-01
  • 1970-01-01
  • 2023-03-19
  • 2020-05-23
  • 1970-01-01
  • 1970-01-01
  • 1970-01-01
  • 1970-01-01
  • 2017-02-14
相关资源
最近更新 更多