CUDA 统一内存页面在 CPU 中访问但未从 GPU 中逐出

Question

我试图了解 CUDA 统一内存的功能。我已经为初学者阅读了关于 CUDA 统一内存的blog。我写了下面给出的代码：

#include <cstdio>
#include <iostream>
#include <fstream>
#include <climits>
#include <vector>

__global__ void transfer(int *X)
{
    X[threadIdx.x] = X[threadIdx.x]+3;
}
using namespace std;
int main()
{
    int *x;
    size_t free_bytes, total_bytes;
    
    cudaMemGetInfo(&free_bytes, &total_bytes);
    std::cout << "Before cudaMallocManaged: " << "free: " << free_bytes << " total: " << total_bytes <<'\n'; 
    cudaMallocManaged(&x,sizeof(int)*512);
    
    cudaMemGetInfo(&free_bytes, &total_bytes);
    std::cout << "After cudaMallocManaged and Before Prefetch to GPU: " << "free: " << free_bytes << " total: " << total_bytes <<'\n';
    std::cout <<  cudaMemPrefetchAsync(x, sizeof(int)*512, 0);
    cudaMemset(x,0,sizeof(int)*512);
    cudaDeviceSynchronize();
    cudaMemGetInfo(&free_bytes, &total_bytes);
    std::cout << "\nAfter Prefetch to GPU Before Kernel call: " << "free: " << free_bytes << " total: " << total_bytes <<'\n'; 
    transfer<<<1,512>>>(x);
    cudaMemGetInfo(&free_bytes, &total_bytes);
    std::cout << "After Kernel call Before memAdvise: " << "free: " << free_bytes << " total: " << total_bytes <<'\n';
    cudaMemAdvise(x,sizeof(int)*512, cudaMemAdviseSetPreferredLocation, cudaCpuDeviceId);
    cudaMemGetInfo(&free_bytes, &total_bytes);
    std::cout << "After memAdvise Before Prefetch to CPU: " << "free: " << free_bytes << " total: " << total_bytes <<'\n';
    std::cout << cudaMemPrefetchAsync(x, sizeof(int)*512, cudaCpuDeviceId);
    cudaDeviceSynchronize();
    cudaMemGetInfo(&free_bytes, &total_bytes);
    std::cout << "\nAfter Prefetch Before processing in CPU: " << "free: " << free_bytes << " total: " << total_bytes <<'\n'; 
    for(int i=0;i<512;i++)
    {
        x[i] = x[i]+1;
        std::cout << x[i];
    }
    cudaMemGetInfo(&free_bytes, &total_bytes);
    std::cout << "\nAfter processing in CPU Before free: " << "free: " << free_bytes << " total: " << total_bytes <<'\n';
    cudaFree(x);
    cudaMemGetInfo(&free_bytes, &total_bytes);
    std::cout << "After free: " << "free: " << free_bytes << " total: " << total_bytes <<'\n';
    return 0;
}

输出：

Before cudaMallocManaged: free: 16804216832 total: 17071734784
After cudaMallocManaged and Before Prefetch to GPU: free: 16804216832 total: 17071734784
0
After Prefetch to GPU Before Kernel call: free: 16669999104 total: 17071734784
After Kernel call Before memAdvise: free: 16669999104 total: 17071734784
After memAdvise Before Prefetch to CPU: free: 16669999104 total: 17071734784
0
After Prefetch Before processing in CPU: free: 16669999104 total: 17071734784
44444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444444
After processing in CPU Before free: free: 16669999104 total: 17071734784
After free: free: 16674193408 total: 17071734784

我在提供 16 GB Tesla P100 PCIe GPU 的 Kaggle 上运行代码。我有一个使用cudaMallocManaged() 分配的整数数组x。首先，我在 GPU 中预取数组并对其进行一些处理，然后将其预取到 CPU 并进行一些处理。在这两者之间，我打印了内存传输前后 GPU 上可用的空闲内存。基于此，我有两个问题：

1. 在cudaMallocManaged() 之后的第一次预取期间，可用内存的减少比我分配的要多得多。为什么？

2. 预取到 CPU 前后的空闲内存是一样的。此外，当我访问和修改 CPU 上的数组时，GPU 上的可用内存在此之前和之后仍然保持不变。我不明白为什么会这样。在预取/处理 CPU 上的统一内存位置时，GPU 上的相应页面不应该被驱逐并移动到 CPU，这不应该释放 GPU 内存吗？

Answer 1

1. 在 GPU 上拥有功能齐全的 CUDA 环境需要相当大的开销。这可能超过 CUDA 开销所需的 100MB 空间，不包括您的数据
2. CUDA 有一个惰性初始化系统。

在 cudaMallocManaged() 之后的第一次预取期间，可用内存减少的比我分配的要多得多。为什么？

因为 CUDA 有一个惰性初始化系统。这意味着当您继续在程序中进行 CUDA 运行时 API 调用时，它可能会构建越来越多的必要环境来运行您的内核代码，以及与之相关的内存开销。在内核启动时，大部分或所有初始化都将完成，除了与新资源使用相关的事情。因此，可用内存的减少是由于您的分配加上 CUDA 本身的额外开销。

预取到 CPU 前后的空闲内存是一样的。此外，当我访问和修改 CPU 上的数组时，GPU 上的可用内存在此之前和之后仍然保持不变。我不明白为什么会这样。

我们所说的内存量大约为 100MB。与此相比，您对512*sizeof(int) 的分配微不足道。此外，CUDA 文档中没有声明由于按需分页，底层分配会发生什么。您似乎认为按需分页会在内容被分页时自动释放底层分配。这在任何地方都没有说明，事实并非如此。此处的确切行为未指定。此外，您设置中的 GPU 具有超额订阅的能力，因此没有特别的理由立即释放分配。

在 CPU 上预取/处理统一内存位置时，GPU 上的相应页面不应该被驱逐并移动到 CPU，这不应该释放 GPU 内存吗？

预取与驱逐不同。但是，是的，预取到 CPU 意味着相应的页面不再驻留在该 GPU 的内存中。不，没有理由认为这会自动/立即释放 GPU 内存。您可以预期，当您对分配的指针执行cudaFree 操作时，内存将被释放，而不是之前。