在扭曲中对齐 12 字节结构 - cuda答案

【问题标题】：Aligning 12-byte struct within a warp - cuda在扭曲中对齐 12 字节结构 - cuda
【发布时间】：2017-03-29 02:39:51
【问题描述】：

假设我们有一个未对齐的 3 个整数的 struct：

struct data {
  int x;
  int y;
  int z;
};

我将这个struct 的数组传递给内核。我知道我应该传递数组结构而不是结构数组，但这对于这个问题并不重要。

warp 中有 32 个线程，以合并的方式（i 到 i + 31）访问内存，这等于 384 字节的总内存。 384 字节是 L1 缓存线（128 字节）的倍数，这意味着三个内存事务，每个 128 字节。

现在如果我们有一个对齐的struct:

struct __align__(16) aligned_data {
  int x;
  int y;
  int z;
};

如果访问模式与前面的示例相同，那么它将获取 512 字节的内存，即 4 个内存事务，每个请求 128 字节。

所以这意味着第一个示例更有效，或者第二个示例仍然更有效，尽管它获取更多内存。

【问题讨论】：

标签： arrays memory struct cuda

【解决方案1】：

回答问题的唯一真正方法是进行基准测试。如果你这样做了，你可能不会得到相同的答案，具体取决于你的硬件。当我运行这个时：

#define NITER (128)

struct data {
  int x;
  int y;
  int z;
};

struct __align__(16) aligned_data {
  int x;
  int y;
  int z;
};

template<typename T, int niter>
__global__
void kernel(T *in, int *out, int dowrite=0)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int nthreads = blockDim.x * gridDim.x;

    int oval = 0;
#pragma unroll
    for(int i=0; i<niter; ++i,tid+=nthreads) {
        T val = in[tid];
        oval += val.x + val.y + val.z;
    }
    if (dowrite) {
        out[tid] = oval;
    }
}

template __global__ void kernel<data,NITER>(data *, int*, int);
template __global__ void kernel<aligned_data,NITER>(aligned_data *, int*, int);

int main()
{
    const int bs = 512;
    const int nb = 32;
    const int nvals = bs * nb * NITER;

    data *d_; cudaMalloc((void **)&d_, sizeof(data) * size_t(nvals));
    aligned_data *ad_; cudaMalloc((void **)&ad_, sizeof(aligned_data) * size_t(nvals));

    for(int i=0; i<10; ++i) {
        kernel<data,NITER><<<nb, bs>>>(d_, (int *)0, 0);
        kernel<aligned_data,NITER><<<nb, bs>>>(ad_, (int *)0, 0);
        cudaDeviceSynchronize();
    }
    cudaDeviceReset();

    return 0;
}

我发现对齐结构版本在具有计算 5.2 功能的设备上提供了更高的整体性能：

Time(%)      Time     Calls       Avg       Min       Max  Name
 52.71%  2.3995ms        10  239.95us  238.10us  241.79us  void kernel<data, int=128>(data*, int*, int)
 47.29%  2.1529ms        10  215.29us  214.91us  215.51us  void kernel<aligned_data, int=128>(aligned_data*, int*, int)

在这种情况下，我假设大约 10% 的改进归因于发出的加载指令数量减少。在未对齐的情况下，编译器发出三个 32 位加载来获取结构，而在对齐的情况下，编译器发出单个 128 位加载来获取结构。指令的减少似乎抵消了 25% 浪费的净内存带宽。在具有不同内存指令吞吐量与带宽比率的其他硬件上，结果可能会有所不同。

【讨论】：