0

我有以下内核： 

__global__ void myKernel(int k, int inc, int width, int* d_Xco, int* d_Xnum, bool* 
        Xvalid, float* d_X)
    {

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

        if( (i < k)  ){
           for(k1 = 0; k1 < inc; k1++){

             int mul = (d_X[i*inc + k1] >= 2e2);
             d_X[i*inc + k1] *= (float)(!mul);
             d_Xco[i*width + k1] = k*mul;
             d_Xnum[i] += mul;
             d_Xvalid[i*inc + k1] = (!mul) ; 

            }
         }// of if

 }

就是这样称呼的：

  int bx = (int)(k/32)+1;
  int by = (int)(inc/32)+1;

  dim3 b(bDim, 1);
  dim3 t(tDim, 1);
  cmyKernel<< b, t >>>( k, inc, width, d_Xco, d_Xnum, d_Xvalid, d_X );

  cudaThreadSynchronize();

k是9000inc左右，5000左右，所以我确定我没有超过块数。如果myKernel在y维度中使用 1thread/1block 调用，内核似乎工作正常，但是，例如，只需将y维度中的线程和块数更改为 10，它会给出错误的输出，即使在内核中我并没有真正使用中的线程和块y。理想情况下，我想摆脱for()使用k = threadIdx.y + blockIdx.y * blockDim.y

4

2 回答 2

2

如果您启动一个 y 维度 = 10 的内核，那么您正在使用它们。只是因为您没有使用线程标识符threadIdx.y，blockIdx.y并不意味着没有启动线程。当您启动一个 y 维度 = 10 的内核时，您将有 10 个 i = 0 的线程、10 个 i = 1 的线程等。

假设您为简单起见启动 2x2 线程。您将拥有线程 (0,0)(0,1) (1,0) (1,1)。在您的代码中，两个线程 (0,0) 和 (0,1) 的 i 变量为 0，但threadIdx.y不同。这意味着两个线程将评估相同 i 变量的代码并导致竞争条件。

您需要解决迭代之间的依赖关系（d_Xnum[i] += mul）。一种方法是使用atomicAdd(..). 取消注释k1，将循环替换为if(k1 < inc)并添加atomicAdd. 那应该给你正确的行为。

于 2012-10-02T14:52:30.007 回答
1

正如评论中已经说明的那样，您当前的解决方案是启动多个线程，每个线程将它们的工作应用于相同的内存空间。这是由于生成的多个线程都具有相同值的 threadIdx.x，而 threadIdx.y 值不同。这意味着您将有多个线程同时读取和写入同一内​​存空间，这有很多潜在的问题，这里是一个简单的描述。

为避免这种情况，您可以采取几个步骤。例如，您可以使用同步数据访问（这将导致速度大幅下降，因为线程等待其他人完成数据访问）。如果您想让每个线程处理一个单元元素，则需要删除 for 循环并像以前一样使用 k1，但是您必须仔细考虑内存读取和写入，因为进程的任何部分都可能在任何之前或之后执行其他在不同的线程！

这里的核心理解是，你永远不能依赖线程之间的操作顺序！

在数据访问方面，将所有数据结构视为网格会有所帮助，其中每个线程只应访问和修改其自身坐标中的数据，例如 (3,2) 用于 threadIdx.x == 3 的线程和 threadIdx.y == 2。这样，您可以轻松地可视化线程的行为和潜在的竞争条件。使用此功能的最简单方法是为输出数据的每个元素创建一个网格条目，因此如果您有一个 9000x5000 元素的矩阵，您可能会从该数量的线程开始，并从那里进行优化。这当然会导致 GPU 必须在其所有单元上执行多次，但这是一个很好的起点。

奥斯陆大学有一个关于这个主题的研究生课程，等等。您可能会发现这些幻灯片 与加深理解高度相关。请特别参阅有关线程批处理、网格和块的部分。

于 2012-10-02T13:16:00.000 回答