连续 1 字节写入全局内存导致多个事务

Question

我正在做一个项目，其中每个线程都将 1 字节的值写入全局内存，并且我正在尝试最小化项目中的全局内存写入延迟。

在第 5.3.2 节中。 CUDA C 编程指南（link）它说：

Global memory resides in device memory and device memory is accessed via 32-, 64-, or 128-byte memory transactions. These memory transactions must be naturally aligned: Only the 32-, 64-, or 128-byte segments of device memory that are aligned to their size (i.e., whose first address is a multiple of their size) can be read or written by memory transactions.

所以我认为应该用一次事务处理对全局内存的连续 1 字节写入，因为它们已正确对齐。

但是当我在 Visual Studio 中使用 Nsight 分析以下最小示例时，尽管 4 个线程访问连续的 1 字节地址需要 1 个事务（如预期的那样），但在 5 个线程的情况下，它需要2 笔交易。

__global__ void copyKernel(const unsigned char* a, unsigned char* b)
{
    int i = threadIdx.x;
    a[i] = b[i];
}

int main()
{
    char *d_a;
    char *d_b;

    // ... (stuffs like cudaMalloc)

    // to check that the address is aligned
    printf("%p\n", d_a); // aligned to 512-Byte
    printf("%p\n", d_b); // aligned to 512-Byte

    // copy 4 elements
    copyKernel<<<1, 4>>>(d_a, d_b);

    // copy 5 elements
    copyKernel<<<1, 5>>>(d_a, d_b);

    // ...
}

分析结果如下。（左 - 4 个线程/右 - 5 个线程）

我在这里错过了什么？我应该如何编写代码以使其在一个事务中执行写入？

环境：Windows10，Visual Studio2015，GeForce GTX 1080 (cc 6.1)

Answer 1

看来我是在看错误的实验结果。 Nsight 为 "Profile CUDA Application" 提供了一些实验，问题中发布的图像来自 "Memory Statistics - Global" 实验的结果。根据 Nsight 用户指南，"Global" 实验报告了以下数据：

The Transactions Per Request chart shows the average number of L1 transactions required per executed global memory instruction, separately for load and store operations.

因此 "Global" 实验中显示的写入事务数实际上是写入到 L1 缓存，而不是写入到 L2。（虽然在Nsight中UI，说是到L2。）

另一方面，"Memory Statistics - Caches" 似乎显示了 L2 交易的数量，其中包含与我正在寻找的内容更相关的数据。那里的数字与 Robert Crovella 评论的相同。

1M线程的测试结果：

更新

L2交易好像是32字节粒度的。查看 4 字节连续存储的分析结果，为 1M 线程报告的 L2 存储事务数为 131,072，等于 1M(#threads) 乘以 4（数据大小）除以 32。

所以我得出结论，我的问题中引用的 "device memory can be accessed via 128-byte transaction" 无法通过 Nsight 验证，因为它似乎不计算 L2 和设备内存之间的事务。 (test code)

连续 1 字节写入全局内存导致多个事务

Consecutive 1 byte write to global memory results in multiple transactions

memory

cuda

nsight