1

我试图了解合并全局内存。
假设我想将一组奇数的浮点数加载到全局内存中。每个线程将处理一组 3 个浮点数。假设这些浮点数是 A、B 和 C。 

A0,  B0,  C0
A1,  B1,  C1
A2,  B2,  C2
..          
A19, B19, C19

所以线程会像这样抓取数据：

Thread 0:  A0,  B0,  C0  
Thread 1:  A1,  B1,  C1  
Thread 2:  A2,  B2,  C2
..
Thread 19:  A19, B19, C19

第一种方法：
我可以加载 3 个数组：float A[20]; float B[20]; floatC[20];我必须 cudaMemcpy() 三个不同的时间才能将数据加载到全局内存中。这种方法可能不会很好地结合起来。

第二种方法：
更好的方法是：

struct {float A, float B, float C} dataPt;
dataPt data[20];

我可以用一个 cudaMemcpy() 加载数据，但我不确定内存访问是否会很好地合并。

第三种方法：

struct {float A, float B, float C, float padding} dataPt2;
dataPt2 data2[20];

或者

struct __align__(16){float A, float B, float C} dataPt3;
dataPt3 data3[20];

我可以使用单个 cudaMemcpy() 将数据加载到全局内存，并且线程对数据的访问将被合并。（以浪费全局内存为代价。）

1）第一种方法不会合并，因为每个线程可能需要 3 个总线周期来加载输入数据。
2) 第二种方法将合并许多线程，但会有一些线程需要两个总线周期来获取输入数据。
3）第三种方法将合并所有线程。

这是准确的吗？第二种和第三种方法之间有显着差异吗？是否有使用 3 个线程维度（threadIdx.x、threadIdx.y、threadIdx.z）的方法？

4

2 回答 2

2

只是放大@talonmies 的答案。假设我们的内核如下所示：

__global__ void kern(float *a, float *b, float *c){

  float local_a, local_b, local_c;
  int idx = threadIdx.x + (blockDim.x * blockIdx.x);

  local_a = a[idx];
  local_b = b[idx];
  local_c = c[idx];
}

忽略优化（这将导致空内核），并假设我们启动 1 个 32 个线程块：

  kern<<<1, 32>>>(d_a, d_b, d_c);

然后我们有 32 个线程（1 个 warp）在锁步中执行。这意味着每个线程将处理以下内核代码行：

  local_a = a[idx];

在同一时间。合并加载（来自全局内存）的定义是当一个扭曲加载一系列数据项时，这些数据项都在全局内存中的单个 128 字节对齐边界内（对于 CC 2.0 设备）。具有 100% 带宽利用率的完美合并负载意味着每个线程在该 128 字节对齐区域内使用一个唯一的 32 位数量。如果线程 0 加载 a[0]，线程 1 加载 a[1]，等等，这可能是合并加载的典型示例。

因此，在您的第一种情况下，由于 a[] 数组都是连续且对齐的，并且 a[0..31] 适合全局内存中的 128 字节对齐区域，因此我们得到了合并负载。线程 0 读取 a[0]，线程 1 读取 a[1] 等等。

在第二种情况下，a[0] 与 a[1] 不连续，而且元素 a[0..31]（它们都在同一代码行中加载）不适合 128 字节对齐的序列全局内存。我将让您解析在第三种情况下发生的情况，但只要说像第二种情况一样，元素 a[0..31] 既不连续也不包含在全局内存中的单个 128 字节对齐区域中. 虽然不需要有连续的数据项来实现某种程度的合并，但 100% 的带宽利用率（“完美”）来自 32 个线程扭曲的合并负载意味着每个线程都使用一个唯一的 32 位项，所有这些都是连续并包含在全局内存中的单个 128 字节对齐序列中。

一个方便的心智模型是对比结构阵列 (AoS)（对应于您的案例 2 和 3）和阵列结构 (SoA)，这基本上是您的第一个案例。SoA 通常比 AoS 提供更好的合并可能性。在nvidia 网络研讨会页面上 ，您可能会发现此演示文稿很有趣，尤其是幻灯片 11-22 左右。

于 2012-10-22T21:33:19.937 回答
0

最佳实践指南中的其他一些相关信息：

对于计算能力为 2.x 的设备，可以很容易地总结出要求：warp 线程的并发访问将合并成多个事务，这些事务等于为 warp 的所有线程提供服务所需的缓存行数. 默认情况下，所有访问都通过 L1 进行缓存，L1 为 128 字节行。对于分散的访问模式，为了减少过度读取，有时只在 L2 中缓存会很有用，它缓存较短的 32 字节段（请参阅 CUDA C 编程指南）。

编译器标志：-Xptxas -dlcm=cg 将禁用 L1 缓存。即仅使用 L2，用于合并不良的数据。

于 2012-10-24T21:05:50.900 回答