cuda - CUDA 内核中的不同线程组

Question

我正在尝试使用 CUDA 内核加速 C 中的互相关函数。现在这就是我所拥有的:

__global__ void xcorr(cuDoubleComplex *temp1, cuDoubleComplex *temp2, cuDoubleComplex *temp3, int Nb, int binM, int Nspb)
{
   for (int k1 = 0; k1 < Nb; k1++)
   {
       int idx = blockIdx.x * blockDim.x + threadIdx.x;
       for (int j1 = 0; j1 < Nspb; j1++)
       {
           if ((j1 + idx) <(Nspb + binM))
           {
               temp3[idx + k1*(binM + 1)].x += (temp1[idx + j1 + (k1*(binM + Nspb))].x*temp2[j1 + (k1*Nspb)].x) + (temp1[idx + j1 + (k1*(binM + Nspb))].y*temp2[j1 + (k1*Nspb)].y);
               temp3[idx + k1*(binM + 1)].y += (-temp1[idx + j1 + (k1*(binM + Nspb))].x*temp2[j1 + (k1*Nspb)].y) + (temp1[idx + j1 + (k1*(binM + Nspb))].y*temp2[j1 + (k1*Nspb)].x);
           }
       }
    }
}

结果符合我的预期，但是运行起来还是需要一段时间，大约50 seconds 。当我调用内核时，我这样做

xcorr << <1, 1000 >> > (cuda_E2, cuda_A2, cuda_temp, Nb, *binM, Nspb);

我的想法是发送 6 个 block ，而不是仅发送一个，每个 block 有 1000 个线程，以避免 j1 的循环(Nspb=5000)。我尝试了不同的方法，但找不到使用两组不同线程的方法，第一个 block 与我使用它的方式相同，其他 5 个 block 替换 j1环形。有人可以告诉我怎么做吗？

任何帮助将不胜感激。

Answer 1

为第一个 block 和其他 block 编写不同的代码，将它们放入 if (blockIdx.x==0) 的分支中并启动 <<<6,1000>>> 网格？

__global__ void xcorr(...)
{
   if (blockIdx.x==0) {
       // do block zero stuff
   }
   else {
       // what the other blocks shall do
   }
}

一般来说，对于一维卷积/互相关:

1. 将内核放入常量内存中
2. 让每个线程 block 执行相同的操作，但针对结果数组的不同图 block
3. 让每个 block /图 block 加载该图 block 所需的输入数组的所有元素到共享内存中(包括边缘的“光环”)