问题标签 [gpu-constant-memory]

我正在尝试将我的模拟参数设置在常量内存中，但没有运气（CUDA.NET）。cudaMemcpyToSymbol 函数返回 cudaErrorInvalidSymbol。cudaMemcpyToSymbol 中的第一个参数是字符串...是符号名称吗？实际上我不明白它是如何解决的。任何帮助表示赞赏。

我的 .cu 文件包含

工作解决方案

我自己无法弄清楚，确保内核中使用的内存保持不变的最佳方法是什么。http://stackoverflow...r-pleasant-way有一个类似的问题。我正在使用 GTX580 并且仅针对 2.0 功能进行编译。我的内核看起来像

我在主机上执行以下代码：

另一种方法是添加

到 .cu 文件，从内核中删除src指针并执行

这两种方式是等价的还是第一种不保证使用常量内存而不是全局内存？大小动态变化，所以第二种方式在我的情况下并不方便。

我想读取一个文本文件并将其存储在一个数组中。然后，我想将数组从主机传输到设备并将其存储在共享内存中。我编写了以下代码，但与使用全局内存相比，执行时间有所增加。我无法理解原因可能是什么？此外，如果有人可以帮助我使用常量内存编写这段代码，那就太好了。

根据“CUDA C Programming Guide”，只有在命中多处理器常量缓存时，常量内存访问才会受益（第 5.3.2.4 节）¹。否则，半扭曲的内存请求可能比合并全局内存读取的情况更多。那么为什么恒定的内存大小限制为 64 KB？

为了不问两次，再问一个问题。据我了解，在 Fermi 架构中，纹理缓存与 L2 缓存相结合。纹理使用是否仍然有意义或全局内存读取以相同的方式缓存？

¹常量内存（第 5.3.2.4 节）

常量内存空间驻留在设备内存中，并缓存在 F.3.1 和 F.4.1 节中提到的常量缓存中。

对于计算能力为 1.x 的设备，对 warp 的恒定内存请求首先被拆分为两个请求，一个用于每个半 warp，它们是独立发出的。

然后，一个请求被拆分为与初始请求中不同的内存地址一样多的单独请求，从而将吞吐量降低等于单独请求数量的因子。

然后在缓存命中的情况下以常量缓存的吞吐量为结果请求提供服务，否则以设备内存的吞吐量提供服务。

我阅读了一些引用本地内存的 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 *。）

当我们使用 xptxas 检查寄存器使用情况时，我们会看到如下内容：

我想知道目前是否有任何文档可以清楚地解释 cmem[x]。将常量内存分成多个bank有什么意义，总共有多少个bank，除了0、2、14、16之外的其他bank有什么用？

作为旁注，@njuffa（特别感谢你）之前在 nvidia 的论坛上解释了什么是银行 0、2、14、16：

使用的常量内存被划分为常量程序“变量”（bank 1），以及编译器生成的常量（bank 14）。

cmem[0]:内核参数

cmem[2]:用户定义的常量对象

cmem[16]：编译器生成的常量（其中一些可能对应源代码中的文字常量）

我正在尝试了解我的每个 CUDA 线程的资源使用情况，以用于手写内核。

我将我kernel.cu的文件编译成一个kernel.o文件nvcc -arch=sm_20 -ptxas-options=-v

我得到以下输出（通过c++filt）：

看上面的输出，这样说对吗

• 每个 CUDA 线程使用 46 个寄存器？
• 没有寄存器溢出到本地内存？

我在理解输出方面也有一些问题。

• 我的内核正在调用很多__device__函数。__global__和__device__函数的堆栈帧的内存总和是 72 字节吗？

• 0 byte spill stores和有什么区别0 bytes spill loads

• 为什么信息cmem（我假设是恒定记忆）用不同的数字重复两次？在内核中，我没有使用任何常量内存。这是否意味着编译器会在后台告诉 GPU 使用一些常量内存？

_{这个问题在以下内容中“继续”：解释 ptxas 的详细输出，第二部分}

有没有办法在二维中分配和复制 CUDA 常量内存区域？cudaMemcpyToSymnbol__constant__ 似乎不是一个选择。

我在常量内存中有一个数组（它是一个全局变量），并通过函数调用 cudaGetSymbolAddress 获得了对它的引用。当我使用这个引用来获取常量数据而不是使用全局变量时，我的内核运行缓慢。这是什么原因？

这是示例代码，warp 中的所有线程同时加载相同的位置。注释代码是通过直接访问常量内存

解释为什么不使用常量内存缓存（由talonmies）

原因是缺少常量缓存。仅当编译器在显式标记为处于常量状态空间中的变量上发出特定的 PTX 指令 (ld.const) 时，才会发生缓存访问。编译器知道这样做的方式是在声明变量时__constant__——它是一个影响代码生成的静态编译时属性。相同的过程不能在运行时发生。

如果您在全局内存中传递一个指针并且编译器无法确定该指针在常量状态空间中，则它不会生成正确的 PTX 以通过常量缓存访问该内存。结果访问速度会变慢。

没有回答的问题

为什么即使将数组g声明为__device__变量，使用引用它时代码也会变慢。通过查看PTX代码，将全局内存加载到寄存器：

• 使用了 2 条指令ld.global.s32，将 4 个字节加载到寄存器中。（在使用参考的代码中）
• 使用1 条指令ld.global.v2.s32，将 8 个字节加载到 2 个寄存器，（在使用全局变量的代码中）

有什么区别，任何文档参考将不胜感激？

问题

我使用常量内存准备了一个示例 CUDA 代码。我可以在 cuda 4.2 中成功运行它，但是 当我使用 CUDA 5 编译时出现“无效的设备符号” 。我已在此处附加示例代码。

编码

我已经尝试过WINDOWS：CUDA 5.0 Production Release，显卡是 GTX 590。
任何帮助将不胜感激。