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

在第二种情况下,有一些介于16和32之间的“读”数。但是我们应该更加小心术语,以便理解。

这个 聚结 过程如下。

1. LD/ST单元接收 请求 . 假设我们讨论的是一个读请求(即一个ld指令)。读请求构成ld指令加上在warp中每个线程生成的地址。

2. 当对照缓存线或内存段查看时,处理请求以确定每个地址相对于其他地址的位置。对于这个讨论,我们假设在任何缓存中都没有命中,因此我们必须针对内存段合理化请求。内存段是全局内存空间的固定细分,对应于DRAM子系统可发行的最小事务大小。在我熟悉的所有CUDAGPU上,内存/DRAM段大小为32字节。根据DRAM段模式,对WARP中的每个线程生成的地址进行映射,将标识必须检索内存中的哪些实际段才能满足此LD请求。

3. 内存控制器将检索这些段。对于DRAM,检索段的每个请求都是 交易 .

4. 检索到的段数据将用于填充缓存线,并满足原始的ld请求,即翘曲宽度。

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

现在,通过上面的描述,让我们看看您的具体示例。

在第一个示例中,您声明“它们都被合并成一个读取”。当然,它们最初只是一个单一的读取请求。但是,满足每个线程16个字节的非重叠全扭曲读取(32个线程)的32字节DRAM事务的最小数目是512字节,或者512/32=16个段。根据测量的位置或方式,它也可以称为4个全局事务,因为全局加载事务的宽度高达128字节。但是,无论我们如何/在何处度量,这都将是一个完全合并的100%最佳事务集,因为生成了满足此类请求所需的最小事务数,并且从内存中检索到的每个字节都被使用,或者至少被翘曲中的线程请求。

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

因此,对于最佳情况模式,第二个示例只向DRAM发出16个事务,与第一个示例在生成的DRAM事务数量和总体效率(100%利用率)方面完全匹配。对于最坏情况的模式(每个线程地址以32字节边界隔开),则需要32个段,因此需要32个DRAM事务来满足翘曲读取请求。

为了给出一个代码示例,以下序列将为每个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个字节)。

以下序列将为第一个扭曲生成16个DRAM事务:

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

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