[amontoison]

Every scalar pointer argument now comes in two forms under the same generic:

• the bind(C) specific takes it as type(c_ptr), value, the C API's own spelling;
• a new <routine>_typed module procedure takes the Fortran variable (integer, real, complex or the enum kind) and passes its address.

Both forms resolve, so either of these calls works:

st = rocblas_daxpy(handle, n, alpha, dx, incx, dy, incy)          ! _typed
st = rocblas_daxpy(handle, n, c_loc(alpha), dx, incx, dy, incy)   ! bind(C)

This covers two kinds of scalar:

• Scalars read per the handle's pointer mode (alpha, beta, nnzTotalDevHostPtr, ...). The type(c_ptr) form is what makes the device pointer mode usable from Fortran.
• Host outputs (HIP runtime queries, getters, buffer sizes, descriptor getters, ...). Until now these were typed in the bind(C) specific itself (Update the classification of the arguments, and cover the query routines #571, Update hip, hipBLAS and rocBLAS bindings #572, 5de6093). They are back to type(c_ptr) there, as in hipfort 0.9.0, and keep their typed form in _typed.

Two more forms cover the calls neither of those can spell:

• <routine>_devptr, for the 42 rocSPARSE/hipSPARSE routines with both a pointer-mode scalar and a host output (alpha next to buffer_size). This is the device-pointer-mode call: the pointer-mode scalars as type(c_ptr), the host outputs typed.
• <routine>_rank_N_devptr, next to the hipSOLVER array overloads. In those overloads devInfo is now an integer, like the device arrays next to it, as in 0.9.0. The _devptr twins take it as a type(c_ptr).

devInfo gets no typed form on the type(c_ptr) interface, because a host integer next to type(c_ptr) buffers would be written by the GPU.

Compatibility with 0.9.0

728 of the 807 test programs of hipfort 0.9.0 now compile against these modules, up from 585 against develop. The 79 that still fail break at compile time, for six reasons; the CHANGELOG update will list them:

• hipSOLVER, 57 programs: an integer devInfo passed with type(c_ptr) buffers. Excluded on purpose.
• hipSOLVER array overloads, 12 programs: tau passed as dTau(1).
• *syevdx/*heevdx, 4 programs: nev passed as a c_ptr while lwork is typed.
• rocsolver_{c,z}hegvdx, 4 programs: nev is now a device pointer (fb8f244).
• hipFFTW guru, 1 program: dims passed as a scalar.
• ROCTx, 1 program: Fortran strings (different library).

One change still compiles but behaves differently: hipsolverRfBatch* now takes its array of host pointers by value.

Contents

• lib/hipfort, regenerated with rocm-fortran 9854af6.
• The bind(C) symbol tests, regenerated against these bindings (on top of Regenerate the bind(C) symbol tests #580, which regenerated them for the previous ones).
• The supported-API tables. Only the generated diff is applied: the variants column gains typed/devptr.

Validation

Against ROCm 10.1, with gfortran and amdflang:

• Compilation: the modules compile with no macro, with USE_ASSUMED_SHAPE, with USE_ASSUMED_RANK, and with USE_ASSUMED_SHAPE + USE_CUDA_NAMES. The current tests build 1245/1245.
• ctest with -DHIPFORT_EXTENDED_TESTS=ON: 1255/1258. The 3 failures are the known WSL ones (hip_mem_advise f2003/f2008, hip_module_kernel).
  • The shared-link tests resolve every symbol against ROCm 10.1.
  • Against CUDA 13.3 they also resolve without libcuda: the _typed forms of the 6 HIP routines mapped to the CUDA driver API exist on ROCm only.
• On a gfx1101 GPU: every 0.9.0 test program that builds behaves the same against the 0.9.0 library and against this one.
• Docs: Doxygen reports no new warning.

Generator fixes found along the way

• gfortran 13 passes a garbage byte when a by-reference character(c_char) is forwarded to a bind(C) value dummy. hipsolverSgesvd_bufferSize then returned INVALID_VALUE. The wrappers now take a C char by value.
• CUDA backend: no _typed form calls the CUDA driver API, so hipfort still needs only the CUDA runtime libraries.