3

在使用 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 是停用。任何帮助，将不胜感激。谢谢。

4

1 回答 1

3

有一个分配粒度。

这意味着如果您要求 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 上的分配粒度不同（更小）。

基于评论中的一个问题，我决定使用以下代码估计分配粒度：

#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 字节。

于 2020-02-17T20:19:27.113 回答