Pointer to fixed size array does not work

When pointer to fixed size array like

#define DIM 3
double (*Position)[DIM];

was accessed after memory allocation, the memory error as

Failing in Thread:1
call to cuStreamSynchronize returned error 700: Illegal address during kernel execution

happened. The specific code is as follows;

#define DIM 3

static double (*Position)[DIM];
#pragma acc declare create(Position)

const int N=10;
const int W=3;

static void allocatePosition( void ){
Position=(double (*)[DIM])malloc(sizeof(double [DIM])*N);
#pragma acc enter data create(Position[0:N][0:DIM])
}

static void calculatePosition( void ){
#pragma acc kernels
#pragma acc loop independent
for(int iP=0;iP<N;++iP){
Position[iP][0]=iP/W;
Position[iP][1]=iP%W;
Position[iP][2]=0;
}
}

static void outputPosition( void ){
#pragma acc update host(Position[0:N][0:DIM])
for(int iP=0;iP<N;++iP){
printf(“%d: %e %e %e\n”,iP,Position[iP][0],Position[iP][1],Position[iP][2]);
}
}

int main(int argc, char *argv){
allocatePosition();
calculatePosition();
outputPosition();
return 0;
}

This was observed at least with the compiler versions of 22.1 and 22.5.
Sorry for missing compiler information.

Hi MasahiroKondo64929,

In this case, you need to attach to the global “Position” since it’s already present on the device.

static void allocatePosition( void ){
Position=(double (*)[DIM])malloc(sizeof(double [DIM])*N);
#pragma acc enter data create(Position[0:N][0:DIM]) attach(Position)
}

Hope this helps,
Mat

Thank you. The program correctly run with the directive “attach”.

Dear Mat Colgrove

About one month ago, the code including global pointer to fixed size array with “attach” directive was tested and it seemed good with HPC-SDK 22.5 (pgc++). However, recently, the code has been tested again but this time with HPC-SDK 21.3 (pgc++), and a strange compiler message was found. In specific, the following message was repeated 6 times in the end of compiling.

invalid subroutine type ref
!9 = !DISubroutineType(types: !10)
!10 = !{!11}
!11 = !{i32 24}

What is this? The test code with “attach” directive is as follows;

using namespace std;

#define DIM 3

static double (*Position)[DIM];
#pragma acc declare create(Position)

const int N=10;
const int W=3;

static void allocatePosition( void ){
	Position=(double (*)[DIM])malloc(sizeof(double [DIM])*N);
	#pragma acc enter data create(Position[0:N][0:DIM]) attach(Position)
}
	
static void calculatePosition( void ){
	#pragma acc kernels
	#pragma acc loop independent
	for(int iP=0;iP<N;++iP){
		Position[iP][0]=iP/W;
		Position[iP][1]=iP%W;
		Position[iP][2]=0;
	}
}

static void outputPosition( void ){
	#pragma acc update host(Position[0:N][0:DIM])
	for(int iP=0;iP<N;++iP){
		printf("%d: %e %e %e\n",iP,Position[iP][0],Position[iP][1],Position[iP][2]);
	}
}

int main(int argc, char *argv[]){
	allocatePosition();
	calculatePosition();
	outputPosition();
	#pragma acc exit data delete(Position[0:N][0:DIM]) detach(Position)
	
	return 0;
}

In upgrading my main code:

HPC-SDK/22.x has another problem: Shared memory is not correctly used in kernels block, but I recently found that it is not emerged in HPC-SDK/21.3. So, firstly, I’m considering to update my code for HPC-SDK/21.3.

regards,

Masahiro Kondo

It was some extraneous output from the LLVM code generation that an engineer accidently left in. It was fixed in 21.5 but should be benign.

For the shared memory issue in 22.5, were you able to use the work around I provided, i.e. hoist the declaration of “previous” and then manually privatize it?

-Mat

Dear Mat

Thank you for replying.
I’ve relieved to know that it’s no problem.

Regarding the shared memory issue. I confirmed that the way you prvided worked in the test code. But, I have not try it in my main code because there are so many such places to be mended. I think it’s better for me to wait for the compiler to overcome the issue.

regards,

Masahiro Kondo