要优化的 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 );

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

那么有没有人有更好的想法如何最大化性能？

Answer 1

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

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

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

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