问题标签 [gpu-local-memory]

我只发现本地内存比寄存器内存慢的说法，即每线程两个类型。

共享内存应该是快的，但它比本地内存（线程的）快吗？

我想做的是一种中值过滤器，但使用给定的百分位数而不是中值。因此，我需要获取列表中的大部分内容，对它们进行排序，然后选择一个合适的。但我无法开始对共享内存列表进行排序，否则会出错。仅复制到本地内存会损失很多性能吗？

假设，我在 CUDA 内核函数中为每个线程声明了一个局部变量：

还假设声明的变量由编译器放置到本地内存（这与全局内存相同，但据我所知，它仅对一个线程可见）。f我的问题是在阅读时会合并访问吗？

我正在尝试用 CUDA 开发一个小程序，但由于它很慢，我做了一些测试并用谷歌搜索了一下。我发现虽然单个变量默认存储在本地线程内存中，但数组通常不是。我想这就是为什么执行需要这么多时间的原因。现在我想知道：由于本地线程内存应该至少为 16KB，而且我的数组只有 52 个字符长，有没有办法（请使用语法 :)）将它们存储在本地内存中？

不应该是这样的：

我阅读了一些引用本地内存的 CUDA 文档。（主要是早期文档。）设备属性报告本地内存大小（每个线程）。“本地”内存是什么意思？什么是“本地”内存？“本地”内存在哪里？如何访问“本地”内存？是__device__记忆，不是吗？

设备属性还报告：全局、共享和恒定的内存大小。这些陈述是否正确： 全局内存就是__device__内存。它具有网格范围和网格（内核）的生命周期。 不变的记忆就是__device__ __constant__记忆。它具有网格范围和网格（内核）的生命周期。 共享内存就是__device__ __shared__内存。它具有单个块范围和该块（线程）的生命周期。

我认为共享内存是 SM 内存。即只有那个单一的SM可以直接访问的内存。相当有限的资源。SM不是一次分配一堆块吗？这是否意味着 SM 可以交错执行不同的块（或不能）？即运行块* A * 线程，直到它们停止。然后运行 ​​block* B * 线程直到它们停止。然后再换回 block* A * 线程。或 SM 是否为 block* A * 运行一组线程，直到它们停止。然后换入另一组块* A * 线程。此交换继续进行，直到块* A * 用尽。然后才开始在块* B上工作*。我问是因为共享内存。如果单个 SM 正在从 2 个不同的块交换代码，那么 SM 如何快速换入/换出共享内存块？（我认为后面的 senerio 是真的，并且没有换入/换出共享内存空间。块 * A * 运行直到完成，然后块 * B * 开始执行。注意：块 * A * 可能是不同的内核比块* B *。）

当谈到 CUDA 中的共享/本地内存时，我仍然有点不确定。目前我有一个内核，在内核中每个线程分配一个列表对象。像这样的东西

根据我目前的理解，每个线程都有自己的dlist存储在本地内存中，这是真的吗？如果是这种情况，在内核执行结束时是否有任何方法来获取每个dlist对象（来自另一个内核），或者我应该使用__shared__由第一个线程分配的动态列表数组？

我想我可能有点过于复杂了，但我从来不需要更改列表，我试图实现的执行是这样的

1. 创建列表（仅在 GPU 上完成）
2. 从每个列表生成输出（在 GPU 上由每个线程完成，只需要为该线程分配的列表中的信息。）
3. 修改/交换列表（仍然在 GPU 上完成）
4. 重复 2 和 3，直到在主机上满足某些中断条件

提前致谢！

我有一个内核，它使用大量寄存器并将它们大量溢出到本地内存中。

由于溢出似乎相当高，我相信它可以通过 L1 甚至 L2 缓存。由于本地内存对每个线程都是私有的，编译器如何合并对本地内存的访问？该内存是否像全局内存一样以 128 字节事务读取？有了这么多的溢出，我的内存带宽利用率很低（50%）。我有类似的内核，没有溢出，可以获得高达 80% 的峰值内存带宽。

编辑我使用该nvprof工具提取了更多指标。如果我很好地理解了这里提到的技术，那么由于寄存器溢出（4 * l1 命中和未命中 / L2 = 的 4 个扇区的所有写入总和），我会产生大量内存流量(4 * (45936 + 4278911)) / (5425005 + 5430832 + 5442361 + 5429185) = 79.6%。有人可以验证我是否在这里吗？

编辑

我意识到这是 128 字节，而不是我想的位事务。

我在 CUDA 上写了一个简单的函数。它将图像大小调整为两倍。对于 1920*1080 的图像，此功能需要 ~20ms 才能完成。我尝试了一些不同的方法来优化该功能。而且我发现可能是本地内存是关键原因。

我尝试了三种不同的方法来获取图像。

• OpenCV中的GPU模块
• 纹理绑定到 OpenCV 中的 GpuMat
• 直接从全局内存中获取 GpuMat

他们都不能给我带来一点进步。

然后我使用nvvp找出原因。在上述所有三种情况下，本地内存开销约为 95%。

所以我转向我的代码来了解 nvcc 如何使用内存。然后我发现一个简单的函数就像这样：

需要 80 字节的堆栈帧（它们在本地内存中）。

还有一个像这样的功能：

还需要 88 字节的堆栈帧。

问题是，为什么我的函数在这个简单的任务中使用了如此多的本地内存和寄存器？为什么OpenCV中的函数可以通过不使用本地内存来执行相同的功能（这是nvvp测试，本地内存负载为零）？

我的代码是在调试模式下编译的。我的卡是 GT650(192 SP/SM, 2 SM)。

我在 NVIDIA 文档（http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications，表 #12）中读到每个线程的本地内存量我的 GPU 为 512 Ko（GTX 580，计算能力 2.0）。

我尝试使用 CUDA 6.5 在 Linux 上检查此限制，但未成功。

这是我使用的代码（它的唯一目的是测试本地内存限制，它不会进行任何有用的计算）：

和编译命令行：

编译正常，并报告：

当我达到每个线程 65000 字节时，我在 GTX 580 上运行时出现“内存不足”错误。这是控制台中程序的确切输出：

我还使用 GTX 770 GPU 进行了测试（在带有 CUDA 6.5 的 Linux 上）。MEMSIZE=200000 运行时没有错误，但 MEMSIZE=250000 在运行时出现“内存不足错误”。

如何解释这种行为？难道我做错了什么 ？

我正在阅读这份演示文档： http: //on-demand.gputechconf.com/gtc-express/2011/presentations/register_spilling.pdf

在演示文稿的第 3 页中，作者指出：

存储总是在加载之前发生——只有 GPU 线程可以访问 LMEM 地址

谁能向我解释为什么？他的意思是什么时候首次初始化本地内存？

我想使用本地/共享内存优化来减少全局内存访问，所以我基本上有这个功能

并对其进行通常的本地/共享内存优化：

现在的困难在于，test_optimized它仅在两个可能的 if/else 分支之一中被内核调用。如果工作组中只有一些线程执行 else-branch，则所有其他线程不得选择 if-branch 以使本地内存优化test_optimized起作用。所以我创建了一个解决方法：将工作组中每个线程的条件atomic_or-ed 转换为一个整数，然后检查对于所有线程都相同的整数是否有分支。这确保了，如果线程块中的 1 个或多个线程选择 else-branch，所有其他线程也会这样做。

上面的代码完美无缺，比test_unoptimized函数快 25%。但是对我来说，从工作组中的所有线程以原子方式将一点干扰到相同的本地内存中似乎有点像 hack，它只对小型工作组大小 ( def_ws) 32、64 或 128 有效运行，而不是 256 或更大。

这个技巧在其他代码中使用过吗？它有名字吗？如果没有：有更好的方法吗？