0

我对 CUDA 很陌生，正在尝试编写一个测试程序。我在 GeForce GT 520 卡上运行应用程序，性能非常差。

该应用程序用于处理一些图像，每一行都由一个单独的线程处理。以下是该应用程序的简化版本。请注意，在实际应用中，所有常量实际上都是变量，只要是调用者。

运行下面的代码时，需要20多秒才能完成执行。

但是与使用malloc/free 不同，当l_SrcIntegral被定义为本地数组（如注释行中所示）时，完成执行所需的时间不到 1 秒。

由于数组的实际大小是动态的（而不是 1700），所以这个本地数组不能在实际应用中使用。

任何关于如何提高这个相当简单的代码性能的建议都将不胜感激。

#include "cuda_runtime.h"
#include <stdio.h>

#define d_MaxParallelRows 320
#define d_MinTreatedRow   5
#define d_MaxTreatedRow   915
#define d_RowsResolution  1
#define k_ThreadsPerBlock 64

__global__ void myKernel(int Xi_FirstTreatedRow)
{
  int l_ThreadIndex = blockDim.x * blockIdx.x + threadIdx.x;
  if (l_ThreadIndex >= d_MaxParallelRows)
    return;
  int l_Row = Xi_FirstTreatedRow + (l_ThreadIndex * d_RowsResolution);
  if (l_Row <= d_MaxTreatedRow) {

    //float l_SrcIntegral[1700];
    float* l_SrcIntegral = (float*)malloc(1700 * sizeof(float));

    for (int x=185; x<1407; x++) {
      for (int i=0; i<1700; i++)
        l_SrcIntegral[i] = i;
    }

    free(l_SrcIntegral);
  }
}

int main()
{
  cudaError_t cudaStatus;

  cudaStatus = cudaSetDevice(0);

  int l_ThreadsPerBlock = k_ThreadsPerBlock;
  int l_BlocksPerGrid = (d_MaxParallelRows + l_ThreadsPerBlock - 1) / l_ThreadsPerBlock;

  int l_FirstRow = d_MinTreatedRow;
  while (l_FirstRow <= d_MaxTreatedRow) {
    printf("CUDA: FirstRow=%d\n", l_FirstRow);
    fflush(stdout);

    myKernel<<<l_BlocksPerGrid, l_ThreadsPerBlock>>>(l_FirstRow);

    cudaDeviceSynchronize();

    l_FirstRow += (d_MaxParallelRows * d_RowsResolution);
  }

  printf("CUDA: Done\n");

  return 0;
}
4

2 回答 2

1

1.

正如@aland 所说，您甚至可能会遇到更差的性能，在每个内核调用中只计算一行。

您必须考虑处理整个输入，只是为了在理论上使用大规模并行处理的力量。

为什么只用 320 个线程启动多个内核来计算一行？如何使用尽可能多的块并让每个块的线程处理一行。

（每块 320 个线程不是一个好的选择，看看如何达到更好的占用率）

2.

如果您作为寄存器和共享内存的快速资源还不够，您必须使用 tile apporach，这是使用 GPGPU 编程的基础之一。

将输入数据分成大小相等的图块，并在线程中循环处理它们。

在这里，我发布了这种平铺方法的示例：

CUDA中的并行化，为每一列分配线程

请注意该瓷砖方法中的范围检查！

给你的想法的例子：

计算任意大小矩阵中列向量中所有元素的总和。

每个块处理一列，并且该块的线程将它们的元素存储在 tile 循环中的共享内存数组中。完成后，他们使用并行归约计算总和，以开始下一次迭代。
最后，每个块计算其向量的总和。

于 2012-05-08T02:11:47.353 回答
0

您仍然可以使用共享内存来使用动态数组大小。<<<...>>>只需在内核调用中传递第三个参数。那就是每个块的共享内存的大小。

到达那里后，只需将所有相关数据放入共享数组（您仍应尝试保持合并访问），为每个线程带来一个或多个（如果与保持合并访问相关）元素。在引入线程后同步线程（仅当您需要停止竞争条件时，以确保在完成任何计算之前整个数组都在共享内存中），您就可以开始了。

另外：您应该使用块和线程进行镶嵌，而不是循环。我知道这只是一个使用本地数组的示例，但是仍然可以通过块/线程进行细分，而不是嵌套 for 循环（这对性能非常不利！）我希望您仅使用 1 个块来运行示例代码和1个线程，否则它没有多大意义。

于 2012-05-03T14:43:56.993 回答