Dynamically allocated jagged 2d array in structure

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

This just looks like a compiler bug to me. I filed FS#25981, and a compiler engineer will look into it.

  • Brent

Thank you for your reply!
By the way, do you have other solution to solve this? Such as using runtime APIs or other way of programming to get round this bug.