用于内循环的 CUDA 线程

Question

我有这个内核

__global__ void kernel1(int keep, int include, int width, int* d_Xco, 
              int* d_Xnum, bool* d_Xvalid, float* d_Xblas)
{

  int i, k;  
  i = threadIdx.x + blockIdx.x * blockDim.x;

  if(i < keep){

    for(k = 0; k < include ; k++){

      int val = (d_Xblas[i*include + k] >= 1e5);
      int aux = d_Xnum[i];

      d_Xblas[i*include + k] *= (!val);
      d_Xco[i*width + aux] = k;
      d_Xnum[i] +=val;
      d_Xvalid[i*include + k] = (!val);
    }
  }
}

发布于

int keep = 9000;
int include = 23000;
int width = 0.2*include;

int threads = 192;
int blocks = keep+threads-1/threads;
kernel1 <<< blocks,threads  >>>( keep, include, width,
                                 d_Xco, d_Xnum, d_Xvalid, d_Xblas );

这个kernel1 工作正常，但显然没有完全优化。我认为消除内部循环k 会很直接，但由于某种原因它不能正常工作。 我的第一个想法是：

__global__ void kernel2(int keep, int include, int width, 
               int* d_Xco, int* d_Xnum, bool* d_Xvalid, 
               float* d_Xblas)
{

  int i, k;  
  i = threadIdx.x + blockIdx.x * blockDim.x;
  k = threadIdx.y + blockIdx.y * blockDim.y;

  if((i < keep)  && (k < include) ) {

      int val = (d_Xblas[i*include + k] >= 1e5);
      int aux = d_Xnum[i];
      d_Xblas[i*include + k] *= (float)(!val);
      d_Xco[i*width + aux] = k;
      atomicAdd(&d_Xnum[i], val);
      d_Xvalid[i*include + k] = (!val);
  }
}

以 2D 网格启动：

int keep = 9000;
int include = 23000;
int width = 0.2*include;

int th = 32;
dim3 threads(th,th);
dim3 blocks ((keep+threads.x-1)/threads.x, (include+threads.y-1)/threads.y);
kernel2 <<< blocks,threads >>>( keep, include, width, d_Xco, d_Xnum, 
                               d_Xvalid, d_Xblas );

虽然我认为这个想法很好，但它不起作用，而且我的想法已经用完了。你能帮帮我吗？我还认为问题可能出在d_Xco 中，它将位置k 存储在一个较小的数组中并将它们推到数组的开头，因此顺序很重要。

d_Xco
-------------------------------
| 2|3 |15 |4 |5 |5 | | | | | | .......
-------------------------------

Answer 1

在原始代码中，你有

for(k = 0; k < include ; k++){
  ...
  int aux = d_Xnum[i];
  ...
  d_Xco[i*width + aux] = k;
  ...
}

d_Xco 数组的索引不依赖于k，因此每次迭代都写入它是多余的。最终值将始终为include-1。因此，将k 循环内的这两行替换为k 循环外的一行：

 d_Xco[i*width + d_Xnum[i]] = include - 1;

一旦你这样做了，当你并行化 k 循环时，当许多 k 线程同时为 d_Xco 中的同一位置分配不同的值时，你将不再有当前的竞争条件（不保证排序)。