c++ - 要优化的 CUDA 内核

Question

您好，我最近有一个 CUDA 内核需要优化。这是原始的 CUDA 内核:

__glboal__ void kernel_base( float *data, int x_dim, int y_dim )
{
  int ix  = blockIdx.x;
  int iy  = blockIdx.y*blockDim.y + threadIdx.y;
  int idx = iy*x_dim + ix;
  float tmp = data[idx];

  if( ix % 2 )
  {
    tmp += sqrtf( sinf(tmp) + 1.f );
  }
  else
  {
     tmp += sqrtf( cosf(tmp) + 1.f );
  }
  data[idx] = tmp;
}


dim3 block( 1, 512 );
dim3 grid( 2048/1, 2048/512 );
kernel<<<grid,block>>>( d_data, 2048, 2048 );

这里的基本问题是内存合并和线程发散的困境。原代码以列为主处理数组，因此具有跨步内存访问模式，但没有发散。我可以将其更改为 row-major，这又存在线程发散的问题。

那么有没有人知道如何最大限度地提高性能？

Answer 1

就性能而言，与跨步内存访问相比，这里的线程分歧不是一个大问题。我会去合并。此外，您的数据存储具有隐式 AoS 排序。如果您可以将数据重新排序到 SoA，则可以解决这两个问题。

所以我会重新排序这个内核，首先以行优先的方式处理事情。这解决了合并问题，但引入了扭曲发散。

如果您无法对数据重新排序，我会考虑通过修改索引方案来消除 warp 发散，以便偶数 warp 处理偶数元素，而奇数 warp 处理奇数元素。

这将消除扭曲发散，但会再次破坏完美合并，但缓存应该有助于解决这个问题。在 Fermi 的情况下，L1 缓存应该可以很好地平滑这种模式。然后我会将这种情况与 warp 发散情况进行比较，看看哪个更快。