CUDA高效划分？

Question

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

Answer 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;
        }
    }
}

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

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

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

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