问题标签 [gpu-shared-memory]

I want to write the following CUDA function:

On the host side, we have a slightly-related facility in the form of cudaPointerGetAttributes, which can tell us whether or not a pointer is to device memory or host memory; perhaps there's some way to distinguish pointers in device code as well, and perhaps it can also discern shared from global pointers. Alternatively, and perhaps even better - maybe there's a compile-time mechanism to do that, since, after all, the device functions are only compiled into kernels and are not freestanding, so nvcc can often know whether they're used with shared memory or not.

我有一个这样的数组：

我想使用 G80 GPU 上的共享内存来计算这个数组的减少。

NVIDIA 文档中引用的内核是这样的：

该论文的作者表示，这种方法存在银行冲突的问题。我试图理解，但我不知道为什么？我知道银行冲突和广播访问的定义，但仍然无法理解。

银行冲突

我有一个扭曲，它将一些数据写入共享内存 - 没有覆盖，并且在从共享内存读取后不久。虽然我的块中可能还有其他扭曲，但它们不会触及共享内存的任何部分或写入我感兴趣的扭曲读取的任何地方。

现在，我记得尽管 warp 以锁步执行，但我们不能保证共享内存写入之后的共享内存读取将返回应该由 warp 先前写入的相应值。（这在理论上可能是由于指令重新排序或 - 正如@RobertCrovella 指出的那样 - 编译器优化了共享内存访问）

所以，我们需要求助于一些显式的同步。显然，块级__syncthreads()工作。这是做什么的：

__syncthreads()用于协调同一块的线程之间的通信。当块中的某些线程访问共享或全局内存中的相同地址时，对于其中一些内存访问，可能存在先读后写、先读后写或先写后写的风险。通过在这些访问之间同步线程可以避免这些数据危害。

这对我的需求来说太强大了：

• 它也适用于全局内存，而不仅仅是共享内存。
• 它执行经线间同步；我只需要intra-warp。
• 它可以防止所有类型的危害R-after-W, W-after-R, W-after-W ; 我只需要R-after-W。
• 它也适用于多个线程执行写入共享内存中同一位置的情况；在我的情况下，所有共享内存写入都是不相交的。

另一方面，类似的东西__threadfence_block() 似乎还不够。这两个强度级别之间有什么“中间”吗？

笔记：

• 相关问题：CUDA__syncthreads()在 warp 中的使用。
• 如果您要建议我改用改组，那么，是的，有时这是可能的——但如果您想对数据进行数组访问，即动态决定您要读取的共享数据的哪个元素，则不行。这可能会溢出到本地内存中，这对我来说似乎很可怕。
• 我在想也许volatile对我有用，但我不确定使用它是否能达到我想要的效果。
• 如果您有一个假设计算机功能至少为 XX.YY 的答案，那就足够了。

我有 5 个大型数组 A(N*5)、B(N*5)、C(N*5)、D(N*5)、E(N*2) 数字 5 和 2 代表这些变量的组成部分在不同的平面/轴上。这就是我以这种方式构建数组的原因，因此我可以在编写代码时可视化数据。N ~ 200^3 ~ 8e06 个节点

例如：这是我的内核最简单的样子，我在全局内存上进行所有计算。

我知道可以消除“for”循环，但我把它留在这里，因为它方便查看代码。这可行，但显然即使在删除“for”循环之后，Tesla K40 卡的效率极低且速度很慢。“for”循环中显示的算术只是为了给出一个想法，实际的计算要长得多，并且与 res1、res2... 也混在一起。

我已经实现了以下改进，但我想通过共享内存的超载来进一步改进它。

这有点帮助，但我想实施其中一种减少方法（没有银行冲突）以提高性能，我可以将所有变量放入共享中（可能是平铺方法），然后进行计算部分。我在 CUDA_Sample 文件夹中看到了归约示例，但该示例仅适用于对共享中的一个向量求和，而无需对共享内存中的多个数组进行任何复杂的算术运算。对于改进我现有的 kernel_shared 方法以包括减少方法的任何帮助或建议，我将不胜感激。

假设我在一个 CUDA 块中有一个完整的线程扭曲，并且这些线程中的每一个都旨在与驻留在共享内存中的 T 类型的 N 个元素一起工作（所以我们总共有 warp_size * N = 32 N 个元素）。不同的线程从不访问彼此的数据。（嗯，他们这样做，但在稍后阶段，我们在这里不关心）。此访问将在循环中发生，如下所示：

现在，不同的线程可能每个都有不同的索引，或者相同 - 我不会以这种或那种方式做出任何假设。但我确实想尽量减少共享内存库冲突。

如果sizeof(T) == 4，那么这很容易：只需将线程 i 的所有数据放在共享内存地址 i、32+i、64+i、96+i 等中。这会将 i 的所有数据放在同一个 bank 中，这也是不同的从另一条车道的银行。伟大的。

但是现在——如果sizeof(T) == 8呢？我应该如何放置和访问我的数据，以尽量减少银行冲突（对指数一无所知）？

注意：假设 T 是普通旧数据。如果这使您的答案更简单，您甚至可以假设它是一个数字。

假设（CUDA 内核网格）块中的许多扭曲正在重复更新相当数量的共享内存位置。

哪些情况下这样的工作会更快完成？：

1. 内部warp访问局部性的情况，例如每个warp访问的内存位置总数很少，其中大部分确实是由多个通道访问的
2. 访问反局部性的情况，所有车道通常都访问不同的位置（并且可能是为了避免银行冲突）？

同样重要的是——这种微架构是依赖于微架构的，还是在所有最近的 NVIDIA 微架构上都基本相同？

我了解银行、银行宽度和银行冲突。但我不明白多个静态分配的共享内存是如何通过银行布局的。

如果我假设： 4 字节的银行宽度 32 银行 sizeof(int) == 4 字节

那么shared1和shared2将如何在整个银行中布局呢？是否会在 shared1 之后立即放置 shared 2 （即 shared1 的最后一个值存储在 bank 0 中，而 shared2 的第一个值存储在 bank 1 中）？或者shared1的末尾会被填充（即shared1的最后一个值存储在bank 0中，bank 1~31被填充，shared 2从下一行的bank 0开始）？

如果 shared2 在嵌套函数中分配，这种行为会有所不同吗？：

最后，声明的顺序会改变共享内存的布局吗？

在计算 Cuda 中的有效带宽时，我是否计算共享内存中的读/写次数。下面给出了一个示例代码。

以下主机代码test.c和设备代码test0.cu旨在提供相同的结果。

test.c

test0.cu

如果我编译并运行它们，它们会给出与我预期相同的结果。

但是，如果我在设备代码中使用共享内存而不是全局内存，如 中test1.cu所示，它会给出不同的结果。

test1.cu

如果我编译并运行它，它会给出与ortest1.cu不同的结果。test0.cutest.c

经纱同步不应该与共享内存一起使用吗？

对此问题的一些调查：

使用 CUDA8.0 时，如果我test1.cu使用-arch=sm_61选项进行编译（我正在使用 GTX 1080 进行测试），它会给出与test0.cuand相同的结果test.c。

但这不适用于较新版本的 CUDA。如果我使用任何比 8.0 更新的版本，即使我给出-arch=sm_61选项，测试结果也会有所不同。

在计算能力 <= 7.2 的设备上，我总是使用

nvprof --events shared_st_bank_conflict

但是当我使用 CUDA10 在 RTX2080ti 上运行它时，它会返回

Warning: Skipping profiling on device 0 since profiling is not supported on devices with compute capability greater than 7.2

那么如何检测此设备上是否存在共享内存库冲突？

我已经安装了 Nvidia Nsight Systems 和 Nsight Compute ，找不到这样的分析报告...

谢谢