CUDA统一内存和Windows10

Question

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

#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。一位朋友试图在一台 PC 上用 Linux 编译相同的代码并且它工作正常，而它在另一台上有同样的问题具有 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 上的分配粒度可能不同（更小）。

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

#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>

请注意，在我的机器上，在 100,000 次分配后，报告的可用内存只剩下 0.5GB。因此，如果出于任何原因您的 8GB GPU 开始时可用内存较少（完全有可能），您可能运行出现内存不足错误，即使我没有。
分配粒度计算如下：
```
7069866393 - 516266393 / 100000 = 65536 bytes per allocation(!)
```
因此，在我的 machine/test 设置中，我之前对每次分配 4096 字节的估计有很大偏差，至少有 1 个数量级。
分配粒度可能因以下因素而异：
- windows 或 linux
- WDDM 或 TCC
- x86 或 Power9
- 托管与普通 cudaMalloc
- 可能还有其他因素（例如 CUDA 版本）
所以我对未来读者的建议是不要假设每次分配总是 65536 字节，最小值。

CUDA统一内存和Windows10

CUDA unified memory and Windows 10

windows

cuda

unified-memory