From 86aebec1daef35939a2197fc0d400ce51bd69e70 Mon Sep 17 00:00:00 2001 From: Alexis Montoison Date: Sat, 7 Nov 2026 03:17:01 -0600 Subject: [PATCH 1/2] Move the hipFFT Xt API to hipfort_hipfftxt CUDA Fortran has a cufft module and a cufftXt module, which includes cufft and holds the Xt and cuFFTMp interfaces. hipfort now mirrors it, so that hipified code maps `use cufftXt` to a module of its own: - hipfort_hipfft and hipfort_hipfft_enums hold the hipfft.h API, without the hipfftXt* and hipfftMp* functions; hipfort_hipfft_types goes away, as hipfft.h binds no type. - hipfort_hipfftxt holds every hipfftXt* and hipfftMp* function, including hipfftXtSetWorkArea (declared in hipfft.h, as cufftXtSetWorkArea is in cufftXt), with hipfort_hipfftxt_enums and hipfort_hipfftxt_types. It uses hipfort_hipfft, so `use hipfort_hipfftxt` alone gives the whole hipFFT API. The Xt API is not released yet (#611), so nothing breaks. The modules are regenerated with the generator's new [library] uses, function_prefixes and skip_function_prefixes keys. A hipfftxt CMake component, linked to hipFFT, holds the Xt tests and the symbol test, and the supported-API table, the hipFFT examples and the changelog follow. --- CHANGELOG.md | 16 +- README.md | 1 + docs/doxygen/input/supported_api_hipfft.md | 42 +- docs/doxygen/input/supported_api_hipfftxt.md | 26 + docs/how-to/using-hipfort.rst | 3 +- docs/reference/index.md | 1 + docs/tutorials/hipfft-examples.rst | 21 +- lib/CMakeLists.txt | 4 +- lib/hipfort-config.cmake.in | 9 +- lib/hipfort/hipfort_hipfft.F90 | 969 ---------------- lib/hipfort/hipfort_hipfft_enums.F90 | 48 - lib/hipfort/hipfort_hipfftxt.F90 | 1007 +++++++++++++++++ lib/hipfort/hipfort_hipfftxt_enums.F90 | 79 ++ ...t_types.F90 => hipfort_hipfftxt_types.F90} | 4 +- test/CMakeLists.txt | 16 +- test/coverage/variant_coverage.py | 1 + test/f2003/hipfft/test_hipfft.F03 | 22 - .../hipfftxt_descriptor_z.f03} | 6 +- .../hipfftxt_makeplanmany_c.f03} | 6 +- test/f2003/hipfftxt/test_hipfftxt.F03 | 65 ++ 20 files changed, 1240 insertions(+), 1106 deletions(-) create mode 100644 docs/doxygen/input/supported_api_hipfftxt.md create mode 100644 lib/hipfort/hipfort_hipfftxt.F90 create mode 100644 lib/hipfort/hipfort_hipfftxt_enums.F90 rename lib/hipfort/{hipfort_hipfft_types.F90 => hipfort_hipfftxt_types.F90} (96%) rename test/f2003/{hipfft/hipfft_xt_descriptor_z.f03 => hipfftxt/hipfftxt_descriptor_z.f03} (97%) rename test/f2003/{hipfft/hipfft_xt_makeplanmany_c.f03 => hipfftxt/hipfftxt_makeplanmany_c.f03} (97%) create mode 100644 test/f2003/hipfftxt/test_hipfftxt.F03 diff --git a/CHANGELOG.md b/CHANGELOG.md index ba3758487..773c1c251 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -32,13 +32,17 @@ and `getrs` routines and `hipsolverDnXtrtri` in hipSOLVER, the `gehd2` and `gehrd` families in rocSOLVER, and `rocsparse_spsort_*`, `rocsparse_dnmat_copy_data` and `rocsparse_dnmat_transpose` in rocSPARSE. -* The hipFFT Xt API in `hipfort_hipfft`, as `bind(C)` interfaces (with `_typed` - forms for `hipfftXtMakePlanMany` and `hipfftXtGetSizeMany`): the multi-GPU - routines (`hipfftXtSetGPUs`, `hipfftXtMalloc`, `hipfftXtMemcpy`, `hipfftXtFree`, - the `hipfftXtExecDescriptor*` family), `hipfftXtMakePlanMany`, +* The hipFFT Xt API in the new `hipfort_hipfftxt` module (`hipfort::hipfftxt` + CMake component), as CUDA Fortran's `cufftXt`: it uses `hipfort_hipfft`, so + `use hipfort_hipfftxt` alone gives the whole hipFFT API, and + `hipfort_hipfft` keeps the `hipfft.h` API only. The routines are `bind(C)` + interfaces (with `_typed` forms for `hipfftXtMakePlanMany` and + `hipfftXtGetSizeMany`): the multi-GPU routines (`hipfftXtSetGPUs`, + `hipfftXtSetWorkArea`, `hipfftXtMalloc`, `hipfftXtMemcpy`, `hipfftXtFree`, the + `hipfftXtExecDescriptor*` family), `hipfftXtMakePlanMany`, `hipfftXtGetSizeMany` and `hipfftXtExec` on `hipDataType`, and the callback - routines. The `hipXtDesc` and `hipLibXtDesc` types are in the new - `hipfort_hipfft_types` module, and the Xt enums in `hipfort_hipfft_enums`. + routines. The `hipXtDesc` and `hipLibXtDesc` types are in + `hipfort_hipfftxt_types`, and the Xt enums in `hipfort_hipfftxt_enums`. On AMD GPUs the descriptor routines need a plan spread over two or more GPUs, and a multi-GPU plan writes one work size per GPU: pass `c_loc` of an `integer(c_size_t)` array as `workSize`. The multi-process routines diff --git a/README.md b/README.md index 3cad38963..ae324a917 100644 --- a/README.md +++ b/README.md @@ -134,6 +134,7 @@ The following tables list the supported API: * [HIP](https://rocm.docs.amd.com/projects/hipfort/en/develop/doxygen/html/md_input_2supported__api__hip.html) * [hipBLAS](https://rocm.docs.amd.com/projects/hipfort/en/develop/doxygen/html/md_input_2supported__api__hipblas.html) * [hipFFT](https://rocm.docs.amd.com/projects/hipfort/en/develop/doxygen/html/md_input_2supported__api__hipfft.html) +* [hipFFT Xt](https://rocm.docs.amd.com/projects/hipfort/en/develop/doxygen/html/md_input_2supported__api__hipfftxt.html) * [hipFFTW](https://rocm.docs.amd.com/projects/hipfort/en/develop/doxygen/html/md_input_2supported__api__hipfftw.html) * [hipRAND](https://rocm.docs.amd.com/projects/hipfort/en/develop/doxygen/html/md_input_2supported__api__hiprand.html) * [hipSOLVER](https://rocm.docs.amd.com/projects/hipfort/en/develop/doxygen/html/md_input_2supported__api__hipsolver.html) diff --git a/docs/doxygen/input/supported_api_hipfft.md b/docs/doxygen/input/supported_api_hipfft.md index fe21a1821..9a279375d 100644 --- a/docs/doxygen/input/supported_api_hipfft.md +++ b/docs/doxygen/input/supported_api_hipfft.md @@ -25,35 +25,13 @@ 21 | [hipfftGetSize](interfacehipfort__hipfft_1_1hipfftgetsize.html "Interface documentation") | C binding, typed 22 | [hipfftSetAutoAllocation](interfacehipfort__hipfft_1_1hipfftsetautoallocation.html "Interface documentation") | C binding 23 | [hipfftSetWorkArea](interfacehipfort__hipfft_1_1hipfftsetworkarea.html "Interface documentation") | C binding -24 | [hipfftXtSetWorkArea](interfacehipfort__hipfft_1_1hipfftxtsetworkarea.html "Interface documentation") | C binding -25 | [hipfftExecC2C](interfacehipfort__hipfft_1_1hipfftexecc2c.html "Interface documentation") | C binding, rank_0, rank_1, rank_2, rank_3, assumed_rank -26 | [hipfftExecR2C](interfacehipfort__hipfft_1_1hipfftexecr2c.html "Interface documentation") | C binding, rank_0, rank_1, rank_2, rank_3, assumed_rank -27 | [hipfftExecC2R](interfacehipfort__hipfft_1_1hipfftexecc2r.html "Interface documentation") | C binding, rank_0, rank_1, rank_2, rank_3, assumed_rank -28 | [hipfftExecZ2Z](interfacehipfort__hipfft_1_1hipfftexecz2z.html "Interface documentation") | C binding, rank_0, rank_1, rank_2, rank_3, assumed_rank -29 | [hipfftExecD2Z](interfacehipfort__hipfft_1_1hipfftexecd2z.html "Interface documentation") | C binding, rank_0, rank_1, rank_2, rank_3, assumed_rank -30 | [hipfftExecZ2D](interfacehipfort__hipfft_1_1hipfftexecz2d.html "Interface documentation") | C binding, rank_0, rank_1, rank_2, rank_3, assumed_rank -31 | [hipfftSetStream](interfacehipfort__hipfft_1_1hipfftsetstream.html "Interface documentation") | C binding -32 | [hipfftDestroy](interfacehipfort__hipfft_1_1hipfftdestroy.html "Interface documentation") | C binding -33 | [hipfftGetVersion](interfacehipfort__hipfft_1_1hipfftgetversion.html "Interface documentation") | C binding, typed -34 | [hipfftGetProperty](interfacehipfort__hipfft_1_1hipfftgetproperty.html "Interface documentation") | C binding, typed -35 | [hipfftXtSetCallback](interfacehipfort__hipfft_1_1hipfftxtsetcallback.html "Interface documentation") | C binding -36 | [hipfftXtClearCallback](interfacehipfort__hipfft_1_1hipfftxtclearcallback.html "Interface documentation") | C binding -37 | [hipfftXtSetCallbackSharedSize](interfacehipfort__hipfft_1_1hipfftxtsetcallbacksharedsize.html "Interface documentation") | C binding -38 | [hipfftXtSetJITCallback](interfacehipfort__hipfft_1_1hipfftxtsetjitcallback.html "Interface documentation") | C binding -39 | [hipfftXtMakePlanMany](interfacehipfort__hipfft_1_1hipfftxtmakeplanmany.html "Interface documentation") | C binding, rank_0, rank_1, assumed_rank, typed -40 | [hipfftXtGetSizeMany](interfacehipfort__hipfft_1_1hipfftxtgetsizemany.html "Interface documentation") | C binding, rank_0, rank_1, assumed_rank, typed -41 | [hipfftXtExec](interfacehipfort__hipfft_1_1hipfftxtexec.html "Interface documentation") | C binding -42 | [hipfftXtSetGPUs](interfacehipfort__hipfft_1_1hipfftxtsetgpus.html "Interface documentation") | C binding, rank_0, rank_1, assumed_rank -43 | [hipfftXtMalloc](interfacehipfort__hipfft_1_1hipfftxtmalloc.html "Interface documentation") | C binding -44 | [hipfftXtMemcpy](interfacehipfort__hipfft_1_1hipfftxtmemcpy.html "Interface documentation") | C binding -45 | [hipfftXtFree](interfacehipfort__hipfft_1_1hipfftxtfree.html "Interface documentation") | C binding -46 | [hipfftXtExecDescriptorC2C](interfacehipfort__hipfft_1_1hipfftxtexecdescriptorc2c.html "Interface documentation") | C binding -47 | [hipfftXtExecDescriptorR2C](interfacehipfort__hipfft_1_1hipfftxtexecdescriptorr2c.html "Interface documentation") | C binding -48 | [hipfftXtExecDescriptorC2R](interfacehipfort__hipfft_1_1hipfftxtexecdescriptorc2r.html "Interface documentation") | C binding -49 | [hipfftXtExecDescriptorZ2Z](interfacehipfort__hipfft_1_1hipfftxtexecdescriptorz2z.html "Interface documentation") | C binding -50 | [hipfftXtExecDescriptorD2Z](interfacehipfort__hipfft_1_1hipfftxtexecdescriptord2z.html "Interface documentation") | C binding -51 | [hipfftXtExecDescriptorZ2D](interfacehipfort__hipfft_1_1hipfftxtexecdescriptorz2d.html "Interface documentation") | C binding -52 | [hipfftXtExecDescriptor](interfacehipfort__hipfft_1_1hipfftxtexecdescriptor.html "Interface documentation") | C binding -53 | [hipfftMpAttachComm](interfacehipfort__hipfft_1_1hipfftmpattachcomm.html "Interface documentation") | C binding -54 | [hipfftXtSetDistribution](interfacehipfort__hipfft_1_1hipfftxtsetdistribution.html "Interface documentation") | C binding -55 | [hipfftXtSetSubformatDefault](interfacehipfort__hipfft_1_1hipfftxtsetsubformatdefault.html "Interface documentation") | C binding +24 | [hipfftExecC2C](interfacehipfort__hipfft_1_1hipfftexecc2c.html "Interface documentation") | C binding, rank_0, rank_1, rank_2, rank_3, assumed_rank +25 | [hipfftExecR2C](interfacehipfort__hipfft_1_1hipfftexecr2c.html "Interface documentation") | C binding, rank_0, rank_1, rank_2, rank_3, assumed_rank +26 | [hipfftExecC2R](interfacehipfort__hipfft_1_1hipfftexecc2r.html "Interface documentation") | C binding, rank_0, rank_1, rank_2, rank_3, assumed_rank +27 | [hipfftExecZ2Z](interfacehipfort__hipfft_1_1hipfftexecz2z.html "Interface documentation") | C binding, rank_0, rank_1, rank_2, rank_3, assumed_rank +28 | [hipfftExecD2Z](interfacehipfort__hipfft_1_1hipfftexecd2z.html "Interface documentation") | C binding, rank_0, rank_1, rank_2, rank_3, assumed_rank +29 | [hipfftExecZ2D](interfacehipfort__hipfft_1_1hipfftexecz2d.html "Interface documentation") | C binding, rank_0, rank_1, rank_2, rank_3, assumed_rank +30 | [hipfftSetStream](interfacehipfort__hipfft_1_1hipfftsetstream.html "Interface documentation") | C binding +31 | [hipfftDestroy](interfacehipfort__hipfft_1_1hipfftdestroy.html "Interface documentation") | C binding +32 | [hipfftGetVersion](interfacehipfort__hipfft_1_1hipfftgetversion.html "Interface documentation") | C binding, typed +33 | [hipfftGetProperty](interfacehipfort__hipfft_1_1hipfftgetproperty.html "Interface documentation") | C binding, typed diff --git a/docs/doxygen/input/supported_api_hipfftxt.md b/docs/doxygen/input/supported_api_hipfftxt.md new file mode 100644 index 000000000..4accddf76 --- /dev/null +++ b/docs/doxygen/input/supported_api_hipfftxt.md @@ -0,0 +1,26 @@ +# hipFFT Xt API Support + +\# | API Name | Variants +----|---------------|--------- +1 | [hipfftXtSetWorkArea](interfacehipfort__hipfftxt_1_1hipfftxtsetworkarea.html "Interface documentation") | C binding +2 | [hipfftXtSetCallback](interfacehipfort__hipfftxt_1_1hipfftxtsetcallback.html "Interface documentation") | C binding +3 | [hipfftXtClearCallback](interfacehipfort__hipfftxt_1_1hipfftxtclearcallback.html "Interface documentation") | C binding +4 | [hipfftXtSetCallbackSharedSize](interfacehipfort__hipfftxt_1_1hipfftxtsetcallbacksharedsize.html "Interface documentation") | C binding +5 | [hipfftXtSetJITCallback](interfacehipfort__hipfftxt_1_1hipfftxtsetjitcallback.html "Interface documentation") | C binding +6 | [hipfftXtMakePlanMany](interfacehipfort__hipfftxt_1_1hipfftxtmakeplanmany.html "Interface documentation") | C binding, rank_0, rank_1, assumed_rank, typed +7 | [hipfftXtGetSizeMany](interfacehipfort__hipfftxt_1_1hipfftxtgetsizemany.html "Interface documentation") | C binding, rank_0, rank_1, assumed_rank, typed +8 | [hipfftXtExec](interfacehipfort__hipfftxt_1_1hipfftxtexec.html "Interface documentation") | C binding +9 | [hipfftXtSetGPUs](interfacehipfort__hipfftxt_1_1hipfftxtsetgpus.html "Interface documentation") | C binding, rank_0, rank_1, assumed_rank +10 | [hipfftXtMalloc](interfacehipfort__hipfftxt_1_1hipfftxtmalloc.html "Interface documentation") | C binding +11 | [hipfftXtMemcpy](interfacehipfort__hipfftxt_1_1hipfftxtmemcpy.html "Interface documentation") | C binding +12 | [hipfftXtFree](interfacehipfort__hipfftxt_1_1hipfftxtfree.html "Interface documentation") | C binding +13 | [hipfftXtExecDescriptorC2C](interfacehipfort__hipfftxt_1_1hipfftxtexecdescriptorc2c.html "Interface documentation") | C binding +14 | [hipfftXtExecDescriptorR2C](interfacehipfort__hipfftxt_1_1hipfftxtexecdescriptorr2c.html "Interface documentation") | C binding +15 | [hipfftXtExecDescriptorC2R](interfacehipfort__hipfftxt_1_1hipfftxtexecdescriptorc2r.html "Interface documentation") | C binding +16 | [hipfftXtExecDescriptorZ2Z](interfacehipfort__hipfftxt_1_1hipfftxtexecdescriptorz2z.html "Interface documentation") | C binding +17 | [hipfftXtExecDescriptorD2Z](interfacehipfort__hipfftxt_1_1hipfftxtexecdescriptord2z.html "Interface documentation") | C binding +18 | [hipfftXtExecDescriptorZ2D](interfacehipfort__hipfftxt_1_1hipfftxtexecdescriptorz2d.html "Interface documentation") | C binding +19 | [hipfftXtExecDescriptor](interfacehipfort__hipfftxt_1_1hipfftxtexecdescriptor.html "Interface documentation") | C binding +20 | [hipfftMpAttachComm](interfacehipfort__hipfftxt_1_1hipfftmpattachcomm.html "Interface documentation") | C binding +21 | [hipfftXtSetDistribution](interfacehipfort__hipfftxt_1_1hipfftxtsetdistribution.html "Interface documentation") | C binding +22 | [hipfftXtSetSubformatDefault](interfacehipfort__hipfftxt_1_1hipfftxtsetsubformatdefault.html "Interface documentation") | C binding diff --git a/docs/how-to/using-hipfort.rst b/docs/how-to/using-hipfort.rst index c9b21c030..6d811bad3 100644 --- a/docs/how-to/using-hipfort.rst +++ b/docs/how-to/using-hipfort.rst @@ -87,7 +87,7 @@ underlying ROCm library it wraps. List the libraries your code uses as ``COMPONENTS`` (``hip``, ``roctx``, ``rocblas``, ``rocfft``, ``rocrand``, ``rocsolver``, ``rocsparse``, -``hipblas``, ``hipfft``, ``hipfftw``, ``hiprand``, ``hipsolver``, +``hipblas``, ``hipfft``, ``hipfftw``, ``hipfftxt``, ``hiprand``, ``hipsolver``, ``hipsparse``) and link the matching ``hipfort::`` targets. A ``hipfort::`` target is only defined when that component is listed, and the Fortran language must be enabled before ``find_package(hipfort)``. @@ -180,6 +180,7 @@ links as long as it does not call them. The following tables list the supported * :doc:`HIP API <../doxygen/html/md_input_2supported__api__hip>` * :doc:`hipBLAS API <../doxygen/html/md_input_2supported__api__hipblas>` * :doc:`hipFFT API <../doxygen/html/md_input_2supported__api__hipfft>` +* :doc:`hipFFT Xt API <../doxygen/html/md_input_2supported__api__hipfftxt>` * :doc:`hipFFTW API <../doxygen/html/md_input_2supported__api__hipfftw>` * :doc:`hipRAND API <../doxygen/html/md_input_2supported__api__hiprand>` * :doc:`hipSOLVER API <../doxygen/html/md_input_2supported__api__hipsolver>` diff --git a/docs/reference/index.md b/docs/reference/index.md index 7170038d1..4b0765ceb 100644 --- a/docs/reference/index.md +++ b/docs/reference/index.md @@ -11,6 +11,7 @@ This section provides technical descriptions and important information about the - {doc}`/doxygen/html/md_input_2supported__api__hip` - {doc}`/doxygen/html/md_input_2supported__api__hipblas` - {doc}`/doxygen/html/md_input_2supported__api__hipfft` +- {doc}`/doxygen/html/md_input_2supported__api__hipfftxt` - {doc}`/doxygen/html/md_input_2supported__api__hipfftw` - {doc}`/doxygen/html/md_input_2supported__api__hiprand` - {doc}`/doxygen/html/md_input_2supported__api__hipsolver` diff --git a/docs/tutorials/hipfft-examples.rst b/docs/tutorials/hipfft-examples.rst index 65018397f..6c58833f7 100644 --- a/docs/tutorials/hipfft-examples.rst +++ b/docs/tutorials/hipfft-examples.rst @@ -8,14 +8,17 @@ hipFFT examples `hipFFT `_ is a thin layer over rocFFT with an API that follows cuFFT. hipFORT exposes it through the -``hipfort_hipfft`` module. +``hipfort_hipfft`` module, and its Xt and multi-process APIs through the +``hipfort_hipfftxt`` module, as CUDA Fortran does with ``cufft`` and ``cufftXt``. +``hipfort_hipfftxt`` uses ``hipfort_hipfft``, so ``use hipfort_hipfftxt`` alone +gives the whole hipFFT API. Every program on this page is complete and self-contained, and is built and run as part of the hipFORT test suite. The Fortran 2008 version of each program lives in ``test/f2008/hipfft``, and the equivalent Fortran 2003 version, which uses ``type(c_ptr)`` device pointers and explicit byte counts instead of Fortran array pointers, lives in ``test/f2003/hipfft``. The Xt programs exist -only in the Fortran 2003 form. +only in the Fortran 2003 form, in ``test/f2003/hipfftxt``. If you want direct access to rocFFT rather than a cuFFT-style interface, see the :doc:`rocFFT examples `. For FFTW3-compatible code, see the @@ -177,7 +180,7 @@ names the input, output and execution data types with ``hipDataType`` values ``hipfftXtExec`` runs the plan whatever its data types, taking the direction for complex transforms. -.. literalinclude:: ../../test/f2003/hipfft/hipfft_xt_makeplanmany_c.f03 +.. literalinclude:: ../../test/f2003/hipfftxt/hipfftxt_makeplanmany_c.f03 :language: fortran Multi-GPU transforms @@ -211,9 +214,9 @@ copy of the type: ``hipfftXtFree`` would then release the wrong address. The ``GPUs``, ``data`` and ``size`` fields say where each part of the data is. The ``hipLibXtDesc`` and ``hipXtDesc`` types are in the -``hipfort_hipfft_types`` module, which ``hipfort_hipfft`` uses, and the Xt -constants (``HIPFFT_XT_FORMAT_INPLACE``, ``HIPFFT_COPY_HOST_TO_DEVICE``, ...) -are in ``hipfort_hipfft_enums``. +``hipfort_hipfftxt_types`` module, and the Xt constants +(``HIPFFT_XT_FORMAT_INPLACE``, ``HIPFFT_COPY_HOST_TO_DEVICE``, ...) in +``hipfort_hipfftxt_enums``; ``hipfort_hipfftxt`` uses both. On AMD GPUs a descriptor needs a plan over two or more GPUs, and unbatched one-dimensional transforms are not supported. This program runs a @@ -221,7 +224,7 @@ two-dimensional transform on two GPUs, forward then inverse, and checks the result against the input; on a machine with fewer than two GPUs it reports so and the test is skipped. -.. literalinclude:: ../../test/f2003/hipfft/hipfft_xt_descriptor_z.f03 +.. literalinclude:: ../../test/f2003/hipfftxt/hipfftxt_descriptor_z.f03 :language: fortran Multi-process transforms with MPI @@ -229,7 +232,7 @@ Multi-process transforms with MPI ``hipfftMpAttachComm``, ``hipfftXtSetDistribution`` and ``hipfftXtSetSubformatDefault`` distribute a transform over the ranks of an MPI -communicator. They are bound in ``hipfort_hipfft``, but exist only in a hipFFT +communicator. They are bound in ``hipfort_hipfftxt``, but exist only in a hipFFT built with MPI support (``HIPFFT_MPI_ENABLE``), which the ROCm packages are not: with those, a program that calls them does not link. @@ -256,7 +259,7 @@ depend on the MPI implementation, so build it in C from the Fortran handle with use iso_c_binding use mpi_f08 use hipfort_check - use hipfort_hipfft + use hipfort_hipfftxt interface function c_comm_from_fortran(fcomm) bind(c, name="c_comm_from_fortran") diff --git a/lib/CMakeLists.txt b/lib/CMakeLists.txt index 8293baac1..50825e673 100644 --- a/lib/CMakeLists.txt +++ b/lib/CMakeLists.txt @@ -322,6 +322,8 @@ if(HIP_PLATFORM STREQUAL "amd" AND TARGET hipfort-amdgcn) else() hipfort_add_component(hipfftw hip::hipfft) endif() + # hipFFT's Xt and Mp APIs (hipfort_hipfftxt), in libhipfft. + hipfort_add_component(hipfftxt hip::hipfft) else() message(STATUS "Skipping hipfort::hipfft target export") endif() @@ -372,7 +374,7 @@ elseif(CUDAToolkit_FOUND AND TARGET hipfort-nvptx) # Only the hip* components exist here; roc* and the FFTW3 API have no CUDA counterpart. set(HIPFORT_BASE_LIB hipfort-nvptx) foreach(pair "hip:cudart" "hipblas:cublas" "hipfft:cufft" - "hiprand:curand" "hipsolver:cusolver" "hipsparse:cusparse") + "hipfftxt:cufft" "hiprand:curand" "hipsolver:cusolver" "hipsparse:cusparse") string(REPLACE ":" ";" pair "${pair}") list(GET pair 0 _comp) list(GET pair 1 _cudalib) diff --git a/lib/hipfort-config.cmake.in b/lib/hipfort-config.cmake.in index deafd8631..1981f41db 100644 --- a/lib/hipfort-config.cmake.in +++ b/lib/hipfort-config.cmake.in @@ -29,19 +29,20 @@ include(CMakeFindDependencyMacro) set(_hipfort_supported_components hip roctx rocblas rocfft rocrand rocsolver rocsparse - hipblas hipfft hipfftw hiprand hipsolver hipsparse) + hipblas hipfft hipfftw hipfftxt hiprand hipsolver hipsparse) # Most components map 1:1 to a ROCm package of the same name. A few do not: -# hipfftw is provided by the hipfft package (there is no standalone hipfftw -# package), and roctx by rocprofiler-sdk-roctx. +# hipfftw and hipfftxt are provided by the hipfft package (there is no standalone +# hipfftw or hipfftxt package), and roctx by rocprofiler-sdk-roctx. set(_hipfort_hipfftw_dependency hipfft) +set(_hipfort_hipfftxt_dependency hipfft) set(_hipfort_roctx_dependency rocprofiler-sdk-roctx) # A build of the CUDA (nvptx) backend exports only the hip* components, and # links them against the CUDA:: targets of CUDAToolkit instead. set(_hipfort_base_arch "@HIPFORT_BASE_ARCH@") if(_hipfort_base_arch STREQUAL "nvptx") - set(_hipfort_supported_components hip hipblas hipfft hiprand hipsolver hipsparse) + set(_hipfort_supported_components hip hipblas hipfft hipfftxt hiprand hipsolver hipsparse) endif() include("${CMAKE_CURRENT_LIST_DIR}/hipfort-${_hipfort_base_arch}-targets.cmake") diff --git a/lib/hipfort/hipfort_hipfft.F90 b/lib/hipfort/hipfort_hipfft.F90 index 44f82910c..132527947 100644 --- a/lib/hipfort/hipfort_hipfft.F90 +++ b/lib/hipfort/hipfort_hipfft.F90 @@ -25,7 +25,6 @@ module hipfort_hipfft use hipfort_hipfft_enums - use hipfort_hipfft_types implicit none !> @brief Create a new one-dimensional FFT plan. @@ -793,26 +792,6 @@ function hipfftSetWorkArea_(plan,workArea) bind(c, name="hipfftSetWorkArea") end function end interface - !> @brief Generalized version of `hipfftSetWorkArea` accepting multi-device plans. - !> - !> @param[in] plan - Pointer to the (possibly multi-device) FFT plan. - !> @param[in] workArea - Array of pointer(s) to the plan's work area(s), which must - !> be accessible by the plan's used device(s) (in the same order as communicated - !> via `hipfftXtSetGPUs`, if used prior). - interface hipfftXtSetWorkArea -#ifdef USE_CUDA_NAMES - function hipfftXtSetWorkArea_(plan,workArea) bind(c, name="cufftXtSetWorkArea") -#else - function hipfftXtSetWorkArea_(plan,workArea) bind(c, name="hipfftXtSetWorkArea") -#endif - import - implicit none - integer(c_int) :: hipfftXtSetWorkArea_ - type(c_ptr),value :: plan - type(c_ptr) :: workArea - end function - end interface - !> @brief Execute a (float) complex-to-complex FFT. !> !> @details If the input and output buffers are equal, an in-place @@ -1102,727 +1081,6 @@ function hipfftGetProperty_(myType,myValue) bind(c, name="hipfftGetProperty") module procedure hipfftGetProperty_typed end interface - interface hipfftXtSetCallback -#ifdef USE_CUDA_NAMES - function hipfftXtSetCallback_(plan,callbacks,cbtype,callbackData) & - bind(c, name="cufftXtSetCallback") -#else - function hipfftXtSetCallback_(plan,callbacks,cbtype,callbackData) & - bind(c, name="hipfftXtSetCallback") -#endif - import - implicit none - integer(c_int) :: hipfftXtSetCallback_ - type(c_ptr),value :: plan - type(c_ptr) :: callbacks - integer(c_int),value :: cbtype - type(c_ptr) :: callbackData - end function - end interface - - interface hipfftXtClearCallback -#ifdef USE_CUDA_NAMES - function hipfftXtClearCallback_(plan,cbtype) bind(c, name="cufftXtClearCallback") -#else - function hipfftXtClearCallback_(plan,cbtype) bind(c, name="hipfftXtClearCallback") -#endif - import - implicit none - integer(c_int) :: hipfftXtClearCallback_ - type(c_ptr),value :: plan - integer(c_int),value :: cbtype - end function - end interface - - !> @brief Set shared memory size for callback. - !> - !> @details Set shared memory required for a callback. The - !> callback of the specified type must have already been set on the - !> plan. - !> - !> @param[in] plan - The FFT plan. - !> @param[in] cbtype - Type of callback being modified. - !> @param[in] sharedSize - Amount of shared memory required, in bytes. - interface hipfftXtSetCallbackSharedSize -#ifdef USE_CUDA_NAMES - function hipfftXtSetCallbackSharedSize_(plan,cbtype,sharedSize) & - bind(c, name="cufftXtSetCallbackSharedSize") -#else - function hipfftXtSetCallbackSharedSize_(plan,cbtype,sharedSize) & - bind(c, name="hipfftXtSetCallbackSharedSize") -#endif - import - implicit none - integer(c_int) :: hipfftXtSetCallbackSharedSize_ - type(c_ptr),value :: plan - integer(c_int),value :: cbtype - integer(c_size_t),value :: sharedSize - end function - end interface - - !> @brief Set a JIT callback on a plan - !> - !> @details Set either a load or store callback to run with a plan. - !> The type of callback is specified with the `cbtype` parameter. - !> The callback is provided as SPIR-V on AMD platforms and as - !> LTO-IR fatbin on CUDA platforms. - !> - !> Note: JIT callbacks cannot currently be used on multi-GPU transforms. - !> This support will be added in a future release of hipFFT. - !> - !> JIT callbacks can be cleared from the plan by passing nullptr - !> for `symbol_name` and `bitcode_data`, and 0 for - !> `bitcode_len_bytes`. - !> - !> `cbdata` is an optional array of pointers of data that is passed - !> to the callback function. As JIT callbacks are currently only - !> supported for single-device transforms, this array must have one - !> pointer element in it if it is specified. The pointer will be - !> passed to the callback function's `cbdata` parameter when it is - !> executed on that device. - !> - !> If specified, the `cbdata` array must remain valid until either - !> the plan is destroyed or callbacks are cleared from the plan. - !> - !> This function must be called after the plan is allocated using - !> `hipfftCreate`, but before the plan is initialized by any of the - !> "MakePlan" functions. Therefore, API functions that combine - !> creation and initialization (`hipfftPlan1d`, `hipfftPlan2d`, - !> `hipfftPlan3d`, and `hipfftPlanMany`) cannot set a JIT callback. - !> - !> @param[in] plan - The FFT plan. - !> @param[in] symbol_name - Name of the symbol in the compiled bitcode. - !> @param[in] bitcode_data - Pointer to bitcode data. - !> @param[in] bitcode_len_bytes - Number of bytes in the bitcode data. - !> @param[in] cbtype - Type of callback function. - !> @param[in] cbdata - Array of callback function data pointers. - interface hipfftXtSetJITCallback -#ifdef USE_CUDA_NAMES - function hipfftXtSetJITCallback_(plan,symbol_name,bitcode_data,bitcode_len_bytes,cbtype, & - cbdata) & - bind(c, name="cufftXtSetJITCallback") -#else - function hipfftXtSetJITCallback_(plan,symbol_name,bitcode_data,bitcode_len_bytes,cbtype, & - cbdata) & - bind(c, name="hipfftXtSetJITCallback") -#endif - import - implicit none - integer(c_int) :: hipfftXtSetJITCallback_ - type(c_ptr),value :: plan - type(c_ptr),value :: symbol_name - type(c_ptr),value :: bitcode_data - integer(c_size_t),value :: bitcode_len_bytes - integer(c_int),value :: cbtype - type(c_ptr) :: cbdata - end function - end interface - - !> @brief Initialize a batched rank-dimensional FFT plan with - !> advanced data layout and specified input, output, execution data - !> types. - !> - !> @details Assumes that the plan has been created already, and - !> modifies the plan associated with the plan handle. The number - !> of elements to transform in each direction of the input data - !> in the FFT plan is specified in n. - !> - !> The batch parameter tells hipFFT how many transforms to perform. - !> The distance between the first elements of two consecutive batches - !> of the input and output data are specified with the idist and odist - !> parameters. - !> - !> The inembed and onembed parameters define the input and output data - !> layouts. The number of elements in the data is assumed to be larger - !> than the number of elements in the transform. Strided data layouts - !> are also supported. Strides along the fastest direction in the input - !> and output data are specified via the istride and ostride parameters. - !> - !> If both inembed and onembed parameters are set to NULL, all the - !> advanced data layout parameters are ignored and reverted to default - !> values, i.e., the batched transform is performed with non-strided data - !> access and the number of data/transform elements are assumed to be - !> equivalent. - !> - !> The inputType, outputType, executionType parameters specify the - !> data types (precision, and whether the data is real-valued or - !> complex-valued) of the transform input, output, and internal - !> representation respectively. Currently, the precision of all - !> three parameters must match, and the execution type must always - !> be complex-valued. At least one of inputType and outputType - !> must be complex. A half-precision transform can be requested - !> by using either the HIP_R_16F or HIP_C_16F types. - !> - !> @param[out] plan - Pointer to the FFT plan handle. - !> @param[in] rank - Dimension of transform (1, 2, or 3). - !> @param[in] n - Number of elements to transform in the x/y/z directions. - !> @param[in] inembed - Number of elements in the input data in the x/y/z directions. - !> @param[in] istride - Distance between two successive elements in the input data. - !> @param[in] idist - Distance between input batches. - !> @param[in] inputType - Format of FFT input. - !> @param[in] onembed - Number of elements in the output data in the x/y/z directions. - !> @param[in] ostride - Distance between two successive elements in the output data. - !> @param[in] odist - Distance between output batches. - !> @param[in] outputType - Format of FFT output. - !> @param[in] batch - Number of batched transforms to perform. - !> @param[out] workSize - Pointer to work area size (returned value). - !> @param[in] executionType - Internal data format used by the library during computation. - interface hipfftXtMakePlanMany -#ifdef USE_CUDA_NAMES - function hipfftXtMakePlanMany_(plan,rank,n,inembed,istride,idist,inputType,onembed,ostride, & - odist,outputType,batch,workSize,executionType) & - bind(c, name="cufftXtMakePlanMany") -#else - function hipfftXtMakePlanMany_(plan,rank,n,inembed,istride,idist,inputType,onembed,ostride, & - odist,outputType,batch,workSize,executionType) & - bind(c, name="hipfftXtMakePlanMany") -#endif - import - implicit none - integer(c_int) :: hipfftXtMakePlanMany_ - type(c_ptr),value :: plan - integer(c_int),value :: rank - type(c_ptr),value :: n - type(c_ptr),value :: inembed - integer(c_int64_t),value :: istride - integer(c_int64_t),value :: idist - integer(c_int),value :: inputType - type(c_ptr),value :: onembed - integer(c_int64_t),value :: ostride - integer(c_int64_t),value :: odist - integer(c_int),value :: outputType - integer(c_int64_t),value :: batch - type(c_ptr),value :: workSize - integer(c_int),value :: executionType - end function - - module procedure hipfftXtMakePlanMany_typed - -#ifdef USE_ASSUMED_RANK - module procedure hipfftXtMakePlanMany_assumed_rank -#else -#ifdef USE_ASSUMED_SHAPE - module procedure & - hipfftXtMakePlanMany_rank_0,& - hipfftXtMakePlanMany_rank_1 -#endif -#endif - end interface - - !> @brief Return size of the work area size required for a - !> rank-dimensional plan, with specified input, output, execution - !> data types. - !> - !> @details See `hipfftXtMakePlanMany` for restrictions on inputType, - !> outputType, executionType parameters. - !> - !> @param[in] plan - Pointer to the FFT plan. - !> @param[in] rank - Dimension of FFT transform (1, 2, or 3). - !> @param[in] n - Number of elements in the x/y/z directions. - !> @param[in] inembed - Number of elements in the input data in the x/y/z directions. - !> @param[in] istride - Distance between two successive elements in the input data. - !> @param[in] idist - Distance between input batches. - !> @param[in] inputType - Format of FFT input. - !> @param[in] onembed - Number of elements in the output data in the x/y/z directions. - !> @param[in] ostride - Distance between two successive elements in the output data. - !> @param[in] odist - Distance between output batches. - !> @param[in] outputType - Format of FFT output. - !> @param[in] batch - Number of batched transforms to perform. - !> @param[out] workSize - Pointer to work area size (returned value). - !> @param[in] executionType - Internal data format used by the library during computation. - interface hipfftXtGetSizeMany -#ifdef USE_CUDA_NAMES - function hipfftXtGetSizeMany_(plan,rank,n,inembed,istride,idist,inputType,onembed,ostride, & - odist,outputType,batch,workSize,executionType) & - bind(c, name="cufftXtGetSizeMany") -#else - function hipfftXtGetSizeMany_(plan,rank,n,inembed,istride,idist,inputType,onembed,ostride, & - odist,outputType,batch,workSize,executionType) & - bind(c, name="hipfftXtGetSizeMany") -#endif - import - implicit none - integer(c_int) :: hipfftXtGetSizeMany_ - type(c_ptr),value :: plan - integer(c_int),value :: rank - type(c_ptr),value :: n - type(c_ptr),value :: inembed - integer(c_int64_t),value :: istride - integer(c_int64_t),value :: idist - integer(c_int),value :: inputType - type(c_ptr),value :: onembed - integer(c_int64_t),value :: ostride - integer(c_int64_t),value :: odist - integer(c_int),value :: outputType - integer(c_int64_t),value :: batch - type(c_ptr),value :: workSize - integer(c_int),value :: executionType - end function - - module procedure hipfftXtGetSizeMany_typed - -#ifdef USE_ASSUMED_RANK - module procedure hipfftXtGetSizeMany_assumed_rank -#else -#ifdef USE_ASSUMED_SHAPE - module procedure & - hipfftXtGetSizeMany_rank_0,& - hipfftXtGetSizeMany_rank_1 -#endif -#endif - end interface - - !> @brief Execute an FFT plan for any precision and type. - !> - !> @details An in-place transform is performed if the input and - !> output pointers have the same value. - !> - !> The direction parameter is ignored if for real-to-complex and - !> complex-to-real transforms, as the direction is already implied by - !> the data types. - !> - !> @param[in] plan - Pointer to the FFT plan. - !> @param[in] input - Pointer to input data for the transform. - !> @param[in] output - Pointer to output data for the transform. - !> @param[in] direction - Either `HIPFFT_FORWARD` or `HIPFFT_BACKWARD`. - interface hipfftXtExec -#ifdef USE_CUDA_NAMES - function hipfftXtExec_(plan,input,output,direction) bind(c, name="cufftXtExec") -#else - function hipfftXtExec_(plan,input,output,direction) bind(c, name="hipfftXtExec") -#endif - import - implicit none - integer(c_int) :: hipfftXtExec_ - type(c_ptr),value :: plan - type(c_ptr),value :: input - type(c_ptr),value :: output - integer(c_int),value :: direction - end function - end interface - - !> @brief Instructs hipFFT to use multiple GPUs for a plan. - !> - !> This function must be called after the plan is allocated using - !> `hipfftCreate`, but before the plan is initialized by any of the - !> "MakePlan" functions. Therefore, API functions that combine - !> creation and initialization (`hipfftPlan1d`, `hipfftPlan2d`, - !> `hipfftPlan3d`, and `hipfftPlanMany`) cannot use multiple GPUs. - !> - !> @note The rocFFT backend currently requires all device IDs in - !> the `gpus` array to be distinct. - !> - !> @param[in, out] plan - The FFT plan. - !> @param[in] count - Number of GPUs (length of the `gpus` array). - !> @param[in] gpus - Array of device IDs. - !> - !> @warning Experimental - interface hipfftXtSetGPUs -#ifdef USE_CUDA_NAMES - function hipfftXtSetGPUs_(plan,count,gpus) bind(c, name="cufftXtSetGPUs") -#else - function hipfftXtSetGPUs_(plan,count,gpus) bind(c, name="hipfftXtSetGPUs") -#endif - import - implicit none - integer(c_int) :: hipfftXtSetGPUs_ - type(c_ptr),value :: plan - integer(c_int),value :: count - type(c_ptr),value :: gpus - end function - -#ifdef USE_ASSUMED_RANK - module procedure hipfftXtSetGPUs_assumed_rank -#else -#ifdef USE_ASSUMED_SHAPE - module procedure & - hipfftXtSetGPUs_rank_0,& - hipfftXtSetGPUs_rank_1 -#endif -#endif - end interface - - !> @brief Allocates a multi-device descriptor (i.e., a `hipLibXtDesc` instance), - !> tailored to a specific multi-device plan's needs. The allocated descriptor - !> must be freed using `hipfftXtFree` when no longer needed. - !> - !> @details The data set of interest that is considered for the definition of - !> the desired descriptor is the plan's expected input data set, unless - !> `format` is `HIPFFT_XT_FORMAT_OUTPUT`. For values of `format` consistent - !> with in-place operations (i.e., `HIPFFT_XT_FORMAT_INPLACE` or - !> `HIPFFT_XT_FORMAT_INPLACE_SHUFFLED`), the created descriptor is usable for - !> both the plan's input and output data sets. - !> - !> @note Multi-device descriptors resulting in a division of real data sets - !> along their fastest dimension are not supported. As a result, - !> `hipfftXtMalloc` does not produce usable descriptors for requests to - !> create a multi-device descriptor of subformat value - !> `HIPFFT_XT_FORMAT_INPLACE_SHUFFLED` (resp. `HIPFFT_XT_FORMAT_INPLACE`) - !> given a plan initialized for an unbatched two-dimensional real forward - !> (resp. inverse) transform. - !> - !> @param[in] plan - Multi-device FFT plan to allocate descriptor memory for. - !> @param[out] desc - Address of the pointer to the created descriptor. - !> @param[in] myFormat - Desired descriptor's subformat, identifying the plan's - !> I/O data set of interest (see `hipfftXtSubFormat` for more details). - !> - !> @warning Experimental - interface hipfftXtMalloc -#ifdef USE_CUDA_NAMES - function hipfftXtMalloc_(plan,desc,myFormat) bind(c, name="cufftXtMalloc") -#else - function hipfftXtMalloc_(plan,desc,myFormat) bind(c, name="hipfftXtMalloc") -#endif - import - implicit none - integer(c_int) :: hipfftXtMalloc_ - type(c_ptr),value :: plan - type(c_ptr) :: desc - integer(c_int),value :: myFormat - end function - end interface - - !> @brief Copy data to/from `hipLibXtDesc` descriptors. - !> - !> @details If `type` is `HIPFFT_COPY_HOST_TO_DEVICE` (resp. - !> `HIPFFT_COPY_DEVICE_TO_HOST`), `src` (resp. `dest`) must point to a - !> host-residing memory buffer and `dest` (resp. `src`) must point to a - !> `hipLibXtDesc` descriptor. That descriptor must be fit for storing the plan's - !> input (resp. output) data set assuming it's distributed consistently with the - !> descriptor's own `subFormat` value (see `hipfftXtSubFormat` for more details). - !> The host buffer must be large enough to hold the plan's input (resp. output) - !> data (accounting for possible padding in real domain for plans and descriptors - !> specific to data for in-place real transforms). - !> `HIPFFT_COPY_DEVICE_TO_DEVICE` copies are supported only from descriptors of - !> format `HIPFFT_XT_FORMAT_INPLACE_SHUFFLED` to descriptors of format - !> `HIPFFT_XT_FORMAT_INPLACE` along with plans that support both formats. - !> - !> @param[in] plan - Multi-device FFT plan of interest, defining the data set of - !> relevance for the desired copy operation. - !> @param[out] dest - Opaque pointer to the copy destination. - !> @param[in] src - Opaque pointer to the copy source. - !> @param[in] myType - Type of copy operation to perform. - !> - !> @warning Experimental - interface hipfftXtMemcpy -#ifdef USE_CUDA_NAMES - function hipfftXtMemcpy_(plan,dest,src,myType) bind(c, name="cufftXtMemcpy") -#else - function hipfftXtMemcpy_(plan,dest,src,myType) bind(c, name="hipfftXtMemcpy") -#endif - import - implicit none - integer(c_int) :: hipfftXtMemcpy_ - type(c_ptr),value :: plan - type(c_ptr),value :: dest - type(c_ptr),value :: src - integer(c_int),value :: myType - end function - end interface - - !> @brief Free memory allocated by `hipfftXtMalloc`. - !> - !> @param[in] desc - Descriptor whose memory will be freed. - !> - !> @warning Experimental - interface hipfftXtFree -#ifdef USE_CUDA_NAMES - function hipfftXtFree_(desc) bind(c, name="cufftXtFree") -#else - function hipfftXtFree_(desc) bind(c, name="hipfftXtFree") -#endif - import - implicit none - integer(c_int) :: hipfftXtFree_ - type(hipLibXtDesc) :: desc - end function - end interface - - !> @ingroup hipfftXtExecDescriptor - !> - !> @param[in] plan - The FFT plan. - !> @param[in] input - Input data. - !> @param[out] output - Output data. - !> @param[in] direction - Either `HIPFFT_FORWARD` or `HIPFFT_BACKWARD`. - !> - !> @brief Execute single-precision complex-to-complex transform. - interface hipfftXtExecDescriptorC2C -#ifdef USE_CUDA_NAMES - function hipfftXtExecDescriptorC2C_(plan,input,output,direction) & - bind(c, name="cufftXtExecDescriptorC2C") -#else - function hipfftXtExecDescriptorC2C_(plan,input,output,direction) & - bind(c, name="hipfftXtExecDescriptorC2C") -#endif - import - implicit none - integer(c_int) :: hipfftXtExecDescriptorC2C_ - type(c_ptr),value :: plan - type(hipLibXtDesc) :: input - type(hipLibXtDesc) :: output - integer(c_int),value :: direction - end function - end interface - - !> @ingroup hipfftXtExecDescriptor - !> - !> @param[in] plan - The FFT plan. - !> @param[in] input - Input data. - !> @param[out] output - Output data. - !> - !> @brief Execute single-precision real forward transform. - interface hipfftXtExecDescriptorR2C -#ifdef USE_CUDA_NAMES - function hipfftXtExecDescriptorR2C_(plan,input,output) bind(c, name="cufftXtExecDescriptorR2C") -#else - function hipfftXtExecDescriptorR2C_(plan,input,output) bind(c, name="hipfftXtExecDescriptorR2C") -#endif - import - implicit none - integer(c_int) :: hipfftXtExecDescriptorR2C_ - type(c_ptr),value :: plan - type(hipLibXtDesc) :: input - type(hipLibXtDesc) :: output - end function - end interface - - !> @ingroup hipfftXtExecDescriptor - !> - !> @param[in] plan - The FFT plan. - !> @param[in] input - Input data. - !> @param[out] output - Output data. - !> - !> @brief Execute single-precision real backward transform. - interface hipfftXtExecDescriptorC2R -#ifdef USE_CUDA_NAMES - function hipfftXtExecDescriptorC2R_(plan,input,output) bind(c, name="cufftXtExecDescriptorC2R") -#else - function hipfftXtExecDescriptorC2R_(plan,input,output) bind(c, name="hipfftXtExecDescriptorC2R") -#endif - import - implicit none - integer(c_int) :: hipfftXtExecDescriptorC2R_ - type(c_ptr),value :: plan - type(hipLibXtDesc) :: input - type(hipLibXtDesc) :: output - end function - end interface - - !> @ingroup hipfftXtExecDescriptor - !> - !> @param[in] plan - The FFT plan. - !> @param[in] input - Input data. - !> @param[out] output - Output data. - !> @param[in] direction - Either `HIPFFT_FORWARD` or `HIPFFT_BACKWARD`. - !> - !> @brief Execute double-precision complex-to-complex transform. - interface hipfftXtExecDescriptorZ2Z -#ifdef USE_CUDA_NAMES - function hipfftXtExecDescriptorZ2Z_(plan,input,output,direction) & - bind(c, name="cufftXtExecDescriptorZ2Z") -#else - function hipfftXtExecDescriptorZ2Z_(plan,input,output,direction) & - bind(c, name="hipfftXtExecDescriptorZ2Z") -#endif - import - implicit none - integer(c_int) :: hipfftXtExecDescriptorZ2Z_ - type(c_ptr),value :: plan - type(hipLibXtDesc) :: input - type(hipLibXtDesc) :: output - integer(c_int),value :: direction - end function - end interface - - !> @ingroup hipfftXtExecDescriptor - !> - !> @param[in] plan - The FFT plan. - !> @param[in] input - Input data. - !> @param[out] output - Output data. - !> - !> @brief Execute double-precision real forward transform. - interface hipfftXtExecDescriptorD2Z -#ifdef USE_CUDA_NAMES - function hipfftXtExecDescriptorD2Z_(plan,input,output) bind(c, name="cufftXtExecDescriptorD2Z") -#else - function hipfftXtExecDescriptorD2Z_(plan,input,output) bind(c, name="hipfftXtExecDescriptorD2Z") -#endif - import - implicit none - integer(c_int) :: hipfftXtExecDescriptorD2Z_ - type(c_ptr),value :: plan - type(hipLibXtDesc) :: input - type(hipLibXtDesc) :: output - end function - end interface - - !> @ingroup hipfftXtExecDescriptor - !> - !> @param[in] plan - The FFT plan. - !> @param[in] input - Input data. - !> @param[out] output - Output data. - !> - !> @brief Execute double-precision real backward transform. - interface hipfftXtExecDescriptorZ2D -#ifdef USE_CUDA_NAMES - function hipfftXtExecDescriptorZ2D_(plan,input,output) bind(c, name="cufftXtExecDescriptorZ2D") -#else - function hipfftXtExecDescriptorZ2D_(plan,input,output) bind(c, name="hipfftXtExecDescriptorZ2D") -#endif - import - implicit none - integer(c_int) :: hipfftXtExecDescriptorZ2D_ - type(c_ptr),value :: plan - type(hipLibXtDesc) :: input - type(hipLibXtDesc) :: output - end function - end interface - - !> @ingroup hipfftXtExecDescriptor - !> - !> @param[in] plan - The FFT plan. - !> @param[in] input - Input data. - !> @param[out] output - Output data. - !> @param[in] direction - Either `HIPFFT_FORWARD` or `HIPFFT_BACKWARD`. - !> - !> @brief Execute general transform - types of the descriptors must match the plan. - interface hipfftXtExecDescriptor -#ifdef USE_CUDA_NAMES - function hipfftXtExecDescriptor_(plan,input,output,direction) & - bind(c, name="cufftXtExecDescriptor") -#else - function hipfftXtExecDescriptor_(plan,input,output,direction) & - bind(c, name="hipfftXtExecDescriptor") -#endif - import - implicit none - integer(c_int) :: hipfftXtExecDescriptor_ - type(c_ptr),value :: plan - type(hipLibXtDesc) :: input - type(hipLibXtDesc) :: output - integer(c_int),value :: direction - end function - end interface - - !> @brief Set a multi-processing communicator on a plan. - !> - !> @details Attach a multi-processing communication handle to a - !> hipFFT plan. 'comm_handle' points to the handle, whose type - !> depends on the multi-processing library being used. With MPI - !> (Message Passing Interface) for example, 'comm_handle' points to - !> an MPI communicator. - !> - !> This function must only be called on a plan that has already - !> been allocated by `hipfftCreate`, but before the plan is - !> initialized with any of the 'hipfftMakePlan' functions. - !> - !> @param[in] plan - The FFT plan. - !> @param[in] comm_type - Type of communication handle. - !> @param[in] comm_handle - Pointer to the communication handle. - !> - !> @warning Experimental -#ifndef USE_CUDA_NAMES - interface hipfftMpAttachComm - function hipfftMpAttachComm_(plan,comm_type,comm_handle) bind(c, name="hipfftMpAttachComm") - import - implicit none - integer(c_int) :: hipfftMpAttachComm_ - type(c_ptr),value :: plan - integer(c_int),value :: comm_type - type(c_ptr),value :: comm_handle - end function - end interface -#endif - - !> @brief Describe the partial input and output for a distributed transform in the current - !> process. - !> - !> @details For the given plan, describe the 'brick' of input and - !> output that is distributed to the current process. A brick is - !> defined with the lower and upper coordinates of the that brick - !> in the global index space of the transform. - !> - !> Strides for both the input and output bricks are also provided, - !> to describe how the brick is laid out in memory. - !> - !> All coordinates and strides are in row-major order, with the - !> slowest-moving dimension specified first. - !> - !> This function must only be called on a plan that has already - !> been allocated by `hipfftCreate`, but before the plan is - !> initialized with any of the 'hipfftMakePlan' functions. - !> - !> @param[in] plan - The FFT plan. - !> @param[in] rank - Dimension of the transform - !> @param[in] input_lower - Array of length rank, specifying the lower index (inclusive) for the - !> brick in the FFT input. - !> @param[in] input_upper - Array of length rank, specifying the upper index (exclusive) for the - !> brick in the FFT input. - !> @param[in] output_lower - Array of length rank, specifying the lower index (inclusive) for the - !> brick in the FFT output. - !> @param[in] output_upper - Array of length rank, specifying the upper index (exclusive) for the - !> brick in the FFT output. - !> @param[in] input_stride - Array of length rank specifying the input brick's stride in memory - !> @param[in] output_stride - Array of length rank specifying the output brick's stride in memory - !> - !> @warning Experimental -#ifndef USE_CUDA_NAMES - interface hipfftXtSetDistribution - function hipfftXtSetDistribution_(plan,rank,input_lower,input_upper,output_lower,output_upper, & - input_stride,output_stride) & - bind(c, name="hipfftXtSetDistribution") - import - implicit none - integer(c_int) :: hipfftXtSetDistribution_ - type(c_ptr),value :: plan - integer(c_int),value :: rank - type(c_ptr),value :: input_lower - type(c_ptr),value :: input_upper - type(c_ptr),value :: output_lower - type(c_ptr),value :: output_upper - type(c_ptr),value :: input_stride - type(c_ptr),value :: output_stride - end function - end interface -#endif - - !> @brief Set the input and output formats for a distributed transform. - !> - !> @details Specifies the distribution of data for each side of a - !> distributed transform. Forward transforms use 'subformat_forward' - !> to describe input and 'subformat_inverse' to describe output. - !> Inverse transforms use 'subformat_inverse' to describe input and - !> 'subformat_forward' to describe output. - !> - !> 'subformat_forward' and 'subformat_inverse' must be set to - !> matching values. One may be HIPFFT_XT_FORMAT_INPLACE and the - !> other HIPFFT_XT_FORMAT_INPLACE_SHUFFLED; Alternatively, one may be - !> HIPFFT_XT_FORMAT_DISTRIBUTED_INPUT and the other - !> HIPFFT_XT_FORMAT_DISTRIBUTED_OUTPUT. - !> - !> This function must only be called on a plan that has already - !> been allocated by `hipfftCreate`, but before the plan is - !> initialized with any of the 'hipfftMakePlan' functions. - !> - !> @param[in] plan - The FFT plan. - !> @param[in] subformat_forward - Format of the input for a forward transform. - !> @param[in] subformat_inverse - Format of the input for an inverse transform. - !> - !> @warning Experimental -#ifndef USE_CUDA_NAMES - interface hipfftXtSetSubformatDefault - function hipfftXtSetSubformatDefault_(plan,subformat_forward,subformat_inverse) & - bind(c, name="hipfftXtSetSubformatDefault") - import - implicit none - integer(c_int) :: hipfftXtSetSubformatDefault_ - type(c_ptr),value :: plan - integer(c_int),value :: subformat_forward - integer(c_int),value :: subformat_inverse - end function - end interface -#endif - contains @@ -2064,52 +1322,6 @@ function hipfftGetProperty_typed(myType,myValue) hipfftGetProperty_typed = hipfftGetProperty_(myType,c_loc(myValue)) end function - function hipfftXtMakePlanMany_typed(plan,rank,n,inembed,istride,idist,inputType,onembed, & - ostride,odist,outputType,batch,workSize,executionType) - implicit none - integer(kind(HIPFFT_SUCCESS)) :: hipfftXtMakePlanMany_typed - type(c_ptr), value :: plan - integer(c_int) :: rank - type(c_ptr) :: n - type(c_ptr) :: inembed - integer(c_int64_t) :: istride - integer(c_int64_t) :: idist - integer(c_int) :: inputType - type(c_ptr) :: onembed - integer(c_int64_t) :: ostride - integer(c_int64_t) :: odist - integer(c_int) :: outputType - integer(c_int64_t) :: batch - integer(c_size_t),target :: workSize - integer(c_int) :: executionType - ! - hipfftXtMakePlanMany_typed = hipfftXtMakePlanMany_(plan,rank,n,inembed,istride,idist, & - inputType,onembed,ostride,odist,outputType,batch,c_loc(workSize),executionType) - end function - - function hipfftXtGetSizeMany_typed(plan,rank,n,inembed,istride,idist,inputType,onembed, & - ostride,odist,outputType,batch,workSize,executionType) - implicit none - integer(kind(HIPFFT_SUCCESS)) :: hipfftXtGetSizeMany_typed - type(c_ptr), value :: plan - integer(c_int) :: rank - type(c_ptr) :: n - type(c_ptr) :: inembed - integer(c_int64_t) :: istride - integer(c_int64_t) :: idist - integer(c_int) :: inputType - type(c_ptr) :: onembed - integer(c_int64_t) :: ostride - integer(c_int64_t) :: odist - integer(c_int) :: outputType - integer(c_int64_t) :: batch - integer(c_size_t),target :: workSize - integer(c_int) :: executionType - ! - hipfftXtGetSizeMany_typed = hipfftXtGetSizeMany_(plan,rank,n,inembed,istride,idist, & - inputType,onembed,ostride,odist,outputType,batch,c_loc(workSize),executionType) - end function - #if defined(USE_ASSUMED_SHAPE) || defined(USE_ASSUMED_RANK) @@ -2830,187 +2042,6 @@ function hipfftExecZ2D_rank_3(plan,idata,odata) hipfftExecZ2D_rank_3 = hipfftExecZ2D_(plan,c_loc(idata),c_loc(odata)) end function -#endif -#ifdef USE_ASSUMED_RANK - function hipfftXtMakePlanMany_assumed_rank(plan,rank,n,inembed,istride,idist,inputType, & - onembed,ostride,odist,outputType,batch,workSize,executionType) - implicit none - integer(kind(HIPFFT_SUCCESS)) :: hipfftXtMakePlanMany_assumed_rank - type(c_ptr), value :: plan - integer(c_int) :: rank - integer(c_int64_t),target,dimension(..) :: n - integer(c_int64_t),target,dimension(..) :: inembed - integer(c_int64_t) :: istride - integer(c_int64_t) :: idist - integer(c_int) :: inputType - integer(c_int64_t),target,dimension(..) :: onembed - integer(c_int64_t) :: ostride - integer(c_int64_t) :: odist - integer(c_int) :: outputType - integer(c_int64_t) :: batch - integer(c_size_t),target :: workSize - integer(c_int) :: executionType - ! - hipfftXtMakePlanMany_assumed_rank = hipfftXtMakePlanMany_(plan,rank,c_loc(n),c_loc(inembed), & - istride,idist,inputType,c_loc(onembed),ostride,odist,outputType,batch,c_loc(workSize), & - executionType) - end function - -#else - function hipfftXtMakePlanMany_rank_0(plan,rank,n,inembed,istride,idist,inputType,onembed, & - ostride,odist,outputType,batch,workSize,executionType) - implicit none - integer(kind(HIPFFT_SUCCESS)) :: hipfftXtMakePlanMany_rank_0 - type(c_ptr), value :: plan - integer(c_int) :: rank - integer(c_int64_t),target :: n - integer(c_int64_t),target :: inembed - integer(c_int64_t) :: istride - integer(c_int64_t) :: idist - integer(c_int) :: inputType - integer(c_int64_t),target :: onembed - integer(c_int64_t) :: ostride - integer(c_int64_t) :: odist - integer(c_int) :: outputType - integer(c_int64_t) :: batch - integer(c_size_t),target :: workSize - integer(c_int) :: executionType - ! - hipfftXtMakePlanMany_rank_0 = hipfftXtMakePlanMany_(plan,rank,c_loc(n),c_loc(inembed), & - istride,idist,inputType,c_loc(onembed),ostride,odist,outputType,batch,c_loc(workSize), & - executionType) - end function - - function hipfftXtMakePlanMany_rank_1(plan,rank,n,inembed,istride,idist,inputType,onembed, & - ostride,odist,outputType,batch,workSize,executionType) - implicit none - integer(kind(HIPFFT_SUCCESS)) :: hipfftXtMakePlanMany_rank_1 - type(c_ptr), value :: plan - integer(c_int) :: rank - integer(c_int64_t),target,dimension(:) :: n - integer(c_int64_t),target,dimension(:) :: inembed - integer(c_int64_t) :: istride - integer(c_int64_t) :: idist - integer(c_int) :: inputType - integer(c_int64_t),target,dimension(:) :: onembed - integer(c_int64_t) :: ostride - integer(c_int64_t) :: odist - integer(c_int) :: outputType - integer(c_int64_t) :: batch - integer(c_size_t),target :: workSize - integer(c_int) :: executionType - ! - hipfftXtMakePlanMany_rank_1 = hipfftXtMakePlanMany_(plan,rank,c_loc(n),c_loc(inembed), & - istride,idist,inputType,c_loc(onembed),ostride,odist,outputType,batch,c_loc(workSize), & - executionType) - end function - -#endif -#ifdef USE_ASSUMED_RANK - function hipfftXtGetSizeMany_assumed_rank(plan,rank,n,inembed,istride,idist,inputType,onembed, & - ostride,odist,outputType,batch,workSize,executionType) - implicit none - integer(kind(HIPFFT_SUCCESS)) :: hipfftXtGetSizeMany_assumed_rank - type(c_ptr), value :: plan - integer(c_int) :: rank - integer(c_int64_t),target,dimension(..) :: n - integer(c_int64_t),target,dimension(..) :: inembed - integer(c_int64_t) :: istride - integer(c_int64_t) :: idist - integer(c_int) :: inputType - integer(c_int64_t),target,dimension(..) :: onembed - integer(c_int64_t) :: ostride - integer(c_int64_t) :: odist - integer(c_int) :: outputType - integer(c_int64_t) :: batch - integer(c_size_t),target :: workSize - integer(c_int) :: executionType - ! - hipfftXtGetSizeMany_assumed_rank = hipfftXtGetSizeMany_(plan,rank,c_loc(n),c_loc(inembed), & - istride,idist,inputType,c_loc(onembed),ostride,odist,outputType,batch,c_loc(workSize), & - executionType) - end function - -#else - function hipfftXtGetSizeMany_rank_0(plan,rank,n,inembed,istride,idist,inputType,onembed, & - ostride,odist,outputType,batch,workSize,executionType) - implicit none - integer(kind(HIPFFT_SUCCESS)) :: hipfftXtGetSizeMany_rank_0 - type(c_ptr), value :: plan - integer(c_int) :: rank - integer(c_int64_t),target :: n - integer(c_int64_t),target :: inembed - integer(c_int64_t) :: istride - integer(c_int64_t) :: idist - integer(c_int) :: inputType - integer(c_int64_t),target :: onembed - integer(c_int64_t) :: ostride - integer(c_int64_t) :: odist - integer(c_int) :: outputType - integer(c_int64_t) :: batch - integer(c_size_t),target :: workSize - integer(c_int) :: executionType - ! - hipfftXtGetSizeMany_rank_0 = hipfftXtGetSizeMany_(plan,rank,c_loc(n),c_loc(inembed),istride, & - idist,inputType,c_loc(onembed),ostride,odist,outputType,batch,c_loc(workSize),executionType) - end function - - function hipfftXtGetSizeMany_rank_1(plan,rank,n,inembed,istride,idist,inputType,onembed, & - ostride,odist,outputType,batch,workSize,executionType) - implicit none - integer(kind(HIPFFT_SUCCESS)) :: hipfftXtGetSizeMany_rank_1 - type(c_ptr), value :: plan - integer(c_int) :: rank - integer(c_int64_t),target,dimension(:) :: n - integer(c_int64_t),target,dimension(:) :: inembed - integer(c_int64_t) :: istride - integer(c_int64_t) :: idist - integer(c_int) :: inputType - integer(c_int64_t),target,dimension(:) :: onembed - integer(c_int64_t) :: ostride - integer(c_int64_t) :: odist - integer(c_int) :: outputType - integer(c_int64_t) :: batch - integer(c_size_t),target :: workSize - integer(c_int) :: executionType - ! - hipfftXtGetSizeMany_rank_1 = hipfftXtGetSizeMany_(plan,rank,c_loc(n),c_loc(inembed),istride, & - idist,inputType,c_loc(onembed),ostride,odist,outputType,batch,c_loc(workSize),executionType) - end function - -#endif -#ifdef USE_ASSUMED_RANK - function hipfftXtSetGPUs_assumed_rank(plan,count,gpus) - implicit none - integer(kind(HIPFFT_SUCCESS)) :: hipfftXtSetGPUs_assumed_rank - type(c_ptr), value :: plan - integer(c_int) :: count - integer(c_int),target,dimension(..) :: gpus - ! - hipfftXtSetGPUs_assumed_rank = hipfftXtSetGPUs_(plan,count,c_loc(gpus)) - end function - -#else - function hipfftXtSetGPUs_rank_0(plan,count,gpus) - implicit none - integer(kind(HIPFFT_SUCCESS)) :: hipfftXtSetGPUs_rank_0 - type(c_ptr), value :: plan - integer(c_int) :: count - integer(c_int),target :: gpus - ! - hipfftXtSetGPUs_rank_0 = hipfftXtSetGPUs_(plan,count,c_loc(gpus)) - end function - - function hipfftXtSetGPUs_rank_1(plan,count,gpus) - implicit none - integer(kind(HIPFFT_SUCCESS)) :: hipfftXtSetGPUs_rank_1 - type(c_ptr), value :: plan - integer(c_int) :: count - integer(c_int),target,dimension(:) :: gpus - ! - hipfftXtSetGPUs_rank_1 = hipfftXtSetGPUs_(plan,count,c_loc(gpus)) - end function - #endif #endif end module hipfort_hipfft diff --git a/lib/hipfort/hipfort_hipfft_enums.F90 b/lib/hipfort/hipfort_hipfft_enums.F90 index a6dc191c4..da1944de7 100644 --- a/lib/hipfort/hipfort_hipfft_enums.F90 +++ b/lib/hipfort/hipfort_hipfft_enums.F90 @@ -64,58 +64,10 @@ module hipfort_hipfft_enums enumerator :: HIPFFT_PATCH_LEVEL = 2 end enum - ! hiplibFormat_t - enum, bind(c) - enumerator :: HIPLIB_FORMAT_HIPFFT = 0 - enumerator :: HIPLIB_FORMAT_UNDEFINED = 1 - end enum - - ! hipfftXtCopyType_t - enum, bind(c) - enumerator :: HIPFFT_COPY_HOST_TO_DEVICE = 0 - enumerator :: HIPFFT_COPY_DEVICE_TO_HOST = 1 - enumerator :: HIPFFT_COPY_DEVICE_TO_DEVICE = 2 - enumerator :: HIPFFT_COPY_UNDEFINED = 3 - end enum - - ! hipfftXtCallbackType_t - enum, bind(c) - enumerator :: HIPFFT_CB_LD_COMPLEX = 0 - enumerator :: HIPFFT_CB_LD_COMPLEX_DOUBLE = 1 - enumerator :: HIPFFT_CB_LD_REAL = 2 - enumerator :: HIPFFT_CB_LD_REAL_DOUBLE = 3 - enumerator :: HIPFFT_CB_ST_COMPLEX = 4 - enumerator :: HIPFFT_CB_ST_COMPLEX_DOUBLE = 5 - enumerator :: HIPFFT_CB_ST_REAL = 6 - enumerator :: HIPFFT_CB_ST_REAL_DOUBLE = 7 - enumerator :: HIPFFT_CB_UNDEFINED = 8 - end enum - - ! hipfftXtSubFormat_t - enum, bind(c) - enumerator :: HIPFFT_XT_FORMAT_INPUT = 0 - enumerator :: HIPFFT_XT_FORMAT_OUTPUT = 1 - enumerator :: HIPFFT_XT_FORMAT_INPLACE = 2 - enumerator :: HIPFFT_XT_FORMAT_INPLACE_SHUFFLED = 3 - enumerator :: HIPFFT_XT_FORMAT_1D_INPUT_SHUFFLED = 4 -#ifdef USE_CUDA_NAMES - enumerator :: HIPFFT_FORMAT_UNDEFINED = 7 -#else - enumerator :: HIPFFT_FORMAT_UNDEFINED = 5 -#endif - end enum - - ! hipfftMpCommType_t - enum, bind(c) - enumerator :: HIPFFT_COMM_MPI = 0 - enumerator :: HIPFFT_COMM_NONE = 1 - end enum - integer(c_int), parameter :: hipfftVersionMajor = 1 integer(c_int), parameter :: hipfftVersionMinor = 0 integer(c_int), parameter :: hipfftVersionPatch = 27 integer(c_int), parameter :: HIPFFT_FORWARD = -1 integer(c_int), parameter :: HIPFFT_BACKWARD = 1 - integer(c_int), parameter :: MAX_HIP_DESCRIPTOR_GPUS = 64 end module hipfort_hipfft_enums diff --git a/lib/hipfort/hipfort_hipfftxt.F90 b/lib/hipfort/hipfort_hipfftxt.F90 new file mode 100644 index 000000000..0d282395e --- /dev/null +++ b/lib/hipfort/hipfort_hipfftxt.F90 @@ -0,0 +1,1007 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +module hipfort_hipfftxt + use hipfort_hipfft + use hipfort_hipfftxt_enums + use hipfort_hipfftxt_types + implicit none + + !> @brief Generalized version of `hipfftSetWorkArea` accepting multi-device plans. + !> + !> @param[in] plan - Pointer to the (possibly multi-device) FFT plan. + !> @param[in] workArea - Array of pointer(s) to the plan's work area(s), which must + !> be accessible by the plan's used device(s) (in the same order as communicated + !> via `hipfftXtSetGPUs`, if used prior). + interface hipfftXtSetWorkArea +#ifdef USE_CUDA_NAMES + function hipfftXtSetWorkArea_(plan,workArea) bind(c, name="cufftXtSetWorkArea") +#else + function hipfftXtSetWorkArea_(plan,workArea) bind(c, name="hipfftXtSetWorkArea") +#endif + import + implicit none + integer(c_int) :: hipfftXtSetWorkArea_ + type(c_ptr),value :: plan + type(c_ptr) :: workArea + end function + end interface + + interface hipfftXtSetCallback +#ifdef USE_CUDA_NAMES + function hipfftXtSetCallback_(plan,callbacks,cbtype,callbackData) & + bind(c, name="cufftXtSetCallback") +#else + function hipfftXtSetCallback_(plan,callbacks,cbtype,callbackData) & + bind(c, name="hipfftXtSetCallback") +#endif + import + implicit none + integer(c_int) :: hipfftXtSetCallback_ + type(c_ptr),value :: plan + type(c_ptr) :: callbacks + integer(c_int),value :: cbtype + type(c_ptr) :: callbackData + end function + end interface + + interface hipfftXtClearCallback +#ifdef USE_CUDA_NAMES + function hipfftXtClearCallback_(plan,cbtype) bind(c, name="cufftXtClearCallback") +#else + function hipfftXtClearCallback_(plan,cbtype) bind(c, name="hipfftXtClearCallback") +#endif + import + implicit none + integer(c_int) :: hipfftXtClearCallback_ + type(c_ptr),value :: plan + integer(c_int),value :: cbtype + end function + end interface + + !> @brief Set shared memory size for callback. + !> + !> @details Set shared memory required for a callback. The + !> callback of the specified type must have already been set on the + !> plan. + !> + !> @param[in] plan - The FFT plan. + !> @param[in] cbtype - Type of callback being modified. + !> @param[in] sharedSize - Amount of shared memory required, in bytes. + interface hipfftXtSetCallbackSharedSize +#ifdef USE_CUDA_NAMES + function hipfftXtSetCallbackSharedSize_(plan,cbtype,sharedSize) & + bind(c, name="cufftXtSetCallbackSharedSize") +#else + function hipfftXtSetCallbackSharedSize_(plan,cbtype,sharedSize) & + bind(c, name="hipfftXtSetCallbackSharedSize") +#endif + import + implicit none + integer(c_int) :: hipfftXtSetCallbackSharedSize_ + type(c_ptr),value :: plan + integer(c_int),value :: cbtype + integer(c_size_t),value :: sharedSize + end function + end interface + + !> @brief Set a JIT callback on a plan + !> + !> @details Set either a load or store callback to run with a plan. + !> The type of callback is specified with the `cbtype` parameter. + !> The callback is provided as SPIR-V on AMD platforms and as + !> LTO-IR fatbin on CUDA platforms. + !> + !> Note: JIT callbacks cannot currently be used on multi-GPU transforms. + !> This support will be added in a future release of hipFFT. + !> + !> JIT callbacks can be cleared from the plan by passing nullptr + !> for `symbol_name` and `bitcode_data`, and 0 for + !> `bitcode_len_bytes`. + !> + !> `cbdata` is an optional array of pointers of data that is passed + !> to the callback function. As JIT callbacks are currently only + !> supported for single-device transforms, this array must have one + !> pointer element in it if it is specified. The pointer will be + !> passed to the callback function's `cbdata` parameter when it is + !> executed on that device. + !> + !> If specified, the `cbdata` array must remain valid until either + !> the plan is destroyed or callbacks are cleared from the plan. + !> + !> This function must be called after the plan is allocated using + !> `hipfftCreate`, but before the plan is initialized by any of the + !> "MakePlan" functions. Therefore, API functions that combine + !> creation and initialization (`hipfftPlan1d`, `hipfftPlan2d`, + !> `hipfftPlan3d`, and `hipfftPlanMany`) cannot set a JIT callback. + !> + !> @param[in] plan - The FFT plan. + !> @param[in] symbol_name - Name of the symbol in the compiled bitcode. + !> @param[in] bitcode_data - Pointer to bitcode data. + !> @param[in] bitcode_len_bytes - Number of bytes in the bitcode data. + !> @param[in] cbtype - Type of callback function. + !> @param[in] cbdata - Array of callback function data pointers. + interface hipfftXtSetJITCallback +#ifdef USE_CUDA_NAMES + function hipfftXtSetJITCallback_(plan,symbol_name,bitcode_data,bitcode_len_bytes,cbtype, & + cbdata) & + bind(c, name="cufftXtSetJITCallback") +#else + function hipfftXtSetJITCallback_(plan,symbol_name,bitcode_data,bitcode_len_bytes,cbtype, & + cbdata) & + bind(c, name="hipfftXtSetJITCallback") +#endif + import + implicit none + integer(c_int) :: hipfftXtSetJITCallback_ + type(c_ptr),value :: plan + type(c_ptr),value :: symbol_name + type(c_ptr),value :: bitcode_data + integer(c_size_t),value :: bitcode_len_bytes + integer(c_int),value :: cbtype + type(c_ptr) :: cbdata + end function + end interface + + !> @brief Initialize a batched rank-dimensional FFT plan with + !> advanced data layout and specified input, output, execution data + !> types. + !> + !> @details Assumes that the plan has been created already, and + !> modifies the plan associated with the plan handle. The number + !> of elements to transform in each direction of the input data + !> in the FFT plan is specified in n. + !> + !> The batch parameter tells hipFFT how many transforms to perform. + !> The distance between the first elements of two consecutive batches + !> of the input and output data are specified with the idist and odist + !> parameters. + !> + !> The inembed and onembed parameters define the input and output data + !> layouts. The number of elements in the data is assumed to be larger + !> than the number of elements in the transform. Strided data layouts + !> are also supported. Strides along the fastest direction in the input + !> and output data are specified via the istride and ostride parameters. + !> + !> If both inembed and onembed parameters are set to NULL, all the + !> advanced data layout parameters are ignored and reverted to default + !> values, i.e., the batched transform is performed with non-strided data + !> access and the number of data/transform elements are assumed to be + !> equivalent. + !> + !> The inputType, outputType, executionType parameters specify the + !> data types (precision, and whether the data is real-valued or + !> complex-valued) of the transform input, output, and internal + !> representation respectively. Currently, the precision of all + !> three parameters must match, and the execution type must always + !> be complex-valued. At least one of inputType and outputType + !> must be complex. A half-precision transform can be requested + !> by using either the HIP_R_16F or HIP_C_16F types. + !> + !> @param[out] plan - Pointer to the FFT plan handle. + !> @param[in] rank - Dimension of transform (1, 2, or 3). + !> @param[in] n - Number of elements to transform in the x/y/z directions. + !> @param[in] inembed - Number of elements in the input data in the x/y/z directions. + !> @param[in] istride - Distance between two successive elements in the input data. + !> @param[in] idist - Distance between input batches. + !> @param[in] inputType - Format of FFT input. + !> @param[in] onembed - Number of elements in the output data in the x/y/z directions. + !> @param[in] ostride - Distance between two successive elements in the output data. + !> @param[in] odist - Distance between output batches. + !> @param[in] outputType - Format of FFT output. + !> @param[in] batch - Number of batched transforms to perform. + !> @param[out] workSize - Pointer to work area size (returned value). + !> @param[in] executionType - Internal data format used by the library during computation. + interface hipfftXtMakePlanMany +#ifdef USE_CUDA_NAMES + function hipfftXtMakePlanMany_(plan,rank,n,inembed,istride,idist,inputType,onembed,ostride, & + odist,outputType,batch,workSize,executionType) & + bind(c, name="cufftXtMakePlanMany") +#else + function hipfftXtMakePlanMany_(plan,rank,n,inembed,istride,idist,inputType,onembed,ostride, & + odist,outputType,batch,workSize,executionType) & + bind(c, name="hipfftXtMakePlanMany") +#endif + import + implicit none + integer(c_int) :: hipfftXtMakePlanMany_ + type(c_ptr),value :: plan + integer(c_int),value :: rank + type(c_ptr),value :: n + type(c_ptr),value :: inembed + integer(c_int64_t),value :: istride + integer(c_int64_t),value :: idist + integer(c_int),value :: inputType + type(c_ptr),value :: onembed + integer(c_int64_t),value :: ostride + integer(c_int64_t),value :: odist + integer(c_int),value :: outputType + integer(c_int64_t),value :: batch + type(c_ptr),value :: workSize + integer(c_int),value :: executionType + end function + + module procedure hipfftXtMakePlanMany_typed + +#ifdef USE_ASSUMED_RANK + module procedure hipfftXtMakePlanMany_assumed_rank +#else +#ifdef USE_ASSUMED_SHAPE + module procedure & + hipfftXtMakePlanMany_rank_0,& + hipfftXtMakePlanMany_rank_1 +#endif +#endif + end interface + + !> @brief Return size of the work area size required for a + !> rank-dimensional plan, with specified input, output, execution + !> data types. + !> + !> @details See `hipfftXtMakePlanMany` for restrictions on inputType, + !> outputType, executionType parameters. + !> + !> @param[in] plan - Pointer to the FFT plan. + !> @param[in] rank - Dimension of FFT transform (1, 2, or 3). + !> @param[in] n - Number of elements in the x/y/z directions. + !> @param[in] inembed - Number of elements in the input data in the x/y/z directions. + !> @param[in] istride - Distance between two successive elements in the input data. + !> @param[in] idist - Distance between input batches. + !> @param[in] inputType - Format of FFT input. + !> @param[in] onembed - Number of elements in the output data in the x/y/z directions. + !> @param[in] ostride - Distance between two successive elements in the output data. + !> @param[in] odist - Distance between output batches. + !> @param[in] outputType - Format of FFT output. + !> @param[in] batch - Number of batched transforms to perform. + !> @param[out] workSize - Pointer to work area size (returned value). + !> @param[in] executionType - Internal data format used by the library during computation. + interface hipfftXtGetSizeMany +#ifdef USE_CUDA_NAMES + function hipfftXtGetSizeMany_(plan,rank,n,inembed,istride,idist,inputType,onembed,ostride, & + odist,outputType,batch,workSize,executionType) & + bind(c, name="cufftXtGetSizeMany") +#else + function hipfftXtGetSizeMany_(plan,rank,n,inembed,istride,idist,inputType,onembed,ostride, & + odist,outputType,batch,workSize,executionType) & + bind(c, name="hipfftXtGetSizeMany") +#endif + import + implicit none + integer(c_int) :: hipfftXtGetSizeMany_ + type(c_ptr),value :: plan + integer(c_int),value :: rank + type(c_ptr),value :: n + type(c_ptr),value :: inembed + integer(c_int64_t),value :: istride + integer(c_int64_t),value :: idist + integer(c_int),value :: inputType + type(c_ptr),value :: onembed + integer(c_int64_t),value :: ostride + integer(c_int64_t),value :: odist + integer(c_int),value :: outputType + integer(c_int64_t),value :: batch + type(c_ptr),value :: workSize + integer(c_int),value :: executionType + end function + + module procedure hipfftXtGetSizeMany_typed + +#ifdef USE_ASSUMED_RANK + module procedure hipfftXtGetSizeMany_assumed_rank +#else +#ifdef USE_ASSUMED_SHAPE + module procedure & + hipfftXtGetSizeMany_rank_0,& + hipfftXtGetSizeMany_rank_1 +#endif +#endif + end interface + + !> @brief Execute an FFT plan for any precision and type. + !> + !> @details An in-place transform is performed if the input and + !> output pointers have the same value. + !> + !> The direction parameter is ignored if for real-to-complex and + !> complex-to-real transforms, as the direction is already implied by + !> the data types. + !> + !> @param[in] plan - Pointer to the FFT plan. + !> @param[in] input - Pointer to input data for the transform. + !> @param[in] output - Pointer to output data for the transform. + !> @param[in] direction - Either `HIPFFT_FORWARD` or `HIPFFT_BACKWARD`. + interface hipfftXtExec +#ifdef USE_CUDA_NAMES + function hipfftXtExec_(plan,input,output,direction) bind(c, name="cufftXtExec") +#else + function hipfftXtExec_(plan,input,output,direction) bind(c, name="hipfftXtExec") +#endif + import + implicit none + integer(c_int) :: hipfftXtExec_ + type(c_ptr),value :: plan + type(c_ptr),value :: input + type(c_ptr),value :: output + integer(c_int),value :: direction + end function + end interface + + !> @brief Instructs hipFFT to use multiple GPUs for a plan. + !> + !> This function must be called after the plan is allocated using + !> `hipfftCreate`, but before the plan is initialized by any of the + !> "MakePlan" functions. Therefore, API functions that combine + !> creation and initialization (`hipfftPlan1d`, `hipfftPlan2d`, + !> `hipfftPlan3d`, and `hipfftPlanMany`) cannot use multiple GPUs. + !> + !> @note The rocFFT backend currently requires all device IDs in + !> the `gpus` array to be distinct. + !> + !> @param[in, out] plan - The FFT plan. + !> @param[in] count - Number of GPUs (length of the `gpus` array). + !> @param[in] gpus - Array of device IDs. + !> + !> @warning Experimental + interface hipfftXtSetGPUs +#ifdef USE_CUDA_NAMES + function hipfftXtSetGPUs_(plan,count,gpus) bind(c, name="cufftXtSetGPUs") +#else + function hipfftXtSetGPUs_(plan,count,gpus) bind(c, name="hipfftXtSetGPUs") +#endif + import + implicit none + integer(c_int) :: hipfftXtSetGPUs_ + type(c_ptr),value :: plan + integer(c_int),value :: count + type(c_ptr),value :: gpus + end function + +#ifdef USE_ASSUMED_RANK + module procedure hipfftXtSetGPUs_assumed_rank +#else +#ifdef USE_ASSUMED_SHAPE + module procedure & + hipfftXtSetGPUs_rank_0,& + hipfftXtSetGPUs_rank_1 +#endif +#endif + end interface + + !> @brief Allocates a multi-device descriptor (i.e., a `hipLibXtDesc` instance), + !> tailored to a specific multi-device plan's needs. The allocated descriptor + !> must be freed using `hipfftXtFree` when no longer needed. + !> + !> @details The data set of interest that is considered for the definition of + !> the desired descriptor is the plan's expected input data set, unless + !> `format` is `HIPFFT_XT_FORMAT_OUTPUT`. For values of `format` consistent + !> with in-place operations (i.e., `HIPFFT_XT_FORMAT_INPLACE` or + !> `HIPFFT_XT_FORMAT_INPLACE_SHUFFLED`), the created descriptor is usable for + !> both the plan's input and output data sets. + !> + !> @note Multi-device descriptors resulting in a division of real data sets + !> along their fastest dimension are not supported. As a result, + !> `hipfftXtMalloc` does not produce usable descriptors for requests to + !> create a multi-device descriptor of subformat value + !> `HIPFFT_XT_FORMAT_INPLACE_SHUFFLED` (resp. `HIPFFT_XT_FORMAT_INPLACE`) + !> given a plan initialized for an unbatched two-dimensional real forward + !> (resp. inverse) transform. + !> + !> @param[in] plan - Multi-device FFT plan to allocate descriptor memory for. + !> @param[out] desc - Address of the pointer to the created descriptor. + !> @param[in] myFormat - Desired descriptor's subformat, identifying the plan's + !> I/O data set of interest (see `hipfftXtSubFormat` for more details). + !> + !> @warning Experimental + interface hipfftXtMalloc +#ifdef USE_CUDA_NAMES + function hipfftXtMalloc_(plan,desc,myFormat) bind(c, name="cufftXtMalloc") +#else + function hipfftXtMalloc_(plan,desc,myFormat) bind(c, name="hipfftXtMalloc") +#endif + import + implicit none + integer(c_int) :: hipfftXtMalloc_ + type(c_ptr),value :: plan + type(c_ptr) :: desc + integer(c_int),value :: myFormat + end function + end interface + + !> @brief Copy data to/from `hipLibXtDesc` descriptors. + !> + !> @details If `type` is `HIPFFT_COPY_HOST_TO_DEVICE` (resp. + !> `HIPFFT_COPY_DEVICE_TO_HOST`), `src` (resp. `dest`) must point to a + !> host-residing memory buffer and `dest` (resp. `src`) must point to a + !> `hipLibXtDesc` descriptor. That descriptor must be fit for storing the plan's + !> input (resp. output) data set assuming it's distributed consistently with the + !> descriptor's own `subFormat` value (see `hipfftXtSubFormat` for more details). + !> The host buffer must be large enough to hold the plan's input (resp. output) + !> data (accounting for possible padding in real domain for plans and descriptors + !> specific to data for in-place real transforms). + !> `HIPFFT_COPY_DEVICE_TO_DEVICE` copies are supported only from descriptors of + !> format `HIPFFT_XT_FORMAT_INPLACE_SHUFFLED` to descriptors of format + !> `HIPFFT_XT_FORMAT_INPLACE` along with plans that support both formats. + !> + !> @param[in] plan - Multi-device FFT plan of interest, defining the data set of + !> relevance for the desired copy operation. + !> @param[out] dest - Opaque pointer to the copy destination. + !> @param[in] src - Opaque pointer to the copy source. + !> @param[in] myType - Type of copy operation to perform. + !> + !> @warning Experimental + interface hipfftXtMemcpy +#ifdef USE_CUDA_NAMES + function hipfftXtMemcpy_(plan,dest,src,myType) bind(c, name="cufftXtMemcpy") +#else + function hipfftXtMemcpy_(plan,dest,src,myType) bind(c, name="hipfftXtMemcpy") +#endif + import + implicit none + integer(c_int) :: hipfftXtMemcpy_ + type(c_ptr),value :: plan + type(c_ptr),value :: dest + type(c_ptr),value :: src + integer(c_int),value :: myType + end function + end interface + + !> @brief Free memory allocated by `hipfftXtMalloc`. + !> + !> @param[in] desc - Descriptor whose memory will be freed. + !> + !> @warning Experimental + interface hipfftXtFree +#ifdef USE_CUDA_NAMES + function hipfftXtFree_(desc) bind(c, name="cufftXtFree") +#else + function hipfftXtFree_(desc) bind(c, name="hipfftXtFree") +#endif + import + implicit none + integer(c_int) :: hipfftXtFree_ + type(hipLibXtDesc) :: desc + end function + end interface + + !> @ingroup hipfftXtExecDescriptor + !> + !> @param[in] plan - The FFT plan. + !> @param[in] input - Input data. + !> @param[out] output - Output data. + !> @param[in] direction - Either `HIPFFT_FORWARD` or `HIPFFT_BACKWARD`. + !> + !> @brief Execute single-precision complex-to-complex transform. + interface hipfftXtExecDescriptorC2C +#ifdef USE_CUDA_NAMES + function hipfftXtExecDescriptorC2C_(plan,input,output,direction) & + bind(c, name="cufftXtExecDescriptorC2C") +#else + function hipfftXtExecDescriptorC2C_(plan,input,output,direction) & + bind(c, name="hipfftXtExecDescriptorC2C") +#endif + import + implicit none + integer(c_int) :: hipfftXtExecDescriptorC2C_ + type(c_ptr),value :: plan + type(hipLibXtDesc) :: input + type(hipLibXtDesc) :: output + integer(c_int),value :: direction + end function + end interface + + !> @ingroup hipfftXtExecDescriptor + !> + !> @param[in] plan - The FFT plan. + !> @param[in] input - Input data. + !> @param[out] output - Output data. + !> + !> @brief Execute single-precision real forward transform. + interface hipfftXtExecDescriptorR2C +#ifdef USE_CUDA_NAMES + function hipfftXtExecDescriptorR2C_(plan,input,output) bind(c, name="cufftXtExecDescriptorR2C") +#else + function hipfftXtExecDescriptorR2C_(plan,input,output) bind(c, name="hipfftXtExecDescriptorR2C") +#endif + import + implicit none + integer(c_int) :: hipfftXtExecDescriptorR2C_ + type(c_ptr),value :: plan + type(hipLibXtDesc) :: input + type(hipLibXtDesc) :: output + end function + end interface + + !> @ingroup hipfftXtExecDescriptor + !> + !> @param[in] plan - The FFT plan. + !> @param[in] input - Input data. + !> @param[out] output - Output data. + !> + !> @brief Execute single-precision real backward transform. + interface hipfftXtExecDescriptorC2R +#ifdef USE_CUDA_NAMES + function hipfftXtExecDescriptorC2R_(plan,input,output) bind(c, name="cufftXtExecDescriptorC2R") +#else + function hipfftXtExecDescriptorC2R_(plan,input,output) bind(c, name="hipfftXtExecDescriptorC2R") +#endif + import + implicit none + integer(c_int) :: hipfftXtExecDescriptorC2R_ + type(c_ptr),value :: plan + type(hipLibXtDesc) :: input + type(hipLibXtDesc) :: output + end function + end interface + + !> @ingroup hipfftXtExecDescriptor + !> + !> @param[in] plan - The FFT plan. + !> @param[in] input - Input data. + !> @param[out] output - Output data. + !> @param[in] direction - Either `HIPFFT_FORWARD` or `HIPFFT_BACKWARD`. + !> + !> @brief Execute double-precision complex-to-complex transform. + interface hipfftXtExecDescriptorZ2Z +#ifdef USE_CUDA_NAMES + function hipfftXtExecDescriptorZ2Z_(plan,input,output,direction) & + bind(c, name="cufftXtExecDescriptorZ2Z") +#else + function hipfftXtExecDescriptorZ2Z_(plan,input,output,direction) & + bind(c, name="hipfftXtExecDescriptorZ2Z") +#endif + import + implicit none + integer(c_int) :: hipfftXtExecDescriptorZ2Z_ + type(c_ptr),value :: plan + type(hipLibXtDesc) :: input + type(hipLibXtDesc) :: output + integer(c_int),value :: direction + end function + end interface + + !> @ingroup hipfftXtExecDescriptor + !> + !> @param[in] plan - The FFT plan. + !> @param[in] input - Input data. + !> @param[out] output - Output data. + !> + !> @brief Execute double-precision real forward transform. + interface hipfftXtExecDescriptorD2Z +#ifdef USE_CUDA_NAMES + function hipfftXtExecDescriptorD2Z_(plan,input,output) bind(c, name="cufftXtExecDescriptorD2Z") +#else + function hipfftXtExecDescriptorD2Z_(plan,input,output) bind(c, name="hipfftXtExecDescriptorD2Z") +#endif + import + implicit none + integer(c_int) :: hipfftXtExecDescriptorD2Z_ + type(c_ptr),value :: plan + type(hipLibXtDesc) :: input + type(hipLibXtDesc) :: output + end function + end interface + + !> @ingroup hipfftXtExecDescriptor + !> + !> @param[in] plan - The FFT plan. + !> @param[in] input - Input data. + !> @param[out] output - Output data. + !> + !> @brief Execute double-precision real backward transform. + interface hipfftXtExecDescriptorZ2D +#ifdef USE_CUDA_NAMES + function hipfftXtExecDescriptorZ2D_(plan,input,output) bind(c, name="cufftXtExecDescriptorZ2D") +#else + function hipfftXtExecDescriptorZ2D_(plan,input,output) bind(c, name="hipfftXtExecDescriptorZ2D") +#endif + import + implicit none + integer(c_int) :: hipfftXtExecDescriptorZ2D_ + type(c_ptr),value :: plan + type(hipLibXtDesc) :: input + type(hipLibXtDesc) :: output + end function + end interface + + !> @ingroup hipfftXtExecDescriptor + !> + !> @param[in] plan - The FFT plan. + !> @param[in] input - Input data. + !> @param[out] output - Output data. + !> @param[in] direction - Either `HIPFFT_FORWARD` or `HIPFFT_BACKWARD`. + !> + !> @brief Execute general transform - types of the descriptors must match the plan. + interface hipfftXtExecDescriptor +#ifdef USE_CUDA_NAMES + function hipfftXtExecDescriptor_(plan,input,output,direction) & + bind(c, name="cufftXtExecDescriptor") +#else + function hipfftXtExecDescriptor_(plan,input,output,direction) & + bind(c, name="hipfftXtExecDescriptor") +#endif + import + implicit none + integer(c_int) :: hipfftXtExecDescriptor_ + type(c_ptr),value :: plan + type(hipLibXtDesc) :: input + type(hipLibXtDesc) :: output + integer(c_int),value :: direction + end function + end interface + + !> @brief Set a multi-processing communicator on a plan. + !> + !> @details Attach a multi-processing communication handle to a + !> hipFFT plan. 'comm_handle' points to the handle, whose type + !> depends on the multi-processing library being used. With MPI + !> (Message Passing Interface) for example, 'comm_handle' points to + !> an MPI communicator. + !> + !> This function must only be called on a plan that has already + !> been allocated by `hipfftCreate`, but before the plan is + !> initialized with any of the 'hipfftMakePlan' functions. + !> + !> @param[in] plan - The FFT plan. + !> @param[in] comm_type - Type of communication handle. + !> @param[in] comm_handle - Pointer to the communication handle. + !> + !> @warning Experimental +#ifndef USE_CUDA_NAMES + interface hipfftMpAttachComm + function hipfftMpAttachComm_(plan,comm_type,comm_handle) bind(c, name="hipfftMpAttachComm") + import + implicit none + integer(c_int) :: hipfftMpAttachComm_ + type(c_ptr),value :: plan + integer(c_int),value :: comm_type + type(c_ptr),value :: comm_handle + end function + end interface +#endif + + !> @brief Describe the partial input and output for a distributed transform in the current + !> process. + !> + !> @details For the given plan, describe the 'brick' of input and + !> output that is distributed to the current process. A brick is + !> defined with the lower and upper coordinates of the that brick + !> in the global index space of the transform. + !> + !> Strides for both the input and output bricks are also provided, + !> to describe how the brick is laid out in memory. + !> + !> All coordinates and strides are in row-major order, with the + !> slowest-moving dimension specified first. + !> + !> This function must only be called on a plan that has already + !> been allocated by `hipfftCreate`, but before the plan is + !> initialized with any of the 'hipfftMakePlan' functions. + !> + !> @param[in] plan - The FFT plan. + !> @param[in] rank - Dimension of the transform + !> @param[in] input_lower - Array of length rank, specifying the lower index (inclusive) for the + !> brick in the FFT input. + !> @param[in] input_upper - Array of length rank, specifying the upper index (exclusive) for the + !> brick in the FFT input. + !> @param[in] output_lower - Array of length rank, specifying the lower index (inclusive) for the + !> brick in the FFT output. + !> @param[in] output_upper - Array of length rank, specifying the upper index (exclusive) for the + !> brick in the FFT output. + !> @param[in] input_stride - Array of length rank specifying the input brick's stride in memory + !> @param[in] output_stride - Array of length rank specifying the output brick's stride in memory + !> + !> @warning Experimental +#ifndef USE_CUDA_NAMES + interface hipfftXtSetDistribution + function hipfftXtSetDistribution_(plan,rank,input_lower,input_upper,output_lower,output_upper, & + input_stride,output_stride) & + bind(c, name="hipfftXtSetDistribution") + import + implicit none + integer(c_int) :: hipfftXtSetDistribution_ + type(c_ptr),value :: plan + integer(c_int),value :: rank + type(c_ptr),value :: input_lower + type(c_ptr),value :: input_upper + type(c_ptr),value :: output_lower + type(c_ptr),value :: output_upper + type(c_ptr),value :: input_stride + type(c_ptr),value :: output_stride + end function + end interface +#endif + + !> @brief Set the input and output formats for a distributed transform. + !> + !> @details Specifies the distribution of data for each side of a + !> distributed transform. Forward transforms use 'subformat_forward' + !> to describe input and 'subformat_inverse' to describe output. + !> Inverse transforms use 'subformat_inverse' to describe input and + !> 'subformat_forward' to describe output. + !> + !> 'subformat_forward' and 'subformat_inverse' must be set to + !> matching values. One may be HIPFFT_XT_FORMAT_INPLACE and the + !> other HIPFFT_XT_FORMAT_INPLACE_SHUFFLED; Alternatively, one may be + !> HIPFFT_XT_FORMAT_DISTRIBUTED_INPUT and the other + !> HIPFFT_XT_FORMAT_DISTRIBUTED_OUTPUT. + !> + !> This function must only be called on a plan that has already + !> been allocated by `hipfftCreate`, but before the plan is + !> initialized with any of the 'hipfftMakePlan' functions. + !> + !> @param[in] plan - The FFT plan. + !> @param[in] subformat_forward - Format of the input for a forward transform. + !> @param[in] subformat_inverse - Format of the input for an inverse transform. + !> + !> @warning Experimental +#ifndef USE_CUDA_NAMES + interface hipfftXtSetSubformatDefault + function hipfftXtSetSubformatDefault_(plan,subformat_forward,subformat_inverse) & + bind(c, name="hipfftXtSetSubformatDefault") + import + implicit none + integer(c_int) :: hipfftXtSetSubformatDefault_ + type(c_ptr),value :: plan + integer(c_int),value :: subformat_forward + integer(c_int),value :: subformat_inverse + end function + end interface +#endif + + + contains + + function hipfftXtMakePlanMany_typed(plan,rank,n,inembed,istride,idist,inputType,onembed, & + ostride,odist,outputType,batch,workSize,executionType) + implicit none + integer(kind(HIPFFT_SUCCESS)) :: hipfftXtMakePlanMany_typed + type(c_ptr), value :: plan + integer(c_int) :: rank + type(c_ptr) :: n + type(c_ptr) :: inembed + integer(c_int64_t) :: istride + integer(c_int64_t) :: idist + integer(c_int) :: inputType + type(c_ptr) :: onembed + integer(c_int64_t) :: ostride + integer(c_int64_t) :: odist + integer(c_int) :: outputType + integer(c_int64_t) :: batch + integer(c_size_t),target :: workSize + integer(c_int) :: executionType + ! + hipfftXtMakePlanMany_typed = hipfftXtMakePlanMany_(plan,rank,n,inembed,istride,idist, & + inputType,onembed,ostride,odist,outputType,batch,c_loc(workSize),executionType) + end function + + function hipfftXtGetSizeMany_typed(plan,rank,n,inembed,istride,idist,inputType,onembed, & + ostride,odist,outputType,batch,workSize,executionType) + implicit none + integer(kind(HIPFFT_SUCCESS)) :: hipfftXtGetSizeMany_typed + type(c_ptr), value :: plan + integer(c_int) :: rank + type(c_ptr) :: n + type(c_ptr) :: inembed + integer(c_int64_t) :: istride + integer(c_int64_t) :: idist + integer(c_int) :: inputType + type(c_ptr) :: onembed + integer(c_int64_t) :: ostride + integer(c_int64_t) :: odist + integer(c_int) :: outputType + integer(c_int64_t) :: batch + integer(c_size_t),target :: workSize + integer(c_int) :: executionType + ! + hipfftXtGetSizeMany_typed = hipfftXtGetSizeMany_(plan,rank,n,inembed,istride,idist, & + inputType,onembed,ostride,odist,outputType,batch,c_loc(workSize),executionType) + end function + + +#if defined(USE_ASSUMED_SHAPE) || defined(USE_ASSUMED_RANK) + +#ifdef USE_ASSUMED_RANK + function hipfftXtMakePlanMany_assumed_rank(plan,rank,n,inembed,istride,idist,inputType, & + onembed,ostride,odist,outputType,batch,workSize,executionType) + implicit none + integer(kind(HIPFFT_SUCCESS)) :: hipfftXtMakePlanMany_assumed_rank + type(c_ptr), value :: plan + integer(c_int) :: rank + integer(c_int64_t),target,dimension(..) :: n + integer(c_int64_t),target,dimension(..) :: inembed + integer(c_int64_t) :: istride + integer(c_int64_t) :: idist + integer(c_int) :: inputType + integer(c_int64_t),target,dimension(..) :: onembed + integer(c_int64_t) :: ostride + integer(c_int64_t) :: odist + integer(c_int) :: outputType + integer(c_int64_t) :: batch + integer(c_size_t),target :: workSize + integer(c_int) :: executionType + ! + hipfftXtMakePlanMany_assumed_rank = hipfftXtMakePlanMany_(plan,rank,c_loc(n),c_loc(inembed), & + istride,idist,inputType,c_loc(onembed),ostride,odist,outputType,batch,c_loc(workSize), & + executionType) + end function + +#else + function hipfftXtMakePlanMany_rank_0(plan,rank,n,inembed,istride,idist,inputType,onembed, & + ostride,odist,outputType,batch,workSize,executionType) + implicit none + integer(kind(HIPFFT_SUCCESS)) :: hipfftXtMakePlanMany_rank_0 + type(c_ptr), value :: plan + integer(c_int) :: rank + integer(c_int64_t),target :: n + integer(c_int64_t),target :: inembed + integer(c_int64_t) :: istride + integer(c_int64_t) :: idist + integer(c_int) :: inputType + integer(c_int64_t),target :: onembed + integer(c_int64_t) :: ostride + integer(c_int64_t) :: odist + integer(c_int) :: outputType + integer(c_int64_t) :: batch + integer(c_size_t),target :: workSize + integer(c_int) :: executionType + ! + hipfftXtMakePlanMany_rank_0 = hipfftXtMakePlanMany_(plan,rank,c_loc(n),c_loc(inembed), & + istride,idist,inputType,c_loc(onembed),ostride,odist,outputType,batch,c_loc(workSize), & + executionType) + end function + + function hipfftXtMakePlanMany_rank_1(plan,rank,n,inembed,istride,idist,inputType,onembed, & + ostride,odist,outputType,batch,workSize,executionType) + implicit none + integer(kind(HIPFFT_SUCCESS)) :: hipfftXtMakePlanMany_rank_1 + type(c_ptr), value :: plan + integer(c_int) :: rank + integer(c_int64_t),target,dimension(:) :: n + integer(c_int64_t),target,dimension(:) :: inembed + integer(c_int64_t) :: istride + integer(c_int64_t) :: idist + integer(c_int) :: inputType + integer(c_int64_t),target,dimension(:) :: onembed + integer(c_int64_t) :: ostride + integer(c_int64_t) :: odist + integer(c_int) :: outputType + integer(c_int64_t) :: batch + integer(c_size_t),target :: workSize + integer(c_int) :: executionType + ! + hipfftXtMakePlanMany_rank_1 = hipfftXtMakePlanMany_(plan,rank,c_loc(n),c_loc(inembed), & + istride,idist,inputType,c_loc(onembed),ostride,odist,outputType,batch,c_loc(workSize), & + executionType) + end function + +#endif +#ifdef USE_ASSUMED_RANK + function hipfftXtGetSizeMany_assumed_rank(plan,rank,n,inembed,istride,idist,inputType,onembed, & + ostride,odist,outputType,batch,workSize,executionType) + implicit none + integer(kind(HIPFFT_SUCCESS)) :: hipfftXtGetSizeMany_assumed_rank + type(c_ptr), value :: plan + integer(c_int) :: rank + integer(c_int64_t),target,dimension(..) :: n + integer(c_int64_t),target,dimension(..) :: inembed + integer(c_int64_t) :: istride + integer(c_int64_t) :: idist + integer(c_int) :: inputType + integer(c_int64_t),target,dimension(..) :: onembed + integer(c_int64_t) :: ostride + integer(c_int64_t) :: odist + integer(c_int) :: outputType + integer(c_int64_t) :: batch + integer(c_size_t),target :: workSize + integer(c_int) :: executionType + ! + hipfftXtGetSizeMany_assumed_rank = hipfftXtGetSizeMany_(plan,rank,c_loc(n),c_loc(inembed), & + istride,idist,inputType,c_loc(onembed),ostride,odist,outputType,batch,c_loc(workSize), & + executionType) + end function + +#else + function hipfftXtGetSizeMany_rank_0(plan,rank,n,inembed,istride,idist,inputType,onembed, & + ostride,odist,outputType,batch,workSize,executionType) + implicit none + integer(kind(HIPFFT_SUCCESS)) :: hipfftXtGetSizeMany_rank_0 + type(c_ptr), value :: plan + integer(c_int) :: rank + integer(c_int64_t),target :: n + integer(c_int64_t),target :: inembed + integer(c_int64_t) :: istride + integer(c_int64_t) :: idist + integer(c_int) :: inputType + integer(c_int64_t),target :: onembed + integer(c_int64_t) :: ostride + integer(c_int64_t) :: odist + integer(c_int) :: outputType + integer(c_int64_t) :: batch + integer(c_size_t),target :: workSize + integer(c_int) :: executionType + ! + hipfftXtGetSizeMany_rank_0 = hipfftXtGetSizeMany_(plan,rank,c_loc(n),c_loc(inembed),istride, & + idist,inputType,c_loc(onembed),ostride,odist,outputType,batch,c_loc(workSize),executionType) + end function + + function hipfftXtGetSizeMany_rank_1(plan,rank,n,inembed,istride,idist,inputType,onembed, & + ostride,odist,outputType,batch,workSize,executionType) + implicit none + integer(kind(HIPFFT_SUCCESS)) :: hipfftXtGetSizeMany_rank_1 + type(c_ptr), value :: plan + integer(c_int) :: rank + integer(c_int64_t),target,dimension(:) :: n + integer(c_int64_t),target,dimension(:) :: inembed + integer(c_int64_t) :: istride + integer(c_int64_t) :: idist + integer(c_int) :: inputType + integer(c_int64_t),target,dimension(:) :: onembed + integer(c_int64_t) :: ostride + integer(c_int64_t) :: odist + integer(c_int) :: outputType + integer(c_int64_t) :: batch + integer(c_size_t),target :: workSize + integer(c_int) :: executionType + ! + hipfftXtGetSizeMany_rank_1 = hipfftXtGetSizeMany_(plan,rank,c_loc(n),c_loc(inembed),istride, & + idist,inputType,c_loc(onembed),ostride,odist,outputType,batch,c_loc(workSize),executionType) + end function + +#endif +#ifdef USE_ASSUMED_RANK + function hipfftXtSetGPUs_assumed_rank(plan,count,gpus) + implicit none + integer(kind(HIPFFT_SUCCESS)) :: hipfftXtSetGPUs_assumed_rank + type(c_ptr), value :: plan + integer(c_int) :: count + integer(c_int),target,dimension(..) :: gpus + ! + hipfftXtSetGPUs_assumed_rank = hipfftXtSetGPUs_(plan,count,c_loc(gpus)) + end function + +#else + function hipfftXtSetGPUs_rank_0(plan,count,gpus) + implicit none + integer(kind(HIPFFT_SUCCESS)) :: hipfftXtSetGPUs_rank_0 + type(c_ptr), value :: plan + integer(c_int) :: count + integer(c_int),target :: gpus + ! + hipfftXtSetGPUs_rank_0 = hipfftXtSetGPUs_(plan,count,c_loc(gpus)) + end function + + function hipfftXtSetGPUs_rank_1(plan,count,gpus) + implicit none + integer(kind(HIPFFT_SUCCESS)) :: hipfftXtSetGPUs_rank_1 + type(c_ptr), value :: plan + integer(c_int) :: count + integer(c_int),target,dimension(:) :: gpus + ! + hipfftXtSetGPUs_rank_1 = hipfftXtSetGPUs_(plan,count,c_loc(gpus)) + end function + +#endif +#endif +end module hipfort_hipfftxt diff --git a/lib/hipfort/hipfort_hipfftxt_enums.F90 b/lib/hipfort/hipfort_hipfftxt_enums.F90 new file mode 100644 index 000000000..9ca1a9880 --- /dev/null +++ b/lib/hipfort/hipfort_hipfftxt_enums.F90 @@ -0,0 +1,79 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +module hipfort_hipfftxt_enums + use, intrinsic :: iso_c_binding + implicit none + + ! hiplibFormat_t + enum, bind(c) + enumerator :: HIPLIB_FORMAT_HIPFFT = 0 + enumerator :: HIPLIB_FORMAT_UNDEFINED = 1 + end enum + + ! hipfftXtCopyType_t + enum, bind(c) + enumerator :: HIPFFT_COPY_HOST_TO_DEVICE = 0 + enumerator :: HIPFFT_COPY_DEVICE_TO_HOST = 1 + enumerator :: HIPFFT_COPY_DEVICE_TO_DEVICE = 2 + enumerator :: HIPFFT_COPY_UNDEFINED = 3 + end enum + + ! hipfftXtCallbackType_t + enum, bind(c) + enumerator :: HIPFFT_CB_LD_COMPLEX = 0 + enumerator :: HIPFFT_CB_LD_COMPLEX_DOUBLE = 1 + enumerator :: HIPFFT_CB_LD_REAL = 2 + enumerator :: HIPFFT_CB_LD_REAL_DOUBLE = 3 + enumerator :: HIPFFT_CB_ST_COMPLEX = 4 + enumerator :: HIPFFT_CB_ST_COMPLEX_DOUBLE = 5 + enumerator :: HIPFFT_CB_ST_REAL = 6 + enumerator :: HIPFFT_CB_ST_REAL_DOUBLE = 7 + enumerator :: HIPFFT_CB_UNDEFINED = 8 + end enum + + ! hipfftXtSubFormat_t + enum, bind(c) + enumerator :: HIPFFT_XT_FORMAT_INPUT = 0 + enumerator :: HIPFFT_XT_FORMAT_OUTPUT = 1 + enumerator :: HIPFFT_XT_FORMAT_INPLACE = 2 + enumerator :: HIPFFT_XT_FORMAT_INPLACE_SHUFFLED = 3 + enumerator :: HIPFFT_XT_FORMAT_1D_INPUT_SHUFFLED = 4 +#ifdef USE_CUDA_NAMES + enumerator :: HIPFFT_FORMAT_UNDEFINED = 7 +#else + enumerator :: HIPFFT_FORMAT_UNDEFINED = 5 +#endif + end enum + + ! hipfftMpCommType_t + enum, bind(c) + enumerator :: HIPFFT_COMM_MPI = 0 + enumerator :: HIPFFT_COMM_NONE = 1 + end enum + + integer(c_int), parameter :: MAX_HIP_DESCRIPTOR_GPUS = 64 + +end module hipfort_hipfftxt_enums diff --git a/lib/hipfort/hipfort_hipfft_types.F90 b/lib/hipfort/hipfort_hipfftxt_types.F90 similarity index 96% rename from lib/hipfort/hipfort_hipfft_types.F90 rename to lib/hipfort/hipfort_hipfftxt_types.F90 index df73c5c46..f32946d5f 100644 --- a/lib/hipfort/hipfort_hipfft_types.F90 +++ b/lib/hipfort/hipfort_hipfftxt_types.F90 @@ -23,7 +23,7 @@ ! !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! -module hipfort_hipfft_types +module hipfort_hipfftxt_types use, intrinsic :: iso_c_binding implicit none @@ -44,4 +44,4 @@ module hipfort_hipfft_types type(c_ptr) :: libDescriptor end type hipLibXtDesc -end module hipfort_hipfft_types +end module hipfort_hipfftxt_types diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 7f00912de..6ddf090cb 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -185,8 +185,8 @@ if(BUILD_TESTING) f2008/hipsparse/hipsparse_zgthr f2008/hipsparse/hipsparse_zsctr # hipFFT Xt API: no cufftXt twin mapped yet, so not bound under USE_CUDA_NAMES. - f2003/hipfft/hipfft_xt_descriptor_z - f2003/hipfft/hipfft_xt_makeplanmany_c + f2003/hipfftxt/hipfftxt_descriptor_z + f2003/hipfftxt/hipfftxt_makeplanmany_c # hipBLAS routines cuBLAS lacks: rotStridedBatched with c and s per the pointer # mode, a non-batched geqrf, getrfStridedBatched; and hipSPARSE size queries # without a cuSPARSE 12 twin (csrsm2, prune). @@ -545,12 +545,16 @@ if(BUILD_TESTING) hipfort_add_test(hipfft hipfft_setstream_z f2008) hipfort_add_test(hipfft hipfft_version f2003) hipfort_add_test(hipfft hipfft_version f2008) + endif() + + if(TARGET hipfort::hipfftxt) + hipfort_add_symbol_test(hipfftxt) # The Xt API: a hipDataType plan on one GPU, and multi-GPU descriptors. - hipfort_add_test(hipfft hipfft_xt_makeplanmany_c f2003) - hipfort_add_test(hipfft hipfft_xt_descriptor_z f2003) + hipfort_add_test(hipfftxt hipfftxt_makeplanmany_c f2003) + hipfort_add_test(hipfftxt hipfftxt_descriptor_z f2003) # Needs two visible GPUs, else exits 77: reported as skipped, not passed. - if(TEST hipfort_test_f2003_hipfft_hipfft_xt_descriptor_z) - set_tests_properties(hipfort_test_f2003_hipfft_hipfft_xt_descriptor_z + if(TEST hipfort_test_f2003_hipfftxt_hipfftxt_descriptor_z) + set_tests_properties(hipfort_test_f2003_hipfftxt_hipfftxt_descriptor_z PROPERTIES SKIP_RETURN_CODE 77) endif() endif() diff --git a/test/coverage/variant_coverage.py b/test/coverage/variant_coverage.py index 5cd5cc494..eb91dacff 100644 --- a/test/coverage/variant_coverage.py +++ b/test/coverage/variant_coverage.py @@ -89,6 +89,7 @@ ("hipsolver", ["hipfort_hipsolver"]), ("hipsparse", ["hipfort_hipsparse"]), ("hipfft", ["hipfort_hipfft"]), + ("hipfftxt", ["hipfort_hipfftxt"]), ("hipfftw", ["hipfort_hipfftw"]), ("hiprand", ["hipfort_hiprand"]), ("roctx", ["hipfort_roctx"]), diff --git a/test/f2003/hipfft/test_hipfft.F03 b/test/f2003/hipfft/test_hipfft.F03 index 9ee0e61e5..7add57aea 100644 --- a/test/f2003/hipfft/test_hipfft.F03 +++ b/test/f2003/hipfft/test_hipfft.F03 @@ -27,15 +27,12 @@ program test_hipfft use iso_c_binding use hipfort_hipfft use hipfort_hipfft_enums - use hipfort_hipfft_types implicit none logical :: never type(c_ptr) :: v0 integer(c_int) :: v1 real(c_double) :: v2 integer(c_int64_t) :: v3 - integer(c_size_t) :: v4 - type(hipLibXtDesc) :: v5 ! Always .false., but not foldable by the compiler, so every guarded ! call is compiled and linked yet never run (no null-pointer segfaults). @@ -66,7 +63,6 @@ program test_hipfft if (never) v1 = hipfftGetSize_(v0, v0) if (never) v1 = hipfftSetAutoAllocation_(v0, v1) if (never) v1 = hipfftSetWorkArea_(v0, v0) - if (never) v1 = hipfftXtSetWorkArea_(v0, v0) if (never) v1 = hipfftExecC2C_(v0, v0, v0, v1) if (never) v1 = hipfftExecR2C_(v0, v0, v0) if (never) v1 = hipfftExecC2R_(v0, v0, v0) @@ -77,24 +73,6 @@ program test_hipfft if (never) v1 = hipfftDestroy_(v0) if (never) v1 = hipfftGetVersion_(v0) if (never) v1 = hipfftGetProperty_(v1, v0) - if (never) v1 = hipfftXtSetCallback_(v0, v0, v1, v0) - if (never) v1 = hipfftXtClearCallback_(v0, v1) - if (never) v1 = hipfftXtSetCallbackSharedSize_(v0, v1, v4) - if (never) v1 = hipfftXtSetJITCallback_(v0, v0, v0, v4, v1, v0) - if (never) v1 = hipfftXtMakePlanMany_(v0, v1, v0, v0, v3, v3, v1, v0, v3, v3, v1, v3, v0, v1) - if (never) v1 = hipfftXtGetSizeMany_(v0, v1, v0, v0, v3, v3, v1, v0, v3, v3, v1, v3, v0, v1) - if (never) v1 = hipfftXtExec_(v0, v0, v0, v1) - if (never) v1 = hipfftXtSetGPUs_(v0, v1, v0) - if (never) v1 = hipfftXtMalloc_(v0, v0, v1) - if (never) v1 = hipfftXtMemcpy_(v0, v0, v0, v1) - if (never) v1 = hipfftXtFree_(v5) - if (never) v1 = hipfftXtExecDescriptorC2C_(v0, v5, v5, v1) - if (never) v1 = hipfftXtExecDescriptorR2C_(v0, v5, v5) - if (never) v1 = hipfftXtExecDescriptorC2R_(v0, v5, v5) - if (never) v1 = hipfftXtExecDescriptorZ2Z_(v0, v5, v5, v1) - if (never) v1 = hipfftXtExecDescriptorD2Z_(v0, v5, v5) - if (never) v1 = hipfftXtExecDescriptorZ2D_(v0, v5, v5) - if (never) v1 = hipfftXtExecDescriptor_(v0, v5, v5, v1) write(*,"(a)") "-- test_hipfft (bind(C) symbols): PASSED (compile+link only)!" end program test_hipfft diff --git a/test/f2003/hipfft/hipfft_xt_descriptor_z.f03 b/test/f2003/hipfftxt/hipfftxt_descriptor_z.f03 similarity index 97% rename from test/f2003/hipfft/hipfft_xt_descriptor_z.f03 rename to test/f2003/hipfftxt/hipfftxt_descriptor_z.f03 index 86fee4895..51066b8ce 100644 --- a/test/f2003/hipfft/hipfft_xt_descriptor_z.f03 +++ b/test/f2003/hipfftxt/hipfftxt_descriptor_z.f03 @@ -23,11 +23,11 @@ ! !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! -program hipfft_xt_descriptor_z +program hipfftxt_descriptor_z use iso_c_binding use hipfort use hipfort_check - use hipfort_hipfft + use hipfort_hipfftxt implicit none @@ -99,4 +99,4 @@ program hipfft_xt_descriptor_z write(*,*) "PASSED!" -end program hipfft_xt_descriptor_z +end program hipfftxt_descriptor_z diff --git a/test/f2003/hipfft/hipfft_xt_makeplanmany_c.f03 b/test/f2003/hipfftxt/hipfftxt_makeplanmany_c.f03 similarity index 97% rename from test/f2003/hipfft/hipfft_xt_makeplanmany_c.f03 rename to test/f2003/hipfftxt/hipfftxt_makeplanmany_c.f03 index 77183f746..9f1bfb075 100644 --- a/test/f2003/hipfft/hipfft_xt_makeplanmany_c.f03 +++ b/test/f2003/hipfftxt/hipfftxt_makeplanmany_c.f03 @@ -23,11 +23,11 @@ ! !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! -program hipfft_xt_makeplanmany_c +program hipfftxt_makeplanmany_c use iso_c_binding use hipfort use hipfort_check - use hipfort_hipfft + use hipfort_hipfftxt implicit none @@ -86,4 +86,4 @@ program hipfft_xt_makeplanmany_c write(*,*) "PASSED!" -end program hipfft_xt_makeplanmany_c +end program hipfftxt_makeplanmany_c diff --git a/test/f2003/hipfftxt/test_hipfftxt.F03 b/test/f2003/hipfftxt/test_hipfftxt.F03 new file mode 100644 index 000000000..9382e4621 --- /dev/null +++ b/test/f2003/hipfftxt/test_hipfftxt.F03 @@ -0,0 +1,65 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! Copyright (C) 2020-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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program test_hipfftxt + use iso_c_binding + use hipfort_hipfftxt + use hipfort_hipfft + use hipfort_hipfftxt_enums + use hipfort_hipfftxt_types + implicit none + logical :: never + type(c_ptr) :: v0 + integer(c_int) :: v1 + integer(c_size_t) :: v2 + integer(c_int64_t) :: v3 + type(hipLibXtDesc) :: v4 + + ! Always .false., but not foldable by the compiler, so every guarded + ! call is compiled and linked yet never run (no null-pointer segfaults). + never = command_argument_count() > 1000000 + + if (never) v1 = hipfftXtSetWorkArea_(v0, v0) + if (never) v1 = hipfftXtSetCallback_(v0, v0, v1, v0) + if (never) v1 = hipfftXtClearCallback_(v0, v1) + if (never) v1 = hipfftXtSetCallbackSharedSize_(v0, v1, v2) + if (never) v1 = hipfftXtSetJITCallback_(v0, v0, v0, v2, v1, v0) + if (never) v1 = hipfftXtMakePlanMany_(v0, v1, v0, v0, v3, v3, v1, v0, v3, v3, v1, v3, v0, v1) + if (never) v1 = hipfftXtGetSizeMany_(v0, v1, v0, v0, v3, v3, v1, v0, v3, v3, v1, v3, v0, v1) + if (never) v1 = hipfftXtExec_(v0, v0, v0, v1) + if (never) v1 = hipfftXtSetGPUs_(v0, v1, v0) + if (never) v1 = hipfftXtMalloc_(v0, v0, v1) + if (never) v1 = hipfftXtMemcpy_(v0, v0, v0, v1) + if (never) v1 = hipfftXtFree_(v4) + if (never) v1 = hipfftXtExecDescriptorC2C_(v0, v4, v4, v1) + if (never) v1 = hipfftXtExecDescriptorR2C_(v0, v4, v4) + if (never) v1 = hipfftXtExecDescriptorC2R_(v0, v4, v4) + if (never) v1 = hipfftXtExecDescriptorZ2Z_(v0, v4, v4, v1) + if (never) v1 = hipfftXtExecDescriptorD2Z_(v0, v4, v4) + if (never) v1 = hipfftXtExecDescriptorZ2D_(v0, v4, v4) + if (never) v1 = hipfftXtExecDescriptor_(v0, v4, v4, v1) + + write(*,"(a)") "-- test_hipfftxt (bind(C) symbols): PASSED (compile+link only)!" +end program test_hipfftxt From a367d4a0cf4f93beed5614011fdbf6e6d41e2da2 Mon Sep 17 00:00:00 2001 From: Alexis Montoison Date: Sat, 7 Nov 2026 03:21:11 -0600 Subject: [PATCH 2/2] Describe the CUDA backend of the hipFFT Xt routines accurately They bind their cufftXt* counterparts; only the multi-process routines have none. --- CHANGELOG.md | 3 ++- test/CMakeLists.txt | 2 +- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 773c1c251..2b1af28a8 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -49,7 +49,8 @@ of `hipfftMp.h` (`hipfftMpAttachComm`, `hipfftXtSetDistribution`, `hipfftXtSetSubformatDefault`) are bound too, but exist only in a hipFFT built with MPI support, which the ROCm packages are not. On the CUDA backend the Xt - and multi-process routines added here are not available. The hipFFT examples document + routines bind their `cufftXt*` counterparts; the multi-process routines have + none. The hipFFT examples document the multi-GPU and MPI workflows. * Array overloads (`rank_N`, `full_rank`, assumed-rank) for the ILP64 (`_64`) routines of rocBLAS, hipBLAS and rocSOLVER (`rocblas_sscal_64`, ...), the same diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 6ddf090cb..5c4ada2c1 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -184,7 +184,7 @@ if(BUILD_TESTING) f2008/hipsparse/hipsparse_zcsrsv2 f2008/hipsparse/hipsparse_zgthr f2008/hipsparse/hipsparse_zsctr - # hipFFT Xt API: no cufftXt twin mapped yet, so not bound under USE_CUDA_NAMES. + # hipFFT Xt tests: not validated on the CUDA backend. f2003/hipfftxt/hipfftxt_descriptor_z f2003/hipfftxt/hipfftxt_makeplanmany_c # hipBLAS routines cuBLAS lacks: rotStridedBatched with c and s per the pointer