OpenAcc Large arrays allocation - Variable partially present issue.

Hi I have a piece of openacc code and running it on a GPU and getting memory error.
Please help me in resolving the same.
Here is my code and the card that I am using.
Also, the gpu card is not running any extra processes. (no gui nothing)

Code:

#define N 1358993817
//This number of chosen such that the array size is approx 5gb
//Such 2 arrays exist making it 10gb and the card size is 12gb

int main (int argc, char *argv[]) {


        float x[N], y[N];
        
        #pragma acc data present_or_create(x[:N]) present_or_copyout(y[:N])
        {

                #pragma acc parallel loop
                for(long long i = 0; i < N; ++i) {
                        y[i] = 0.0;
                        x[i] = (float) (i+1);
                }

                #pragma acc parallel loop
                for(long long i = 0; i < N; ++i) {
                        y[i] = 2.0f * x[i] + y[i];
                }
        }

        return 0;
}



Card Memory Info:
==============NVSMI LOG==============

Timestamp                           : Tue Jun 18 12:09:56 2019
Driver Version                      : 418.67
CUDA Version                        : 10.1

Attached GPUs                       : 1
GPU 00000000:05:00.0
    FB Memory Usage
        Total                       : 12198 MiB
        Used                        : 0 MiB
        Free                        : 12198 MiB
    BAR1 Memory Usage
        Total                       : 16384 MiB
        Used                        : 2 MiB
        Free                        : 16382 MiB

Compiled using :

pgcc -fast -acc -Minfo=all -ta:tesla data.c
main:
     23, Generating create(x[:])
         Generating copyout(y[:])
     26, Generating Tesla code
         27, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     32, Generating Tesla code
         33, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */

         
pgcc --version

pgcc 19.4-0 LLVM 64-bit target on x86-64 Linux -tp haswell
PGI Compilers and Tools
Copyright (c) 2019, NVIDIA CORPORATION.  All rights reserved.

Error is :

y lives at 0x7ffaf2e39bd8 size 5435975268 partially present
Present table dump for device[1]: NVIDIA Tesla GPU 0, compute capability 6.0, threadid=1
host:0x7ffb36e6023c device:0x2af91e000000 size:5435975268 presentcount:1+0 line:23 name:x
allocated block device:0x2af91e000000 size:5435975680 thread:1
FATAL ERROR: variable in data clause is partially present on the device: name=y
 file:/home/vineetm/openacc/other/data.c main line:23

Thanks

Hi ramdev,

While I’m not sure if this is a compiler issue or a limitation of the allowable size of static arrays, I went ahead and filed a problem report, TPR#27318, so our engineers can investigate.

The easy workaround is to use dynamically allocated arrays.

Thanks,
Mat

% cat test.c
#include <stdlib.h>

#ifdef TWOGB
#define N  536870912LL
#else  // 5GB
#define N 1358993817LL
#endif
//This number of chosen such that the array size is approx 5gb
//Such 2 arrays exist making it 10gb and the card size is 12gb
int main (int argc, char *argv[]) {

#ifdef STATIC
        float x[N], y[N];
#else
        float *x, *y;
        x = (float*) malloc(sizeof(float)*N);
        y = (float*) malloc(sizeof(float)*N);
#endif
        #pragma acc data create(x[0:N]) copyout(y[0:N])
        {
                #pragma acc parallel loop
                for(long long i = 0; i < N; ++i) {
                        y[i] = 0.0;
                        x[i] = (float) (i+1);
                }
                #pragma acc parallel loop
                for(long long i = 0; i < N; ++i) {
                        y[i] = 2.0f * x[i] + y[i];
                }
        }
#ifndef STATIC
        free(x);
        free(y);
#endif
        return 0;
}

% pgcc -ta=tesla -Minfo=accel test.c -Mlarge_arrays ; a.out
main:
     19, Generating create(x[:1358993817])
         Generating copyout(y[:1358993817])
     21, Generating Tesla code
         22, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     26, Generating Tesla code
         27, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
% pgcc -ta=tesla -Minfo=accel test.c -DSTATIC -mcmodel=medium ; a.out
main:
     19, Generating create(x[:])
         Generating copyout(y[:])
     21, Generating Tesla code
         22, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     26, Generating Tesla code
         27, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
y lives at 0x7ffc0a00fe00 size 5435975268 partially present
Present table dump for device[1]: NVIDIA Tesla GPU 0, compute capability 7.0, threadid=1
host:0x7ffc4e036464 device:0x2b0e78000000 size:5435975268 presentcount:1+0 line:19 name:x
allocated block device:0x2b0e78000000 size:5435975680 thread:1
FATAL ERROR: variable in data clause is partially present on the device: name=y
 file:/local/home/colgrove/test.c main line:19

Thanks for the reply.
I will continue using dynamic arrays for now. But looking forward for the large static arrays working as well.
Thanks