2

我正在运行一些 GPU 基准测试来了解如何最大化来自/到全局内存的内存带宽。我有一个 128 MB（32*1024*1024 单精度浮点数）的数组，对齐到 128 字节的边距，在实际数据之前和之后有三个光环值。因此，数组中的第一个元素与 128 字节边界对齐。

在下文中，n指的是我的数组中的元素个数（不包括光环）：n = 32*1024*1024. m指的是数组中的 128 字节字：m = 1024*1024 = 1048576.

*array     // Aligned to a 128-bytes boundary
*(array-3) // Start of the (unaligned) halo region

我也有一个类似的输出数组，它对齐到相同的边界并且不包含光环。

我有几个内核，它们使用不同的访问模式实现了所有类似的计算：

P1: *(output+i) = *(array+i) // for i in 0..n
P2: *(output+i) = *(array+i) + *(array+i+1)
P3: *(output+i) = *(array+i-1) + *(array+i+1)

所有这些计算显然都受带宽限制。我正在尝试优化全局内存事务。我使用的代码非常简单：

__global__ void P1(const float* input, float* output)
{
    const int i = threadIdx.x + blockDim.x*blockIdx.x;
    *(output+i) = *(input+i);
}

__global__ void P2(const float* input, float* output)
{
    const int i = threadIdx.x + blockDim.x*blockIdx.x;
    *(output+i) = *(input+i) + *(input+i+1);
}

__global__ void P3(const float* input, float* output)
{
    const int i = threadIdx.x + blockDim.x*blockIdx.x;
    *(output+i) = *(input+i-1) + *(input+i+1);
}

我每个块有 1024 个线程和正确数量的块，这样每个线程都被分配了一个输出数组的值。

我使用缓存和非缓存选项 ( -Xptxas -dclm={ca,cg}) 进行了编译，并使用 nvprof 进行了基准测试，提取了以下指标：

  • ldst_issued：发布加载/存储指令
  • ldst_executed：执行的加载/存储指令
  • gld_transactions: 全局加载事务
  • gst_transactions: 全球店铺交易
  • dram_read_throughput：设备内存读取吞吐量
  • dram_write_throughput：设备内存写入吞吐量

我使用的 GPU 是 Nvidia K20X。

我希望ldst_executed是(k+1) * m，其中 k 对于 P1 是 1，对于 P2 是 2，对于 P3 是 3，并且表示每个线程读取的值的数量。我还希望gst_transactions（m合并访问：按 128 字节字写入）对于 P1，介于P2m和2m P2 之间m以及3mP3 之间的某个地方，因为每个 warp 都必须像 P1 一样访问其“分配”的内存部分，加上以下内容P2 的 128 个字节，加上 P3 的前 128 个字节，但我不确定 warp 是否是正确的单位。我期望一些线程能够避免全局内存访问，因为数据已经被前一个线程提取到 L1 缓存中。

这些是结果：

P1：

     gld_transactions   1048576
     gst_transactions   1048576
          ldst_issued   2097152
        ldst_executed   2097152
 dram_read_throughput   92.552 GB/s
dram_write_throughput   93.067 GB/s

P2：

     gld_transactions   3145728
     gst_transactions   1048576
          ldst_issued   5242880
        ldst_executed   3145728
 dram_read_throughput   80.748 GB/s
dram_write_throughput   79.878 GB/s

P3：

     gld_transactions   5242880
     gst_transactions   1048576
          ldst_issued   8052318
        ldst_executed   4194304
 dram_read_throughput   79.693 GB/s
dram_write_throughput   78.510 GB/s

我已经看到了一些差异：

  • 负载事务的数量从 P1 大幅增加到 P2 和 P3。
  • P2 和 P3 中发出的加载/存储指令的数量也非常高，超出了我的解释范围。我不确定我是否理解这个数字代表什么。

当我转向非缓存测试时，这些是结果

P1：

     gld_transactions   1048576
     gst_transactions   1048576
          ldst_issued   2097152
        ldst_executed   2097152
 dram_read_throughput   92.577 GB/s
dram_write_throughput   93.079 GB/s

P2：

     gld_transactions   3145728
     gst_transactions   1048576
          ldst_issued   5242880
        ldst_executed   3145728
 dram_read_throughput   80.857 GB/s
dram_write_throughput   79.959 GB/s

P3：

     gld_transactions   5242880
     gst_transactions   1048576
          ldst_issued   8053556
        ldst_executed   4194304
 dram_read_throughput   79.661 GB/s
dram_write_throughput   78.484 GB/s

如您所见，没有任何变化。我期待看到一些差异，因为在非缓存情况下，L1 缓存被丢弃，但事务以 32 字节字的形式发生。

问题：

  • 我的方法完全正确吗？
  • 共享内存可以帮助我减少传输量吗？
  • 为什么我看不到缓存和非缓存案例之间的实质性区别？
  • 为什么 P3 不比 P2 慢，就像 P2 比 P1 慢一样？
  • 还有哪些其他指标可以帮助我了解实际发生的情况？
4

0 回答 0