CUDA高效划分？答案

【问题标题】：CUDA efficient division?CUDA高效划分？
【发布时间】：2012-10-04 14:13:56
【问题描述】：

我想知道是否有一种有效的方法来划分数组的元素。我正在使用矩阵值 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 也会给出错误的结果？这里缺少什么？

【问题讨论】：

那么在这个内核中，对于 1000x1000 的情况 N=1000？
@talonmies：是的，对不起。更新
您现在有 N=10000。那应该是 N=1000 吗？
不，其实是10000
所以您正在启动 1 亿 个线程，然后故意让其中一半只执行几次 IOP，您想知道为什么内核运行缓慢？我认为您在错误的地方寻找性能改进。慢除法不是你的问题......

标签： cuda gpu gpgpu nvidia

【解决方案1】：

您的基本问题是您正在启动一个难以置信的巨大网格（对于您的 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;
        }
    }
}

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

这里使用一维网格，每个块一次计算一行。块中的线程沿行移动，因此内存访问被合并。在 cmets 中，您提到您的 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);

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

【讨论】：

@talonmies：它看起来是一个更好的方法，但它不会加速代码。谢谢你，它真的帮助我更多地了解了 CUDA。
@Manolete：自从我发布答案以来，我已经对它进行了基准测试，它肯定会加快代码速度。在 GTX670 上，我的速度比原来的内核快了大约 7 倍。你确定你的时间安排正确吗？
@talonmies：让我仔细检查一下，但整个应用程序需要相同的时间
@talonmies: N=8999 -> mykernel: 1.30893 sec, yourKernel:1.28179 sec
如果你启动一个包含大量块（即 >> 65535）的内核会怎样？可能是结果错误的原因？

【解决方案2】：

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

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

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

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

【讨论】：

每个块的线程数少？块大小为 1024 (32x32)，这是最新硬件上的最大值.....
@talonmies：哎呀。答案固定。
是的，这是我可以在 Fermi 2050 中使用的最大数量
怀疑内核没有使用太多寄存器，我也会尝试使用 192x1 线程块来为每个 SM 提供可用的最大线程数。

【解决方案3】：

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

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

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

【讨论】：