几年前，NVIDIA 的 Mark Harris 发布了这个：

CUDA C/C++ 中的高效矩阵转置

他在其中描述了如何使用共享内存而不是简单的方法更快地执行矩阵转置。出于方法论的目的，他还实现了基于共享内存块的简单矩阵副本版本。

有点令人惊讶的是，通过共享内存块进行复制的速度比“原始”复制（使用 2D 网格）要快：原始复制为 136 GB/秒，基于共享内存块的复制为 152.3 GB/秒。那是在 Kepler 微架构卡 Tesla K20c 上。

我的问题：为什么这有意义？也就是说，当所有所做的都是合并读取和写入时，为什么有效带宽没有降低？具体来说，它是否与 __restrict未使用（因此__ldg()可能未使用）的事实有关？

注意：这个问题不是关于换位的。这篇文章是关于换位的，它的教训很好。它没有讨论涉及简单、非转置复制的奇怪现象。

我有一个内存带宽有限的问题——我需要从 RAM 中顺序读取大量（许多 GB）数据，进行一些快速处理并将其顺序写入 RAM 中的不同位置。内存延迟不是问题。

在不同 NUMA 区域中的两个或多个内核之间分配工作有什么好处？同样，跨区域工作是否会减少可用带宽？

我想将一个小次数的多项式 (2-5) 应用于长度可以在 50 到 3000 之间的向量，并尽可能高效地执行此操作。示例：例如，我们可以取函数：(1+x^2)^3，当x>3时，当x<=3时为0。对于双元素向量，这样的函数将执行 100k 次。每个向量的大小可以在 50 到 3000 之间。

一种想法是使用 Eigen： Eigen::ArrayXd v; 然后简单地应用一个仿函数： v.unaryExpr([&](double x) {return x>3 ? std::pow((1+x*x), 3.00) : 0.00;});

尝试使用 GCC 9 和 GCC 10，我发现这个循环没有被矢量化。我确实手动对其进行了矢量化，结果发现增益比我预期的要小得多（1.5 倍）。我还用逻辑 AND 指令替换了条件，基本上执行两个分支并在 x<=3 时将结果归零。我认为收益主要来自缺乏分支错误预测。

一些考虑 因素有多种因素在起作用。首先，我的代码中有 RAW 依赖项（使用内在函数）。我不确定这如何影响计算。我用 AVX2 编写了我的代码，所以我期待 4 倍的增益。我认为这起到了一定的作用，但我不能确定，因为 CPU 有无序处理。另一个问题是我不确定我尝试编写的循环的性能是否受内存带宽的限制。

问题 如何确定内存带宽或管道危害是否影响此循环的实现？我在哪里可以学习更好地矢量化这个循环的技术？Eigenr MSVC 或 Linux 中是否有用于此目的的好工具？我使用的是 AMD CPU，而不是 Intel。

概括：

我正在尝试编写一个内存绑定 OpenCL 程序，该程序接近我 GPU 上宣传的内存带宽。实际上，我偏离了约 50 倍。

设置：

我只有一张比较旧的 Polaris Card (RX580)，所以我不能使用 CUDA，现在只能选择 OpenCL。我知道这是次优的，我无法让任何调试/性能计数器工作，但这就是我所拥有的。

我是 GPU 计算的新手，想感受一下我可以从 GPU 与 CPU 中获得的一些性能。对我来说首先要做的是内存带宽。

我编写了一个非常小的 OpenCL 内核，它从跨步内存位置读取，我希望波前中的所有工作人员一起在一个大内存段上执行连续内存访问，合并访问。然后内核对加载的数据所做的所有事情就是将这些值相加，并在最后将和写回另一个内存位置。代码（大部分是我从各种来源无耻地复制在一起的）非常简单

我生成iNrOfWorkers这些内核并测量它们完成处理所需的时间。对于我的测试，我设置iNrOfWorkers = 1024和iNrOfIterations = 64*1024. 从处理时间和iMemorySize = iNrOfWorkers * iNrOfIterations * sizeof(float)我计算出大约 5GByte/s 的内存带宽。

期望：

我的问题是内存访问似乎比我被认为可以使用的 256GByte/s 慢一到两个数量级。

GCN ISA 手册 [1] 让我假设我有 36 个 CU，每个 CU 包含 4 个 SIMD 单元，每个单元处理 16 个元素的向量。因此，我应该有 36 4 16 = 2304 个可用的处理元素。

我产生的数量少于这个数量，即 1024 个全局工作单元（“线程”）。线程按顺序访问内存位置，相隔 1024 个位置，因此在循环的每次迭代中，整个波前访问 1024 个连续元素。因此，我相信 GPU 应该能够产生连续的内存地址访问，并且中间没有中断。

我的猜测是，它只产生很少的线程而不是 1024，可能每个 CU 一个？这样一来，它就必须一遍又一遍地重新读取数据。不过，我不知道如何验证这一点。

[1] http://developer.amd.com/wordpress/media/2013/12/AMD_GCN3_Instruction_Set_Architecture_rev1.1.pdf

我正在尝试使用 SIMD 指令（_mm256_add_pd、存储、加载等）优化 C 中的二维矩阵加法。但是，我根本没有看到很大的加速。使用一些时序代码，我看到在 0.8x-1.5x 范围内的加速是天真的解决方案）。我想知道这是否是典型的？我在想这可能是一个内存瓶颈，因为在这种情况下计算似乎很少。我相信这应该会给我大约 4 倍的速度提升，因为我一次要进行 4 次加法，所以我不完全确定瓶颈是什么。

我编写了一些代码来演示我在做什么（测试并行 + SIMD 与仅 SIMD）：

我注意到结果似乎完全取决于 LOOP_COUNT。高循环数的并行/SIMD 版本做得很好，但低循环数的幼稚解决方案往往会做得更好。

我创建了一个简单的测试，该测试使用 SIMD 一次添加一个数组的 4 个元素（求和），而不是用 4 个求和变量累加它并在最后将它们相加。这是我的测试用例代码：

我预计 SIMD 会明显更快（一次添加 4 个），但事实并非如此。我想知道是否有人可以指出我为什么会这样？我运行代码得到的结果是：

A: 1249999975000000.000000, B: 1249999975000000.000000

时间一：0.317978，时间二：0.207965

展开速度快 1.528996 倍

我正在编译没有优化，gcc 选项是 gcc -fopenmp -mavx -mfma -pthread

我一直在查看一些 stackoverflow 帖子（为什么 std::fill(0) 比 std::fill(1) 慢？和增强的 REP MOVSB for memcpy）和一个决定最佳 memcpy/memset 策略的因素似乎是操作是否会受到延迟或 DRAM 带宽限制。其中一点是rep movsb切换延迟比我不理解的正常写入要长。

为什么ERMSB rep movsb切换延迟比movapsmemcpy/memset 的（或任何其他正常写入）循环更长

BeeOnRope 评论写道：

上面描述的 rep movsb 与跨各种缓冲区大小的单个核心上的 movaps 显式循环的行为与我们之前在服务器核心上看到的非常一致。正如您所指出的，竞争是在非 RFO 协议 [Read For Ownership] 和 RFO 协议之间进行的。前者在所有缓存级别之间使用较少的带宽，但特别是在服务器芯片上，一直到内存的切换延迟很长。由于单核通常具有并发限制，因此延迟很重要，并且非 RFO 协议胜出，这就是您在 30 MB L3 之外的区域中看到的情况

然而，在增强的 REP MOVSB for memcpy中，BeeOnRope 说

但是，如果您受到并发限制，那么情况会平衡，有时甚至会逆转。您有空闲的 DRAM 带宽，因此 NT 存储无济于事，它们甚至会受到伤害，因为它们可能会增加延迟，因为线缓冲区的切换时间可能比预取将 RFO 线带入 LLC（甚至L2) 然后存储在 LLC 中完成以有效降低延迟。最后，服务器非内核的 NT 存储往往比客户端的（和高带宽）慢得多，这突出了这种效果。

我无法理解非 RFO 方法rep movsb（动态随机存取存储器。

用于 memcpy的增强型 REP MOVSB帖子讨论了rep movsbon of 的优点：

立即准确地发出预取​​请求。硬件预取在检测类似 memcpy 的模式方面做得很好，但它仍然需要几次读取才能启动，并且会“过度预取”复制区域末尾之外的许多缓存行。rep movsb 准确地知道区域大小并且可以准确地预取。

鉴于这rep movsb是预取（比movaps循环更有效），与循环相比，您不会期望 LFB 能够以至少 L2/LLC 切换到线路的可能性更高（或相等）movaps。如果是这种情况，我不明白：

前者在所有缓存级别之间使用较少的带宽，但特别是在服务器芯片上，一直到内存的切换延迟很长

特别是长延迟切换是否来自。

所以我的问题是

1. LFB 上的额外切换延迟从何rep movsb而来？
2. 更一般地说，是什么导致了rep movsbmemcpy/memset 中的延迟界限？

我在大步访问时读过

两个循环的执行应该相似，因为内存访问的顺序高于乘法。

我正在玩谷歌基准测试，在测试类似的缓存行为时，我得到了我不理解的结果。

我希望访问字节数组比整数数组更快，因为更多元素适合缓存行，但事实并非如此。

以下是启用优化的结果：

任何人都可以澄清这一点吗？谢谢 ：）

更新 1：

我已经阅读了旧文章“程序员应该了解的内存”，现在一切都更加清楚了。但是，我尝试了以下基准：

当工作大小不适合缓存时，我预计随机访问的性能会更差。然而，这些是结果：

我错过了什么？

更新 2：

我现在正在使用您建议的 (linear_congruential_engine) 来生成随机数，而且我只使用静态数组，但结果现在让我更加困惑。

这是更新的代码：

以下是结果（已启用优化）：

(L1d < workingSet < L2) 结果与 (workingSet < L1d) 的结果有何不同？L2 的吞吐量和延迟仍然​​非常高，但是通过随机访问，我试图防止预取和强制缓存未命中.. 那么，为什么我什至没有注意到最小增量？

即使在尝试从主内存（workingSet > L3）中获取数据时，我的性能也没有大幅下降。你提到最新的架构可以保持高达每时钟 8 字节的带宽，但我知道他们必须复制一个保持缓存行，并且如果没有使用可预测的线性模式预取，延迟应该在我的测试中更加明显......为什么是不是这样？

我怀疑页面错误和 tlb 也可能有问题。

（我已经下载了 vtune 分析器来尝试更好地理解所有这些东西，但它挂在我的机器上，我一直在等待支持）

我真的很感谢你的帮助 Peter Cordes :)

我只是一个游戏程序员，试图向我的队友展示在我们的代码中使用某些整数类型是否会（或不会）对我们的游戏性能产生影响。例如，我们是否应该担心使用快速类型（例如 int_fast16_t）或在变量中使用尽可能少的字节以便更好地打包（例如 int8_t）。

如何使用perfand测量应用程序的内存带宽mpirun？我想知道这个应用程序是否受内存带宽限制。

我正在创建一个具有 1024 * 1024 * 1024 元素的 int（32 位）向量，如下所示：

那时它拥有 4 GB 的随机数据。然后我只是简单地总结了向量中的所有元素，如下所示：

这大约需要 0.18 秒，这意味着数据的处理速度约为 22.2 GB/s。我在 M1 上运行它，内存带宽更高，约为 60GB/s。有没有办法让上面的代码在单核上运行得更快？

编辑：手动 SIMD 版本：

SIMD 版本与非手动 SIMD 版本具有相同的性能。

编辑 2：好的，所以我将向量元素更改为 uint32_t 并将结果类型更改为 uint32_t（如@Peter Cordes 所建议）：

这运行得更快（~45 GB/s）。这是反汇编：

我还重写了 Manual-SIMD 版本：

它仍然比非手动 SIMD 版本慢 2 倍，并导致以下反汇编：

为了达到与自动矢量化版本相同的速度，我们可以在手动 SIMD 版本中使用 uint32x4x2 而不是 uint32x4：

为了获得更快的速度，我们可以利用 uint32x4x4（这让我们获得大约 ~53 GB/s）：

这让我们得到以下反汇编：

疯狂的事情