Reputation: 89
I'm trying to handle a global variable declared in this way in head.h:
extern int NX1_TOT;
#pragma acc declare create(NX1_TOT)
While in globals.h:
int NX1_TOT;
in main.c:
#include "head.h"
#include "globals.h"
In the file initialize.c I set the NX1_TOT value:
NX1_TOT=ntot[IDIR];
What happens is that if I try to print the value of NX1_TOT
in different kernels, I get different and sometimes wrong results.
I tried with #pragma acc update device
but it doesn't seem to make a difference.
I just want to be able to use a variable (that is set at the beginning of the application and never modified) in every device kernel without having to pass it as an argument to every routine (seems a very simple problem to me but the documentation is null).
Upvotes: 0
Views: 46
Reputation: 5646
Can you please provide a minimal reproducing test case? This should work as expected, but without an example we can't tell what the code is doing.
Here's a simple working example you can start with and then modify to show how your code is doing the operations.
% cat head.h
#include <stdlib.h>
#include <stdio.h>
extern int NX1_TOT;
#pragma acc declare create(NX1_TOT)
% cat globals.h
int NX1_TOT;
% cat test.c
#include "head.h"
#include "globals.h"
int main() {
int *A;
A=(int*) malloc(sizeof(int)*1024);
NX1_TOT=2;
#pragma acc update device(NX1_TOT)
#pragma acc parallel loop copyout(A[:1024])
for (int i=0; i<1024; ++i) {
A[i] = NX1_TOT;
}
printf("%d %d\n", A[10],A[100]);
free(A);
}
% nvc test.c -acc -Minfo=accel; a.out
main:
10, Generating update device(NX1_TOT)
Generating copyout(A[:1024]) [if not already present]
Generating NVIDIA GPU code
14, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
2 2
Upvotes: 0