-ta=tesla:managed:cuda8 但 cuMemAllocManaged 返回错误 2：内存不足

Question

我是 OpenACC 的新手。就我对OpenMP的熟悉程度而言，我非常喜欢它。

我有 2 张 1080Ti 卡，每张 9GB，内存为 128GB。我正在尝试一个非常基本的测试来分配一个数组，初始化它，然后并行地总结它。这适用于 8 GB，但当我增加到 10 GB 时，出现内存不足错误。我的理解是，使用 Pascal（这些卡是）和 CUDA 8 的统一内存，我可以分配一个大于 GPU 内存的数组，并且硬件将按需调入调出。

这是我的完整 C 代码测试：

$ cat firstAcc.c 

#include <stdio.h>
#include <openacc.h>
#include <stdlib.h>

#define GB 10

int main()
{
  float *a;
  size_t n = GB*1024*1024*1024/sizeof(float);
  size_t s = n * sizeof(float);
  a = (float *)malloc(s);
  if (!a) { printf("Failed to malloc.\n"); return 1; }
  printf("Initializing ... ");
  for (int i = 0; i < n; ++i) {
    a[i] = 0.1f;
  }
  printf("done\n");
  float sum=0.0;
  #pragma acc loop reduction (+:sum)
  for (int i = 0; i < n; ++i) {
    sum+=a[i];
  }
  printf("Sum is %f\n", sum);
  free(a);
  return 0;
}

根据 this article 的 "Enable Unified Memory" 部分，我用 :

编译它

$ pgcc -acc -fast -ta=tesla:managed:cuda8 -Minfo firstAcc.c
main:
 20, Loop not fused: function call before adjacent loop
     Generated vector simd code for the loop
 28, Loop not fused: function call before adjacent loop
     Generated vector simd code for the loop containing reductions
     Generated a prefetch instruction for the loop

我需要了解这些消息，但目前我认为它们不相关。那我运行吧:

$ ./a.out
malloc: call to cuMemAllocManaged returned error 2: Out of memory
Aborted (core dumped)

如果我将 GB 更改为 8，则可以正常工作。由于 Pascal 1080Ti 和 CUDA 8，我希望 10GB 可以工作（尽管 GPU 卡有 9GB）。

我是不是理解错了，或者我做错了什么？提前致谢。

$ pgcc -V
pgcc 17.4-0 64-bit target on x86-64 Linux -tp haswell 
PGI Compilers and Tools
Copyright (c) 2017, NVIDIA CORPORATION.  All rights reserved.

$ cat /usr/local/cuda-8.0/version.txt 
CUDA Version 8.0.61

Answer 1

我认为这里有问题：

size_t n = GB*1024*1024*1024/sizeof(float);

当我用 g++ 编译那行代码时，我收到了关于整数溢出的警告。出于某种原因，PGI 编译器没有警告，但同样的错误正在发生。在 s 和 n 的声明之后，如果我添加这样的打印输出：

  size_t n = GB*1024*1024*1024/sizeof(float);
  size_t s = n * sizeof(float);
  printf("n = %lu, s = %lu\n", n, s);  // add this line

并使用 PGI 17.04 和运行（在 P100 上，16GB）编译，我得到这样的输出：

$ pgcc -acc -fast -ta=tesla:managed:cuda8 -Minfo m1.c
main:
     16, Loop not fused: function call before adjacent loop
         Generated vector simd code for the loop
     22, Loop not fused: function call before adjacent loop
         Generated vector simd code for the loop containing reductions
         Generated a prefetch instruction for the loop
$ ./a.out
n = 4611686017890516992, s = 18446744071562067968
malloc: call to cuMemAllocManaged returned error 2: Out of memory
Aborted
$

所以很明显 n 和 s 不是您想要的。

我们可以通过用 ULL 标记所有这些常量来解决这个问题，然后事情似乎对我来说是正确的：

$ cat m1.c
#include <stdio.h>
#include <openacc.h>
#include <stdlib.h>

#define GB 20ULL

int main()
{
  float *a;
  size_t n = GB*1024ULL*1024ULL*1024ULL/sizeof(float);
  size_t s = n * sizeof(float);
  printf("n = %lu, s = %lu\n", n, s);
  a = (float *)malloc(s);
  if (!a) { printf("Failed to malloc.\n"); return 1; }
  printf("Initializing ... ");
  for (int i = 0; i < n; ++i) {
    a[i] = 0.1f;
  }
  printf("done\n");
  double sum=0.0;
  #pragma acc loop reduction (+:sum)
  for (int i = 0; i < n; ++i) {
    sum+=a[i];
  }
  printf("Sum is %f\n", sum);
  free(a);
  return 0;
}
$ pgcc -acc -fast -ta=tesla:managed:cuda8 -Minfo m1.c
main:
     16, Loop not fused: function call before adjacent loop
         Generated vector simd code for the loop
     22, Loop not fused: function call before adjacent loop
         Generated vector simd code for the loop containing reductions
         Generated a prefetch instruction for the loop
$ ./a.out
n = 5368709120, s = 21474836480
Initializing ... done
Sum is 536870920.000000
$

请注意，我还在上面进行了另一项更改。我将 sum 累积变量从 float 更改为 double。当对非常小的数量进行非常大的减少时，这对于保留一些 "sensible" 结果是必要的。

而且，正如@MatColgrove 在他的回答中指出的那样，我还错过了其他一些事情。

Answer 2

除了 Bob 提到的，我还做了一些修复。

首先，您实际上并没有生成 OpenACC 计算区域，因为您只有一个“#pragma acc loop”指令。这应该是“#pragma acc parallel loop”。您可以在仅显示主机代码优化的编译器反馈消息中看到这一点。

其次，"i" 索引应声明为 "long"。否则，您将溢出索引。

最后，您需要将 "cc60" 添加到您的目标加速器选项中，以告知编译器以基于 Pascal 的 GPU 为目标。

% cat mi.c  
#include <stdio.h>
#include <openacc.h>
#include <stdlib.h>

#define GB 20ULL

int main()
{
  float *a;
  size_t n = GB*1024ULL*1024ULL*1024ULL/sizeof(float);
  size_t s = n * sizeof(float);
  printf("n = %lu, s = %lu\n", n, s);
  a = (float *)malloc(s);
  if (!a) { printf("Failed to malloc.\n"); return 1; }
  printf("Initializing ... ");
  for (int i = 0; i < n; ++i) {
    a[i] = 0.1f;
  }
  printf("done\n");
  double sum=0.0;
  #pragma acc parallel loop reduction (+:sum)
  for (long i = 0; i < n; ++i) {
    sum+=a[i];
  }
  printf("Sum is %f\n", sum);
  free(a);
  return 0;
}

% pgcc -fast -acc -ta=tesla:managed,cuda8.0,cc60 -Minfo=accel mi.c
main:
     21, Accelerator kernel generated
         Generating Tesla code
         21, Generating reduction(+:sum)
         22, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     21, Generating implicit copyin(a[:5368709120])
% ./a.out
n = 5368709120, s = 21474836480
Initializing ... done
Sum is 536870920.000000

-ta=tesla:managed:cuda8 但 cuMemAllocManaged 返回错误 2：内存不足

-ta=tesla:managed:cuda8 but cuMemAllocManaged returned error 2: Out of memory

cuda

openacc