On Mon, Nov 02, 2015 at 06:33:39PM +0000, Julian Brown wrote: > Firstly, on trunk at least, use_device_ptr variables are restricted to > pointer or array types: that restriction doesn't exist in OpenACC, nor > actually could I find it in the OpenMP 4.1 document (my guess is the > standards are supposed to match in this regard). I think that a program > such as this should work:
So, after talking about this on omp-lang, it seems there is agreement that only arrays and pointer types (or reference to arrays or pointers) should be allowed in use_device_ptr clause and that for pointers/reference to pointers it should probably act the way I've coded it up, i.e. that for them it translates the pointer to point to corresponding object to the one to which it points on the host. It is too late to change the standard now, but will be changed soon, and hopefully clarified in examples. > void target_fn (int *targ_data); > > int > main (int argc, char *argv[]) > { > char out; > int myvar; > #pragma omp target enter data map(to: myvar) > > #pragma omp target data use_device_ptr(myvar) map(from:out) > { > target_fn (&myvar); > out = 5; > } > > return 0; > } That would make the above non-conforming for OpenMP. > Secondly, attempts to use use_device_ptr on (e.g. > dynamically-allocated) arrays accessed through a pointer cause an ICE > with the existing trunk OpenMP code: > > #include <stdlib.h> > > void target_fn (char *targ_data); > > int > main (int argc, char *argv[]) > { > char *myarr, out; > > myarr = malloc (1024); > > #pragma omp target data map(to: myarr[0:1024]) > { > #pragma omp target data use_device_ptr(myarr) map(from:out) > { > target_fn (myarr); > out = 5; > } > } > > return 0; > } Can't reproduce this ICE (at least not on gomp-4_5-branch, but there aren't significant changes from the trunk there). > Furthermore, this looks strange to me (006t.omplower): > > .omp_data_arr.5.out = &out; > myarr.8 = myarr; > .omp_data_arr.5.myarr = myarr.8; > #pragma omp target data map(from:out [len: 1]) use_device_ptr(myarr) > { > D.2436 = .omp_data_arr.5.myarr; > myarr = D.2436; > > That's clobbering the original myarr variable, right? Just use -fdump-tree-omplower-uid to see that it is a different variable. Basically, for OpenMP use_device_ptr creates a private copy of the pointer for the body of the target data construct, and that pointer is assigned the target device's address. For arrays the implementation creates an artificial pointer variable (holding the start of the array initially) and replaces all references to the array in the target data body with dereference of the pointer. Jakub