indexed constant和immediate constant的区别

  • 5 replies
  • 1134 views
indexed constant和immediate constant的区别
« 于: 十一月 29, 2021, 06:57:07 pm »
请问这两个有什么区别呢,目前我所知道的是,constant memory是off-chip memory,每个SM都有constant cache,是on-chip memory,这个图里面的 Immediate Constant Cache和 Indexed Constant Cache和上面提到的每个SM上的constant cache有什么关系呢?什么是Immediate Constant Cache?什么是 Indexed Constant Cache?

Re: indexed constant和immediate constant的区别
« 回复 #1 于: 十一月 29, 2021, 07:08:20 pm »
请问这两个有什么区别呢,目前我所知道的是,constant memory是off-chip memory,每个SM都有constant cache,是on-chip memory,这个图里面的 Immediate Constant Cache和 Indexed Constant Cache和上面提到的每个SM上的constant cache有什么关系呢?什么是Immediate Constant Cache?什么是 Indexed Constant Cache?
再细化下问题:
问题1:什么是Immediate Constant Cache?什么是 Indexed Constant Cache?能否举个例子,
问题2:这个图里面的 Immediate Constant Cache和 Indexed Constant Cache和上面提到的每个SM上的constant cache有什么关系呢?
问题3:核函数传输的参数存储在constant cache中嘛?

Re: indexed constant和immediate constant的区别
« 回复 #2 于: 十二月 11, 2021, 08:16:16 pm »
请问这两个有什么区别呢,目前我所知道的是,constant memory是off-chip memory,每个SM都有constant cache,是on-chip memory,这个图里面的 Immediate Constant Cache和 Indexed Constant Cache和上面提到的每个SM上的constant cache有什么关系呢?什么是Immediate Constant Cache?什么是 Indexed Constant Cache?

光看你的这个图,而不看其中的文字定义的话,则是常见的CUDA的SASS中的,constant memory中的内容作为操作数的用法。

一种是直接作为通用指令的操作数,例如FADD指令,其中的某个加数可能直接以c[0x4][0x100]类似这种的形式出现,看上去就类似立即数的效果。此时无需单独的从constant memory中加载,就可以直接使用。一般通用的指令仅限于1个操作数是来自于constant。

另外一种是作为专用指令的操作数,例如LDC指令,单独的从c[0x4][0x100]这种加载到某个寄存器,加载完毕后,再将这个寄存器作为操作数给通用的指令使用。后者往往见于某些通用指令的多个操作数来自constant,但是只能指定1个,另外一个提前用LDC指令加载到寄存器。

关于你的文字中的两种文字定义,我暂时不清楚由来。

Re: indexed constant和immediate constant的区别
« 回复 #3 于: 十二月 14, 2021, 09:59:15 am »
根据您的说法,FADD指令,其中的某个加数可能直接以c[0x4][0x100]类似这种的形式出现 是直接访问constant memory,这种是不是会速度比较慢?而LDC指令,单独的从c[0x4][0x100]这种加载到某个寄存器,加载完毕后,再将这个寄存器作为操作数给通用的指令使用这样子的话应该会比较快。
另外,我还看到《CUDA编程基础与实践》一书中的76页说到pascal架构开始编译器能够判断一个全局内存变量在这个核函数范围内都是只读的时候则会自动用函数__ldg()读取全局内存,从而对数据读取进行缓存,缓解非合并访问带来的影响。书后面还给出了使用该函数进行矩阵转置时的测试。,我没太想明白使用__ldg()读取全局内存时,对数据进行缓存是为什么可以缓解非合并访问带来的影响,矩阵转置是只读取一次数据的,都是要从全局内存加载进来,进行缓存了后面也没有用到,为什么会快呢?
当调用__ldg()读取全局内存,数据的传输途径和方法是什么样子的呢?和正常的直接读取全局内存的差别到底是在哪里呢?
希望您能解答一下

Re: indexed constant和immediate constant的区别
« 回复 #4 于: 十二月 14, 2021, 01:22:29 pm »
根据您的说法,FADD指令,其中的某个加数可能直接以c[0x4][0x100]类似这种的形式出现 是直接访问constant memory,这种是不是会速度比较慢?而LDC指令,单独的从c[0x4][0x100]这种加载到某个寄存器,加载完毕后,再将这个寄存器作为操作数给通用的指令使用这样子的话应该会比较快。
另外,我还看到《CUDA编程基础与实践》一书中的76页说到pascal架构开始编译器能够判断一个全局内存变量在这个核函数范围内都是只读的时候则会自动用函数__ldg()读取全局内存,从而对数据读取进行缓存,缓解非合并访问带来的影响。书后面还给出了使用该函数进行矩阵转置时的测试。,我没太想明白使用__ldg()读取全局内存时,对数据进行缓存是为什么可以缓解非合并访问带来的影响,矩阵转置是只读取一次数据的,都是要从全局内存加载进来,进行缓存了后面也没有用到,为什么会快呢?
当调用__ldg()读取全局内存,数据的传输途径和方法是什么样子的呢?和正常的直接读取全局内存的差别到底是在哪里呢?
希望您能解答一下

你的两个问题:

(1)如果某指令中的某操作数,使用了c[XX][XX] 的形式,硬件执行上的策略区别,这个无具体资料。但是考虑到我们的GPU是一个latency processor,设计于掩盖延迟的:所以c[XX][XX] 作为操作数,可能读取上的确比直接读取寄存器,延迟要高一点,但这种高一点点,可能并不会影响GPU的任何性能。
(《dissecting XXX via microbenchmarking》, 其中XXX是GPU架构名,这系列文章里面有讲述constant memory的各种情况下的延迟,感兴趣可以看一下)。

(2)关于Maxwell/Pascal(5.X & 6.X)上的L1/Tex cache的特性。这里简单说一下:很多时候,5.0,5.2和6.1的硬件,无法使用L1 cache缓存读写的数据,但是用__ldg()进行只读加载,可以有效的被L1缓冲,从而越过了很多场合的必须从L2加载的弊端,往往可以提升性能。在这两代的硬件上,使用const __restrict__修饰的指针,可以提示编译器自动进行这一点,但我不知道编译器能否像老樊那样说的,全自动的发现这点(因为你这里没有贴出老樊的书的原文)。

扩展阅读:分别查找 __ldg, const __restrict__, LDG.E.CI以获取更多信息。阅读手册的计算能力的5.X和6.X的章节,和自带的pascal tuning/compatiblity guide手册,获取Pascal在特定情况下,L1被自动禁用的章节(或者自动修复occupancy为0的BUG的章节),以获取更多信息。
« 最后编辑时间: 十二月 14, 2021, 01:23:45 pm 作者 屠戮人神 »

Re: indexed constant和immediate constant的区别
« 回复 #5 于: 十二月 14, 2021, 03:06:47 pm »
你的两个问题:

(1)如果某指令中的某操作数,使用了c[XX][XX] 的形式,硬件执行上的策略区别,这个无具体资料。但是考虑到我们的GPU是一个latency processor,设计于掩盖延迟的:所以c[XX][XX] 作为操作数,可能读取上的确比直接读取寄存器,延迟要高一点,但这种高一点点,可能并不会影响GPU的任何性能。
(《dissecting XXX via microbenchmarking》, 其中XXX是GPU架构名,这系列文章里面有讲述constant memory的各种情况下的延迟,感兴趣可以看一下)。

(2)关于Maxwell/Pascal(5.X & 6.X)上的L1/Tex cache的特性。这里简单说一下:很多时候,5.0,5.2和6.1的硬件,无法使用L1 cache缓存读写的数据,但是用__ldg()进行只读加载,可以有效的被L1缓冲,从而越过了很多场合的必须从L2加载的弊端,往往可以提升性能。在这两代的硬件上,使用const __restrict__修饰的指针,可以提示编译器自动进行这一点,但我不知道编译器能否像老樊那样说的,全自动的发现这点(因为你这里没有贴出老樊的书的原文)。

扩展阅读:分别查找 __ldg, const __restrict__, LDG.E.CI以获取更多信息。阅读手册的计算能力的5.X和6.X的章节,和自带的pascal tuning/compatiblity guide手册,获取Pascal在特定情况下,L1被自动禁用的章节(或者自动修复occupancy为0的BUG的章节),以获取更多信息。

在官方论坛https://forums.developer.nvidia.com/t/ldg-versus-textures/31383找到了ldg指令的一些解释:
The LDG instruction is a global memory load that uses the texture path. It has the advantage that it does not require the explicit use of textures. Explicit uses of textures causes a certain amount of code clutter and overhead (e.g. for API calls to bind textures), and textures are objects unfamiliar to many programmers new to CUDA. The introduction of LDG therefore increases the ease of use.

Whether the use of LDG results in higher or lower performance compared to the use of classical textures depends on the use case, I am not aware of any hard-and-fast rule about that. When comparing the performance with LDG vs regular global loads, I have found that use of LDG results in higher performance in most cases, but recall at least one real-life use case where this was not the case.

As eyalhir74 points out, declaring pointers as “const restrict” facilitates, but does not guarantee, the generation of LDG instructions on Kepler-class GPUs. The reason for this is that the use of “const restrict” pointers makes assertions about local read-only behavior, whereas use of LDG requires the data to be read-only for the lifetime of a kernel. Only the use of the __ldg() device function ensures that LDG is generated.