Repository navigation
Support scalars on host and device (when supported) - #564
Merged
Merged
Conversation
amontoison
force-pushed
the
fixes-10.1
branch
from
September 30, 2026 05:46
c38f822 to
da86ae9
Compare
amontoison
added a commit
to amontoison/hipfort
that referenced
this pull request
Sep 30, 2026
ROCm#564 turned scalar and out-pointer arguments into type(c_ptr),value in the raw C bindings, but the test updates did not follow the two calling conventions the generics now expose: * the raw binding takes every pointer as type(c_ptr), so device-pointer call sites must wrap host scalars in c_loc() and the target must be a variable with the TARGET attribute, not a named constant; * the _rank_N/_full_rank/_assumed_rank wrappers still take host scalars and call c_loc() themselves, so Fortran-array call sites must keep passing the scalar directly. There is no overload mixing a c_ptr scalar with a Fortran array, so the f2008/f2018 tests, which pass Fortran pointer arrays, keep host scalars. * f2008/f2018 hipblas, rocblas, hipsparse, rocsparse: drop the c_loc() added around alpha/beta; keep them as TARGET variables since the wrapper now takes their address. * f2003 hipsparse/rocsparse: add the missing TARGET attribute the new c_loc() call sites require, and wrap nnz_C for xcsrgemmNnz. * hipsolver (f2003 + f2008), untouched by ROCm#564 although tau, devInfo and nev changed: wrap the device scalars in c_loc(), pass tau as a whole array where the array wrapper is selected, and leave nev alone in the *_bufferSize calls, which still bind it by reference. * rocblas dgemv_batched: c_ptr batch arrays select the raw binding, so alpha/beta need c_loc() there. Verified against ROCm 10.2 amdflang on gfx1101, with HIPFORT_ASSUMED_RANK both OFF and ON: the suite builds clean and the remaining 19 ctest failures are identical to develop's.
amontoison
added a commit
to amontoison/hipfort
that referenced
this pull request
Sep 30, 2026
The coverage tests under HIPFORT_EXTENDED_TESTS call every raw binding directly with one placeholder per type (v0 = c_ptr, and one each for real/complex/integer). ROCm#564 turned the scalar dummies into type(c_ptr),value, so 2681 of those call sites had to switch to v0. Two arguments needed the opposite move. hipsolverDsyevdx_bufferSize's nev and hipsolverSpDcsrlsvcholHost's singularity were the only members of their families still declared type(c_ptr),value; ROCm#564 normalises them to integer(c_int), matching every sibling. Both forms pass an address, so this is a source-form change only, and the two call sites now use the integer placeholder. Verified with ROCm 10.2 amdflang on gfx1101 and with gfortran 13, both with HIPFORT_EXTENDED_TESTS=ON: the suite builds clean and the 19 remaining ctest failures out of 1012 are identical to develop's.
#564 turned scalar and out-pointer arguments into type(c_ptr),value in the raw C bindings, but the test updates did not follow the two calling conventions the generics now expose: * the raw binding takes every pointer as type(c_ptr), so device-pointer call sites must wrap host scalars in c_loc() and the target must be a variable with the TARGET attribute, not a named constant; * the _rank_N/_full_rank/_assumed_rank wrappers still take host scalars and call c_loc() themselves, so Fortran-array call sites must keep passing the scalar directly. There is no overload mixing a c_ptr scalar with a Fortran array, so the f2008/f2018 tests, which pass Fortran pointer arrays, keep host scalars. * f2008/f2018 hipblas, rocblas, hipsparse, rocsparse: drop the c_loc() added around alpha/beta; keep them as TARGET variables since the wrapper now takes their address. * f2003 hipsparse/rocsparse: add the missing TARGET attribute the new c_loc() call sites require, and wrap nnz_C for xcsrgemmNnz. * hipsolver (f2003 + f2008), untouched by #564 although tau, devInfo and nev changed: wrap the device scalars in c_loc(), pass tau as a whole array where the array wrapper is selected, and leave nev alone in the *_bufferSize calls, which still bind it by reference. * rocblas dgemv_batched: c_ptr batch arrays select the raw binding, so alpha/beta need c_loc() there. Verified against ROCm 10.2 amdflang on gfx1101, with HIPFORT_ASSUMED_RANK both OFF and ON: the suite builds clean and the remaining 19 ctest failures are identical to develop's.
The coverage tests under HIPFORT_EXTENDED_TESTS call every raw binding directly with one placeholder per type (v0 = c_ptr, and one each for real/complex/integer). #564 turned the scalar dummies into type(c_ptr),value, so 2681 of those call sites had to switch to v0. Two arguments needed the opposite move. hipsolverDsyevdx_bufferSize's nev and hipsolverSpDcsrlsvcholHost's singularity were the only members of their families still declared type(c_ptr),value; #564 normalises them to integer(c_int), matching every sibling. Both forms pass an address, so this is a source-form change only, and the two call sites now use the integer placeholder. Verified with ROCm 10.2 amdflang on gfx1101 and with gfortran 13, both with HIPFORT_EXTENDED_TESTS=ON: the suite builds clean and the 19 remaining ctest failures out of 1012 are identical to develop's.
This was referenced Oct 1, 2026
amontoison
added a commit
that referenced
this pull request
Oct 6, 2026
The symbol tests were last regenerated before #553, #564, #571, #572 and 5de6093, so they no longer matched the bindings: eight of them failed to compile with -DHIPFORT_EXTENDED_TESTS=ON, passing type(c_ptr) to arguments those changes typed (hipblasGetProperty_, hipsolver *gels niters, the hipSPARSE/rocSPARSE descriptor getters, rocfft_cache_serialize, the hipFFTW guru dims, ...). Regenerated with gen_symbol_tests.jl from rocm-fortran main, against the bindings on develop. Besides the signature fixes, the tests now reference the routines added by the ROCm 10.1 regeneration (rocsolver cholqr and larft_64, rocblas gemm_grouped_batched, ...). Most of the remaining churn is the scratch variables being renumbered around the new types. Built with gfortran against ROCm 10.1 with -DHIPFORT_EXTENDED_TESTS=ON: all 13 symbol tests compile, link and pass, and the full ctest run shows only the three known WSL failures (hip_mem_advise f2003/f2008, hip_module_kernel).
amontoison
added a commit
that referenced
this pull request
Oct 7, 2026
#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.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.