3

We have a project that is written in Fortran.

Now I know this can be done using PGI compilers but I don't want to get stuck with licenses.

I am trying to see whether we could use OpenACC in our project. I got gcc5.2 installed using instructions here.

https://github.com/olcf/OLCFHack15

Now I want to do something similar to what is stated here.

https://gcc.gnu.org/onlinedocs/libgomp/OpenACC-Library-Interoperability.html

More specifically what is stated in section 8.3. I am trying to exactly reproduce it using gfortran. Unfortunately I don't see how I can do it in fortran. In the example, 

d_X = acc_copyin(&h_X[0], N * sizeof (float));

This allows d_X to be directly used in 

s = cublasSaxpy(h, N, &alpha, d_X, 1, d_Y, 1);

But in fortran, the acc_copyin does not return anything.

So how would I replicate the case in Fortran?

Vladimir F Героям слава
  • 57,977
  • 4
  • 76
  • 119
Vikram
  • 308
  • 1
  • 5
  • Looking at the specification it seems that the Fortran interface of OpenACC is pretty limited. It avoids any pointer stuff, even when it would have been possible using `type(c_ptr)`. I am not sure if it is possible to call the C function using `bind(C)`? – Vladimir F Героям слава Apr 12 '16 at 16:23

1 Answers1

1

Are you looking to interface with cuBLAS or is this more general? cuBLAS does provide a F77 style interface (See: http://docs.nvidia.com/cuda/cublas/#appendix-b-cublas-fortran-bindings)

The OpenACC solution is to manage your data as you normally would using the "data" directive, but then call the CUDA C routine from within a "host_data" region. "host_data" specifies that the device pointer should be used with this region. Hence when passing "d_X" to cublasSaxpy, the device pointer will be passed in.

On caveat with cuBLAS is that the F77 interface mentioned above is expecting host arrays and will manage the data movement for you. Hence you'll need to write CUDA C wrapper functions to call the correct device routines. (CUDA Fortran does provide a cublas module for this but is PGI only)

Though, GNU 5.2 doesn't support "host_data" and I just looked on their status page (https://gcc.gnu.org/wiki/OpenACC) and it doesn't look like it will be supported in Fortran in 6.0 either. Which is unfortunate since "host_data" is the best solution for you.

Note that NVIDIA does give free PGI license to students and academics for teaching purposes as part of the OpenACC Toolkit (See: https://developer.nvidia.com/openacc).

Mat Colgrove
  • 5,441
  • 1
  • 10
  • 11
  • Maybe I interpret the question wrong, but I see it as: Can I get the device pointer for the memory I allocated (and data copied) by the Fortran OpenACC routine `acc_copyin`? As I go through the OpenACC specification they basically ignored anything which has something to do with pointers and memory allocation in the Fortran interface, thus crippling it considerably. I would try writing a custom Fortran interface to the more potent C `acc_` routines, but I don't know if they can be just called or there is some more magic going on behind the scenes. – Vladimir F Героям слава Apr 13 '16 at 18:07
  • I don't think it's an oversight since handling raw C pointers isn't easily done in Fortran. It's easier in CUDA Fortran since the language was extended to include the notion of device data, but in regular Fortran it's a bit difficult. Do you have a use case? If so, my suggestion would be to send a note directly to the OpenACC committee (feedback@openacc.org) asking for this to be added. They are always looking for good suggestions and ideas. – Mat Colgrove Apr 13 '16 at 22:25
  • If you don't do pointer arithmetics, but just store and pass an address `(void*)` you can use C pointers directly as `type(c_ptr)` in Fortran. Another option is `integer(c_intptr_t)`, but is worse. Both available for more than 10 years in the language. If the device pointer is not compatible with `void*`, that would be a different story. CUDA Fortran... isn't the point of OpenACC to avoid directly programming in CUDA? As I wrote one can call the C `acc_malloc` from Fortran if it is just a function, but I don't know if there is not something more going on behind the scenes. – Vladimir F Героям слава Apr 13 '16 at 22:39
  • So it just confirms what I initially thought. This is not going to be possible with gfortran unless I start writing wrapper code. For now, I'll try to do it using PGI compilers. – Vikram Apr 14 '16 at 10:06