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

docs/how-to/fortran-interfaces.rst

@@ -1,6 +1,6 @@
 .. meta::
   :description: The Fortran interface variants hipFORT generates and how to call them
-  :keywords: hipFORT, ROCm, Fortran, interfaces, hipMalloc, assumed-rank, f2003, f2008, f2018
+  :keywords: hipFORT, ROCm, Fortran, interfaces, hipMalloc, assumed-rank, f2003, f2008, f2018, c_loc, pointer mode
 
 **************************
 Fortran interface variants
@@ -33,6 +33,10 @@ By convention, application and test sources that rely on the array overloads use
 the ``.f08`` file extension (see the ``test/f2008`` examples), while Fortran 2003
 sources use ``.f03``.
 
+Independently of these variants, a routine that takes a scalar through a pointer
+accepts either the Fortran variable or its address; see
+`Scalar arguments passed by pointer`_.
+
 Allocating and copying device memory
 ====================================
 
@@ -83,6 +87,87 @@ Unlike the array interfaces of the math libraries, these ``hipMalloc`` and
 ``hipMemcpy`` overloads are unconditional, so they are available in every
 hipFORT build.
 
+Scalar arguments passed by pointer
+==================================
+
+Many C routines take a scalar through a pointer: ``alpha`` and ``beta`` in the
+BLAS and sparse libraries, or an output such as a buffer size, a version or a
+count. hipFORT offers two forms of each such routine under the same name, and the
+compiler picks one from the type of the argument:
+
+* **The Fortran variable.** hipFORT passes its address to the library. This is
+  the usual call, for a scalar held on the host:
+
+  .. code-block:: fortran
+
+     real(c_double)    :: alpha, beta
+     integer(c_size_t) :: buffer_size
+     !
+     ierr = rocblas_dgemm(handle, transa, transb, m, n, k, alpha, &
+                          dA, lda, dB, ldb, beta, dC, ldc)
+     ierr = rocsparse_dgemvi_buffer_size(handle, trans, m, n, nnz, buffer_size)
+
+* **Its address, a** ``type(c_ptr)``. This is the C API's own spelling. Use it
+  to pass a device address, to pass ``c_null_ptr`` for an output the C API lets
+  you skip, or to keep calling with ``c_loc``:
+
+  .. code-block:: fortran
+
+     real(c_double), target :: alpha, beta
+     !
+     ierr = rocblas_dgemm(handle, transa, transb, m, n, k, c_loc(alpha), &
+                          dA, lda, dB, ldb, c_loc(beta), dC, ldc)
+
+An output the C API lets you skip is an ``optional`` argument of the first form.
+
+Device pointer mode
+-------------------
+
+``alpha``, ``beta`` and a few other scalars are read from host or device memory
+depending on the pointer mode of the handle (``rocblas_set_pointer_mode``,
+``hipblasSetPointerMode``, ``rocsparse_set_pointer_mode``,
+``hipsparseSetPointerMode``). The Fortran-variable form passes the address of the
+variable, which is normally host memory, so it is for the host pointer mode, the
+default. In device pointer mode, pass the device address as a ``type(c_ptr)``:
+
+.. code-block:: fortran
+
+   type(c_ptr) :: d_alpha, d_beta   ! device memory, from hipMalloc
+   !
+   ierr = rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device)
+   ierr = rocblas_dgemm(handle, transa, transb, m, n, k, d_alpha, &
+                        dA, lda, dB, ldb, d_beta, dC, ldc)
+
+A few rocSPARSE and hipSPARSE routines take both kinds of scalar, for example
+``rocsparse_dcsrgemm_buffer_size``, which reads ``alpha`` and ``beta`` per the
+pointer mode but always writes ``buffer_size`` on the host. In device pointer
+mode, pass ``alpha`` and ``beta`` as ``type(c_ptr)`` and ``buffer_size`` as the
+variable.
+
+The Fortran 2008 array overloads take ``alpha`` and ``beta`` as Fortran variables
+only. To use them in device pointer mode, the variables themselves must live on
+the device, for example a ``real(c_double), pointer`` allocated with
+``hipMalloc``; otherwise, call the ``type(c_ptr)`` form.
+
+Device-only scalars
+-------------------
+
+hipSOLVER's ``devInfo`` is written by the GPU, so it is always device memory.
+
+* With ``type(c_ptr)`` buffers, pass it as a ``type(c_ptr)``, for example the
+  one ``hipMalloc`` returned. There is no Fortran-variable form here: a host
+  integer would be written by the GPU.
+* With the Fortran 2008 array overloads, pass either an integer that lives on
+  the device, such as an ``integer(c_int), pointer`` allocated with
+  ``hipMalloc``, like the arrays beside it, or a ``type(c_ptr)``.
+
+.. note::
+
+   In the API reference and the supported-API tables, the Fortran-variable form
+   of a routine ``X`` appears as the module procedure ``X_typed``, and the
+   device-pointer-mode form of the routines that mix both kinds of scalar as
+   ``X_devptr``. Call them through the generic name ``X``.
+
 Assumed-rank interfaces (Fortran 2018)
 ======================================
 

test/CMakeLists.txt

@@ -151,6 +151,7 @@ if(BUILD_TESTING)
     hipfort_add_test(hipblas cgemm f2003)
     hipfort_add_test(hipblas dgemm f2003)
     hipfort_add_test(hipblas sgemm f2003)
+    hipfort_add_test(hipblas sgemm_typed f2003)
     hipfort_add_test(hipblas zgemm f2003)
     hipfort_add_test(hipblas sgemm_batched f2003)
     hipfort_add_test(hipblas dgemm_batched f2003)
@@ -166,6 +167,7 @@ if(BUILD_TESTING)
     hipfort_add_test(hipblas daxpy f2003)
     hipfort_add_test(hipblas caxpy f2003)
     hipfort_add_test(hipblas zaxpy f2003)
+    hipfort_add_test(hipblas zaxpy_typed f2003)
     hipfort_add_test(hipblas scopy f2003)
     hipfort_add_test(hipblas sgemv f2003)
     hipfort_add_test(hipblas dgemv f2003)
@@ -573,6 +575,7 @@ if(BUILD_TESTING)
     hipfort_add_test(hipsparse hipsparse_sgemvi f2003)
     hipfort_add_test(hipsparse hipsparse_sgemvi f2008)
     hipfort_add_test(hipsparse hipsparse_dgemvi f2003)
+    hipfort_add_test(hipsparse hipsparse_dgemvi_typed f2003)
     hipfort_add_test(hipsparse hipsparse_dgemvi f2008)
     hipfort_add_test(hipsparse hipsparse_cgemvi f2003)
     hipfort_add_test(hipsparse hipsparse_cgemvi f2008)
@@ -585,6 +588,7 @@ if(BUILD_TESTING)
     hipfort_add_test(hipsparse hipsparse_ccsrsv2 f2003)
     hipfort_add_test(hipsparse hipsparse_ccsrsv2 f2008)
     hipfort_add_test(hipsparse hipsparse_zcsrsv2 f2003)
+    hipfort_add_test(hipsparse hipsparse_zcsrsv2_typed f2003)
     hipfort_add_test(hipsparse hipsparse_zcsrsv2 f2008)
     hipfort_add_test(hipsparse hipsparse_scsrilu02 f2003)
     hipfort_add_test(hipsparse hipsparse_scsrilu02 f2008)
@@ -693,6 +697,7 @@ if(BUILD_TESTING)
     hipfort_add_test(rocblas saxpy f2003)
     hipfort_add_test(rocblas daxpy f2003)
     hipfort_add_test(rocblas caxpy f2003)
+    hipfort_add_test(rocblas caxpy_typed f2003)
     hipfort_add_test(rocblas zaxpy f2003)
     hipfort_add_test(rocblas dtrmm f2003)
     hipfort_add_test(rocblas saxpy f2008)
@@ -703,6 +708,7 @@ if(BUILD_TESTING)
     hipfort_add_test(rocblas dtrmm f2008)
     hipfort_add_test(rocblas sgemm f2003)
     hipfort_add_test(rocblas dgemm f2003)
+    hipfort_add_test(rocblas dgemm_typed f2003)
     hipfort_add_test(rocblas cgemm f2003)
     hipfort_add_test(rocblas zgemm f2003)
     hipfort_add_test(rocblas sgemm f2008)
@@ -1285,6 +1291,7 @@ if(BUILD_TESTING)
     hipfort_add_test(rocsparse rocsparse_sgemvi f2003)
     hipfort_add_test(rocsparse rocsparse_sgemvi f2008)
     hipfort_add_test(rocsparse rocsparse_dgemvi f2003)
+    hipfort_add_test(rocsparse rocsparse_dgemvi_typed f2003)
     hipfort_add_test(rocsparse rocsparse_dgemvi f2008)
     hipfort_add_test(rocsparse rocsparse_cgemvi f2003)
     hipfort_add_test(rocsparse rocsparse_cgemvi f2008)
@@ -1383,6 +1390,7 @@ if(BUILD_TESTING)
     hipfort_add_test(rocsparse rocsparse_zspmv f2003)
     hipfort_add_test(rocsparse rocsparse_scsrgemm f2003)
     hipfort_add_test(rocsparse rocsparse_dcsrgemm f2003)
+    hipfort_add_test(rocsparse rocsparse_dcsrgemm_typed f2003)
     hipfort_add_test(rocsparse rocsparse_ccsrgemm f2003)
     hipfort_add_test(rocsparse rocsparse_zcsrgemm f2003)
     hipfort_add_test(rocsparse rocsparse_ssptrsv f2003)

test/f2003/hipblas/sgemm_typed.f03

@@ -0,0 +1,126 @@
+!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!
+! Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved.
+!
+! SPDX-License-Identifier: MIT
+!
+! Permission is hereby granted, free of charge, to any person obtaining a copy
+! of this software and associated documentation files (the "Software"), to deal
+! in the Software without restriction, including without limitation the rights
+! to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+! copies of the Software, and to permit persons to whom the Software is
+! furnished to do so, subject to the following conditions:
+!
+! The above copyright notice and this permission notice shall be included in
+! all copies or substantial portions of the Software.
+!
+! THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+! IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+! FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+! AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+! LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+! OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+! THE SOFTWARE.
+!
+!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!
+
+! The same check as sgemm.f03, through the _typed form: alpha and beta are passed as
+! Fortran variables, here named constants, rather than with c_loc(), so the
+! generic resolves to the _typed module procedure, which hands the C API their
+! address. The handle is in host pointer mode, as that form requires.
+program hip_sgemm_typed
+
+  use iso_c_binding
+  use hipfort
+  use hipfort_check
+  use hipfort_hipblas
+
+  implicit none
+
+  integer(kind(HIPBLAS_OP_N)), parameter :: transa = HIPBLAS_OP_N, transb = HIPBLAS_OP_N;
+  real(c_float), parameter :: alpha = 1.1, beta = 0.9;
+
+  integer, parameter ::  m = 1024, n = 1024, k = 1024;
+  integer :: lda, ldb, ldc, size_a, size_b, size_c;
+
+  real(c_float), allocatable, target, dimension(:) :: ha, hb, hc
+  real(c_float), allocatable, dimension(:) :: hc_exact
+
+  type(c_ptr) :: da = c_null_ptr, db = c_null_ptr, dc = c_null_ptr
+  type(c_ptr) :: handle = c_null_ptr
+
+  integer, parameter :: bytes_per_element = 4 !single precision
+  integer(c_size_t) :: Nabytes, Nbbytes, Ncbytes
+
+  integer :: i
+  real :: error
+  real, parameter :: error_max = 10*epsilon(error)
+
+  write(*,"(a)",advance="no") "-- Running test 'SGEMM typed' (Fortran 2003 interfaces) - "
+
+  call hipblasCheck(hipblasCreate(handle))
+
+  lda        = m;
+  size_a     = k * lda;
+  Nabytes = size_a*bytes_per_element
+
+  ldb        = k;
+  size_b     = n * ldb;
+  Nbbytes = size_b*bytes_per_element
+
+  ldc    = m;
+  size_c = n * ldc;
+  Ncbytes = size_c*bytes_per_element
+
+  allocate(ha(size_a))
+  allocate(hb(size_b))
+  allocate(hc(size_c))
+  allocate(hc_exact(size_c))
+
+  ! Use these constant matrices so the exact answer is also a
+  ! constant matrix and therefore easy to check
+  ha(:) = 1.0
+  hb(:) = 2.0
+  hc(:) = 3.0
+  hc_exact = alpha*k*2.0 + beta*3.0
+
+  ! Allocate device memory
+  call hipCheck(hipMalloc(da,Nabytes))
+  call hipCheck(hipMalloc(db,Nbbytes))
+  call hipCheck(hipMalloc(dc,Ncbytes))
+
+  !Transfer from host to device
+  call hipCheck(hipMemcpy(da, c_loc(ha(1)), Nabytes, hipMemcpyHostToDevice))
+  call hipCheck(hipMemcpy(db, c_loc(hb(1)), Nbbytes, hipMemcpyHostToDevice))
+  call hipCheck(hipMemcpy(dc, c_loc(hc(1)), Ncbytes, hipMemcpyHostToDevice))
+
+
+  call hipblasCheck(hipblasSgemm(handle,transa,transb,m,n,k,alpha,da,lda,db,ldb,beta,dc,ldc))
+
+  call hipCheck(hipDeviceSynchronize())
+
+  ! Transfer data back to host memory
+  call hipCheck(hipMemcpy(c_loc(hc(1)), dc, Ncbytes, hipMemcpyDeviceToHost))
+
+
+  do i = 1,size_c
+     error = abs((hc_exact(i) - hc(i))/hc_exact(i))
+     if( error > error_max )then
+        write(*,*) "FAILED! Error bigger than max! Error = ", error
+        call exit(1)
+     end if
+  end do
+
+  call hipCheck(hipFree(da))
+  call hipCheck(hipFree(db))
+  call hipCheck(hipFree(dc))
+
+  call hipblasCheck(hipblasDestroy(handle))
+
+  deallocate(ha)
+  deallocate(hb)
+  deallocate(hc)
+  deallocate(hc_exact)
+
+  write(*,*) "PASSED!"
+
+end program hip_sgemm_typed