在不移动数据的情况下在 CUDA 中实现 realloc

Posted 2023-02-21

【中文标题】在不移动数据的情况下在 CUDA 中实现 realloc

【英文标题】：Implementing realloc in CUDA without moving data

【发布时间】：2014-10-15 16:03:38

【问题描述】：

根据this question和引用NVIDIA CUDA Programming Guide，realloc功能没有实现：

内核中的 CUDA 函数 malloc() 至少分配了 size 字节 从设备堆中返回一个指向已分配内存的指针或 如果没有足够的内存来满足请求，则为 NULL。这 返回的指针保证与 16 字节边界对齐。

内核中的 CUDA free() 函数释放指向的内存 由ptr 提供，它必须是由先前对malloc() 的调用返回的。 如果ptr 为NULL，则忽略对free() 的调用。反复调用 free() 与相同的 ptr 具有未定义的行为。

我现在被 GMP 库的某些部分困住了（或者更严格地说，我尝试将它移植到 CUDA 上），它依赖于这个功能：

__host__ __device__ static void * // generate this function for both CPU and GPU
gmp_default_realloc (void *old, size_t old_size, size_t new_size)

    mp_ptr p;

#if __CUDA_ARCH__ // this directive separates device and host code
    /* ? */
#else
    p = (mp_ptr) realloc (old, new_size); /* host code has realloc from glibc */
#endif

    if (!p)
        gmp_die("gmp_default_realoc: Virtual memory exhausted.");

    return p;

基本上我可以只需简单地用new_size 调用malloc，然后调用memcpy（或者可能是memmove），然后调用free 上一个块，但这需要强制移动数据（大型数组），我想避免。

是否有任何有效 有效方式来实现（标准 C 或 C++）realloc 函数（即内核内部）？假设我有一些动态分配的大数据数组（已经由 malloc 分配），然后在其他地方调用 realloc 以便为 那个 块请求更多的内存.简而言之，我想避免将整个数据数组复制到新位置，我特别询问如何做到这一点（当然，如果可能的话）。

我对 PTX ISA 或内核堆函数的底层实现并不是特别熟悉，但也许值得研究一下这个方向？

【问题讨论】：

realloc 将copy data in some cases。如果您的问题是我如何在没有数据副本的情况下实现realloc（任何地方），对于一般情况，我认为它无法完成。你的问题到底是什么？ 有效这个词并没有真正告诉我。换句话说，您的问题标题是：“”我会问您可以在主机上执行此操作吗？因为realloc 不保证这一点。

我的意思是当我已经有一些大数据块（更准确地说是数组，例如unsigned long 对象），然后realloc 用于获取更多内存。这只是artibratry精度数字的情况，无法确定需要多少内存。我知道 C99/C11 标准不保证数据会被保留，但通常情况下是这样。

如果您要求更大的内存，realloc 经常需要进行数据复制。我认为这种说法值得怀疑：“我知道 C99/C11 标准不能保证数据被保存，但通常情况下是这样的”。即使这是真的，在某些情况下，也不是所有情况都可以处理（即使在主机代码中）而不需要数据副本。因此，无论您是在谈论主机还是 GPU，我都怀疑您的问题是可能的（在没有数据副本的情况下实现 realloc）。

我再说一遍：在 c 或 c++ 标签上问这个问题：“如何在不移动数据的情况下实现 realloc？”无论您想出什么答案，都可能对创建 GPU 版本具有指导意义。我认为没有人可以在那里（或者）给你答案，但我可能是错的。

我认为你的观点是正确的。从malloc、calloc 或realloc 分配的块都必须连续，并且没有什么真正保证更大的块将“适合”到可用的可用空间中（这以同样的方式影响主机和设备内存）。

【参考方案1】：

大多数 malloc 实现过度分配，这就是 realloc 有时可以避免复制字节的原因 - 旧块可能足够大以容纳新大小。但显然在您的环境中系统 malloc 不会这样做，所以我认为您唯一的选择是在系统提供的 malloc/free 之上重新实现所有 3 个原语 gmp_default_alloc,realloc,free。

那里有许多开源 malloc 实现，glibc 有一个您可能能够适应。

我不熟悉 CUDA 或 GMP，但我不知道：

gmp_malloc() 后跟普通 free() 可能适用于“普通”平台，但如果继续这样做可能会导致堆损坏

1234563 .你甚至不需要一个成熟的堆实现。

您的实现可能需要使用互斥锁或类似的东西来保护您的堆免受并发修改

如果您从不（或不经常）将 malloc()ed 块从自定义堆中返回给操作系统，则可以进一步提高性能，即保留 gmp_free()ed 块以供后续重用，而不是立即对它们调用系统 free()

想一想，一个更好的主意是在您的 GMP 库之外向该平台引入一个健全的 malloc 实现，以便其他程序和库可以从同一个池中提取它们的内存，而不是GMP 做一件事，其他一切都做另一件事。这应该有助于整体内存消耗 w.r.t 前一点。也许你应该先移植 glibc :)

【讨论】：

"最接近 2 的倍数" 我想你的意思是最接近 2 的 power

你说得对，谢谢（将“倍数”改为“幂”）

啊，我想我错了。实际上，它是分配单位的最接近 2 的幂的倍数