DAY56：阅读Dynamic Global Memory Allocation and Operations

B.20. Dynamic Global Memory Allocation and Operations

Dynamic global memory allocation and operations are only supported by devices of compute capability 2.x and higher.

void* malloc(size_t size);void free(void* ptr);

allocate and free memory dynamically from a fixed-size heap in global memory.

void* memcpy(void* dest, const void* src, size_t size);

copy size bytes from the memory location pointed by src to the memory location pointed by dest.

void* memset(void* ptr, int value, size_t size);

set size bytes of memory block pointed by ptr to value (interpreted as an unsigned char).

The CUDA in-kernel malloc() function allocates at least size bytes from the device heap and returns a pointer to the allocated memory or NULL if insufficient memory exists to fulfill the request. The returned pointer is guaranteed to be aligned to a 16-byte boundary.

The CUDA in-kernel free() function deallocates the memory pointed to by ptr, which must have been returned by a previous call to malloc(). If ptr is NULL, the call to free() is ignored. Repeated calls to free() with the same ptr has undefined behavior.

The memory allocated by a given CUDA thread via malloc() remains allocated for the lifetime of the CUDA context, or until it is explicitly released by a call to free(). It can be used by any other CUDA threads even from subsequent kernel launches. Any CUDA thread may free memory allocated by another thread, but care should be taken to ensure that the same pointer is not freed more than once.

B.20.1. Heap Memory Allocation

The device memory heap has a fixed size that must be specified before any program using malloc() or free() is loaded into the context. A default heap of eight megabytes is allocated if any program uses malloc() without explicitly specifying the heap size.

The following API functions get and set the heap size:

• cudaDeviceGetLimit(size_t* size, cudaLimitMallocHeapSize)
• cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)

The heap size granted will be at least size bytes. cuCtxGetLimit()and cudaDeviceGetLimit() return the currently requested heap size.

The actual memory allocation for the heap occurs when a module is loaded into the context, either explicitly via the CUDA driver API (see Module), or implicitly via the CUDA runtime API (see CUDA C Runtime). If the memory allocation fails, the module load will generate a CUDA_ERROR_SHARED_OBJECT_INIT_FAILED error.

Heap size cannot be changed once a module load has occurred and it does not resize dynamically according to need.

Memory reserved for the device heap is in addition to memory allocated through host-side CUDA API calls such as cudaMalloc().

B.20.2. Interoperability with Host Memory API

Memory allocated via device malloc() cannot be freed using the runtime (i.e., by calling any of the free memory functions from Device Memory).

Similarly, memory allocated via the runtime (i.e., by calling any of the memory allocation functions from Device Memory) cannot be freed via free().

In addition, device malloc() memory cannot be used in any runtime or driver API calls (i.e. cudaMemcpy, cudaMemset, etc).

B.20.3. Examples

B.20.3.1. Per Thread Allocation

The following code sample:

will output:

Notice how each thread encounters the malloc() and memset() commands and so receives and initializes its own allocation. (Exact pointer values will vary: these are illustrative.)

B.20.3.2. Per Thread Block Allocation

B.20.3.3. Allocation Persisting Between Kernel Launches

本文备注/经验分享：

本章节讲的说, GPU设备端调用的Heap管理函数(malloc()/free()), 以及, 两个辅助的memcpy和memset函数. 我们分别来看一下它们.

以前显存的分配, 只能从CPU端进行, 也就是之前我们说过的cudaMalloc*()之类的函数, 但是很多人感觉不方便---例如假设一些人需要从GPU端构造一个链表(Linked List),每个节点数据的存储需要动态的分配和插入. 此时能在GPU端进行直接分配, 从而能就地构造这个链标, 很多人感觉很方便很多.基于类似这种要求, NV分别两次引入了设备端的动态显存分配能力.(请注意, 这个例子不是一个恰当的例子, 以后说) 第一次的设备端动态分配能力的引入, 是在V2 API的时候, 也就是在CUDA 3.2时期(CUDA 3.2和CUDA 9.0一样, 也是一次重大的更新, 引入了大量的不兼容的改变. 这个之前说过).设备端的malloc()/free()就是那次引入的, 为当时的fermi架构.而第二次的设备端动态分配能力的引入, 则是为了后来的Kepler二代, 也就是计算能力3.5, 这时不仅仅有了malloc()在设备端的支持, 还能直接使用cudaMalloc(需要Device RT, 设备端运行时支持),这主要是为了能在动态并行的时候, 动态的分配global memory, 从而能在两次kernel启动之间传递数据.当然此时你也依然可以继续保持使用第一次引入的malloc()名字, 而不是cudaMalloc(), 但用后者能给你更多原本在CPU端启动kernel的感觉.需要说明的是, 这种用法也不是一个很好的例子. 以后再说(虽然不好, 但能用与否是个特性问题, 而用的好不好则是一个使用问题. 就像哪怕暴雨了, 城市里面的下水道虽然排水效果不好,但却不能缺少它, 本章节的动态分配是类似的). 不好的原因则是, 如同前几天章节里面的assert(), printf()一样, malloc()/free()也是被实现成所谓的"设备端系统调用",性能较差.而性能较差的原因则可能是(NV没说, 根据多年资料猜测): 每个系统调用均不是在GPU上执行的(N卡上面没有类似CPU之类的, 或者管理单元之类的结构), 而是通知Host端, 在CPU上完成的---例如这里的堆管理(Heap Management)相关函数.如果你不能知道heap是什么, 请阅读一下基本的<操作系统>或者<数据结构>课程, 这些都是从事本行业所必须的基本知识. 这里就不普及了.

除了缓慢的设备端的malloc()/free()之外, 本章节还提供了memset()函数, 和memcpy()函数.后面两个函数可以完成两个常见操作: 清空一段存储器, 或者从复制一段内容.幸运的是, 后面两个函数不是系统调用, 会被就地展开为一系列的指令操作(例如memcpy复制操作, 会被编译器一系列的读取--写入操作, 从而完成复制)，因此开销比前两个要好一些. 但又不幸的是, 这两个函数的实现均为每个线程各自为战的串行版本, 很多时候,性能远比你自己手工写的并行复制或者清0之类的代码要惨很多.(不仅仅是串行的原因, 还因为串行导致的访存问题--例如不合并或者shared memory上的bank conflict之类)，所以实际应用中, 总是推荐客户手工实现memcpy和memset，进行手工展开.实际上, 这个测试是当年在Fermi上进行的. 为了防止今天的CUDA 9.2有变化，例如做了优化, 我们进行了测试,测试结果很遗憾的表明, 都快8年过去了, 花儿还是那样的不红.所以你依然应当手工展开的, 而不应当调用它们(memset/memcpy)。

回到malloc上, 这个函数使用的是单独的一个设备端的堆(heap), 默认很小的, 本章节说明了如何扩大它.如果你不扩大它, 会导致一系列的问题, 实际上, 论坛的客户们遇到多次了: http://bbs.gpuworld.cn/thread-10601-1-1.html 还有 http://bbs.gpuworld.cn/thread-10708-1-1.html

这种问题论坛很多. 用户一定要仔细看本章节的手册.除了默认heap较小的问题外,设备端的malloc还存在一个问题, 它的分配粒度较小(粒度就是指的最小分配的大小)，这个大小在设备端和CPU端分配为80B(为何是这么一个奇特的数字, 用户可以自己想想)和512B(Pascal上），也就是同样你要分配1B, 从CPU上调用cudaMalloc()将总是会分配512B, 浪费掉511B，而从GPU端调用则会分配80B, 只浪费79B.所以你看, 设备端的malloc()是为了较小的分配而设计的.其实本章节说了, 设备端的分配是对齐到16B的边界的(忘记对齐到边界是什么意思了? 看看之前的文章里面的, 小区住房和巨人)，这里实际上是80B, 这80B = 16 * 5, 而为何不是喜闻乐见的64B这种2的整数次幂,则是因为很多情况下, malloc()在kernel里面被并行的调用, 每个线程都正好用64B这种而不是80B的边界,会造成一系列的问题, 例如对Cache的低效利用, 或者卡显存的某个bank之类的(NV的手册从来不讲这种问题,但是会默默的为你避免它, 例如这里的16 * 5;AMD的会很认真的讲解显存bank conflict的问题, 但却不会为你自动的从小细节默默避免它，所以这个各有利弊吧, CUDA毕竟是使用简易的, 这个细节也可见一斑)，以及, 实际上的生活中, 正常的CUDA C程序员都不会大量的使用本章的函数的,例如本章节的2个我举出的例子, 之前说它们不好, 完全可以规避这两个函数的使用的.例如前面说过的例子1中的Linked List的Node的分配和插入, 完全可以自行实现一个高效的分配(读者自己想)，例如前面的说过的例子2中的动态并行时候的kernel间的数据分配, 也完全可以提前准备好空间的, 而不是从GPU端现场分配.

关于本章节手册说的, 和Host端的CUDA Runtime API的问题, 请参考本章节手册说法.例如你不能从Host上释放一段Device分配的缓冲区.用户可以直接理解成, 存在2个不同的heap, 这两方面的函数, 分别是对这2个不同的heap操作的, 不具有互换性.此外, 用户需要注意的是, 当常见的并行malloc()或者new的时候,一个warp可能正好分配出来的N段缓冲区, 或者N个实例的首地址, 相差了规律的倍数关系(例如都相差80B)，但这个不能保证总是如此.如果用户需要这个特性, 应当直接warp整体分配一大段缓冲区, 然后再手工在warp内部规律的分配.

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

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