超量分配统一内存的上限

  • 6 replies
  • 197 views
超量分配统一内存的上限
« 于: 十二月 15, 2019, 03:19:33 am »
我尝试用下面的程序测试统一内存的超量分配:

#include "error.cuh"
#include <stdio.h>

int main(void)
{
    long long GB = 1 << 30;
    for (long long n = 1; n <= 1000L; ++n)
    {
        char *d_x;
#ifdef UNIFIED
        CHECK(cudaMallocManaged(&d_x, GB * n));
        printf("Can allocate %lld GB unified memory.\n", n);
#else
        CHECK(cudaMalloc(&d_x, GB * n));
        printf("Can allocate %lld GB device memory.\n", n);
#endif
        CHECK(cudaFree(d_x));
    }
    return 0;
}


我的显卡是显存略小于8G的RTX 2070,主机是16G内存的笔记本。

如果用如下命令编译
nvcc -arch=sm_75  xxx.cu

则得到如下输出:
Can allocate 1 GB device memory.
Can allocate 2 GB device memory.
Can allocate 3 GB device memory.
Can allocate 4 GB device memory.
Can allocate 5 GB device memory.
Can allocate 6 GB device memory.
Can allocate 7 GB device memory.
CUDA Error:
    File:       test.cu
    Line:       14
    Error code: 2
    Error text: out of memory
这是我可以理解的。

如果用如下命令编译:
nvcc -arch=sm_75  -DUNIFIED xxx.cu
则得到如下输出:

...
Can allocate 1000 GB unified memory.

问题:如何理解此时程序的输出?



PS: error.cuh文件内容如下:

#pragma once
#include <stdio.h>

#define CHECK(call)                                   \
do                                                    \
{                                                     \
    const cudaError_t error_code = call;              \
    if (error_code != cudaSuccess)                    \
    {                                                 \
        printf("CUDA Error:\n");                      \
        printf("    File:       %s\n", __FILE__);     \
        printf("    Line:       %d\n", __LINE__);     \
        printf("    Error code: %d\n", error_code);   \
        printf("    Error text: %s\n",                \
            cudaGetErrorString(error_code));          \
        exit(1);                                      \
    }                                                 \
} while (0)



Re: 超量分配统一内存的上限
« 回复 #1 于: 十二月 16, 2019, 02:03:57 pm »
Hello bruce,

cudaMallocManaged()的调用成功,在Linux, amd64下的时候,只代表地址空间(address space)保留成功,此时将返回到你的结果指针中,该段保留(reserved)的地址空间的起始值。

而实际的分配,则发生于first touch的时刻(很多分配器都这样,不仅仅是UVM),具体的说,我们可以分成在CPU上和在GPU上直接访问两种情况。

假设你已经分配了128GB的Unified Memory。显卡大小8GB,物理内存大小为64GB(不算swap,unified memory在非power平台上,不能利用swap).

(A)请参考如下函数模拟CPU的First Touch过程:
程序代码: [选择]
#include <stdint.h>
void CPU_Touch(uint64_t *p, size_t size)
{
    for (size_t i = 0; i < size/sizeof(uint64_t); i++) p[i] = 0;   
}
然后在你的cudaFree前面加上:
CPU_Touch((uint64_t *)d_x, (size_t)n * 1024 * 1024 * 1024);

然后nvcc -arch sm_75 -DUNIFIED 你的cu文件 回车
然后./a.out 回车

然后你能观察到,你目前系统(64GB物理内存)的Unified Memory使用极限应该在略小于64GB(因为所有的unified memory的内存部分,都是pinned memory)。

同时请观察,使用超过64GB的时候,你的进程将挂掉,此时检查dmesg, 搜索最后的OOM(out of memory)字样。

(B)进一步的,请参考如下代码,模拟GPU进行first touch的过程:
程序代码: [选择]
__global__ void _touch(uint64_t *p)
{
    size_t index = blockIdx.x * blockDim.x + threadIdx.x;
    p[index] = 0;
}
void touch_gpu(uint64_t *p, size_t size)
{
    _touch<<<size / sizeof(uint64_t) / 1024, 1024>>>(p);
    CHECK(cudaGetLastError());
    CHECK(cudaDeviceSynchronize());
}
然后:nvcc -arch sm_75 -DUNIFIED 你的cu文件 回车
然后:./a.out 回车

运行程序,注意观察如下现象:
(B1)当kernel使用的存储器越来越大后,用nvidia-smi可以观察到,设备的存储器使用显示了XXXX MB / YYYY MB字样,其中XXXX和YYYY都接近于8GB(你的显存容量)。然而虽然kernel使用满了显存,却依然在正常运行没有挂掉(因为此时,超量分配机制让显存变成了L3 cache). 然后用nvidia-smi -a观察到,存在的从显卡方向往CPU方向的TX传输(超过显存部分的淘汰数据将自动被传输到后后备的内存中)。

(B2)观察当分配的Unified Memory越大,这个kernel运行的越慢。当kernel使用到的缓冲区大小,从[8GB, 略小于64GB)的区间内,kernel能正常运行。

(B3)观察kernel使用超过64GB的缓冲区的时候,挂掉。此时检查CUDA错误代码,和dmesg中查找UVM动态分配失败字样。

感谢来访。
« 最后编辑时间: 十二月 16, 2019, 02:05:20 pm 作者 屠戮人神 »

Re: 超量分配统一内存的上限
« 回复 #2 于: 十二月 25, 2019, 12:40:19 pm »
谢谢解答。非常专业,我还需要消化消化。 ;D

Re: 超量分配统一内存的上限
« 回复 #3 于: 十二月 25, 2019, 01:58:17 pm »
谢谢解答。非常专业,我还需要消化消化。 ;D

留个思考题:
在情况(B)的时候,几乎总是能观察到,从GPU端能使用的Unified Memory,即kernel能使用的缓冲区大小,总是要大小CPU端所能够使用的缓冲区大小。请解释这个现象。

Re: 超量分配统一内存的上限
« 回复 #4 于: 十二月 25, 2019, 05:17:38 pm »
还有思考题!嗯,我会好好思考的。

Re: 超量分配统一内存的上限
« 回复 #5 于: 十二月 30, 2019, 12:03:20 am »
留个思考题:
在情况(B)的时候,几乎总是能观察到,从GPU端能使用的Unified Memory,即kernel能使用的缓冲区大小,总是要大小CPU端所能够使用的缓冲区大小。请解释这个现象。

我不知道我的理解是否正确。我今天测试的结果是:
1) 使用CPU touch时,只能分配13GB的统一内存,小于主机内存16GB。
2) 使用GPU touch时,能分配20 GB的统一内存,大于GPU或CPU的内存,但小于它们的和24GB。

我的理是:
1)操作系统和其它应用程序可能使用了大约3G内存。
2)如果仅仅用 CPU touch,则无法将 GPU 内存纳入统一内存中。但我不知为何?

还请继续解惑。



Re: 超量分配统一内存的上限
« 回复 #6 于: 十二月 30, 2019, 11:57:28 am »
我不知道我的理解是否正确。我今天测试的结果是:
1) 使用CPU touch时,只能分配13GB的统一内存,小于主机内存16GB。
2) 使用GPU touch时,能分配20 GB的统一内存,大于GPU或CPU的内存,但小于它们的和24GB。

我的理是:
1)操作系统和其它应用程序可能使用了大约3G内存。
2)如果仅仅用 CPU touch,则无法将 GPU 内存纳入统一内存中。但我不知为何?

还请继续解惑。

的确是这样的。

(1)被操作系统内核,或者其他应用进程所使用的page-locked memory, 不能被使用。

(2)单独在GPU上初次使用的话,当该GPU上被填充满了,然后才会被淘汰传输往内存。此时大约等于:你单独在CPU上能使用的大小 加上GPU显存大小(不太到)。也就是内存大小 + 显存大小 - (1)中被占用的,最终的大小可能接近物理内存总容量或者轻微超出。

(3)的确不会CPU在使用满内存后,自动往显存淘汰。

但是你可以尝试使用cudaMemPrefetchAsync来辅助迁移到显存(只有当没有设定只读副本(cudaMemAdviceSetReadMostly)的时候有效,设定只读副本后,迁移了,但对应的内存还在,会同时出现相同内容的内存和显存副本)。该函数需要增强版本版本的Unified Memory支持(请检测并发CPU/GPU Managed Memory访问能力,即设备属性的ConcurrentManagedAccess。)该属性实际上更多的表示是否具有增强的Unified Memory支持,而不仅仅是能CPU/GPU同时访问。

此外,如果没有增强的Unified Memory支持,则可以考虑cudaStreamAttachMemAsync()到单一GPU上的单一逻辑Stream范围内。该方法不需要增强版本的Unified Memory。注意论坛没有测试过这种对内存使用的情况的变化,不过欢迎bruce测试并报告。

« 最后编辑时间: 十二月 30, 2019, 11:59:38 am 作者 屠戮人神 »