Skip to content

[Sandbox] Move the Fortran bindings to one module per library (ROCm only) - #540

Open
amontoison wants to merge 5 commits into
developfrom
rocm-fortran
Open

amontoison wants to merge 5 commits into
developfrom
rocm-fortran

Conversation

@amontoison

@amontoison amontoison commented Sep 18, 2026

Copy link
Copy Markdown
Collaborator

A sandbox preview, not for merge.
It shows what the Fortran bindings look like once they are co-located with the C headers they wrap, so the shape can be discussed before anything is committed to.

The layout

lib/hipfort is a single flat directory of 33 files today.
This replaces it with one directory per library, grouped by upstream repository:

lib/rocm-systems/hip
lib/rocm-systems/roctx
lib/rocm-libraries/hipblas
lib/rocm-libraries/rocsolver
lib/rocm-libraries/hipsolver
lib/rocm-libraries/rocsparse
lib/rocm-libraries/hipsparse
lib/rocm-libraries/rocfft
lib/rocm-libraries/hipfft
lib/rocm-libraries/hipfftw
lib/rocm-libraries/rocrand
lib/rocm-libraries/hiprand

Each library is now a single self-contained module.
What used to be scattered across sibling modules (the enums, the derived types, the opaque handles, the check helper) is folded in.
--> 12 files replace 33.

@amontoison
amontoison force-pushed the rocm-fortran branch 2 times, most recently from 7794245 to ab3686f Compare September 20, 2026 16:04
@amontoison amontoison changed the title Split the Fortran bindings into one directory per library [Sandbox] Move the Fortran bindings to one module per library (ROCm only) Sep 20, 2026
@amontoison
amontoison force-pushed the rocm-fortran branch 2 times, most recently from bf7a1d1 to fd2fee6 Compare September 20, 2026 22:55
@amontoison
amontoison requested a review from a team as a code owner September 20, 2026 22:55
Replace the single flat lib/hipfort/ with the layout the bindings would have
once they are co-located with the C headers they wrap, one directory per
library grouped by upstream repository:

  lib/rocm-systems/{hip,roctx}/
  lib/rocm-libraries/{rocblas,hipblas,rocsolver,hipsolver,rocsparse,hipsparse,
                      rocfft,hipfft,hipfftw,rocrand,hiprand}/

The module names, the generic names and the bind(C) symbols are unchanged, so
no user code and no test needs to be touched. Each library now ships its module
plus two included fragments (<mod>_interfaces.f90, <mod>_wrappers.f90), which
keeps a single library's surface reviewable on its own.

lib/CMakeLists.txt changes accordingly:

  * the source list is built per directory instead of from one flat glob, and
    excludes *_interfaces.f90 / *_wrappers.f90. Those are INCLUDEd by their
    module, not compilation units: compiling them standalone fails, since they
    are bare interface and contains bodies with no enclosing module.
  * every library directory is added to the include path. gfortran resolves an
    INCLUDE relative to the including file, but that is not guaranteed by the
    standard, so the directories are named explicitly for other compilers.
  * hipfort_cuda_errors.f90 is hand-written rather than generated; it moves to
    rocm-systems/hip/ with the rest of the HIP runtime.

Validated on gfortran with -DHIPFORT_ASSUMED_RANK=ON: 901 Fortran objects build
with no error, and ctest reports 843 of 863 passing on a gfx1101 GPU. The 20
failures are identical to those the flat bindings produce, and each was checked
to reproduce against upstream's own unmodified lib/hipfort: 16 rocSPARSE
csrgeam/csrgemm SIGSEGV inside librocsparse's nnz phase, 1 sddmm_batched
returning rocsparse_status 2 (not implemented), 2 hipMemRangeGetAttribute
ReadMostly unsupported on this device, 1 hipModuleLaunchKernel returning 719.

hipBLAS is skipped at configure time on this machine because its config package
cannot find hipblas-common, so its 79 tests are not covered by the run above.
Replace the flat lib/hipfort/ with the shape the bindings take once they are
co-located with the C headers they wrap: one self-contained module per library,
named after the library, grouped by upstream repository.

  lib/rocm-systems/hip/
  lib/rocm-libraries/{rocblas,hipblas,rocsolver,hipsolver,rocsparse,hipsparse,
                      rocfft,hipfft,hipfftw,rocrand,hiprand}/

Each module folds in what used to be scattered across sibling modules: the
enums, the derived types, the opaque handles and the <lib>Check helper. Twelve
files replace sixty-three, and a library's whole Fortran surface is one file.
The only cross-library dependency left is rocsolver naming six rocBLAS entities
in an `only` clause, which mirrors rocsolver.h including rocblas.h.

This renames the modules: hipfort_hipblas and its _enums/_types become hipblas,
the HIP runtime modules all become hip. Every test's `use` lines were rewritten
accordingly (3728 rewritten, 1543 now-duplicate uses collapsed). No procedure,
constant or bind(C) symbol name changes, so only the `use` lines move.

ROCm only. The generator loads the CUDA symbol maps solely for the flat layout,
so these sources carry no USE_CUDA_NAMES branch at all. Rather than let an
"nvptx" build silently produce an archive still bound to hip*/roc* symbols,
HIPFORT_BUILD_NVPTX now raises a fatal error, and hipfort_cuda_errors.f90 goes
with the rest of the CUDA support. Teaching the per-library layout to emit the
CUDA bindings is a separate piece of work.

rocTX is left out. It exists as two incompatible APIs: roctracer/roctx.h with
libroctx64 exports versioned symbols (roctxMarkA@@ROCTX_4.1) and declares
roctx_version_major/minor, while rocprofiler-sdk-roctx/roctx.h with
librocprofiler-sdk-roctx exports plain symbols, drops those two and adds
roctxProfilerPause/Resume, roctxGetThreadId and the roctxName* family.
hipfort::roctx links the latter while the bindings are generated from the
former. Choosing a side is an API decision of its own and does not belong in a
layout change, so no rocTX module ships here and its tests are guarded off with
their sources left in place.

Validated with gfortran, -DHIPFORT_ASSUMED_RANK=ON -DHIPFORT_EXTENDED_TESTS=ON:
every target builds, and ctest reports 855 of 875 passing on a gfx1101 GPU. The
20 failures are identical to those the flat bindings produce, and each was
checked to reproduce against upstream's own unmodified lib/hipfort: 16 rocSPARSE
csrgeam/csrgemm SIGSEGV inside librocsparse's nnz phase, 1 sddmm_batched
returning rocsparse_status 2 (not implemented), 2 hipMemRangeGetAttribute
ReadMostly unsupported on this device, 1 hipModuleLaunchKernel returning 719.
The same configuration under amdflang gives the same 20.

hipBLAS is skipped at configure time on this machine because its config package
cannot find hipblas-common; its tests were compiled and linked by hand instead
(137 of 137 clean).
The prose still told readers to `use hipfort_rocblas` and to reach for
`hipfort_check`, neither of which exists in this layout. Updating it makes
the proof-of-concept complete: a reader can follow the tutorials against
these bindings without hitting a module that is not there.

Most of it is the rename, but three groups of sentences described a
structure that no longer exists and had to be rewritten rather than
substituted:

* The `<lib>Check` helpers are no longer in a shared `hipfort_check`
  module; each one travels with the library it checks. `rocblasCheck` now
  comes from `rocblas`, `hipsparseCheck` from `hipsparse`, and only
  `hipCheck` from `hip`. A blind rename would have pointed every one of
  them at `hip`.
* The `<lib>_enums` modules are folded in, so sentences of the form "the
  enumerations come from `hipfort_rocblas_enums`, which `hipfort_rocblas`
  re-exports" degenerated into a module re-exporting itself.
* The HIP page listed where each piece lived: enumerators in
  `hipfort_enums`, types in `hipfort_types`, `hipGetDeviceProperties` in
  `hipfort_auxiliary`, and so on. That whole inventory collapses to one
  sentence now that `hip` holds all of it.

Two `hipfort_` strings are deliberately untouched: `-Dhipfort_ROOT` is a
CMake variable, and `hipfort_test_f2008_hipblas_dgemm` is a ctest target
name. Neither is a Fortran module.

The generated supported-API tables under docs/doxygen/input/ still carry
`hipfort_`-prefixed anchors. They are produced from the bindings, so they
follow whenever the generator is pointed at this layout.
Three globs and two snippets still named `lib/hipfort/`, which this branch
deletes. None of them fail loudly: a glob that matches nothing produces an
empty file list, and an empty file list builds fine.

* `docs/conf.py` preprocessed `../lib/hipfort/*.[fF]90` into
  `docs/doxygen/input/`, which the Doxyfile then reads. The glob now matched
  zero files, so the whole API reference would have been generated from no
  sources at all.
* `test/openmp/CMakeLists.txt` compiled the bindings straight from the same
  path. That project is standalone, not reached by `add_subdirectory`, so the
  main build never noticed that `test_rocsolver_zhegvdx` had lost every module
  it depends on.
* `README.md` still opened both of its Fortran examples with `use hipfort`,
  and `docs/README.md` documented the old preprocessing path.

Also adds the changelog entry for the restructure itself, which had none, and
stops the ROCTx entries from naming a `hipfort_roctx` module that no longer
exists under that name.

An empty `lib/hipfort/` may linger in a working tree that predates the move;
git does not track directories, so it is only a local artefact.
The tree was generated with the convenience specifics on: native array overloads,
typed handles, character(len=*) forms. That is 803k lines of Fortran, and it makes
this branch hard to compare against what hipfort ships.

Regenerated with NATIVE/TYPED/CSTR=0, the same flags regenerate.sh uses for the
compatibility track, so each module now exposes the plain bind(C) surface and
nothing else. 264k lines. rocblas.F90 alone drops from 179k to 67k.

The derived types go with them, and that is not a loss: emit_types only fires when
some signature references a type, and with handles as type(c_ptr) and buffers as
type(c_ptr) none does. hipfort ships three *_types modules for the same reason.

Adding a convenience back later is a small change; removing one from an API people
already call is not. So this track starts at parity with hipfort and grows from
there.

The generated symbol tests are regenerated to match: without a generic, a routine
is a bare bind(C) named after the C symbol rather than a <name>_ specific, so the
old tests referenced names that no longer exist.

Built with -DHIPFORT_EXTENDED_TESTS=ON and run on a gfx1101 GPU: 992 of 1011
tests pass, and every failure is in the environment set this branch already
documents.

This branch has not been deployed

No deployments
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