CUDA统一内存和Windows 10

Question

在使用 CudaMallocManaged() 分配内部包含数组的结构数组时，即使我有足够的可用内存，我也会收到“内存不足”错误。这是一些复制我的问题的代码：

#include <iostream>
#include <cuda.h>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

#define N 100000
#define ARR_SZ 100

struct Struct
{
    float* arr;
};

int main()
{
    Struct* struct_arr;

    gpuErrchk( cudaMallocManaged((void**)&struct_arr, sizeof(Struct)*N) );
    for(int i = 0; i < N; ++i)
        gpuErrchk( cudaMallocManaged((void**)&(struct_arr[i].arr), sizeof(float)*ARR_SZ) ); //out of memory...

    for(int i = 0; i < N; ++i)
        cudaFree(struct_arr[i].arr);
    cudaFree(struct_arr);

    /*float* f;
    gpuErrchk( cudaMallocManaged((void**)&f, sizeof(float)*N*ARR_SZ) ); //this works ok
    cudaFree(f);*/

    return 0;
}

当我调用 cudaMallocManaged() 一次来分配一块内存时，似乎没有问题，正如我在最后一段注释代码中所展示的那样。 我有一个 GeForce GTX 1070 Ti，我使用的是 Windows 10。一个朋友试图在一台装有 Linux 的 PC 上编译相同的代码，它工作正常，而在另一台装有 Windows 10 的 PC 中它也有同样的问题。WDDM TDR 是停用。 任何帮助，将不胜感激。谢谢。

Answer 1

有一个分配粒度。

这意味着如果你要求 1 个字节，或者 400 个字节，实际用完的就是 4096 65536 个字节。因此，一堆非常小的分配实际上会以比您根据请求的分配大小预测的速度更快的速度耗尽内存。解决方案是不要进行非常小的分配，而是分配更大的块。

这里的另一种策略也是扁平化您的分配，并从中为您的每个数组划分出部分：

#include <iostream>
#include <cstdio>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess)
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

#define N 100000
#define ARR_SZ 100

struct Struct
{
    float* arr;
};

int main()
{
    Struct* struct_arr;
    float* f;

    gpuErrchk( cudaMallocManaged((void**)&struct_arr, sizeof(Struct)*N) );
    gpuErrchk( cudaMallocManaged((void**)&f, sizeof(float)*N*ARR_SZ) );
    for(int i = 0; i < N; ++i)
        struct_arr[i].arr = f+i*ARR_SZ;
    cudaFree(struct_arr);
    cudaFree(f);

    return 0;
}

ARR_SZ 可被 4 整除意味着各种创建的指针也可以向上转换为更大的向量类型，例如float2 或 float4，如果您的用户有任何这样做的意图。

原始代码在 linux 上运行的一个可能原因是，在适当的设置下，linux 上的托管内存可能会超额订阅 GPU 物理内存。结果是实际分配限制远高于 GPU 板载内存的建议。也可能是linux的情况下空闲内存多一点，或者linux上的分配粒度不同（更小）。

基于cmets中的一个问题，我决定估计分配粒度，使用这段代码：

#include <iostream>
#include <cstdio>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char* file, int line, bool abort = true)
{
    if (code != cudaSuccess)
    {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

#define N 100000
#define ARR_SZ 100

struct Struct
{
    float* arr;
};

int main()
{
    Struct* struct_arr;
    //float* f;

    gpuErrchk(cudaMallocManaged((void**)& struct_arr, sizeof(Struct) * N));
#if 0
    gpuErrchk(cudaMallocManaged((void**)& f, sizeof(float) * N * ARR_SZ));
    for (int i = 0; i < N; ++i)
        struct_arr[i].arr = f + i * ARR_SZ;
#else
    size_t fre, tot;
    gpuErrchk(cudaMemGetInfo(&fre, &tot));
    std::cout << "Free: " << fre << " total: " << tot << std::endl;
    for (int i = 0; i < N; ++i)
        gpuErrchk(cudaMallocManaged((void**) & (struct_arr[i].arr), sizeof(float) * ARR_SZ)); 
    gpuErrchk(cudaMemGetInfo(&fre, &tot));
    std::cout << "Free: " << fre << " total: " << tot << std::endl;

    for (int i = 0; i < N; ++i)
        cudaFree(struct_arr[i].arr);
#endif
    cudaFree(struct_arr);
    //cudaFree(f);

    return 0;
}

当我使用该代码编译调试项目并在具有 RTX 2070 GPU（8GB 内存，与 GTX 1070 Ti 相同）的 Windows 10 桌面上运行时，我得到以下输出：

Microsoft Windows [Version 10.0.17763.973]
(c) 2018 Microsoft Corporation. All rights reserved.

C:\Users\Robert Crovella>cd C:\Users\Robert Crovella\source\repos\test12\x64\Debug

C:\Users\Robert Crovella\source\repos\test12\x64\Debug>test12
Free: 7069866393 total: 8589934592
Free: 516266393 total: 8589934592

C:\Users\Robert Crovella\source\repos\test12\x64\Debug>test12
Free: 7069866393 total: 8589934592
Free: 516266393 total: 8589934592

C:\Users\Robert Crovella\source\repos\test12\x64\Debug>

1. 请注意，在我的机器上，在 100,000 次分配后报告的可用内存仅剩 0.5GB。因此，如果出于某种原因，您的 8GB GPU 开始时可用内存较少（完全有可能），您可能会遇到内存不足错误，即使我没有这样做。

2. 分配粒度的计算如下：

   7069866393 - 516266393 / 100000 = 65536 bytes per allocation(!)
   

   因此，在我的机器/测试设置上，我之前估计的每次分配 4096 字节至少相差 1 个数量级。

3. 分配粒度可能因以下因素而异：

   • windows 或 linux
   • WDDM 或 TCC
   • x86 或 Power9
   • 托管与普通cudaMalloc
   • 可能是其他因素（例如 CUDA 版本）

   所以我对未来读者的建议是不要假设每次分配总是至少 65536 字节。