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

Question

I'm new to OpenACC. I like it very much so far as I'm familiar with OpenMP.

I have 2 1080Ti cards each with 9GB and I've 128GB of RAM. I'm trying a very basic test to allocate an array, initialize it, then sum it up in parallel. This works for 8 GB but when I increase to 10 GB I get out-of-memory error. My understanding was that with unified memory of Pascal (which these card are) and CUDA 8, I could allocate an array larger than the GPU's memory and the hardware will page in and page out on demand.

Here's my full C code test :

$ 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;
}

As per the "Enable Unified Memory" section of this article I compile it with :

$ 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

I need to understand those messages but for now I don't think they are relevant. Then I run it :

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

This works fine if I change GB to 8. I expected 10GB to work (despite the GPU card having 9GB) thanks to Pascal 1080Ti and CUDA 8.

Have I misunderstand, or what am I doing wrong? Thanks in advance.

$ 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

Besides what Bob mentioned, I made a few more fixes.

First, you're not actually generating an OpenACC compute region since you only have a "#pragma acc loop" directive. This should be "#pragma acc parallel loop". You can see this in the compiler feedback messages where it's only showing host code optimizations.

Second, the "i" index should be declared as a "long". Otherwise, you'll overflow the index.

Finally, you need to add "cc60" to your target accelerator options to tell the compiler to target a Pascal based 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

Answer 2

I believe a problem is here:

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

when I compile that line of code with g++, I get a warning about integer overflow. For some reason the PGI compiler is not warning, but the same badness is occurring under the hood. After the declarations of s, and n, if I add a printout like this:

  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

and compile with PGI 17.04, and run (on a P100, with 16GB) I get output like this:

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

so it's evident that n and s are not what you intended.

We can fix this by marking all of those constants with ULL, and then things seem to work correctly for me:

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

Note that I've made another change above as well. I changed the sum accumulation variable from float to double. This is necessary to preserve somewhat "sensible" results when doing a very large reduction across very small quantities.

And, as @MatColgrove pointed out in his answer, I missed a few other things as well.