0

我需要使用一个整数 (0-99) 作为参数，使用相同的数据运行我的 GPU 内核（ALEA 库）100 次。我试图在内核中实现这个循环，但我得到了奇怪的结果。我必须将循环从内核中取出并围绕 GPULaunch 函数，如下所示：

var lp = new LaunchParam(GridDim, BlockDim);
for (int i= 0; i < 100; i++)
{
   GPULaunch(TestKernel, lp, Data, i);
}

代码的 CPU 版本经过高度优化，高效使用 4 个内核 (%100)。根据合并的内存访问原则重新组织内存中的数据后，我可以获得 %92 的占用率和 %96 的全局负载效率。但是，GPU 版本仅比 CPU 版本快 %50。我怀疑以这种方式循环 GPULaunch 是否有效。

正如您在下图中看到的，我没有在 NVIDIA Visual Profiler 中看到重复的内存传输。一旦我将数据加载到 GPU（图中看不到，但对我来说并不重要），我会得到 100 个循环输出的短暂内存传输，如右端所示。所以我的问题是：

  1. 这种在循环中调用 GPULaunch 的方法是否对相同数据进行了看不见的内存传输？
  2. 如果有这样的开销，我需要在内核中有这个循环。我该怎么做。我试过但结果不稳定，认为这种方法不适合 GPU 并行编程架构。

提前致谢

NVIdia Visual Profiler 结果

4

1 回答 1

0

我再次尝试在内核中实现循环并且它起作用了。我不确定这次有什么不同。这是明显的代码（只是模板不是工作代码）：

public class GPUModule : ILGPUModule
{

  public GPUModule (GPUModuleTarget target) : base(target)
  {
  }

  [Kernel]
  Public MyKernel(deviceptr<int> Data)
  {
    var start = blockIdx.x * blockDim.x + threadIdx.x;
    int ind = threadIdx.x;

    for (int i=0;i<100;i++)
    {
      //Kernel Code here
    }
  }

  public void Dilimle_VerilerB(deviceptr<int> Data
  {
    ...
    var lp = new LaunchParam(GridDim, BlockDim);
    GPULaunch(TestKernel, lp, Data, HOIndex);
    ...
  }
}

对内核的唯一补充是整数“i”上的循环。不幸的是，它导致寄存器/线程计数从 26 跳到 42，导致占用率从 %100 下降到 %50，这将执行时间从 2.1 秒略微增加到 2.3 秒。因此，如果保持 %100 Occupancy，将循环放入内核将提高性能，从而大大消除 GPULaunch 开销。

GPULaunch 周围循环的 %100 Occupancy 使用 1024 个线程/块。更改为内部内核循环后，我将其更改为 128 个线程/块。这将占用率增加到 %62，并导致执行时间为 1.1 秒。总而言之，在内核中采用这样的循环可将 GPU 性能提高 2 倍。

所以问题是为什么内核的寄存器/线程数从 26 增加到 42 并在内核中添加一个整数循环。我想如果寄存器数量在 30 左右，占用率仍然可能接近 %100。

于 2017-01-30T17:25:23.000 回答