DAY68:阅读 Memory Declarations

我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第68天，我们正在讲解CUDA C语法，希望在接下来的32天里，您可以学习到原汁原味的CUDA，同时能养成英文阅读的习惯。

关注微信公众号，查看历史信息，可以看到之前的阅读

本文共计448字，阅读时间12分钟

D.3.1.6. Memory Declarations

D.3.1.6.1. Device and Constant Memory

Memory declared at file scope with __device__ or __constant__ memory space specifiers behaves identically when using the device runtime. All kernels may read or write device variables, whether the kernel was initially launched by the host or device runtime. Equivalently, all kernels will have the same view of __constant__s as declared at the module scope.

D.3.1.6.2. Textures & Surfaces

CUDA supports dynamically created texture and surface objects1, where a texture reference may be created on the host, passed to a kernel, used by that kernel, and then destroyed from the host. The device runtime does not allow creation or destruction of texture or surface objects from within device code, but texture and surface objects created from the host may be used and passed around freely on the device. Regardless of where they are created, dynamically created texture objects are always valid and may be passed to child kernels from a parent.

Note: The device runtime does not support legacy module-scope (i.e., Fermi-style) textures and surfaces within a kernel launched from the device. Module-scope (legacy) textures may be created from the host and used in device code as for any kernel, but may only be used by a top-level kernel (i.e., the one which is launched from the host).

D.3.1.6.3. Shared Memory Variable Declarations

In CUDA C/C++ shared memory can be declared either as a statically sized file-scope or function-scoped variable, or as an extern variable with the size determined at runtime by the kernel's caller via a launch configuration argument. Both types of declarations are valid under the device runtime.

D.3.1.6.4. Symbol Addresses

Device-side symbols (i.e., those marked __device__) may be referenced from within a kernel simply via the & operator, as all global-scope device variables are in the kernel's visible address space. This also applies to __constant__ symbols, although in this case the pointer will reference read-only data.

Given that device-side symbols can be referenced directly, those CUDA runtime APIs which reference symbols (e.g., cudaMemcpyToSymbol() or cudaGetSymbolAddress()) are redundant and hence not supported by the device runtime. Note this implies that constant data cannot be altered from within a running kernel, even ahead of a child kernel launch, as references to __constant__ space are read-only.

本文备注/经验分享：

今天的章节主要是说, 各种存储器的具体使用注意事项. 请注意本章节存在诸多问题, 主要是历史原因.(这章节从CUDA 5到现在就没改过)。 首先是说, 用__device__声明的global memory上的变量和数组的使用.通过我们之前的global memory章节, 你知道global memory有两种, 一种是静态分配的, 一种是动态分配的.而__device__分配的属于静态分配的, 在CUDA Runtime API初始化的时候, 会自动为这种变量/数组分配显存.不需要手工的cudaMalloc*()的过程.这种静态分配的global memory上的变量和数组, 第一段落说明, 使用起来和普通的Host上cudaMalloc*()动态分配毫无区别(但需要注意一致性的问题, 一致性的问题在上次的章节中说过.). 因为实际上随着CUDA的演进, 动态分配的global memory也有过变化, 这里一并说一下: (A)动态分配global memory: (1)从Host上调用cudaMalloc*()系列函数. (2)从设备端调用malloc(), 需要计算能力2.0+, 不需要动态并行支持. (3)从设备端调用cudaMalloc(), 类似(2), 但需要动态并行支持. (B)静态分配, 通过__device__, 或者__constant__(后者将通过constant cache访问) 请注意因为这本手册基本讲述的是runtime api, 但有时候会不经意的引入driver api的概念, 例如需要注意__device__实际上是每模块的(driver api需要考虑同一个进程, 多个context多个模块的问题. 这个再说) (C)Unified memory, 因为本章节的历史原因, 这里没有提到. 实际上这个在动态并行里面, 和普通的__device__, 以及, cudaMalloc/malloc出来的一样.这是动态并行的时候说道__device__和__constant__需要注意的, 以及, 不建议任何时候在父kernel和任何它的子kernel中修改__constant__的内容(通过获取对应的后备显存指针进行修改). 这点和texture还是不同的. 然后这里来说到texture和surface,这里从动态并行启动的子kernel中可以使用它们, 但需要注意的是, 只能使用texture object和surface object,不能使用texture和surface reference。 精确一点讲, 动态并行的时候, 只能在最外层的kernel里使用(从CPU启动的那层), 使用texture和surface引用. 再深入一层或更多层(从GPU启动的后续层次), 则不能使用它们. 如果你真的要使用, 能正常通过编译. 但所有的纹理和表面读取出来的结果都是错误的, 表面写入的结果也是错误的. 如果想在非最外层次的父kernel中使用它们, 则必须使用texture和surface object对像， 这个可以在任意层次使用(包括最外层和里面的任意多层). 使用的时候, 将这两种object, 直接当成参数, 传递给多层的kernel即可（ 或者你认为的其他传递方式).

需要注意的是, 本章节说到texture和surface的部分存在若干问题.这主要是因为历史原因造成的。在当初推出动态并行的时候, 是CUDA 5，一并推出了surface和texture object(CUDA 5当初是Kepler时代, 对应计算能力3.X)。因为是刚出来的, 名字还不稳定, 人们也不熟悉object这种名字.所以这时候文档使用了多种叫法(见本章节)而不是直接叫object。 例如:Fermi Style Texture/Surface这指的是纹理引用和表面引用.再例如本章节还叫做: module scope或者file scope的texture/surface，这个实际上也是纹理和表面引用. 前者是用driver api的角度叫的. 后者是从编译的时候, 做为文件里的全局变量的角度叫的。再类似的, 这里的动态创建的纹理和表面, 实际上则是指的新的纹理和表面对像，这种才能再动态并行的子kernel里中.

本章节的主要问题在于没有直接的指出, 再动态并行的时候, 它们的准确指对. 会对读者造成干扰.特别是没有从老CUDA时代走过来的人, 往往会不知所云.但是实际上, 直接写点代码试验就知道指的是什么了.类似的, 这里的段落里的角标1,后面说的, 请参考CUDA Progamming Guide, 难道我们现在读的不是CUDA编程指南手册么?实际上这也是历史原因.最初的时候, 动态并行是单独的一本pdf，和手册是独立开的.所以我们这章实际上之前是在那本"动态并行手册"里,所以类似这里的让你看CUDA手册, 实际上现在你直接无视即可. 已经在一起了. 明白是reference不能用, object能用即可. 需要注意的是, NV的很多文档有这个问题, 一并说一下吧: 1)见到需要手工绑定, 解绑的都是指的reference; 2)见到不需要绑定(bindless)的, 指的是object；很多时候不直接出现reference和object字样的. 以及, (1)见到有限数量的纹理也是指的reference； (2)现在的纹理对像之类的, 并不存在数量限制. 看到一些资料说道纹理有限制之类的, 直接改成object就没有了. 类似这样的. 总之本章节都是CUDA 5时代的历史残留, 用户知道就好. 至于shared memory, 这个读者之前知道有静态分配的(通过__shared__)和动态分配的(通过第三个<<<>>>参数),这两种均可以无障碍的在动态并行启动的kernel里使用.并不存在特别的问题.但需要注意的是, 用户不能跨kernel传递shared memory指针,但是却可以传值.这点需要注意. 有一些技巧允许你这样传递使用, 但使用的时候一定要小心.对于新用户来说, 应当直接理解成不能这样做. 本章节也重新给了一个<<<>>>的例子, 用来说明动态分配的参数如何使用的问题. 读者也可以直接看之前的章节, Host上的"启动配置"章节, 这里的动态shared memory大小参数是一样的, 可以参考.最后是说明了如何对symbol进行地址获取,获取地址后往往可以用来复制或者填充一些初始化的值,以前在host上可以通过cudaMemcpyToSymbol进行直接复制到一个符号上,这里需要有几点注意的: (1)以前的cudaMemcpyToSymbol实际上等于分步的cudaGetSymbolAddress获取地址后, 再用普通的cudaMemcpy*()传输. 只是直接cudaMemcpyToSymbol用起来合并了步骤而已. (2)动态并行在设备端的时候, 因为可以直接用&符号取地址, 非常简单. 于是设备端的cuda Runtime api, 就取消了相关的函数,进行了简化, 但实际上并不妨碍你使用的.注意这里对__constant__的说明, 说是只读的, 但实际上依然是可以有技巧写入的, 但只是需要下次"从Host上"启动的kernel才能生效而已.

注意之前章节曾经提到过设备端的cudaMemcpy*Async(),这个实际上你总是可以配合&符号, 进行之前的cudaMemcpyToSymbol之类的操作的.只是在动态并行的时候, cudaMemcpy*Async()在设备端的调用, 很有可能被实现成一个子kernel(隐形的)，这是因为实际上从Host上进行显存到显存的传输复制, 也有可能会启动一个隐形的kernel(看不见的, profiler里也看不见).这是因为实际的显存传输, 有DMA和PIO的模式.直接以隐形的kernel启动一次显存到显存传输, 可以节省硬件设计成本,这种隐形的kernel等于是在PIO了.

有不明白的地方，请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖