Skip to content

Document and test the typed and type(c_ptr) forms of scalar arguments - #583

Merged
amontoison merged 2 commits into
developfrom
typed_forms_docs_tests
Oct 7, 2026
Merged

amontoison merged 2 commits into
developfrom
typed_forms_docs_tests

Conversation

@amontoison

@amontoison amontoison commented Oct 7, 2026 •

Copy link
Copy Markdown
Collaborator

Follow-up to #581: documents and tests the typed and type(c_ptr) forms it introduced.

Docs

docs/how-to/fortran-interfaces.rst gains a section Scalar arguments passed by pointer that explains:

  • when to pass the Fortran variable (_typed) and when to pass its address (type(c_ptr)): device address, c_null_ptr for an output that can be skipped, or existing c_loc code;
  • the device pointer mode:
    • the routines that mix a pointer-mode scalar with a host output (X_devptr);
    • the array overloads, whose alpha/beta are typed, so the variable itself has to live on the device;
  • hipSOLVER devInfo: type(c_ptr) with type(c_ptr) buffers, and an integer(c_int), pointer or a type(c_ptr) with the array overloads;
  • how the X_typed/X_devptr names that appear in the reference relate to the generic.

The examples compile against the #581 bindings with gfortran and amdflang. The RST is checked with docutils (beyond the Sphinx :doc: role).

Tests

Until now, alpha/beta only went through the c_loc(alpha) form in CTest. Eight tests now take an existing check through the _typed form, with alpha/beta passed as named constants:

Library Tests
rocBLAS dgemm_typed, caxpy_typed
hipBLAS sgemm_typed, zaxpy_typed
rocSPARSE rocsparse_dgemvi_typed, rocsparse_dcsrgemm_typed (alpha, beta and buffer_size all typed)
hipSPARSE hipsparse_dgemvi_typed, hipsparse_zcsrsv2_typed
  • Each object references exactly the _typed procedures it targets (checked with nm).
  • The two tests that copied an alpha of 1 use 2 and (2,1) instead, so a lost or misread alpha fails the check.
  • The _devptr form is already covered by the *_devptr tests (rocSPARSE dcsrgemm/csrgeam...).
  • The four f2008 hipsolver_?getrf tests now pass dInfo itself (an integer(c_int), pointer from hipMalloc) instead of c_loc(dInfo), which puts the typed devInfo of the array overloads under CTest and fixes a comment that said "by reference". The other f2008 hipSOLVER tests keep c_loc(dInfo) and cover _rank_N_devptr. hipsolver_dgetrf.f08 is the program the hipSOLVER tutorial shows.

Validation

  • gfortran, ROCm 10.1, gfx1101: ctest -R '_typed$' gives 8/8.
  • amdflang: all 8 compile.

@amontoison
amontoison requested a review from a team as a code owner October 7, 2026 03:37
Base automatically changed from fortran_typed_variants to develop October 7, 2026 04:16
#581 gives every scalar pointer argument two forms under one generic: the
bind(C) specific takes a type(c_ptr), the _typed module procedure the
Fortran variable. Nothing in the user documentation said so, and the test
suite only called the type(c_ptr) form for alpha/beta (with c_loc, since
#564), so the _typed path was not covered by CTest.

docs/how-to/fortran-interfaces.rst gains a section "Scalar arguments
passed by pointer": when to pass the variable and when its address, the
device pointer mode (including the routines that mix a pointer-mode scalar
with a host output, and the array overloads, whose alpha/beta are typed),
and hipSOLVER's device-only devInfo. Its examples compile against these
bindings with gfortran and amdflang.

Eight tests take an existing check through the _typed form, with alpha and
beta passed as named constants instead of c_loc(): rocBLAS dgemm and
caxpy, hipBLAS sgemm and zaxpy, rocSPARSE dgemvi and dcsrgemm (alpha, beta
and buffer_size all typed), hipSPARSE dgemvi and zcsrsv2. Each object file
references the _typed procedures it is meant to (checked with nm), and the
two copies whose alpha was 1 use 2 and (2,1) instead, so a lost or
misread alpha fails the check.
dInfo is an integer(c_int), pointer allocated with hipMalloc, so it lives
on the device like the arrays beside it, and the array overloads take it
as is. The four tests passed c_loc(dInfo), which resolves to the _devptr
twin, under a comment saying "devInfo by reference". They now pass dInfo,
which also puts the typed devInfo of the array overloads under CTest; the
other f2008 hipSOLVER tests still pass c_loc(dInfo) and cover _devptr.
hipsolver_dgetrf.f08 is the program the hipSOLVER tutorial shows.
@amontoison
amontoison force-pushed the typed_forms_docs_tests branch from 6365348 to 6d2910b Compare October 7, 2026 04:23
@amontoison
amontoison merged commit fea0a22 into develop Oct 7, 2026
2 of 3 checks passed
@amontoison
amontoison deleted the typed_forms_docs_tests branch October 7, 2026 13:12
amontoison added a commit that referenced this pull request Oct 7, 2026
Regenerated with rocm-fortran 3b65cdb.

Typed forms for scalar pointers that still needed c_loc(): the BLAS
dot/nrm2/asum/iamax/iamin results and rotg/rotmg scalars (rocBLAS and
hipBLAS) and hipBLAS rot's c/s; the sparse doti/dotci results, nnz
counts, pivot positions and csr2ell_width's ell_width; hipBLAS
getrs/geqrf/gels info (host); hipChooseDevice. The array overloads take
the BLAS results and hipBLAS info typed and gain _rank_N_devptr twins
that keep the type(c_ptr), so existing c_loc() calls still compile.

Fixes:
- rocSOLVER larf/larfg alpha/tau are device scalars and lose their
  _typed form, which would have handed the kernel a host address;
- hipsolverRfBatchZeroPivot position (an array) and
  rocblas_set_solution_fitness_query fitness (a retained pointer) lose
  theirs too;
- csrcolor's coloring/reordering are arrays again in the array
  overloads;
- rocSPARSE's nullable `error` arguments are optional;
- hipMalloc(..., source=) and its siblings report an allocation failure
  instead of the copy's status;
- the compat hipsolver*gels/gesvd_bufferSize are no longer bound to
  cusolverDn symbols with different arguments under USE_CUDA_NAMES;
- hipChooseDevice is a generic over hipChooseDeviceR0600's specifics
  instead of a second interface on the same binding label.

Three tests take a result through the new typed forms: rocBLAS ddot,
hipBLAS snrm2 and rocSPARSE ddoti (host pointer mode).

Also regenerates the supported-API tables (only the generated diff: the
variants column gains typed and devptr), and updates the changelog with
the Scalar arguments section and tests of #583, the typed forms above and
the fixes.
amontoison added a commit that referenced this pull request Oct 9, 2026
Regenerated with rocm-fortran 3b65cdb.

Typed forms for scalar pointers that still needed c_loc(): the BLAS
dot/nrm2/asum/iamax/iamin results and rotg/rotmg scalars (rocBLAS and
hipBLAS) and hipBLAS rot's c/s; the sparse doti/dotci results, nnz
counts, pivot positions and csr2ell_width's ell_width; hipBLAS
getrs/geqrf/gels info (host); hipChooseDevice. The array overloads take
the BLAS results and hipBLAS info typed and gain _rank_N_devptr twins
that keep the type(c_ptr), so existing c_loc() calls still compile.

Fixes:
- rocSOLVER larf/larfg alpha/tau are device scalars and lose their
  _typed form, which would have handed the kernel a host address;
- hipsolverRfBatchZeroPivot position (an array) and
  rocblas_set_solution_fitness_query fitness (a retained pointer) lose
  theirs too;
- csrcolor's coloring/reordering are arrays again in the array
  overloads;
- rocSPARSE's nullable `error` arguments are optional;
- hipMalloc(..., source=) and its siblings report an allocation failure
  instead of the copy's status;
- the compat hipsolver*gels/gesvd_bufferSize are no longer bound to
  cusolverDn symbols with different arguments under USE_CUDA_NAMES;
- hipChooseDevice is a generic over hipChooseDeviceR0600's specifics
  instead of a second interface on the same binding label.

Three tests take a result through the new typed forms: rocBLAS ddot,
hipBLAS snrm2 and rocSPARSE ddoti (host pointer mode).

Also regenerates the supported-API tables (only the generated diff: the
variants column gains typed and devptr), and updates the changelog with
the Scalar arguments section and tests of #583, the typed forms above and
the fixes.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant