cuda - 在缓冲区的不同部分合并读写

在第二种情况下，会有 16 到 32 个“读取”之间的一些数字。但是我们应该更加小心术语，以便理解。

聚结过程如下工作。

LD/ST 单元接收请求。假设我们正在讨论读取请求（即 LD 指令）。读请求构成 LD 指令加上 warp 中每个线程生成的地址。
处理请求以确定每个地址相对于其他地址的位置，当针对高速缓存行或内存段查看时。对于这个讨论，我们假设在任何缓存中都没有命中，因此我们必须合理化针对内存段的请求。内存段是全局内存空间的固定细分，对应于可发布给 DRAM 子系统的最小事务大小。在我熟悉的所有 CUDA GPU 上，内存/DRAM 段大小为 32 字节。扭曲中每个线程生成的地址与 DRAM 段模式的映射将识别必须检索内存中的哪些实际段以满足此 LD 请求。
内存控制器将检索这些段。对于 DRAM，检索段的每个请求都是一个事务。
检索到的段数据将用于适当地填充高速缓存行，并满足原始 LD 请求，warp-wide。

合并基本上发生在第 2 步。由于跨 warp 发出的地址被映射到 DRAM 段的底层模式，如果多个地址落入单个段，则不会多次请求该段。它只会被请求一次。这就是合并的中心思想。

现在，通过以上描述，让我们看一下您的具体示例。

在第一个示例中，您声明“它们都合并为一个读取”。好吧，它们当然是从单个读取请求开始的。但是，满足每个线程 16 字节的非重叠 full-warp 读取（32 个线程）的 32 字节 DRAM 事务的最小数量是 512 字节，或 512/32 = 16 个段。根据您测量的位置或方式，它也可以称为 4 个全局事务，因为全局加载事务最多 128 个字节宽。但无论我们如何/在何处衡量这一点，这将是一个完全合并的、100% 优化的事务集，因为生成满足此类请求所需的最小事务数，并且使用或至少使用从内存中检索到的每个字节经线中的线程请求。

在第二个示例中，如果不知道 warp 中的线程生成的实际地址模式，就无法确定精确的活动。对于读取 0 到 512 之间位置的线程，此范围内最多有 512/32 = 16 个段。并且有16个线程。因此，您可能处于每个线程都需要自己的段的最坏情况（对于这种特殊安排）。或者，如果线程地址在 32 字节边界处没有完美间隔，则前 16 个线程所需的 DRAM 事务数可能低于 16，可能低至 8 个。同样，对于第二组 16 个线程，以及第二组 512 字节在内存中。

因此，对于最佳情况模式，第二个示例仅向 DRAM 发出 16 个事务，就生成的 DRAM 事务数量以及整体效率（100% 利用率）而言，与第一个示例完全匹配。对于最坏的情况模式（每个线程地址间隔为 32 字节边界），则需要 32 个段，因此需要 32 个 DRAM 事务来满足 warp 读取请求。

举一个代码示例，以下序列将在每个 warp 中生成 32 个 DRAM 事务：

__global__ void k(float4 *d){
  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  float4 temp = d[idx*2];
  ...
  }

在上面的示例中，每个线程生成的底层字节地址将完美地间隔在 32 字节边界处。前 16 个线程将从内存中的第一个 512 字节区域请求数据，后 16 个线程将从内存中的第二个 512 字节区域请求数据。该请求的总体效率为 50%（将从内存中请求 1024 个字节，但 warp 中的线程只需要 512 个字节）。

以下序列将为第一个 warp 生成 16 个 DRAM 事务：

__global__ void k(float4 *d){
  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  float4 temp = d[idx + (idx/16)*16];
  ...
  }

在上面的示例中，对于前 16 个线程（在第一个 warp 中），每个线程生成的底层字节地址将是 0,16,32,48...,252。对于第二个 16 个线程（在第一个线程中），地址将是 512,528,544....,764。前 16 个线程将从内存中的第一个 512 字节区域请求数据，后 16 个线程将从内存中的第二个 512 字节区域请求数据。然而，前 16 个线程只需要 8 个 DRAM 事务，而后 16 个线程只需要 8 个 DRAM 事务。该请求的整体效率将是 100%（将从内存中请求 512 个字节，对于 warp 中的线程所需的 512 个字节）。

cuda - 在缓冲区的不同部分合并读写

问题描述

解决方案

推荐阅读