3

我想知道是否有一种有效的方法来划分数组的元素。我正在使用矩阵值 10000x10000 运行，与其他内核相比，它需要相当长的时间。除法是昂贵的操作，我看不出如何改进它。

__global__ void division(int N, float* A, int* B){

  int row = blockIdx.x * blockDim.x + threadIdx.x;
  int col = blockIdx.y * blockDim.y + threadIdx.y;

  if((row < N) && (col <= row) ){
    if( B[row*N+col] >0 )
      A[row*N+col] /= (float)B[row*N+col];
  }

}

内核启动

  int N = 10000;
  int threads = 32
  int blocks = (N+threads-1)/threads
  dim3 t(threads,threads);
  dim3 b(blocks, blocks);
  division<<< b, t >>>(N, A, B);
  cudaThreadSynchronize();

选项 B：

__global__ void division(int N, float* A, int* B){
  int k =  blockIdx.x * blockDim.x + threadIdx.x;
  int kmax = N*(N+1)/2 
  int i,j;
  if(k< kmax){
    row = (int)(sqrt(0.25+2.0*k)-0.5); 
    col = k - (row*(row+1))>>1;
    if( B[row*N+col] >0 )
      A[row*N+col] /= (float)B[row*N+col];
  }
}

推出与

  int threads =192;
  int totalThreadsNeeded = (N*(N+1)/2;
  int blocks = ( threads + (totalThreadsNeeded)-1 )/threads;
  division<<<blocks, threads >>>(N, A, B);

即使 threadId 是正确的，为什么选项 B 会给出错误的结果？这里缺少什么？

4

3 回答 3

4

您的基本问题是您正在启动一个难以置信的巨大网格（对于您的 10000x10000 数组示例，有超过 1 亿个线程），然后由于内核中访问模式的三角形性质，这些线程中有一半从来没有做任何有效的事情。因此，大量的 GPU 周期被无缘无故地浪费了。此外，您使用的访问模式不允许合并内存访问，这将进一步降低实际执行有用工作的线程的性能。

如果我正确理解了您的问题，则内核仅对正方形数组的下三角形执行逐元素除法。如果是这种情况，同样可以使用以下方法完成：

__global__ 
void division(int N, float* A, int* B)
{
    for(int row=blockIdx.x; row<N; row+=gridDim.x) {
        for(int col=threadIdx.x; col<=row; col+=blockDim.x) {
            int val = max(1,B[row*N+col]);
            A[row*N+col] /= (float)val;
        }
    }
}

[免责声明：用浏览器编写，从未编译，从未测试，使用风险自负]

这里使用一维网格，每个块一次计算一行。块中的线程沿行移动，因此内存访问被合并。在评论中您提到您的 GPU 是 Tesla C2050。该设备只需要 112 个块，每个块 192 个线程，就可以完全“填充”14 个 SM 中的每一个，每个块有 8 个块和每个 SM 的最大并发线程数。所以启动参数可能是这样的：

int N = 10000;
int threads = 192;
int blocks = min(8*14, N);
division<<<blocks, threads>>>(N, A, B);

我希望这比您当前的方法运行得快得多。如果数值精度不是那么重要，您可以通过用近似倒数内在和浮点乘法替换除法来进一步加快速度。

于 2012-10-04T17:27:53.197 回答
3

因为线程是以 32 个一组的形式执行的，称为 warp，如果两个if条件都只true针对其中一个线程，那么您需要为一个 warp 中的所有 32 个线程的划分付费。如果条件是false多线程，看看是否可以在单独的内核中过滤掉不需要除法的值。

从 int 到 float 的转换本身可能很慢。如果是这样，您也许可以在前面的步骤中直接生成浮点数，并将 B 作为浮点数数组传入。

您可以在前面生成 B 数组的步骤中生成倒数。如果是这样，您可以在此内核中使用乘法而不是除法。(a / b == a * 1 / b).

根据您的算法，也许您可​​以使用较低精度的除法。有一个内在的, __fdividef(x, y), 你可以试试。还有一个编译器标志，-prec-div=false.

于 2012-10-04T14:55:05.377 回答
2

首先要看的应该是合并的内存访问。这里没有非合并模式的原因，只是交换行和列以避免浪费大量内存带宽：

int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
...
A[row*N+col] ...

即使这是在计算能力 2.0 或更高版本上运行的，缓存也不足以弥补这种次优模式。

于 2012-10-04T17:30:34.110 回答