CUDA知识点之Bank Conflict(3)-向量化访问

#01 背景

在前文的讨论中，我们分析了一个warp内的每个线程访问4个Byte的情况，这样所有线程发起一次读写请求，可以访问到128 Byte的 Shared Memory。在涉及到向量化访存时，访问模式就变了，因为通过一个 LDS.64 或 LDS.128 指令就可以让一个线程一次性访问 8 个或 16 个字节（相当于 2 个或 4 个元素）。

#02 基础知识

放在 shared memory 中的数据是以 4 bytes（即 32 bits）作为 1 个 word，依次放在 32 个 banks 中。所以，第 i 个 word，就存放在第 ( i mod 32 ) 个 bank 上。每个 bank 在每个 cycle 的 bandwidth 为 32 bits，所以 shared memory 在每个 cycle 的 bandwidth 为 32 * 32 bits = 32 * 4 bytes = 128 bytes，也就是说，每次 memory transaction 最多访问 128 bytes 的数据。

当使用 float2 进行访存时，每个 thread 同时读取 8 bytes 数据，若每个 thread 读取不同的数据，则一个 warp 共访问 256 bytes 数据，不产生 bank conflicts 时会拆分为 2 个 memory transactions 进行，且拆分顺序是按照 0-15，16-31 这样 16 个连续 threads 为一组进行的，我们将每一组称为一个 half-warp，即 1/2 warp。

当使用 float4 进行访存时，每个 thread 同时读取 16 bytes 数据，若每个 thread 读取不同的数据，则一个 warp 共访问 512 bytes 数据，不产生 bank conflicts 时会拆分为 4 个 memory transactions 进行，且拆分顺序是按照 0-7，8-15，16-23，24-31 这样 8 个连续 threads 为一组进行的，我们将每一组称为一个 quarter-warp，即 1/4 warp。

当每个 thread 读取一个 float 数据时，如果一个 warp 中不同的 thread 读取相同的数据，会触发 broadcast （广播）机制，合并为一次访存。

实际上，Bank Conflict是以每个 Wrap 内所产生的 Memory Transaction作为分析对象，而非以内存读写请求来划分粒度。如果单次 memory transaction 需要访问的 128 bytes 中有多个 word 属于同一个 bank，就产生了 bank conflict，从而需要拆分为多次 transaction。

一个 warp 中所有线程在同时执行一条 Shared Memory 访存指令时会对应到 1 个或多个 Memory Transaction，一个 Memory Transaction 最长是 128 字节。如果一个 warp 内在同一时刻所需要的访存超过了 128 字节，那么会则被拆成多个 Transaction 进行。因为一个 warp 同一时刻执行的访存指令的位宽应该是一样的（即例如不存在线程 0 执行 LDS.32 而线程 1 执行 LDS.128），因此我们只需要对 64 位宽和 128 位宽的访存指令分别讨论即可。

#03 64位宽访问

使用 LDS.64 指令（或者通过 float2、uint2 等类型）取数据时，每个 thread 请求 64 bits（即 8 bytes）数据，那么每 16 个 thread 就需要请求 128 bytes 的数据。所以 CUDA 会默认将一个 warp 拆分为两个 half warp，每个 half warp 产生一次 memory transaction，即一共两次 transaction。

只有以下两个广播条件之一满足时，这两个 half warp 的访问才会合并成一次 memory transaction：

  • 对于 warp 内所有活跃的第 i 号线程，第 i xor 1 号线程不活跃或者访存地址和其一致；(i.e. T0==T1, T2==T3, T4==T5, T6==T7, T8 == T9, ......, T30 == T31, etc.)
  • 对于 warp 内所有活跃的第 i 号线程，第 i xor 2 号线程不活跃或者访存地址和其一致；(i.e. T0==T2, T1==T3, T4==T6, T5==T7 etc.)

一个 Half-warp 活跃的定义是这个 Half-warp 内有任意一个线程活跃。
memory transaction合并，可以视为是从SMEM中取128bytes数据，然后再在warp内部分发。

第一个例子：


image.png

在这个例子中，活跃线程全部在第 1 个 Half-warp 内，每个线程依次访问连续的 uint2，需要 1 个 Memory Transaction，没有 Bank Conflict，需要 1 个 Wavefront。

第二个例子：


image.png

活跃线程分散在了 2 个 Half-warp 内，每个线程一次访问连续的uint2，需要 2 个 Memory Transaction，没有 Bank Conflict，需要 2 个 Wavefront（注意第 15 号和第 16 号线程）。

第三个例子：


image.png

活跃线程分散在了 2 个 Half-warp 内，但因为触发了广播机制中的第一条，因此仍然只需要 1 个 Memory Transaction，没有 Bank Conflict，需要 1 个 Wavefront。

第四个例子：


image.png

活跃线程分散在了 2 个 Half-warp 内，看似好像触发了广播机制，但其实并没有，因为第一个 Half-warp 触发的是第一条，第二个 Half-warp 触发的是第二条，因此仍然需要 2 个 Memory Transaction，没有 Bank Conflict，需要 2 个 Wavefront。

第五个例子：


image.png

活跃线程分散在了 2 个 Half-warp 内，没有触发广播机制，需要 2 个 Memory Transaction，没有 Bank Conflict，需要 2 个 Wavefront

可以通过下面的代码来进行验证：

#include <cstdint>

__global__ void smem_1(uint32_t *a) {
  __shared__ uint32_t smem[128];
  uint32_t tid = threadIdx.x;
  for (int i = 0; i < 4; i++) {
    smem[i * 32 + tid] = tid;
  }
  __syncthreads();
  if (tid < 16) {
    reinterpret_cast<uint2 *>(a)[tid] =
        reinterpret_cast<const uint2 *>(smem)[tid];
  }
}

__global__ void smem_2(uint32_t *a) {
  __shared__ uint32_t smem[128];
  uint32_t tid = threadIdx.x;
  for (int i = 0; i < 4; i++) {
    smem[i * 32 + tid] = tid;
  }
  __syncthreads();
  if (tid < 15 || tid == 16) {
    reinterpret_cast<uint2 *>(a)[tid] =
        reinterpret_cast<const uint2 *>(smem)[tid == 16 ? 15 : tid];
  }
}

__global__ void smem_3(uint32_t *a) {
  __shared__ uint32_t smem[128];
  uint32_t tid = threadIdx.x;
  for (int i = 0; i < 4; i++) {
    smem[i * 32 + tid] = tid;
  }
  __syncthreads();
  reinterpret_cast<uint2 *>(a)[tid] =
      reinterpret_cast<const uint2 *>(smem)[tid / 2];
}

__global__ void smem_4(uint32_t *a) {
  __shared__ uint32_t smem[128];
  uint32_t tid = threadIdx.x;
  for (int i = 0; i < 4; i++) {
    smem[i * 32 + tid] = tid;
  }
  __syncthreads();
  uint32_t addr;
  if (tid < 16) {
    addr = tid / 2;
  } else {
    addr = (tid / 4) * 4 + (tid % 4) % 2;
  }
  reinterpret_cast<uint2 *>(a)[tid] =
      reinterpret_cast<const uint2 *>(smem)[addr];
}

__global__ void smem_5(uint32_t *a) {
  __shared__ uint32_t smem[128];
  uint32_t tid = threadIdx.x;
  for (int i = 0; i < 4; i++) {
    smem[i * 32 + tid] = tid;
  }
  __syncthreads();
  reinterpret_cast<uint2 *>(a)[tid] =
      reinterpret_cast<const uint2 *>(smem)[tid % 16];
}

int main() {
  uint32_t *d_a;
  cudaMalloc(&d_a, sizeof(uint32_t) * 128);
  smem_1<<<1, 32>>>(d_a);
  smem_2<<<1, 32>>>(d_a);
  smem_3<<<1, 32>>>(d_a);
  smem_4<<<1, 32>>>(d_a);
  smem_5<<<1, 32>>>(d_a);
  cudaFree(d_a);
  cudaDeviceSynchronize();
  return 0;
}

#04 128位宽访问

使用 LDS.128 指令（或者通过 float4、uint4 等类型）取数据时，每个 thread 请求 128 bits（即 16 bytes）数据，那么每 8 个 thread 就需要请求 128 bytes 的数据。

CUDA 会默认把每个 half warp 进一步切分成两个 quarter warp，每个包含 8 个 thread。每个 quarter warp 产生一次 memory transaction。所以每个 warp 每次请求，默认会有 4 次 memory transaction。（没有 bank conflict 的情况下）。

类似 64 位宽的情况，当满足特定条件时，一个 half warp 内的两个 quarter warp 的访存请求会合并为 1 次 memory transaction，但是两个 half warp 不会再进一步合并了。所以，假设一个 warp 中 32 个线程都活跃，即使它们的访存地址都一样，也需要 2 个 Memory transaction。

128位的广播条件与64位是一样的。当满足特定条件时，一个 half warp 内的两个 quarter warp 的访存请求会合并为 1 次 memory transaction，但是两个 half warp 不会再进一步合并了。

合并原理：

  • 对于 warp 内所有活跃的第 i 号线程，第 i xor 1 号线程不活跃或者访存地址和其一致；(i.e. T0==T1, T2==T3, T4==T5, T6==T7, T8 == T9, ......, T30 == T31, etc.)
  • 对于 warp 内所有活跃的第 i 号线程，第 i xor 2 号线程不活跃或者访存地址和其一致；(i.e. T0==T2, T1==T3, T4==T6, T5==T7 etc.)

第一个例子：


image.png

活跃线程分散在了 2 个 Half-warp 和 2 个 Quarter-warp 内，每个 Half-warp 需要 1 个 Memory Transaction，总共需要 2 个 Memory Transaction，没有 Bank Conflict，需要 2 个 Wavefront

第二个例子：


image.png

活跃线程分散在了 1 个 Half-warp 和 2 个 Quarter-warp 内，需要 1 个 Memory Transaction，没有 Bank Conflict，需要 1 个 Wavefront。

第三个例子：


image.png

活跃线程分散在了 2 个 Half-warp 和 4 个 Quarter-warp 内，但触发了广播机制（第一条），每个 Half-warp 需要 1 个 Memory Transaction，总共需要 2 个 Memory Transaction，没有 Bank Conflict，需要 2 个 Wavefront。

第四个例子：


image.png

活跃线程分散在了 2 个 Half-warp 和 4 个 Quarter-warp 内，没有触发广播机制，每个 Half-warp 需要 2 个 Memory Transaction，总共需要 4 个 Memory Transaction，没有 Bank Conflict，需要 4 个 Wavefront

第五个例子：


image.png

活跃线程分散在了 2 个 Half-warp 和 4 个 Quarter-warp 内，但触发了广播机制（第一条），所以前两个和后两个 quarter warp 分别合并，每个 Half-warp 需要 1 个 Memory Transaction，总共需要 2 个 Memory Transaction。但因为每个Half-warp 中产生了2-way Bank Conflict，因此会拆分成 4 个 Memory Transaction，对应需要 4 个 Wavefront。

第六个例子：


image.png

活跃线程分散在了 2 个 Half-warp 和 4 个 Quarter-warp 内，没有触发广播机制，每个 Half-warp 需要 2 个 Memory Transaction，总共需要 4 个 Memory Transaction，没有 Bank Conflict，需要 4 个 Wavefront

可以通过下面的代码来进行验证：

#include <cstdint>

__global__ void smem_1(uint32_t *a) {
  __shared__ uint32_t smem[128];
  uint32_t tid = threadIdx.x;
  for (int i = 0; i < 4; i++) {
    smem[i * 32 + tid] = tid;
  }
  __syncthreads();
  if (tid == 15 || tid == 16) {
    reinterpret_cast<uint4 *>(a)[tid] =
        reinterpret_cast<const uint4 *>(smem)[4];
  }
}

__global__ void smem_2(uint32_t *a) {
  __shared__ uint32_t smem[128];
  uint32_t tid = threadIdx.x;
  for (int i = 0; i < 4; i++) {
    smem[i * 32 + tid] = tid;
  }
  __syncthreads();
  if (tid == 0 || tid == 15) {
    reinterpret_cast<uint4 *>(a)[tid] =
        reinterpret_cast<const uint4 *>(smem)[4];
  }
}

__global__ void smem_3(uint32_t *a) {
  __shared__ uint32_t smem[128];
  uint32_t tid = threadIdx.x;
  for (int i = 0; i < 4; i++) {
    smem[i * 32 + tid] = tid;
  }
  __syncthreads();
  reinterpret_cast<uint4 *>(a)[tid] = reinterpret_cast<const uint4 *>(
      smem)[(tid / 8) * 2 + ((tid % 8) / 2) % 2];
}

__global__ void smem_4(uint32_t *a) {
  __shared__ uint32_t smem[128];
  uint32_t tid = threadIdx.x;
  for (int i = 0; i < 4; i++) {
    smem[i * 32 + tid] = tid;
  }
  __syncthreads();
  uint32_t addr;
  if (tid < 16) {
    addr = (tid / 8) * 2 + ((tid % 8) / 2) % 2;
  } else {
    addr = (tid / 8) * 2 + ((tid % 8) % 2);
  }
  reinterpret_cast<uint4 *>(a)[tid] =
      reinterpret_cast<const uint4 *>(smem)[addr];
}

__global__ void smem_5(uint32_t *a) {
  __shared__ uint32_t smem[128];
  uint32_t tid = threadIdx.x;
  for (int i = 0; i < 4; i++) {
    smem[i * 32 + tid] = tid;
  }
  __syncthreads();
  reinterpret_cast<uint4 *>(a)[tid] =
      reinterpret_cast<const uint4 *>(smem)[(tid / 16) * 4 + (tid % 16) / 8 + (tid % 8) / 4 * 8];
}

__global__ void smem_6(uint32_t *a) {
  __shared__ uint32_t smem[128];
  uint32_t tid = threadIdx.x;
  for (int i = 0; i < 4; i++) {
    smem[i * 32 + tid] = tid;
  }
  __syncthreads();
  uint32_t addr = (tid / 16) * 4 + (tid % 16 / 8) * 8;
  if (tid < 16) {
    addr += (tid % 4 / 2) * 2;
  } else {
    addr += (tid % 4 % 2) * 2;
  }
  reinterpret_cast<uint4 *>(a)[tid] =
      reinterpret_cast<const uint4 *>(smem)[addr];
}

int main() {
  uint32_t *d_a;
  cudaMalloc(&d_a, sizeof(uint32_t) * 128);
  smem_1<<<1, 32>>>(d_a);
  smem_2<<<1, 32>>>(d_a);
  smem_3<<<1, 32>>>(d_a);
  smem_4<<<1, 32>>>(d_a);
  smem_5<<<1, 32>>>(d_a);
  smem_6<<<1, 32>>>(d_a);
  cudaFree(d_a);
  cudaDeviceSynchronize();
  return 0;
}

#05 总结

可以参考NV官方的表述来理解Bank Conflict与memory transaction的关系：


image.png

从上面的分析过程可以看出，64/128位宽的访存指令对应的广播机制与32位宽情况下，Bank Conflict的情况是不一样的。在32位宽情况下，我们主要关心统一Bank下的并发度，而在64/128位宽情况下，我们不仅要关心并发度，还需要结合 Memory Transaction与warp的分拆情况，最终我们可以看到，Bank Conflict实际上是在memory transaction中分析出来的，而不是由shared memory访问请求决定的。

最后编辑于 ：
©著作权归作者所有,转载或内容合作请联系作者
平台声明：文章内容（如有图片或视频亦包括在内）由作者上传并发布，文章内容仅代表作者本人观点，简书系信息发布平台，仅提供信息存储服务。

推荐阅读更多精彩内容