Purpose: Testing of deep copy of user-defined data structure with dynamically allocated jagged 2d array component
Compiler: PGI Community Edition 18.4
Platform: Windows 7 64-bit
GPU: Geforce GTX 980M
Command: pgcc -acc -Minfo acc_deepcopy_jagged.c
The source code of acc_deepcopy_jagged.c is:
#include <stdlib.h>
#include <stdio.h>
#include <openacc.h>
int ROW = 128;
struct A
{
int n; //number of elements in x
int **data2D;
int *data1D;
};
double sum1D(const struct A *z, const int n, int usegpu)
{
double sum1D=0.0;
/* The disjoint data structure is expected to be present */
#pragma acc parallel loop reduction(+:sum1D) present(z[0:n]) if(usegpu)
for (int i = 0; i < n; i++)
{
for (int r = 0; r < ROW; r++)
{
sum1D += z[i].data1D[r];
}
}
return sum1D;
}
double sum2D(const struct A *z, const int n, int usegpu)
{
double sum2D = 0.0;
/* The disjoint data structure is expected to be present */
#pragma acc parallel loop reduction(+:sum2D) present(z[0:n]) if(usegpu)
for (int i = 0; i < n; i++)
{
for (int r = 0; r < ROW; r++)
{
if (r > 0)
{
for (int c = 0; c < r; c++)
{
sum2D += z[i].data2D[r][c];
}
}
}
}
return sum2D;
}
int main(int argc, char* argv[])
{
int n = 4, cnt1D, cnt2D; // cnt1D and cnt2D are counters
acc_set_device_num(1, acc_device_nvidia); // use the second GPU
//Allocate and initalize a disjoint data structure
struct A *z;
//The top level is an array of struct A objects
z = (struct A*) malloc(n*sizeof(struct A));
#pragma acc enter data create(z[0:n])
for (int i = 0; i < n; i++)
{
z[i].n = n;
#pragma acc update device(z[i].n)
z[i].data1D = (int*)malloc(ROW*sizeof(int));
#pragma acc enter data create(z[i].data1D[0:ROW])
z[i].data2D = (int**)malloc(ROW*sizeof(int*));
#pragma acc enter data create(z[i].data2D[0:ROW][0:1])
for (int r = 0; r < ROW; r++)
{
if (r > 0)
{
z[i].data2D[r] = (int*)malloc(r*sizeof(int));
#pragma acc enter data create(z[i].data2D[r:1][0:r])
}
}
cnt1D = 0, cnt2D=0;
for (int r = 0; r < ROW; r++)
{
z[i].data1D[r] = ++cnt1D;
if (r > 0)
{
for (int c = 0; c < r; c++)
{
z[i].data2D[r][c] = ++cnt2D; //OR *(*(z[i].data2D+r)+c) = ++cnt2D
}
#pragma acc update device(z[i].data2D[r:1][0:r])
}
}
#pragma acc update device(z[i].data1D[0:ROW])
}
//Compute and print host sum
printf("host_sum1D=%g\n", sum1D(z, n, 0));
printf("host_sum2D=%g\n", sum2D(z, n, 0));
//compute and print accelerator sum
printf("acc_sum1D=%g\n", sum1D(z, n, 1));
printf("acc_sum2D=%g\n", sum2D(z, n, 1));
return 0;
}
The compiler response is:
sum1D:
18, Generating present(z[:n])
Generating implicit copy(sum1D)
Accelerator kernel generated
Generating Tesla code
19, #pragma acc loop gang /* blockIdx.x /
Generating reduction(+:sum1D)
21, #pragma acc loop vector(128) / threadIdx.x /
21, Loop is parallelizable
sum2D:
33, Generating present(z[:n])
Generating implicit copy(sum2D)
Accelerator kernel generated
Generating Tesla code
34, #pragma acc loop gang / blockIdx.x /
Generating reduction(+:sum2D)
36, #pragma acc loop seq
40, #pragma acc loop vector(128) / threadIdx.x */
36, Loop is parallelizable
40, Loop is parallelizable
main:
60, Generating enter data create(z[:n])
66, Generating update device(z->n)
67, Generating enter data create(z->data1D[:ROW])
70, Generating enter data create(z->data2D[:ROW][:1])
76, Generating enter data create(z->data2D[r][:r])
93, Generating update device(z->data2D[r][:r])
95, Generating update device(z->data1D[:ROW])
When run the exe, error information is:
F:\Files\SourceCode\OpenaccCExamples\acc_deepcopy_jagged>acc_deepcopy_jagged
(null) lives at 00000000058F6350 size 180 partially present
Present table dump for device[2]: NVIDIA Tesla GPU 1, compute capability 5.2, th
readid=1
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:0000000000000000 device:0000000000000000 size:0 presentcount:0+1 line:76 na
me:(null)
host:00000000024ADC20 device:00000007036E4A00 size:108 presentcount:0+1 line:76
name:(null)
host:00000000024ADCA0 device:00000007036E4C00 size:112 presentcount:0+1 line:76
name:(null)
host:00000000024ADD20 device:00000007036E4E00 size:116 presentcount:0+1 line:76
name:(null)
host:00000000024ADDA0 device:00000007036E5000 size:120 presentcount:0+1 line:76
name:(null)
host:00000000024AFC80 device:00000007036E5A00 size:140 presentcount:0+1 line:76
name:(null)
host:00000000024AFD20 device:00000007036E5C00 size:144 presentcount:0+1 line:76
name:(null)
host:00000000024AFDC0 device:00000007036E5E00 size:148 presentcount:0+1 line:76
name:(null)
host:00000000024AFE60 device:00000007036E6000 size:152 presentcount:0+1 line:76
name:(null)
host:00000000028E31E0 device:00000007036E2A00 size:44 presentcount:0+1 line:76 n
ame:(null)
host:00000000028E3220 device:00000007036E2C00 size:48 presentcount:0+1 line:76 n
ame:(null)
host:00000000028E3260 device:00000007036E2E00 size:52 presentcount:0+1 line:76 n
ame:(null)
host:00000000028E32A0 device:00000007036E3000 size:56 presentcount:0+1 line:76 n
ame:(null)
host:0000000002922FE0 device:00000007036E6A00 size:172 presentcount:0+1 line:76
name:(null)
host:0000000005873290 device:00000007036E6200 size:156 presentcount:0+1 line:76
name:(null)
host:0000000005873340 device:00000007036E6400 size:160 presentcount:0+1 line:76
name:(null)
host:00000000058733F0 device:00000007036E6600 size:164 presentcount:0+1 line:76
name:(null)
host:00000000058734A0 device:00000007036E6800 size:168 presentcount:0+1 line:76
name:(null)
host:00000000058F29A0 device:00000007036E3200 size:60 presentcount:0+1 line:76 n
ame:(null)
host:00000000058F29F0 device:00000007036E3400 size:64 presentcount:0+1 line:76 n
ame:(null)
host:00000000058F2A40 device:00000007036E3600 size:68 presentcount:0+1 line:76 n
ame:(null)
host:00000000058F2A90 device:00000007036E3800 size:72 presentcount:0+1 line:76 n
ame:(null)
host:00000000058F6350 device:00000007036E0A00 size:4 presentcount:0+1 line:70 na
me:(null)
host:000000000592D2B0 device:00000007036E0000 size:96 presentcount:0+1 line:60 n
ame:z
host:000000000592D320 device:00000007036E4200 size:92 presentcount:0+1 line:76 n
ame:(null)
host:000000000592D390 device:00000007036E4400 size:96 presentcount:0+1 line:76 n
ame:(null)
host:000000000592D400 device:00000007036E4600 size:100 presentcount:0+1 line:76
name:(null)
host:000000000592D470 device:00000007036E4800 size:104 presentcount:0+1 line:76
name:(null)
host:0000000005A23240 device:00000007036E1600 size:4 presentcount:0+1 line:76 na
me:(null)
host:0000000005A23250 device:00000007036E1800 size:8 presentcount:0+1 line:76 na
me:(null)
host:0000000005A238C0 device:00000007036E3C00 size:80 presentcount:0+1 line:76 n
ame:(null)
host:0000000005A23980 device:00000007036E3A00 size:76 presentcount:0+1 line:76 n
ame:(null)
host:0000000005A239E0 device:00000007036E3E00 size:84 presentcount:0+1 line:76 n
ame:(null)
host:0000000005A23A40 device:00000007036E4000 size:88 presentcount:0+1 line:76 n
ame:(null)
host:0000000005A252E0 device:00000007036E2200 size:28 presentcount:0+1 line:76 n
ame:(null)
host:0000000005A25820 device:00000007036E2600 size:36 presentcount:0+1 line:76 n
ame:(null)
host:0000000005A25850 device:00000007036E2400 size:32 presentcount:0+1 line:76 n
ame:(null)
host:0000000005A25AF0 device:00000007036E2800 size:40 presentcount:0+1 line:76 n
ame:(null)
host:0000000005A3FA70 device:00000007036E0200 size:512 presentcount:0+1 line:67
name:(null)
host:0000000005A45FA0 device:00000007036E1A00 size:12 presentcount:0+1 line:76 n
ame:(null)
host:0000000005A45FE0 device:00000007036E1C00 size:16 presentcount:0+1 line:76 n
ame:(null)
host:0000000005A46020 device:00000007036E1E00 size:20 presentcount:0+1 line:76 n
ame:(null)
host:0000000005A46060 device:00000007036E2000 size:24 presentcount:0+1 line:76 n
ame:(null)
host:0000000005A7ADD0 device:00000007036E1200 size:4 presentcount:0+1 line:70 na
me:(null)
host:0000000005A7B180 device:00000007036E0400 size:1024 presentcount:0+1 line:70
name:(null)
host:0000000005A7B590 device:00000007036E6C00 size:176 presentcount:0+1 line:76
name:(null)
host:0000000005A9A290 device:00000007036E0800 size:4 presentcount:0+1 line:70 na
me:(null)
host:0000000005AA21D0 device:00000007036E5200 size:124 presentcount:0+1 line:76
name:(null)
host:0000000005AA24A0 device:00000007036E5400 size:128 presentcount:0+1 line:76
name:(null)
host:0000000005AA2770 device:00000007036E5600 size:132 presentcount:0+1 line:76
name:(null)
host:0000000005AA2A40 device:00000007036E5800 size:136 presentcount:0+1 line:76
name:(null)
host:00000000FFFFFFFF device:00000007036E0E00 size:4 presentcount:0+1 line:70 na
me:(null)
host:00000001A1695385 device:00000007036E0C00 size:4 presentcount:0+1 line:70 na
me:(null)
host:00000002B98D70E8 device:00000007036E1400 size:4 presentcount:0+1 line:70 na
me:(null)
host:000000200000000D device:00000007036E1000 size:4 presentcount:0+1 line:70 na
me:(null)
allocated block device:00000007036E0000 size:512 thread:1
allocated block device:00000007036E0200 size:512 thread:1
allocated block device:00000007036E0400 size:1024 thread:1
allocated block device:00000007036E0800 size:512 thread:1
allocated block device:00000007036E0A00 size:512 thread:1
allocated block device:00000007036E0C00 size:512 thread:1
allocated block device:00000007036E0E00 size:512 thread:1
allocated block device:00000007036E1000 size:512 thread:1
allocated block device:00000007036E1200 size:512 thread:1
allocated block device:00000007036E1400 size:512 thread:1
allocated block device:00000007036E1600 size:512 thread:1
allocated block device:00000007036E1800 size:512 thread:1
allocated block device:00000007036E1A00 size:512 thread:1
allocated block device:00000007036E1C00 size:512 thread:1
allocated block device:00000007036E1E00 size:512 thread:1
allocated block device:00000007036E2000 size:512 thread:1
allocated block device:00000007036E2200 size:512 thread:1
allocated block device:00000007036E2400 size:512 thread:1
allocated block device:00000007036E2600 size:512 thread:1
allocated block device:00000007036E2800 size:512 thread:1
allocated block device:00000007036E2A00 size:512 thread:1
allocated block device:00000007036E2C00 size:512 thread:1
allocated block device:00000007036E2E00 size:512 thread:1
allocated block device:00000007036E3000 size:512 thread:1
allocated block device:00000007036E3200 size:512 thread:1
allocated block device:00000007036E3400 size:512 thread:1
allocated block device:00000007036E3600 size:512 thread:1
allocated block device:00000007036E3800 size:512 thread:1
allocated block device:00000007036E3A00 size:512 thread:1
allocated block device:00000007036E3C00 size:512 thread:1
allocated block device:00000007036E3E00 size:512 thread:1
allocated block device:00000007036E4000 size:512 thread:1
allocated block device:00000007036E4200 size:512 thread:1
allocated block device:00000007036E4400 size:512 thread:1
allocated block device:00000007036E4600 size:512 thread:1
allocated block device:00000007036E4800 size:512 thread:1
allocated block device:00000007036E4A00 size:512 thread:1
allocated block device:00000007036E4C00 size:512 thread:1
allocated block device:00000007036E4E00 size:512 thread:1
allocated block device:00000007036E5000 size:512 thread:1
allocated block device:00000007036E5200 size:512 thread:1
allocated block device:00000007036E5400 size:512 thread:1
allocated block device:00000007036E5600 size:512 thread:1
allocated block device:00000007036E5800 size:512 thread:1
allocated block device:00000007036E5A00 size:512 thread:1
allocated block device:00000007036E5C00 size:512 thread:1
allocated block device:00000007036E5E00 size:512 thread:1
allocated block device:00000007036E6000 size:512 thread:1
allocated block device:00000007036E6200 size:512 thread:1
allocated block device:00000007036E6400 size:512 thread:1
allocated block device:00000007036E6600 size:512 thread:1
allocated block device:00000007036E6800 size:512 thread:1
allocated block device:00000007036E6A00 size:512 thread:1
allocated block device:00000007036E6C00 size:512 thread:1
FATAL ERROR: variable in data clause is partially present on the device: name=(u
nknown)
file:F:\Files\SourceCode\OpenaccCExamples\acc_deepcopy_jagged\acc_deepcopy_jagg
ed.c main line:76
The utilization of the jagged 2d array follows the instructions from the chapter 5 of the book Parallel Programming with OpenACC. When testing the jagged 2d array separately, it is correct. When combine the 2d jagged array into user-defined structure, the exe can not run. How to solve this? Please help me. Thank you!
Guangyuan Kan