CUDA高效除法？

Question

CUDA高效除法？

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);

为什么即使线程ID正确，选项B仍然给出错误的结果？这里缺少了什么？

- Manolete

所以在这个内核中，对于1000x1000的情况，N=1000？ - talonmies

@talonmies：是的，抱歉。已更新。 - Manolete

你现在有N=10000。是否应该为N=1000？ - talonmies

不，实际上是10000。 - Manolete

5

所以你正在启动1亿个线程，然后故意让其中一半的线程仅执行少量IOPs操作，然后你想知道为什么内核运行缓慢？我认为你在寻找提高性能的错误地方。缓慢的除法不是你的问题...... - talonmies

你说得对，我想我在错误的地方寻找。然而，我已经尝试使用不同的内核启动了 N*(N+1)/2 个线程，但是我无法得到正确的结果。 - Manolete

3个回答

3

由于线程是以32个为一组，称为warp的方式执行的，如果对于其中一个线程，两个if条件都为true，则您需要为warp中的所有32个线程进行分割。如果对于许多线程，条件都是false，请尝试在单独的内核中筛选出不需要进行除法的值。

将int转换为float本身可能很慢。如果是这样，您可以在较早的步骤中直接生成浮点数，并将B作为浮点数组传递。

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

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

- Roger Dahl

1

每个块的线程数很少？块大小为1024（32x32），这是现代硬件上的最大值..... - talonmies

@talonmies：糟糕，答案已修正。 - Roger Dahl

是的，这是我可以在Fermi 2050中使用的最大数字。 - Manolete

怀疑内核没有使用太多的寄存器，我也会尝试使用192x1线程块来为每个SM提供最大数量的线程。 - pQB

2

首先要考虑的是合并内存访问。这里没有非合并模式的理由，只需交换行和列以避免浪费大量内存带宽：

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

即使在计算能力为2.0或更高的情况下运行，缓存也不足以弥补这种次优模式。

- tera

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- talonmies · Accepted Answer

你的基本问题在于你正在启动一个不可思议巨大的网格（对于你的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，每个SM都具有8个块和最大数量的并发线程。因此，启动参数可以是以下内容：

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

我希望这种方法的运行速度比你目前的方法快得多。如果数值精度不是很重要，你可以用近似倒数内置函数和浮点乘法来取代除法，从而进一步提高速度。