Using create() with void * blocks of memory

I have some code that I’m trying to get right on the data movement side.

I have an allocator that fills a field in a struct with a chunk of memory based on a computed size. I don’t have the type information of the chunk of memory at this point, and I’m not sure what to put into the create() clause of a #pragma acc enter data create(v) statement.

PGCC can’t figure out the bounds, so I am wondering whether I use the complete computed size or whether this is not supported? In all other code examples I have seen, you specify the count of the array that has a type associated with it, not a void *. So, as an example:

void aa(A*a,I s){
        void *v=malloc(s);
        #pragma acc enter data create(v[:s])
        a->v=v;
}

Is this the correct way to do this? Later on I expect to be able to do something like this:

int *v = a->v;
#pragma acc kernels loop present(v[:n])
for (...) { ... }

But I’ve been seeing some errors there about v not being present on the device? In this case the value n would be the number of elements in V, not the same as the s value used to allocate space.

Hi Aaron W. Hsu,

A “void” type has no size so when you create it, it will have a size of zero on the device. How about converting it to a char?

  • Mat
% cat test.c
#include <stdlib.h>
#include <stdio.h>

struct A {
   void * v;
   int size;
};

void aa(struct A*a,int s){
        void *v=(char*)malloc(s);
#ifdef _OPENACC
        char *vc=(char*)v;
        #pragma acc enter data create(vc[:s])
#endif
        a->v=v;
        printf("a->V = %p, size=%d \n",v,s);
}

int main () {

   struct A* a;
   int n = 1024;
   a = (struct A*) malloc(sizeof(struct A));
   a->size = n*sizeof(int);
   aa(a,a->size);

int *v = (int*) a->v;
printf("v = %p \n",v);
#pragma acc kernels loop present(v[:n])
for (int i=0; i < n; ++i) {
   v[i] = i;
}

#pragma acc update host(v[0:n])
  printf ("v[1] = %d  v[n-1]=%d\n", v[1], v[n-1]);

  free(a->v);
  free(a);
  exit(0);



}
% pgcc test.c -Minfo=accel -acc; a.out
aa:
     13, Generating enter data create(vc[:s])
main:
     29, Generating present(v[:n])
     30, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         30, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     35, Generating update host(v[:n])
a->V = 0x604170, size=4096
v = 0x604170
v[1] = 1  v[n-1]=1023