PGF90-S-0148-Reference to TYPE(C_PTR) expression required

I am looking at a code which was programmed several years ago with GPU support and it throws up the above warning.
I can reproduce it with the example at:
https://www.olcf.ornl.gov/tutorials/mixing-openacc-with-gpu-libraries/
$ pgf90 -acc -c fft.f90
PGF90-S-0148-Reference to TYPE(C_PTR) expression required (fft.f90: 41)
0 inform, 0 warnings, 1 severes, 0 fatal for fft
It looks like something has changed in between PGI versions?
I am currently trying PGI 16.10 and CUDA 8.0.
(The C version is fine)

Hi RikaK,

The code has a slight error in that “acc_get_cuda_stream” is expecting “stream” to be a “integer(acc_handle_kind)” type, not a C_PTR.

Hope this helps,
Mat

% cat fft.f90
module cufft
  INTERFACE
    subroutine launchcufft(data, n, stream) BIND (C, NAME='launchCUFFT')
      USE ISO_C_BINDING
      USE openacc
      implicit none
      type (C_PTR), value :: data
      integer (C_INT), value :: n
      integer(acc_handle_kind), value :: stream
    end subroutine
  END INTERFACE
end module cufft

program fft
    USE ISO_C_BINDING
    USE cufft
    USE openacc
    IMPLICIT NONE

    INTEGER, PARAMETER :: n = 256
    COMPLEX (C_FLOAT_COMPLEX) :: data(n)
    INTEGER (C_INT):: i
    INTEGER :: max_id,istat
    integer(acc_handle_kind) :: stream

    ! Initialize interleaved input data on host
    REAL :: w = 7.0
    REAL :: x
    REAL, PARAMETER :: PI = 3.1415927
    do i=1,n
        x = (i-1.0)/(n-1.0);
        data(i) = CMPLX(COS(2.0*PI*w*x),0.0)
    enddo

    ! Copy data to device at start of region and back to host and end of region
    !$acc data copy(data)

        ! Inside this region the device data pointer will be used
        !$acc host_data use_device(data)
        stream = acc_get_cuda_stream(acc_async_sync)
        call launchcufft(C_LOC(data), n, stream)
        !$acc end host_data

    !$acc end data

    ! Find the frequency
    max_id = 1
    do i=1,n/2
        if (REAL(data(i)) .gt. REAL(data(max_id))) then
            max_id = i-1
        endif
    enddo
    print *, "frequency:", max_id

end program fft


% nvcc -c cufft.cu
% pgfortran -Mcudalib=cufft -acc fft.f90 cufft.o -V16.5
fft.f90:
% a.out
 frequency:            7

Hi Mat,
Thanks. Got past that in situ in the original code but it appears to have knock-on effects…
PGF90-S-0188-Argument number 2 to cublassetstream_v2: type mismatch (/short/z00/rxk900/DALTON2016.2-Source/lsdalton/src/deccc/cc_tools.F90: 664)
0 inform, 0 warnings, 1 severes, 0 fatal for get_a22_and_prepb22_terms_ex
coming from
dummy47 = acc_get_cuda_stream(acc_h(curr_id))
stat = cublasSetStream_v2(cub_h(curr_id), dummy47)

I’m assuming this code must have worked at some point to give:
https://developer.nvidia.com/openacc-toolkit
Might have to see if there is a more up-to-date source (this is attempting to work from the latest official release).
Rika

Hi Rika,

I’m thinking that LSDALTON has it’s own interface to cublasSetStream_v2. Can you look to see what that interface is expecting?

PGI does provide a interface module to cuBLAS as well, “use cublas”, but with ours you’d call cublasSetStream, which in turn binds to the CUDA cublasSetStream_v2 routine. These interfaces are documented in: http://www.pgroup.com/doc/pgi17cudaint.pdf.

-Mat

Hi Mat,
Yep, you’re right:
src/cuda/gpu_interfaces.F90: integer (C_INT) function cublasSetStream_v2(handle,stream) bind(C,name=“cublasSetStream_v2”)

I should have looked/thought harder. This looks like it is going to require more effort than I had originally planned.
What happened to the good old days of “config make”.
Rika