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

Question

根据 this question和引用 NVIDIA CUDA Programming Guide realloc 函数未实现:

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.

我目前坚持使用 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 前一个 block ，但这需要强制移动数据(大数组)，我想避免这种情况。

是否有任何effective efficient 方法来实现(标准 C 或 C++)realloc 函数(即内部核心) ？假设我有一些动态分配数据的大数组(已经由 malloc 分配)，然后在其他地方调用 realloc 以请求更大的内存量对于that block 。简而言之，我想避免将整个数据数组复制到新位置，我会具体询问如何执行此操作(当然，如果可能的话)。

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

Answer 1

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

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

我不熟悉 CUDA 或 GMP，但想不通:

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

• 如果您想要的只是一个更高效的 realloc，您可以简单地在您的自定义 malloc 中进行过度分配(达到某个大小，比如说最接近的 2 的幂)，这样您就可以避免在 subseauent 重新分配中进行复制.为此，您甚至不需要成熟的堆实现。

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

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

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