diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS index 3df545d74867..6696234dbb46 100644 --- a/.github/CODEOWNERS +++ b/.github/CODEOWNERS @@ -244,6 +244,7 @@ /projects/rocfft/**/.readthedocs.yaml @ROCm/fft-reviewers @ROCm/fft-docs-reviewers /projects/rocfft/docs/ @ROCm/fft-reviewers @ROCm/fft-docs-reviewers /projects/rocfft/library/include/ @ROCm/fft-reviewers @ROCm/fft-docs-reviewers +/bindings/fortran/rocfft/ @ROCm/rocm-fortran-maintainers @ROCm/fft-reviewers /projects/rocprim/**/*.md @ROCm/prims-rands-reviewers @ROCm/prims-docs-reviewers /projects/rocprim/**/*.rst @ROCm/prims-rands-reviewers @ROCm/prims-docs-reviewers diff --git a/.github/repos-config.json b/.github/repos-config.json index 218164737f74..8f85df84a817 100644 --- a/.github/repos-config.json +++ b/.github/repos-config.json @@ -305,6 +305,15 @@ "auto_subtree_pull": false, "auto_subtree_push": false, "monorepo_source_of_truth": true + }, + { + "name": "rocfft", + "url": "ROCm/rocm-libraries", + "branch": "develop", + "category": "bindings/fortran", + "auto_subtree_pull": false, + "auto_subtree_push": false, + "monorepo_source_of_truth": true } ] } diff --git a/bindings/fortran/rocfft/.gitattributes b/bindings/fortran/rocfft/.gitattributes new file mode 100644 index 000000000000..c59665de42a8 --- /dev/null +++ b/bindings/fortran/rocfft/.gitattributes @@ -0,0 +1,2 @@ +# Generated by the rocm-fortran generator +/rocfft.F90 linguist-generated=true diff --git a/bindings/fortran/rocfft/CHANGELOG.md b/bindings/fortran/rocfft/CHANGELOG.md new file mode 100644 index 000000000000..de7247a6e7e1 --- /dev/null +++ b/bindings/fortran/rocfft/CHANGELOG.md @@ -0,0 +1,13 @@ +# Changelog for the rocFFT Fortran binding + +## Unreleased + +### Added + +* Generated Fortran bindings, as a single `rocfft` module: `use rocfft`, link `roc::rocfft_fortran`. + Controlled by `BUILD_FORTRAN_BINDINGS` (on by default on Linux, off on Windows), + `BUILD_FORTRAN_CLIENTS` and `FORTRAN_ARRAY_INTERFACES` (`assumed-shape` by default; `none` or `assumed-rank`). + Found with `find_package(rocfft-fortran)`; the archive and the `.mod` files install per compiler, + under `/fortran/` and `/fortran/`. + The archive has one section per procedure: link with `-Wl,--gc-sections` to keep only the wrappers + you call. diff --git a/bindings/fortran/rocfft/CMakeLists.txt b/bindings/fortran/rocfft/CMakeLists.txt new file mode 100644 index 000000000000..0bc6e9e65752 --- /dev/null +++ b/bindings/fortran/rocfft/CMakeLists.txt @@ -0,0 +1,208 @@ +# ############################################################################# +# 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. +# ############################################################################# + +# Fortran bindings for rocFFT. rocfft.F90 is GENERATED by the rocm-fortran +# generator: regenerate it, do not edit it. + +cmake_minimum_required(VERSION 3.19) + +project(rocfft-fortran LANGUAGES NONE) + +# Not ON: ROCm ships no Fortran compiler on Windows. +option(BUILD_FORTRAN_BINDINGS "Build the rocFFT Fortran bindings" "${UNIX}") +option(BUILD_FORTRAN_CLIENTS "Build the rocFFT Fortran binding tests" ON) + +set(FORTRAN_ARRAY_INTERFACES "assumed-shape" CACHE STRING + "Array-argument overloads to expose: none, assumed-shape, or assumed-rank") +set_property(CACHE FORTRAN_ARRAY_INTERFACES PROPERTY STRINGS + none assumed-shape assumed-rank) + +if(NOT BUILD_FORTRAN_BINDINGS) + return() +endif() + +include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/ROCmFortranCompiler.cmake") +if(NOT CMAKE_Fortran_COMPILER) + message(STATUS "rocFFT: no Fortran compiler found, skipping the Fortran bindings") + return() +endif() + +enable_language(Fortran) + +include(GNUInstallDirs) + +if(NOT FORTRAN_ARRAY_INTERFACES MATCHES "^(none|assumed-shape|assumed-rank)$") + message(FATAL_ERROR + "FORTRAN_ARRAY_INTERFACES must be none, assumed-shape or assumed-rank, " + "not '${FORTRAN_ARRAY_INTERFACES}'.") +endif() + +add_library(rocfft_fortran STATIC rocfft.F90) +add_library(roc::rocfft_fortran ALIAS rocfft_fortran) + +# Each module compiles to one object, and a linker extracts whole objects from +# an archive, so a program that calls one wrapper links all of them. One section +# per procedure lets a program linked with -Wl,--gc-sections keep only the +# wrappers it calls, and still link when another wrapper names a C function the +# installed library does not export. Only these compilers split the sections +# (the Intel ones accept the flags but do not). +target_compile_options(rocfft_fortran PRIVATE + "$<$:-ffunction-sections;-fdata-sections>") + +# A .mod is compiler-specific, so artifacts install per compiler. This mapping +# MUST match the one in rocfft-fortran-config.cmake.in. +if(CMAKE_Fortran_COMPILER_ID STREQUAL "GNU") + set(_fc_name "gfortran") +elseif(CMAKE_Fortran_COMPILER_ID MATCHES "Flang") # Flang, LLVMFlang + set(_fc_name "amdflang") +elseif(CMAKE_Fortran_COMPILER_ID STREQUAL "Cray") + set(_fc_name "ftn") +elseif(CMAKE_Fortran_COMPILER_ID STREQUAL "IntelLLVM") + set(_fc_name "ifx") +elseif(CMAKE_Fortran_COMPILER_ID STREQUAL "Intel") + set(_fc_name "ifort") +elseif(CMAKE_Fortran_COMPILER_ID STREQUAL "NVHPC") + set(_fc_name "nvfortran") +else() + get_filename_component(_fc_name "${CMAKE_Fortran_COMPILER}" NAME_WE) + string(TOLOWER "${_fc_name}" _fc_name) +endif() + +set(ROCFFT_FORTRAN_COMPILER_DIR "${_fc_name}" CACHE STRING + "Per-compiler subdirectory the .mod and the archive install under") + +set(_moddir "${CMAKE_CURRENT_BINARY_DIR}/include") +set_target_properties(rocfft_fortran PROPERTIES + Fortran_MODULE_DIRECTORY "${_moddir}" + Fortran_PREPROCESS ON + Fortran_FORMAT FREE + POSITION_INDEPENDENT_CODE ON) + +if(FORTRAN_ARRAY_INTERFACES STREQUAL "assumed-shape") + target_compile_definitions(rocfft_fortran PRIVATE USE_ASSUMED_SHAPE=1) +elseif(FORTRAN_ARRAY_INTERFACES STREQUAL "assumed-rank") + target_compile_definitions(rocfft_fortran PRIVATE USE_ASSUMED_RANK=1) +endif() + +target_include_directories(rocfft_fortran PUBLIC + "$" + "$") + +# Link interface only. C is enabled because hip-config.cmake reaches FindThreads. +enable_language(C) +find_package(rocfft QUIET) + +# Always TRUE or FALSE: the config template substitutes it into a bare if(). +if(TARGET roc::rocfft) + target_link_libraries(rocfft_fortran INTERFACE roc::rocfft) + set(ROCFFT_FORTRAN_NEEDS_C_PACKAGE TRUE) +else() + set(ROCFFT_FORTRAN_NEEDS_C_PACKAGE FALSE) + message(STATUS + "rocFFT: librocfft was not found, so the Fortran bindings carry no link " + "interface; a consumer will have to link it (-lrocfft) itself.") +endif() + +# The uninstrumented Fortran executables cannot link an ASan librocfft. +if(BUILD_FORTRAN_CLIENTS AND BUILD_ADDRESS_SANITIZER) + message(STATUS + "rocFFT: address sanitizer is enabled, skipping the Fortran binding tests.") + set(BUILD_FORTRAN_CLIENTS OFF) +endif() + +if(BUILD_FORTRAN_CLIENTS) + enable_testing() + add_subdirectory(test) +endif() + +include(CMakePackageConfigHelpers) + +set(_fortran_pkgdir "${CMAKE_INSTALL_LIBDIR}/cmake/rocfft-fortran") + +install(TARGETS rocfft_fortran + EXPORT rocfft-fortran-targets + ARCHIVE DESTINATION "${CMAKE_INSTALL_LIBDIR}/fortran/${ROCFFT_FORTRAN_COMPILER_DIR}" + COMPONENT devel) +install(DIRECTORY "${_moddir}/" + DESTINATION "${CMAKE_INSTALL_INCLUDEDIR}/fortran/${ROCFFT_FORTRAN_COMPILER_DIR}" + COMPONENT devel + FILES_MATCHING PATTERN "*.mod") + +# Under share/: the config reads every entry of include/fortran/ as a compiler. +install(FILES "${CMAKE_CURRENT_SOURCE_DIR}/rocfft.F90" + DESTINATION "${CMAKE_INSTALL_DATADIR}/rocfft/fortran" + COMPONENT devel) +install(EXPORT rocfft-fortran-targets + FILE rocfft-fortran-targets.cmake + NAMESPACE roc:: + DESTINATION "${_fortran_pkgdir}" + COMPONENT devel) + +set(ROCFFT_FORTRAN_LIBDIR "${CMAKE_INSTALL_LIBDIR}/fortran") +set(ROCFFT_FORTRAN_INCLUDEDIR "${CMAKE_INSTALL_INCLUDEDIR}/fortran") + +configure_package_config_file( + "${CMAKE_CURRENT_SOURCE_DIR}/rocfft-fortran-config.cmake.in" + "${CMAKE_CURRENT_BINARY_DIR}/rocfft-fortran-config.cmake" + INSTALL_DESTINATION "${_fortran_pkgdir}" + PATH_VARS ROCFFT_FORTRAN_LIBDIR ROCFFT_FORTRAN_INCLUDEDIR) +install(FILES "${CMAKE_CURRENT_BINARY_DIR}/rocfft-fortran-config.cmake" + DESTINATION "${_fortran_pkgdir}" + COMPONENT devel) + +# Read from this tree, not rocfft_VERSION, which is the INSTALLED rocFFT's. Judged on the whole match: the minor component "0" is falsy. +set(_fortran_version "") +if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/../../../projects/rocfft/CMakeLists.txt") + file(READ "${CMAKE_CURRENT_SOURCE_DIR}/../../../projects/rocfft/CMakeLists.txt" _rocfft_root) + set(_rocfft_parts "") + foreach(_part MAJOR MINOR PATCH) + string(REGEX MATCH "set[ \t]*\\([ \t]*ROCFFT_VERSION_${_part}[ \t]+\"?([0-9]+)" + _m "${_rocfft_root}") + if(_m STREQUAL "") + set(_rocfft_parts "") + break() + endif() + list(APPEND _rocfft_parts "${CMAKE_MATCH_1}") + endforeach() + if(_rocfft_parts) + list(JOIN _rocfft_parts "." _fortran_version) + endif() + unset(_rocfft_parts) + unset(_rocfft_root) +endif() + +if(_fortran_version) + write_basic_package_version_file( + "${CMAKE_CURRENT_BINARY_DIR}/rocfft-fortran-config-version.cmake" + VERSION "${_fortran_version}" + COMPATIBILITY SameMajorVersion) + install(FILES "${CMAKE_CURRENT_BINARY_DIR}/rocfft-fortran-config-version.cmake" + DESTINATION "${_fortran_pkgdir}" + COMPONENT devel) +else() + message(STATUS + "rocFFT: could not determine the version for the Fortran package, so no " + "rocfft-fortran-config-version.cmake is installed; " + "find_package(rocfft-fortran ) will not match.") +endif() diff --git a/bindings/fortran/rocfft/cmake/ROCmFortranCompiler.cmake b/bindings/fortran/rocfft/cmake/ROCmFortranCompiler.cmake new file mode 100644 index 000000000000..616f4fa03c68 --- /dev/null +++ b/bindings/fortran/rocfft/cmake/ROCmFortranCompiler.cmake @@ -0,0 +1,59 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + +# Runs check_language(Fortran), preferring ROCm's amdflang over the first +# Fortran compiler on PATH (often gfortran). An explicit CMAKE_Fortran_COMPILER +# or FC wins. CMAKE_Fortran_COMPILER is left false when none is found. + +include(CheckLanguage) + +# rocm-cmake's ROCMChecks watches CMAKE_Fortran_COMPILER and fails under +# ROCM_ERROR_TOOLCHAIN_VAR=ON; silence it across the writes. +set(_rocm_fc_warn_toolchain_var "${ROCM_WARN_TOOLCHAIN_VAR}") +set(_rocm_fc_error_toolchain_var "${ROCM_ERROR_TOOLCHAIN_VAR}") +set(ROCM_WARN_TOOLCHAIN_VAR OFF) +set(ROCM_ERROR_TOOLCHAIN_VAR OFF) + +if(NOT CMAKE_Fortran_COMPILER AND NOT DEFINED ENV{FC}) + set(_rocm_fc_hints) + foreach(_rocm_fc_cc "${CMAKE_CXX_COMPILER}" "${CMAKE_C_COMPILER}" + "${CMAKE_HIP_COMPILER}") + if(_rocm_fc_cc) + get_filename_component(_rocm_fc_dir "${_rocm_fc_cc}" DIRECTORY) + list(APPEND _rocm_fc_hints "${_rocm_fc_dir}") + endif() + endforeach() + foreach(_rocm_fc_root "${ROCM_PATH}" "$ENV{ROCM_PATH}" "/opt/rocm") + if(_rocm_fc_root) + list(APPEND _rocm_fc_hints "${_rocm_fc_root}/bin" + "${_rocm_fc_root}/lib/llvm/bin") + endif() + endforeach() + # Uncached, so a changed ROCM_PATH is honoured on reconfigure. + unset(_rocm_amdflang CACHE) + find_program(_rocm_amdflang NAMES amdflang HINTS ${_rocm_fc_hints} + NO_DEFAULT_PATH) + # A ROCm tree can carry a dangling amdflang symlink. + if(_rocm_amdflang) + execute_process(COMMAND "${_rocm_amdflang}" --version + RESULT_VARIABLE _rocm_fc_rc + OUTPUT_QUIET ERROR_QUIET) + if(_rocm_fc_rc EQUAL 0) + set(CMAKE_Fortran_COMPILER "${_rocm_amdflang}" CACHE FILEPATH + "Fortran compiler used for the ROCm Fortran bindings" FORCE) + endif() + endif() + unset(_rocm_amdflang CACHE) +endif() + +check_language(Fortran) + +set(ROCM_WARN_TOOLCHAIN_VAR "${_rocm_fc_warn_toolchain_var}") +set(ROCM_ERROR_TOOLCHAIN_VAR "${_rocm_fc_error_toolchain_var}") +unset(_rocm_fc_warn_toolchain_var) +unset(_rocm_fc_error_toolchain_var) +unset(_rocm_fc_hints) +unset(_rocm_fc_cc) +unset(_rocm_fc_dir) +unset(_rocm_fc_root) +unset(_rocm_fc_rc) diff --git a/bindings/fortran/rocfft/cmake/toolchains/amdflang.cmake b/bindings/fortran/rocfft/cmake/toolchains/amdflang.cmake new file mode 100644 index 000000000000..94f565807650 --- /dev/null +++ b/bindings/fortran/rocfft/cmake/toolchains/amdflang.cmake @@ -0,0 +1,30 @@ +# ############################################################################# +# 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. +# ############################################################################# + +# AMD ROCm toolchain (amdflang). +# Usage, from bindings/fortran/rocfft: +# cmake -S . -B build -DCMAKE_TOOLCHAIN_FILE=cmake/toolchains/amdflang.cmake + +set(CMAKE_Fortran_COMPILER amdflang CACHE FILEPATH "Fortran compiler") +set(CMAKE_C_COMPILER amdclang CACHE FILEPATH "C compiler") diff --git a/bindings/fortran/rocfft/cmake/toolchains/cray.cmake b/bindings/fortran/rocfft/cmake/toolchains/cray.cmake new file mode 100644 index 000000000000..b26b38ce6420 --- /dev/null +++ b/bindings/fortran/rocfft/cmake/toolchains/cray.cmake @@ -0,0 +1,30 @@ +# ############################################################################# +# 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. +# ############################################################################# + +# Cray toolchain (ftn and cc wrappers; the loaded PrgEnv module picks the compiler). +# Usage, from bindings/fortran/rocfft: +# cmake -S . -B build -DCMAKE_TOOLCHAIN_FILE=cmake/toolchains/cray.cmake + +set(CMAKE_Fortran_COMPILER ftn CACHE FILEPATH "Fortran compiler") +set(CMAKE_C_COMPILER cc CACHE FILEPATH "C compiler") diff --git a/bindings/fortran/rocfft/cmake/toolchains/gnu.cmake b/bindings/fortran/rocfft/cmake/toolchains/gnu.cmake new file mode 100644 index 000000000000..0f181295bf3a --- /dev/null +++ b/bindings/fortran/rocfft/cmake/toolchains/gnu.cmake @@ -0,0 +1,30 @@ +# ############################################################################# +# 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. +# ############################################################################# + +# GNU toolchain (gfortran). +# Usage, from bindings/fortran/rocfft: +# cmake -S . -B build -DCMAKE_TOOLCHAIN_FILE=cmake/toolchains/gnu.cmake + +set(CMAKE_Fortran_COMPILER gfortran CACHE FILEPATH "Fortran compiler") +set(CMAKE_C_COMPILER gcc CACHE FILEPATH "C compiler") diff --git a/bindings/fortran/rocfft/cmake/toolchains/intel-classic.cmake b/bindings/fortran/rocfft/cmake/toolchains/intel-classic.cmake new file mode 100644 index 000000000000..6d69ec14ca90 --- /dev/null +++ b/bindings/fortran/rocfft/cmake/toolchains/intel-classic.cmake @@ -0,0 +1,30 @@ +# ############################################################################# +# 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. +# ############################################################################# + +# Classic Intel toolchain (ifort, deprecated by Intel in favour of ifx). +# Usage, from bindings/fortran/rocfft: +# cmake -S . -B build -DCMAKE_TOOLCHAIN_FILE=cmake/toolchains/intel-classic.cmake + +set(CMAKE_Fortran_COMPILER ifort CACHE FILEPATH "Fortran compiler") +set(CMAKE_C_COMPILER icx CACHE FILEPATH "C compiler") diff --git a/bindings/fortran/rocfft/cmake/toolchains/intel.cmake b/bindings/fortran/rocfft/cmake/toolchains/intel.cmake new file mode 100644 index 000000000000..0ca38cecbeed --- /dev/null +++ b/bindings/fortran/rocfft/cmake/toolchains/intel.cmake @@ -0,0 +1,30 @@ +# ############################################################################# +# 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. +# ############################################################################# + +# Intel oneAPI toolchain (ifx, the LLVM-based compiler). +# Usage, from bindings/fortran/rocfft: +# cmake -S . -B build -DCMAKE_TOOLCHAIN_FILE=cmake/toolchains/intel.cmake + +set(CMAKE_Fortran_COMPILER ifx CACHE FILEPATH "Fortran compiler") +set(CMAKE_C_COMPILER icx CACHE FILEPATH "C compiler") diff --git a/bindings/fortran/rocfft/cmake/toolchains/nvhpc.cmake b/bindings/fortran/rocfft/cmake/toolchains/nvhpc.cmake new file mode 100644 index 000000000000..e62774b10e0f --- /dev/null +++ b/bindings/fortran/rocfft/cmake/toolchains/nvhpc.cmake @@ -0,0 +1,30 @@ +# ############################################################################# +# 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. +# ############################################################################# + +# NVIDIA HPC SDK toolchain (nvfortran). +# Usage, from bindings/fortran/rocfft: +# cmake -S . -B build -DCMAKE_TOOLCHAIN_FILE=cmake/toolchains/nvhpc.cmake + +set(CMAKE_Fortran_COMPILER nvfortran CACHE FILEPATH "Fortran compiler") +set(CMAKE_C_COMPILER nvc CACHE FILEPATH "C compiler") diff --git a/bindings/fortran/rocfft/metadata/rocfft.toml b/bindings/fortran/rocfft/metadata/rocfft.toml new file mode 100644 index 000000000000..e8c6a9a2bdb6 --- /dev/null +++ b/bindings/fortran/rocfft/metadata/rocfft.toml @@ -0,0 +1,37 @@ +# Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. +# SPDX-License-Identifier: MIT +# +# What the rocfft C headers cannot express, beyond the defaults the generator +# derives from them: one table per C function, and a key with braces names +# several. A _64 or _spt variant without a table takes its base's. See "Fixing a +# binding" in the generator README; `gen.sh --explain ` shows what +# applies. + +format = 1 +# by-value: +# The callback `void**` arrays are documented nullable. By reference, a variable +# holding c_null_ptr arrives as its own non-NULL address, so a callback could +# never be cleared. + +[rocfft_cache_serialize] +buffer_len_bytes = "scalar-host" + +# The buffer arrays are on the host and hold one pointer per buffer: by +# reference, a caller passes one device pointer variable, or c_null_ptr. +[rocfft_execute] +in_buffer = "by-reference" +out_buffer = "by-reference" + +[rocfft_execution_info_set_load_callback] +cb_functions = "by-value" +cb_data = "by-value" + +[rocfft_execution_info_set_load_callback_data] +cb_data = "by-value" + +[rocfft_execution_info_set_store_callback] +cb_functions = "by-value" +cb_data = "by-value" + +[rocfft_execution_info_set_store_callback_data] +cb_data = "by-value" diff --git a/bindings/fortran/rocfft/rocfft-fortran-config.cmake.in b/bindings/fortran/rocfft/rocfft-fortran-config.cmake.in new file mode 100644 index 000000000000..08e0102c6761 --- /dev/null +++ b/bindings/fortran/rocfft/rocfft-fortran-config.cmake.in @@ -0,0 +1,157 @@ +# ############################################################################# +# 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. +# ############################################################################# + +# Config package for the rocFFT Fortran bindings. +# +# find_package(rocfft-fortran REQUIRED) +# target_link_libraries(my_app PRIVATE roc::rocfft_fortran) +# +# One archive and .mod set is installed per Fortran compiler. Set +# ROCFFT_FORTRAN_COMPILER_DIR before find_package to pick one explicitly. + +@PACKAGE_INIT@ + +set(_rocfft_built_dir "@ROCFFT_FORTRAN_COMPILER_DIR@") + +# Absolute, relocatable paths written by PATH_VARS (see PACKAGE_INIT above). +set(_rocfft_libdir "@PACKAGE_ROCFFT_FORTRAN_LIBDIR@") +set(_rocfft_incdir "@PACKAGE_ROCFFT_FORTRAN_INCLUDEDIR@") + +set(rocfft_FORTRAN_FOUND FALSE) + +if(NOT CMAKE_Fortran_COMPILER_LOADED) + set(rocfft-fortran_FOUND FALSE) + string(CONCAT rocfft-fortran_NOT_FOUND_MESSAGE + "The rocFFT Fortran bindings need Fortran enabled before find_package: " + "add Fortran to your project() languages, or call enable_language(Fortran).") + return() +endif() + +# Accepts FC itself or FC plus a separator and suffix (gfortran-13). Not a regex: +# FC may be an arbitrary basename, and CMake has no regex-escape. +function(_rocfft_dir_matches DIR FC OUT) + set(${OUT} FALSE PARENT_SCOPE) + if(DIR STREQUAL FC) + set(${OUT} TRUE PARENT_SCOPE) + return() + endif() + string(LENGTH "${FC}" _n) + string(LENGTH "${DIR}" _dn) + if(_dn GREATER _n) + string(SUBSTRING "${DIR}" 0 ${_n} _head) + string(SUBSTRING "${DIR}" ${_n} 1 _sep) + if(_head STREQUAL FC AND _sep MATCHES "^[-._]$") + set(${OUT} TRUE PARENT_SCOPE) + endif() + endif() +endfunction() + +file(GLOB _rocfft_avail RELATIVE "${_rocfft_incdir}" "${_rocfft_incdir}/*") + +set(_rocfft_dir "") +if(DEFINED ROCFFT_FORTRAN_COMPILER_DIR) + if(NOT EXISTS "${_rocfft_incdir}/${ROCFFT_FORTRAN_COMPILER_DIR}/rocfft.mod" OR + NOT EXISTS "${_rocfft_libdir}/${ROCFFT_FORTRAN_COMPILER_DIR}/librocfft_fortran.a") + set(rocfft-fortran_FOUND FALSE) + string(CONCAT rocfft-fortran_NOT_FOUND_MESSAGE + "ROCFFT_FORTRAN_COMPILER_DIR is '${ROCFFT_FORTRAN_COMPILER_DIR}', but " + "no rocfft.mod and librocfft_fortran.a are installed under " + "${_rocfft_incdir}/${ROCFFT_FORTRAN_COMPILER_DIR} and " + "${_rocfft_libdir}/${ROCFFT_FORTRAN_COMPILER_DIR}. Available: " + "${_rocfft_avail}.") + return() + endif() + set(_rocfft_dir "${ROCFFT_FORTRAN_COMPILER_DIR}") +else() + # Must mirror the mapping in bindings/fortran/rocfft/CMakeLists.txt. + if(CMAKE_Fortran_COMPILER_ID STREQUAL "GNU") + set(_rocfft_fc "gfortran") + elseif(CMAKE_Fortran_COMPILER_ID MATCHES "Flang") + set(_rocfft_fc "amdflang") + elseif(CMAKE_Fortran_COMPILER_ID STREQUAL "Cray") + set(_rocfft_fc "ftn") + elseif(CMAKE_Fortran_COMPILER_ID STREQUAL "IntelLLVM") + set(_rocfft_fc "ifx") + elseif(CMAKE_Fortran_COMPILER_ID STREQUAL "Intel") + set(_rocfft_fc "ifort") + elseif(CMAKE_Fortran_COMPILER_ID STREQUAL "NVHPC") + set(_rocfft_fc "nvfortran") + else() + get_filename_component(_rocfft_fc "${CMAKE_Fortran_COMPILER}" NAME_WE) + string(TOLOWER "${_rocfft_fc}" _rocfft_fc) + endif() + + # A prefix can hold several compilers' artifacts; the recorded one is a fallback. + if(EXISTS "${_rocfft_incdir}/${_rocfft_fc}/rocfft.mod") + set(_rocfft_dir "${_rocfft_fc}") + else() + _rocfft_dir_matches("${_rocfft_built_dir}" "${_rocfft_fc}" _rocfft_ok) + if(_rocfft_ok) + set(_rocfft_dir "${_rocfft_built_dir}") + endif() + endif() + + if(NOT _rocfft_dir) + set(rocfft-fortran_FOUND FALSE) + string(CONCAT rocfft-fortran_NOT_FOUND_MESSAGE + "This project builds with ${_rocfft_fc}, but no rocfft.mod for it is " + "installed under ${_rocfft_incdir}. Available: " + "${_rocfft_avail}. A Fortran .mod cannot be shared across compilers; " + "rebuild the bindings with ${_rocfft_fc}, or set " + "ROCFFT_FORTRAN_COMPILER_DIR to the directory to use.") + return() + endif() +endif() + +# Written TRUE only when the archive links roc::rocfft, which must be imported +# before the targets file. hip's config needs C or CXX enabled (FindThreads). +if(@ROCFFT_FORTRAN_NEEDS_C_PACKAGE@) + # Not from rocfft-config.cmake, which includes the Fortran package after its targets. + string(TOLOWER "${CMAKE_FIND_PACKAGE_NAME}" _rocfft_caller) + if(NOT TARGET roc::rocfft AND NOT _rocfft_caller STREQUAL "rocfft") + if(NOT CMAKE_C_COMPILER_LOADED AND NOT CMAKE_CXX_COMPILER_LOADED) + enable_language(C) + endif() + include(CMakeFindDependencyMacro) + find_dependency(rocfft) + endif() +endif() + +include("${CMAKE_CURRENT_LIST_DIR}/rocfft-fortran-targets.cmake") + +# The targets file hardcodes the directory the package was built with. +if(NOT _rocfft_dir STREQUAL "${_rocfft_built_dir}") + set_target_properties(roc::rocfft_fortran PROPERTIES + INTERFACE_INCLUDE_DIRECTORIES "${_rocfft_incdir}/${_rocfft_dir}") + get_target_property(_rocfft_cfgs roc::rocfft_fortran IMPORTED_CONFIGURATIONS) + foreach(_rocfft_c IN LISTS _rocfft_cfgs) + set_target_properties(roc::rocfft_fortran PROPERTIES + IMPORTED_LOCATION_${_rocfft_c} + "${_rocfft_libdir}/${_rocfft_dir}/librocfft_fortran.a") + endforeach() +endif() + +set(rocfft_FORTRAN_FOUND TRUE) + +check_required_components(rocfft-fortran) diff --git a/bindings/fortran/rocfft/rocfft.F90 b/bindings/fortran/rocfft/rocfft.F90 new file mode 100644 index 000000000000..7f0f1bcc2b59 --- /dev/null +++ b/bindings/fortran/rocfft/rocfft.F90 @@ -0,0 +1,721 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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 rocfft + use, intrinsic :: iso_c_binding + implicit none + + ! rocfft_status_e + enum, bind(c) + enumerator :: rocfft_status_success = 0 + enumerator :: rocfft_status_failure = 1 + enumerator :: rocfft_status_invalid_arg_value = 2 + enumerator :: rocfft_status_invalid_dimensions = 3 + enumerator :: rocfft_status_invalid_array_type = 4 + enumerator :: rocfft_status_invalid_strides = 5 + enumerator :: rocfft_status_invalid_distance = 6 + enumerator :: rocfft_status_invalid_offset = 7 + enumerator :: rocfft_status_invalid_work_buffer = 8 + end enum + + ! rocfft_transform_type_e + enum, bind(c) + enumerator :: rocfft_transform_type_complex_forward = 0 + enumerator :: rocfft_transform_type_complex_inverse = 1 + enumerator :: rocfft_transform_type_real_forward = 2 + enumerator :: rocfft_transform_type_real_inverse = 3 + end enum + + ! rocfft_precision_e + enum, bind(c) + enumerator :: rocfft_precision_single = 0 + enumerator :: rocfft_precision_double = 1 + enumerator :: rocfft_precision_half = 2 + end enum + + ! rocfft_result_placement_e + enum, bind(c) + enumerator :: rocfft_placement_inplace = 0 + enumerator :: rocfft_placement_notinplace = 1 + end enum + + ! rocfft_array_type_e + enum, bind(c) + enumerator :: rocfft_array_type_complex_interleaved = 0 + enumerator :: rocfft_array_type_complex_planar = 1 + enumerator :: rocfft_array_type_real = 2 + enumerator :: rocfft_array_type_hermitian_interleaved = 3 + enumerator :: rocfft_array_type_hermitian_planar = 4 + enumerator :: rocfft_array_type_unset = 5 + end enum + + ! rocfft_comm_type_e + enum, bind(c) + enumerator :: rocfft_comm_none = 0 + enumerator :: rocfft_comm_mpi = 1 + end enum + + + interface rocfft_setup + function rocfft_setup_() bind(c, name="rocfft_setup") + import + implicit none + integer(c_int) :: rocfft_setup_ + end function + end interface + + interface rocfft_cleanup + function rocfft_cleanup_() bind(c, name="rocfft_cleanup") + import + implicit none + integer(c_int) :: rocfft_cleanup_ + end function + end interface + + interface rocfft_plan_create + function rocfft_plan_create_(plan,placement,transform_type,myPrecision,dimensions,lengths, & + number_of_transforms,description) & + bind(c, name="rocfft_plan_create") + import + implicit none + integer(c_int) :: rocfft_plan_create_ + type(c_ptr) :: plan + integer(c_int),value :: placement + integer(c_int),value :: transform_type + integer(c_int),value :: myPrecision + integer(c_size_t),value :: dimensions + type(c_ptr),value :: lengths + integer(c_size_t),value :: number_of_transforms + type(c_ptr),value :: description + end function + +#ifdef USE_ASSUMED_RANK + module procedure rocfft_plan_create_assumed_rank +#else +#ifdef USE_ASSUMED_SHAPE + module procedure & + rocfft_plan_create_rank_0,& + rocfft_plan_create_rank_1 +#endif +#endif + end interface + + interface rocfft_execute + function rocfft_execute_(plan,in_buffer,out_buffer,myInfo) bind(c, name="rocfft_execute") + import + implicit none + integer(c_int) :: rocfft_execute_ + type(c_ptr),value :: plan + type(c_ptr) :: in_buffer + type(c_ptr) :: out_buffer + type(c_ptr),value :: myInfo + end function + end interface + + interface rocfft_plan_destroy + function rocfft_plan_destroy_(plan) bind(c, name="rocfft_plan_destroy") + import + implicit none + integer(c_int) :: rocfft_plan_destroy_ + type(c_ptr),value :: plan + end function + end interface + + interface rocfft_plan_description_set_scale_factor + function rocfft_plan_description_set_scale_factor_(description,scale_factor) & + bind(c, name="rocfft_plan_description_set_scale_factor") + import + implicit none + integer(c_int) :: rocfft_plan_description_set_scale_factor_ + type(c_ptr),value :: description + real(c_double),value :: scale_factor + end function + end interface + + interface rocfft_plan_description_set_data_layout + function rocfft_plan_description_set_data_layout_(description,in_array_type,out_array_type, & + in_offsets,out_offsets,in_strides_size,in_strides,in_distance,out_strides_size, & + out_strides,out_distance) & + bind(c, name="rocfft_plan_description_set_data_layout") + import + implicit none + integer(c_int) :: rocfft_plan_description_set_data_layout_ + type(c_ptr),value :: description + integer(c_int),value :: in_array_type + integer(c_int),value :: out_array_type + type(c_ptr),value :: in_offsets + type(c_ptr),value :: out_offsets + integer(c_size_t),value :: in_strides_size + type(c_ptr),value :: in_strides + integer(c_size_t),value :: in_distance + integer(c_size_t),value :: out_strides_size + type(c_ptr),value :: out_strides + integer(c_size_t),value :: out_distance + end function + +#ifdef USE_ASSUMED_RANK + module procedure rocfft_plan_description_set_data_layout_assumed_rank +#else +#ifdef USE_ASSUMED_SHAPE + module procedure & + rocfft_plan_description_set_data_layout_rank_0,& + rocfft_plan_description_set_data_layout_rank_1 +#endif +#endif + end interface + + interface rocfft_field_create + function rocfft_field_create_(field) bind(c, name="rocfft_field_create") + import + implicit none + integer(c_int) :: rocfft_field_create_ + type(c_ptr) :: field + end function + end interface + + interface rocfft_field_destroy + function rocfft_field_destroy_(field) bind(c, name="rocfft_field_destroy") + import + implicit none + integer(c_int) :: rocfft_field_destroy_ + type(c_ptr),value :: field + end function + end interface + + interface rocfft_get_version_string + function rocfft_get_version_string_(buf,len) bind(c, name="rocfft_get_version_string") + import + implicit none + integer(c_int) :: rocfft_get_version_string_ + type(c_ptr),value :: buf + integer(c_size_t),value :: len + end function + end interface + + interface rocfft_plan_description_set_comm + function rocfft_plan_description_set_comm_(description,comm_type,comm_handle) & + bind(c, name="rocfft_plan_description_set_comm") + import + implicit none + integer(c_int) :: rocfft_plan_description_set_comm_ + type(c_ptr),value :: description + integer(c_int),value :: comm_type + type(c_ptr),value :: comm_handle + end function + end interface + + interface rocfft_plan_description_set_load_callback + function rocfft_plan_description_set_load_callback_(description,symbol_name,bitcode_data, & + bitcode_len_bytes,shared_mem_bytes) & + bind(c, name="rocfft_plan_description_set_load_callback") + import + implicit none + integer(c_int) :: rocfft_plan_description_set_load_callback_ + type(c_ptr),value :: description + type(c_ptr),value :: symbol_name + type(c_ptr),value :: bitcode_data + integer(c_size_t),value :: bitcode_len_bytes + integer(c_size_t),value :: shared_mem_bytes + end function + end interface + + interface rocfft_plan_description_set_store_callback + function rocfft_plan_description_set_store_callback_(description,symbol_name,bitcode_data, & + bitcode_len_bytes,shared_mem_bytes) & + bind(c, name="rocfft_plan_description_set_store_callback") + import + implicit none + integer(c_int) :: rocfft_plan_description_set_store_callback_ + type(c_ptr),value :: description + type(c_ptr),value :: symbol_name + type(c_ptr),value :: bitcode_data + integer(c_size_t),value :: bitcode_len_bytes + integer(c_size_t),value :: shared_mem_bytes + end function + end interface + + interface rocfft_brick_create + function rocfft_brick_create_(brick,field_lower,field_upper,brick_stride,dim_with_batch, & + deviceID) & + bind(c, name="rocfft_brick_create") + import + implicit none + integer(c_int) :: rocfft_brick_create_ + type(c_ptr) :: brick + type(c_ptr),value :: field_lower + type(c_ptr),value :: field_upper + type(c_ptr),value :: brick_stride + integer(c_size_t),value :: dim_with_batch + integer(c_int),value :: deviceID + end function + +#ifdef USE_ASSUMED_RANK + module procedure rocfft_brick_create_assumed_rank +#else +#ifdef USE_ASSUMED_SHAPE + module procedure & + rocfft_brick_create_rank_0,& + rocfft_brick_create_rank_1 +#endif +#endif + end interface + + interface rocfft_brick_destroy + function rocfft_brick_destroy_(brick) bind(c, name="rocfft_brick_destroy") + import + implicit none + integer(c_int) :: rocfft_brick_destroy_ + type(c_ptr),value :: brick + end function + end interface + + interface rocfft_field_add_brick + function rocfft_field_add_brick_(field,brick) bind(c, name="rocfft_field_add_brick") + import + implicit none + integer(c_int) :: rocfft_field_add_brick_ + type(c_ptr),value :: field + type(c_ptr),value :: brick + end function + end interface + + interface rocfft_plan_description_add_infield + function rocfft_plan_description_add_infield_(description,field) & + bind(c, name="rocfft_plan_description_add_infield") + import + implicit none + integer(c_int) :: rocfft_plan_description_add_infield_ + type(c_ptr),value :: description + type(c_ptr),value :: field + end function + end interface + + interface rocfft_plan_description_add_outfield + function rocfft_plan_description_add_outfield_(description,field) & + bind(c, name="rocfft_plan_description_add_outfield") + import + implicit none + integer(c_int) :: rocfft_plan_description_add_outfield_ + type(c_ptr),value :: description + type(c_ptr),value :: field + end function + end interface + + interface rocfft_plan_get_work_buffer_size + function rocfft_plan_get_work_buffer_size_(plan,size_in_bytes) & + bind(c, name="rocfft_plan_get_work_buffer_size") + import + implicit none + integer(c_int) :: rocfft_plan_get_work_buffer_size_ + type(c_ptr),value :: plan + type(c_ptr),value :: size_in_bytes + end function + + module procedure rocfft_plan_get_work_buffer_size_typed + end interface + + interface rocfft_plan_get_print + function rocfft_plan_get_print_(plan) bind(c, name="rocfft_plan_get_print") + import + implicit none + integer(c_int) :: rocfft_plan_get_print_ + type(c_ptr),value :: plan + end function + end interface + + interface rocfft_plan_description_create + function rocfft_plan_description_create_(description) & + bind(c, name="rocfft_plan_description_create") + import + implicit none + integer(c_int) :: rocfft_plan_description_create_ + type(c_ptr) :: description + end function + end interface + + interface rocfft_plan_description_destroy + function rocfft_plan_description_destroy_(description) & + bind(c, name="rocfft_plan_description_destroy") + import + implicit none + integer(c_int) :: rocfft_plan_description_destroy_ + type(c_ptr),value :: description + end function + end interface + + interface rocfft_execution_info_create + function rocfft_execution_info_create_(myInfo) bind(c, name="rocfft_execution_info_create") + import + implicit none + integer(c_int) :: rocfft_execution_info_create_ + type(c_ptr) :: myInfo + end function + end interface + + interface rocfft_execution_info_destroy + function rocfft_execution_info_destroy_(myInfo) bind(c, name="rocfft_execution_info_destroy") + import + implicit none + integer(c_int) :: rocfft_execution_info_destroy_ + type(c_ptr),value :: myInfo + end function + end interface + + interface rocfft_execution_info_set_work_buffer + function rocfft_execution_info_set_work_buffer_(myInfo,work_buffer,size_in_bytes) & + bind(c, name="rocfft_execution_info_set_work_buffer") + import + implicit none + integer(c_int) :: rocfft_execution_info_set_work_buffer_ + type(c_ptr),value :: myInfo + type(c_ptr),value :: work_buffer + integer(c_size_t),value :: size_in_bytes + end function + end interface + + interface rocfft_execution_info_set_stream + function rocfft_execution_info_set_stream_(myInfo,stream) & + bind(c, name="rocfft_execution_info_set_stream") + import + implicit none + integer(c_int) :: rocfft_execution_info_set_stream_ + type(c_ptr),value :: myInfo + type(c_ptr),value :: stream + end function + end interface + + interface rocfft_execution_info_set_load_callback + function rocfft_execution_info_set_load_callback_(myInfo,cb_functions,cb_data, & + shared_mem_bytes) & + bind(c, name="rocfft_execution_info_set_load_callback") + import + implicit none + integer(c_int) :: rocfft_execution_info_set_load_callback_ + type(c_ptr),value :: myInfo + type(c_ptr),value :: cb_functions + type(c_ptr),value :: cb_data + integer(c_size_t),value :: shared_mem_bytes + end function + end interface + + interface rocfft_execution_info_set_load_callback_data + function rocfft_execution_info_set_load_callback_data_(myInfo,cb_data,count) & + bind(c, name="rocfft_execution_info_set_load_callback_data") + import + implicit none + integer(c_int) :: rocfft_execution_info_set_load_callback_data_ + type(c_ptr),value :: myInfo + type(c_ptr),value :: cb_data + integer(c_size_t),value :: count + end function + end interface + + interface rocfft_execution_info_set_store_callback + function rocfft_execution_info_set_store_callback_(myInfo,cb_functions,cb_data, & + shared_mem_bytes) & + bind(c, name="rocfft_execution_info_set_store_callback") + import + implicit none + integer(c_int) :: rocfft_execution_info_set_store_callback_ + type(c_ptr),value :: myInfo + type(c_ptr),value :: cb_functions + type(c_ptr),value :: cb_data + integer(c_size_t),value :: shared_mem_bytes + end function + end interface + + interface rocfft_execution_info_set_store_callback_data + function rocfft_execution_info_set_store_callback_data_(myInfo,cb_data,count) & + bind(c, name="rocfft_execution_info_set_store_callback_data") + import + implicit none + integer(c_int) :: rocfft_execution_info_set_store_callback_data_ + type(c_ptr),value :: myInfo + type(c_ptr),value :: cb_data + integer(c_size_t),value :: count + end function + end interface + + interface rocfft_cache_serialize + function rocfft_cache_serialize_(buffer,buffer_len_bytes) bind(c, name="rocfft_cache_serialize") + import + implicit none + integer(c_int) :: rocfft_cache_serialize_ + type(c_ptr) :: buffer + type(c_ptr),value :: buffer_len_bytes + end function + + module procedure rocfft_cache_serialize_typed + end interface + + interface rocfft_cache_buffer_free + function rocfft_cache_buffer_free_(buffer) bind(c, name="rocfft_cache_buffer_free") + import + implicit none + integer(c_int) :: rocfft_cache_buffer_free_ + type(c_ptr),value :: buffer + end function + end interface + + interface rocfft_cache_deserialize + function rocfft_cache_deserialize_(buffer,buffer_len_bytes) & + bind(c, name="rocfft_cache_deserialize") + import + implicit none + integer(c_int) :: rocfft_cache_deserialize_ + type(c_ptr),value :: buffer + integer(c_size_t),value :: buffer_len_bytes + end function + end interface + + + contains + + subroutine rocfftCheck(status) + implicit none + integer(kind(rocfft_status_success)) :: status + character(len=33) :: name + if (status /= rocfft_status_success) then + name = "unknown status" + select case (status) + case (rocfft_status_failure) + name = "rocfft_status_failure" + case (rocfft_status_invalid_arg_value) + name = "rocfft_status_invalid_arg_value" + case (rocfft_status_invalid_dimensions) + name = "rocfft_status_invalid_dimensions" + case (rocfft_status_invalid_array_type) + name = "rocfft_status_invalid_array_type" + case (rocfft_status_invalid_strides) + name = "rocfft_status_invalid_strides" + case (rocfft_status_invalid_distance) + name = "rocfft_status_invalid_distance" + case (rocfft_status_invalid_offset) + name = "rocfft_status_invalid_offset" + case (rocfft_status_invalid_work_buffer) + name = "rocfft_status_invalid_work_buffer" + end select + write (*, "(a,i0,a)") "ROCFFT ERROR: " // trim(name) // " (code ", status, ")" + stop 1 + end if + end subroutine rocfftCheck + function rocfft_plan_get_work_buffer_size_typed(plan,size_in_bytes) + implicit none + integer(kind(rocfft_status_success)) :: rocfft_plan_get_work_buffer_size_typed + type(c_ptr), value :: plan + integer(c_size_t),target :: size_in_bytes + ! + rocfft_plan_get_work_buffer_size_typed = rocfft_plan_get_work_buffer_size_(plan, & + c_loc(size_in_bytes)) + end function + + function rocfft_cache_serialize_typed(buffer,buffer_len_bytes) + implicit none + integer(kind(rocfft_status_success)) :: rocfft_cache_serialize_typed + type(c_ptr) :: buffer + integer(c_size_t),target :: buffer_len_bytes + ! + rocfft_cache_serialize_typed = rocfft_cache_serialize_(buffer,c_loc(buffer_len_bytes)) + end function + + +#if defined(USE_ASSUMED_SHAPE) || defined(USE_ASSUMED_RANK) + +#ifdef USE_ASSUMED_RANK + function rocfft_plan_create_assumed_rank(plan,placement,transform_type,myPrecision,dimensions, & + lengths,number_of_transforms,description) + implicit none + integer(kind(rocfft_status_success)) :: rocfft_plan_create_assumed_rank + type(c_ptr) :: plan + integer(kind(rocfft_placement_inplace)) :: placement + integer(kind(rocfft_transform_type_complex_forward)) :: transform_type + integer(kind(rocfft_precision_single)) :: myPrecision + integer(c_size_t) :: dimensions + integer(c_size_t),target,dimension(..) :: lengths + integer(c_size_t) :: number_of_transforms + type(c_ptr), value :: description + ! + rocfft_plan_create_assumed_rank = rocfft_plan_create_(plan,placement,transform_type, & + myPrecision,dimensions,c_loc(lengths),number_of_transforms,description) + end function + +#else + function rocfft_plan_create_rank_0(plan,placement,transform_type,myPrecision,dimensions, & + lengths,number_of_transforms,description) + implicit none + integer(kind(rocfft_status_success)) :: rocfft_plan_create_rank_0 + type(c_ptr) :: plan + integer(kind(rocfft_placement_inplace)) :: placement + integer(kind(rocfft_transform_type_complex_forward)) :: transform_type + integer(kind(rocfft_precision_single)) :: myPrecision + integer(c_size_t) :: dimensions + integer(c_size_t),target :: lengths + integer(c_size_t) :: number_of_transforms + type(c_ptr), value :: description + ! + rocfft_plan_create_rank_0 = rocfft_plan_create_(plan,placement,transform_type,myPrecision, & + dimensions,c_loc(lengths),number_of_transforms,description) + end function + + function rocfft_plan_create_rank_1(plan,placement,transform_type,myPrecision,dimensions, & + lengths,number_of_transforms,description) + implicit none + integer(kind(rocfft_status_success)) :: rocfft_plan_create_rank_1 + type(c_ptr) :: plan + integer(kind(rocfft_placement_inplace)) :: placement + integer(kind(rocfft_transform_type_complex_forward)) :: transform_type + integer(kind(rocfft_precision_single)) :: myPrecision + integer(c_size_t) :: dimensions + integer(c_size_t),target,dimension(:) :: lengths + integer(c_size_t) :: number_of_transforms + type(c_ptr), value :: description + ! + rocfft_plan_create_rank_1 = rocfft_plan_create_(plan,placement,transform_type,myPrecision, & + dimensions,c_loc(lengths),number_of_transforms,description) + end function + +#endif +#ifdef USE_ASSUMED_RANK + function rocfft_plan_description_set_data_layout_assumed_rank(description,in_array_type, & + out_array_type,in_offsets,out_offsets,in_strides_size,in_strides,in_distance, & + out_strides_size,out_strides,out_distance) + implicit none + integer(kind(rocfft_status_success)) :: rocfft_plan_description_set_data_layout_assumed_rank + type(c_ptr), value :: description + integer(kind(rocfft_array_type_complex_interleaved)) :: in_array_type + integer(kind(rocfft_array_type_complex_interleaved)) :: out_array_type + integer(c_size_t),target,dimension(..) :: in_offsets + integer(c_size_t),target,dimension(..) :: out_offsets + integer(c_size_t) :: in_strides_size + integer(c_size_t),target,dimension(..) :: in_strides + integer(c_size_t) :: in_distance + integer(c_size_t) :: out_strides_size + integer(c_size_t),target,dimension(..) :: out_strides + integer(c_size_t) :: out_distance + ! + rocfft_plan_description_set_data_layout_assumed_rank = & + rocfft_plan_description_set_data_layout_(description,in_array_type,out_array_type, & + c_loc(in_offsets),c_loc(out_offsets),in_strides_size,c_loc(in_strides),in_distance, & + out_strides_size,c_loc(out_strides),out_distance) + end function + +#else + function rocfft_plan_description_set_data_layout_rank_0(description,in_array_type, & + out_array_type,in_offsets,out_offsets,in_strides_size,in_strides,in_distance, & + out_strides_size,out_strides,out_distance) + implicit none + integer(kind(rocfft_status_success)) :: rocfft_plan_description_set_data_layout_rank_0 + type(c_ptr), value :: description + integer(kind(rocfft_array_type_complex_interleaved)) :: in_array_type + integer(kind(rocfft_array_type_complex_interleaved)) :: out_array_type + integer(c_size_t),target :: in_offsets + integer(c_size_t),target :: out_offsets + integer(c_size_t) :: in_strides_size + integer(c_size_t),target :: in_strides + integer(c_size_t) :: in_distance + integer(c_size_t) :: out_strides_size + integer(c_size_t),target :: out_strides + integer(c_size_t) :: out_distance + ! + rocfft_plan_description_set_data_layout_rank_0 = rocfft_plan_description_set_data_layout_( & + description,in_array_type,out_array_type,c_loc(in_offsets),c_loc(out_offsets), & + in_strides_size,c_loc(in_strides),in_distance,out_strides_size,c_loc(out_strides), & + out_distance) + end function + + function rocfft_plan_description_set_data_layout_rank_1(description,in_array_type, & + out_array_type,in_offsets,out_offsets,in_strides_size,in_strides,in_distance, & + out_strides_size,out_strides,out_distance) + implicit none + integer(kind(rocfft_status_success)) :: rocfft_plan_description_set_data_layout_rank_1 + type(c_ptr), value :: description + integer(kind(rocfft_array_type_complex_interleaved)) :: in_array_type + integer(kind(rocfft_array_type_complex_interleaved)) :: out_array_type + integer(c_size_t),target,dimension(:) :: in_offsets + integer(c_size_t),target,dimension(:) :: out_offsets + integer(c_size_t) :: in_strides_size + integer(c_size_t),target,dimension(:) :: in_strides + integer(c_size_t) :: in_distance + integer(c_size_t) :: out_strides_size + integer(c_size_t),target,dimension(:) :: out_strides + integer(c_size_t) :: out_distance + ! + rocfft_plan_description_set_data_layout_rank_1 = rocfft_plan_description_set_data_layout_( & + description,in_array_type,out_array_type,c_loc(in_offsets),c_loc(out_offsets), & + in_strides_size,c_loc(in_strides),in_distance,out_strides_size,c_loc(out_strides), & + out_distance) + end function + +#endif +#ifdef USE_ASSUMED_RANK + function rocfft_brick_create_assumed_rank(brick,field_lower,field_upper,brick_stride, & + dim_with_batch,deviceID) + implicit none + integer(kind(rocfft_status_success)) :: rocfft_brick_create_assumed_rank + type(c_ptr) :: brick + integer(c_size_t),target,dimension(..) :: field_lower + integer(c_size_t),target,dimension(..) :: field_upper + integer(c_size_t),target,dimension(..) :: brick_stride + integer(c_size_t) :: dim_with_batch + integer(c_int) :: deviceID + ! + rocfft_brick_create_assumed_rank = rocfft_brick_create_(brick,c_loc(field_lower), & + c_loc(field_upper),c_loc(brick_stride),dim_with_batch,deviceID) + end function + +#else + function rocfft_brick_create_rank_0(brick,field_lower,field_upper,brick_stride,dim_with_batch, & + deviceID) + implicit none + integer(kind(rocfft_status_success)) :: rocfft_brick_create_rank_0 + type(c_ptr) :: brick + integer(c_size_t),target :: field_lower + integer(c_size_t),target :: field_upper + integer(c_size_t),target :: brick_stride + integer(c_size_t) :: dim_with_batch + integer(c_int) :: deviceID + ! + rocfft_brick_create_rank_0 = rocfft_brick_create_(brick,c_loc(field_lower), & + c_loc(field_upper),c_loc(brick_stride),dim_with_batch,deviceID) + end function + + function rocfft_brick_create_rank_1(brick,field_lower,field_upper,brick_stride,dim_with_batch, & + deviceID) + implicit none + integer(kind(rocfft_status_success)) :: rocfft_brick_create_rank_1 + type(c_ptr) :: brick + integer(c_size_t),target,dimension(:) :: field_lower + integer(c_size_t),target,dimension(:) :: field_upper + integer(c_size_t),target,dimension(:) :: brick_stride + integer(c_size_t) :: dim_with_batch + integer(c_int) :: deviceID + ! + rocfft_brick_create_rank_1 = rocfft_brick_create_(brick,c_loc(field_lower), & + c_loc(field_upper),c_loc(brick_stride),dim_with_batch,deviceID) + end function + +#endif +#endif +end module rocfft diff --git a/bindings/fortran/rocfft/test/CMakeLists.txt b/bindings/fortran/rocfft/test/CMakeLists.txt new file mode 100644 index 000000000000..aa4209a3544e --- /dev/null +++ b/bindings/fortran/rocfft/test/CMakeLists.txt @@ -0,0 +1,158 @@ +# ############################################################################# +# 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. +# ############################################################################# + +# rocFFT Fortran binding tests: one GPU executable per test. + +# ifx/ifort skip .f03/.f08 files unless -Tf precedes the source. +if(CMAKE_Fortran_COMPILER_ID MATCHES "Intel") + string(REPLACE "" "-Tf " + CMAKE_Fortran_COMPILE_OBJECT "${CMAKE_Fortran_COMPILE_OBJECT}") +endif() + +# Keep only the binding wrappers a test calls (the archive has one section per +# procedure), so a test links even when another wrapper names a C symbol the +# installed library does not export. Spelled -Wl, rather than LINKER:, which +# CMake 3.27 passes to LLVMFlang as a bare --gc-sections. +if(NOT APPLE AND NOT WIN32) + add_link_options(-Wl,--gc-sections) +endif() + +# hip-config.cmake reaches FindThreads, which needs C or CXX enabled. +if(NOT CMAKE_C_COMPILER_LOADED AND NOT CMAKE_CXX_COMPILER_LOADED) + enable_language(C) +endif() + +# The tests `use hip`, so hip.mod (rocm-systems) is a compile-time dependency. +if(NOT TARGET hip::hip_fortran) + find_package(hip-fortran QUIET) +endif() +set(_have_hip_binding TRUE) +if(NOT TARGET hip::hip_fortran) + set(_have_hip_binding FALSE) + message(STATUS + "rocFFT: no hip::hip_fortran target, skipping the Fortran binding tests " + "(they `use hip`). Point CMAKE_PREFIX_PATH at a ROCm whose HIP Fortran " + "bindings were built with this same compiler to build them.") +endif() + +if(NOT TARGET hip::host) + find_package(hip QUIET) +endif() +if(NOT TARGET roc::rocfft OR NOT TARGET hip::host) + message(STATUS + "rocFFT: no roc::rocfft or hip::host target, skipping the Fortran binding " + "tests (they link librocfft and the HIP runtime). Point CMAKE_PREFIX_PATH " + "at an installed ROCm.") + return() +endif() + +# rocfft_fortran_add_test( [F2008]): F2008 builds .f08, which +# exercises the array overloads, as test rocfft_fortran__f2008. +function(rocfft_fortran_add_test name) + if("F2008" IN_LIST ARGN) + set(src "${CMAKE_CURRENT_SOURCE_DIR}/${name}.f08") + set(test_name "${name}_f2008") + else() + set(src "${CMAKE_CURRENT_SOURCE_DIR}/${name}.f03") + set(test_name "${name}") + endif() + if(NOT EXISTS "${src}") + message(FATAL_ERROR "Missing test source: ${src}") + endif() + set_source_files_properties("${src}" PROPERTIES LANGUAGE Fortran) + add_executable(rocfft_fortran_test_${test_name} "${src}") + set_target_properties(rocfft_fortran_test_${test_name} PROPERTIES LINKER_LANGUAGE Fortran) + target_link_libraries(rocfft_fortran_test_${test_name} PRIVATE + rocfft_fortran hip::hip_fortran roc::rocfft hip::host) + # EXIT is a GNU extension. + target_compile_options(rocfft_fortran_test_${test_name} PRIVATE + $<$:-std=gnu>) + add_test(NAME rocfft_fortran_${test_name} COMMAND rocfft_fortran_test_${test_name}) + # ctest -LE gpu skips them on a GPU-less machine. + set_tests_properties(rocfft_fortran_${test_name} PROPERTIES LABELS "gpu") +endfunction() + +set(ROCFFT_FORTRAN_TESTS + rocfft + rocfft_c2c_1d_batched_z + rocfft_c2c_1d_c + rocfft_c2c_1d_notinplace_z + rocfft_c2c_1d_z + rocfft_c2c_2d_z + rocfft_c2c_3d_z + rocfft_cache_z + rocfft_callback_z + rocfft_field_brick_z + rocfft_plan_print_z + rocfft_r2c_c2r_1d_d + rocfft_r2c_c2r_1d_inplace_d + rocfft_r2c_c2r_1d_s + rocfft_scale_factor_z + rocfft_stream_z + rocfft_version + rocfft_work_buffer_z) + +if(_have_hip_binding) + foreach(_t IN LISTS ROCFFT_FORTRAN_TESTS) + rocfft_fortran_add_test(${_t}) + endforeach() + # Exits 77 when fewer than two GPUs are present. + set_tests_properties(rocfft_fortran_rocfft_field_brick_z PROPERTIES + SKIP_RETURN_CODE 77) + # The .f08 variants use the array overloads, present in both the + # assumed-shape and the assumed-rank tier. + if(NOT FORTRAN_ARRAY_INTERFACES STREQUAL "none") + foreach(_t IN LISTS ROCFFT_FORTRAN_TESTS) + rocfft_fortran_add_test(${_t} F2008) + endforeach() + set_tests_properties(rocfft_fortran_rocfft_field_brick_z_f2008 PROPERTIES + SKIP_RETURN_CODE 77) + endif() +endif() + +# Installed tests: the executables go to bin/, and a CTestTestfile.cmake with +# the same tests and properties to bin/rocfft-fortran/, for ctest --test-dir. +include(GNUInstallDirs) +get_property(_rocfft_fortran_tests DIRECTORY PROPERTY TESTS) +if(_rocfft_fortran_tests) + set(_ctest "${CMAKE_CURRENT_BINARY_DIR}/install-CTestTestfile.cmake") + file(WRITE "${_ctest}" "") + foreach(_test IN LISTS _rocfft_fortran_tests) + string(REPLACE "_fortran_" "_fortran_test_" _target "${_test}") + install(TARGETS ${_target} RUNTIME DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT tests) + file(APPEND "${_ctest}" "add_test(${_test} \"../${_target}\")\n") + foreach(_prop LABELS SKIP_RETURN_CODE ENVIRONMENT PASS_REGULAR_EXPRESSION FAIL_REGULAR_EXPRESSION) + get_test_property(${_test} ${_prop} _value) + if(_value) + file(APPEND "${_ctest}" "set_tests_properties(${_test} PROPERTIES ${_prop} [=[${_value}]=])\n") + endif() + endforeach() + endforeach() + install( + FILES "${_ctest}" + DESTINATION "${CMAKE_INSTALL_BINDIR}/rocfft-fortran" + COMPONENT tests + RENAME CTestTestfile.cmake + ) +endif() diff --git a/bindings/fortran/rocfft/test/rocfft.f03 b/bindings/fortran/rocfft/test/rocfft.f03 new file mode 100644 index 000000000000..8210dbe24f7d --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft.f03 @@ -0,0 +1,111 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! Copyright (C) 2021 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 rocfft_example + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N=16 + integer(c_size_t), parameter :: Nbytes = N*8*2 + + type double2 + double precision :: x + double precision :: y + end type double2 + + type(double2), allocatable, target, dimension(:) :: hx + complex(8), allocatable, dimension(:) :: x_in, x_ref + type(c_ptr) :: dx = c_null_ptr + type(c_ptr) :: plan = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i, j + double precision :: error + double precision, parameter :: error_max = 100*epsilon(error) + double precision, parameter :: two_pi = 2*acos(-1.0d0) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT' (Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(3)) + lengths(1) = N + + ! A non-constant input, so that a transform that leaves the data untouched + ! cannot pass. The reference is a direct O(N^2) DFT on the host; the + ! twiddle exponent is reduced modulo N so the angle stays in [0, 2*pi). + allocate(hx(N), x_in(N), x_ref(N)) + do i = 1, int(N) + x_in(i) = cmplx(i, mod(3*i, 5) - 2, kind=8) + hx(i)%x = real(x_in(i), 8) + hx(i)%y = aimag(x_in(i)) + end do + do i = 1, int(N) + x_ref(i) = (0.d0, 0.d0) + do j = 1, int(N) + x_ref(i) = x_ref(i) + x_in(j) * exp(cmplx(0.d0, -two_pi*mod((i-1)*(j-1), int(N))/N, kind=8)) + end do + end do + + call hipCheck(hipMalloc(dx,Nbytes)) + call hipCheck(hipMemcpy(dx,c_loc(hx(1)),Nbytes,hipMemcpyHostToDevice)) + + call rocfftCheck(rocfft_plan_create(plan,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + + call rocfftCheck(rocfft_execute(plan,dx,c_null_ptr,c_null_ptr)) + + call hipCheck(hipDeviceSynchronize()) + + call rocfftCheck(rocfft_plan_destroy(plan)) + + call hipCheck(hipMemcpy(c_loc(hx(1)),dx,Nbytes,hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + do i = 1, int(N) + error = abs(cmplx(hx(i)%x, hx(i)%y, kind=8) - x_ref(i)) + if (.not. (error <= error_max * sum(abs(x_in)))) then + write(*,*) "FAILED! i = ", i, " error = ", error, " hx(i) = ", hx(i)%x, hx(i)%y + call exit(1) + end if + end do + + deallocate(hx, x_in, x_ref) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_example diff --git a/bindings/fortran/rocfft/test/rocfft.f08 b/bindings/fortran/rocfft/test/rocfft.f08 new file mode 100644 index 000000000000..cec45722baf0 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft.f08 @@ -0,0 +1,101 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_example + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N=16 + + complex(8), allocatable, target, dimension(:) :: hx + complex(8), pointer, dimension(:) :: dx => null() + type(c_ptr) :: plan = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + complex(8), allocatable, dimension(:) :: x_in, x_ref + integer :: i, j + double precision :: error + double precision, parameter :: error_max = 100*epsilon(error) + double precision, parameter :: two_pi = 2*acos(-1.0d0) + + + allocate(lengths(3)) + lengths(1) = N + ! A non-constant input, so that a transform that leaves the data untouched + ! cannot pass. The reference is a direct O(N^2) DFT on the host. + allocate(hx(N), x_in(N), x_ref(N)) + do i = 1, int(N) + x_in(i) = cmplx(i, mod(3*i, 5) - 2, kind=8) + end do + hx(:) = x_in(:) + do i = 1, int(N) + x_ref(i) = (0.d0, 0.d0) + do j = 1, int(N) + x_ref(i) = x_ref(i) + x_in(j) * exp(cmplx(0.d0, -two_pi*mod((i-1)*(j-1), int(N))/N, kind=8)) + end do + end do + + call hipCheck(hipMalloc(dx,source=hx)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT' (Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + call rocfftCheck(rocfft_plan_create(plan,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + + call rocfftCheck(rocfft_execute(plan,c_loc(dx),c_null_ptr,c_null_ptr)) + + call hipCheck(hipDeviceSynchronize()) + + call rocfftCheck(rocfft_plan_destroy(plan)) + + call hipCheck(hipMemcpy(hx,dx,hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + do i = 1, int(N) + error = abs(hx(i) - x_ref(i)) + if (.not. (error <= error_max * sum(abs(x_in)))) then + write(*,*) "FAILED! Error = ", error, "hx(i)=",hx(i) + STOP 1 + end if + end do + + deallocate(hx, x_in, x_ref) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_example diff --git a/bindings/fortran/rocfft/test/rocfft_c2c_1d_batched_z.f03 b/bindings/fortran/rocfft/test/rocfft_c2c_1d_batched_z.f03 new file mode 100644 index 000000000000..31b544baeaf7 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_c2c_1d_batched_z.f03 @@ -0,0 +1,132 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_c2c_1d_batched_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Nbatch = 4 + integer(c_size_t), parameter :: Ntot = N * Nbatch + integer(c_size_t), parameter :: Nbytes = Ntot*8*2 + + type double2 + double precision :: x + double precision :: y + end type double2 + + type(double2), allocatable, target, dimension(:) :: hx, hx_input + type(c_ptr) :: dx = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: desc = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), target :: strides(1) + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 1D batched double (z)' (Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(1)) + lengths(1) = N + strides(1) = 1 + + allocate(hx(Ntot)) + allocate(hx_input(Ntot)) + do i = 1, Ntot + hx(i)%x = dble(i) + hx(i)%y = dble(Ntot - i) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, Nbytes)) + call hipCheck(hipMemcpy(dx, c_loc(hx(1)), Nbytes, hipMemcpyHostToDevice)) + + ! Describe a batched layout: contiguous, stride 1, distance N between transforms. + call rocfftCheck(rocfft_plan_description_create(desc)) + call rocfftCheck(rocfft_plan_description_set_data_layout(desc,& + rocfft_array_type_complex_interleaved,& + rocfft_array_type_complex_interleaved,& + c_null_ptr, c_null_ptr,& + one, c_loc(strides), N,& + one, c_loc(strides), N)) + + ! Forward batched transform (in-place). + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + Nbatch,& + desc)) + call rocfftCheck(rocfft_execute(plan_fwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse batched transform (in-place). rocFFT unnormalized -> N*input per transform. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + Nbatch,& + desc)) + call rocfftCheck(rocfft_execute(plan_bwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call rocfftCheck(rocfft_plan_description_destroy(desc)) + + call hipCheck(hipMemcpy(c_loc(hx(1)), dx, Nbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! After forward+inverse each batch element should equal N times the original input. + do i = 1, Ntot + error = abs(hx(i)%x - N * hx_input(i)%x) + abs(hx(i)%y - N * hx_input(i)%y) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_c2c_1d_batched_z diff --git a/bindings/fortran/rocfft/test/rocfft_c2c_1d_batched_z.f08 b/bindings/fortran/rocfft/test/rocfft_c2c_1d_batched_z.f08 new file mode 100644 index 000000000000..588a590a8f05 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_c2c_1d_batched_z.f08 @@ -0,0 +1,124 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_c2c_1d_batched_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Nbatch = 4 + integer(c_size_t), parameter :: Ntot = N * Nbatch + + complex(8), allocatable, target, dimension(:) :: hx, hx_input + complex(8), pointer, dimension(:) :: dx => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: desc = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), target :: strides(1) + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + allocate(lengths(1)) + lengths(1) = N + strides(1) = 1 + + allocate(hx(Ntot)) + allocate(hx_input(Ntot)) + do i = 1, Ntot + hx(i) = cmplx(dble(i), dble(Ntot - i), kind=8) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, source=hx)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 1D batched double (z)' (Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! Describe a batched layout: contiguous, stride 1, distance N between transforms. + call rocfftCheck(rocfft_plan_description_create(desc)) + call rocfftCheck(rocfft_plan_description_set_data_layout(desc,& + rocfft_array_type_complex_interleaved,& + rocfft_array_type_complex_interleaved,& + c_null_ptr, c_null_ptr,& + one, c_loc(strides), N,& + one, c_loc(strides), N)) + + ! Forward batched transform (in-place). + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + Nbatch,& + desc)) + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse batched transform (in-place). rocFFT unnormalized -> N*input per transform. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + Nbatch,& + desc)) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call rocfftCheck(rocfft_plan_description_destroy(desc)) + + call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! After forward+inverse each batch element should equal N times the original input. + do i = 1, Ntot + error = abs(hx(i) - N * hx_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), " expected=", N * hx_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_c2c_1d_batched_z diff --git a/bindings/fortran/rocfft/test/rocfft_c2c_1d_c.f03 b/bindings/fortran/rocfft/test/rocfft_c2c_1d_c.f03 new file mode 100644 index 000000000000..187e582e6242 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_c2c_1d_c.f03 @@ -0,0 +1,116 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_c2c_1d_c + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Nbytes = N*4*2 + + type float2 + real :: x + real :: y + end type float2 + + type(float2), allocatable, target, dimension(:) :: hx, hx_input + type(c_ptr) :: dx = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + real :: error + real, parameter :: error_max = 1.0e-4 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 1D single (c)' (Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i)%x = real(i) + hx(i)%y = real(N - i) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, Nbytes)) + call hipCheck(hipMemcpy(dx, c_loc(hx(1)), Nbytes, hipMemcpyHostToDevice)) + + ! Forward transform (in-place). + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_single,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse transform (in-place). rocFFT is unnormalized, so this yields N*input. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_single,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(c_loc(hx(1)), dx, Nbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! After forward+inverse the data should equal N times the original input. + do i = 1, N + error = abs(hx(i)%x - N * hx_input(i)%x) + abs(hx(i)%y - N * hx_input(i)%y) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_c2c_1d_c diff --git a/bindings/fortran/rocfft/test/rocfft_c2c_1d_c.f08 b/bindings/fortran/rocfft/test/rocfft_c2c_1d_c.f08 new file mode 100644 index 000000000000..639c8e1c79f6 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_c2c_1d_c.f08 @@ -0,0 +1,108 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_c2c_1d_c + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + + complex(4), allocatable, target, dimension(:) :: hx, hx_input + complex(4), pointer, dimension(:) :: dx => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + real :: error + real, parameter :: error_max = 1.0e-4 + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i) = cmplx(real(i), real(N - i), kind=4) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, source=hx)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 1D single (c)' (Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! Forward transform (in-place). + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_single,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse transform (in-place). rocFFT is unnormalized, so this yields N*input. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_single,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! After forward+inverse the data should equal N times the original input. + do i = 1, N + error = abs(hx(i) - N * hx_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), " expected=", N * hx_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_c2c_1d_c diff --git a/bindings/fortran/rocfft/test/rocfft_c2c_1d_notinplace_z.f03 b/bindings/fortran/rocfft/test/rocfft_c2c_1d_notinplace_z.f03 new file mode 100644 index 000000000000..a1febacc39ca --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_c2c_1d_notinplace_z.f03 @@ -0,0 +1,124 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Fortran counterpart of the out-of-place path of the rocFFT sample +! clients/samples/rocfft/rocfft_example_complexcomplex.cpp (--outofplace). +program rocfft_c2c_1d_notinplace_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Nbytes = N*8*2 + + type double2 + double precision :: x + double precision :: y + end type double2 + + type(double2), allocatable, target, dimension(:) :: hx, hx_input + type(c_ptr) :: dx = c_null_ptr ! device input + type(c_ptr) :: dy = c_null_ptr ! device output + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 1D out-of-place double (z)' & + &(Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i)%x = dble(i) + hx(i)%y = dble(N - i) + end do + hx_input(:) = hx(:) + + ! Out-of-place transforms need a separate output buffer of the same size. + call hipCheck(hipMalloc(dx, Nbytes)) + call hipCheck(hipMalloc(dy, Nbytes)) + call hipCheck(hipMemcpy(dx, c_loc(hx(1)), Nbytes, hipMemcpyHostToDevice)) + + ! Forward transform, out-of-place: dx -> dy. + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_notinplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, dx, dy, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse transform, out-of-place: dy -> dx. rocFFT is unnormalized, so the + ! round trip yields N*input. Note that rocFFT may overwrite the input buffer + ! of an out-of-place transform, so dy must be treated as clobbered afterwards. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_notinplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, dy, dx, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(c_loc(hx(1)), dx, Nbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + call hipCheck(hipFree(dy)) + + do i = 1, N + error = abs(hx(i)%x - N * hx_input(i)%x) + abs(hx(i)%y - N * hx_input(i)%y) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_c2c_1d_notinplace_z diff --git a/bindings/fortran/rocfft/test/rocfft_c2c_1d_notinplace_z.f08 b/bindings/fortran/rocfft/test/rocfft_c2c_1d_notinplace_z.f08 new file mode 100644 index 000000000000..2dec2fee2eb2 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_c2c_1d_notinplace_z.f08 @@ -0,0 +1,117 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Fortran counterpart of the out-of-place path of the rocFFT sample +! clients/samples/rocfft/rocfft_example_complexcomplex.cpp (--outofplace). +program rocfft_c2c_1d_notinplace_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + + complex(8), allocatable, target, dimension(:) :: hx, hx_input + complex(8), pointer, dimension(:) :: dx => null() ! device input + complex(8), pointer, dimension(:) :: dy => null() ! device output + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i) = cmplx(dble(i), dble(N - i), kind=8) + end do + hx_input(:) = hx(:) + + ! Out-of-place transforms need a separate output buffer of the same size. + call hipCheck(hipMalloc(dx, source=hx)) + call hipCheck(hipMalloc(dy, N)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 1D out-of-place double (z)' & + &(Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! Forward transform, out-of-place: dx -> dy. + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_notinplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_loc(dy), c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse transform, out-of-place: dy -> dx. rocFFT is unnormalized, so the + ! round trip yields N*input. Note that rocFFT may overwrite the input buffer + ! of an out-of-place transform, so dy must be treated as clobbered afterwards. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_notinplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dy), c_loc(dx), c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + call hipCheck(hipFree(dy)) + + do i = 1, N + error = abs(hx(i) - N * hx_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), & + " expected=", N * hx_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_c2c_1d_notinplace_z diff --git a/bindings/fortran/rocfft/test/rocfft_c2c_1d_z.f03 b/bindings/fortran/rocfft/test/rocfft_c2c_1d_z.f03 new file mode 100644 index 000000000000..3468f78ee148 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_c2c_1d_z.f03 @@ -0,0 +1,116 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_c2c_1d_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Nbytes = N*8*2 + + type double2 + double precision :: x + double precision :: y + end type double2 + + type(double2), allocatable, target, dimension(:) :: hx, hx_input + type(c_ptr) :: dx = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 1D double (z)' (Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i)%x = dble(i) + hx(i)%y = dble(N - i) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, Nbytes)) + call hipCheck(hipMemcpy(dx, c_loc(hx(1)), Nbytes, hipMemcpyHostToDevice)) + + ! Forward transform (in-place). + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse transform (in-place). rocFFT is unnormalized, so this yields N*input. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(c_loc(hx(1)), dx, Nbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! After forward+inverse the data should equal N times the original input. + do i = 1, N + error = abs(hx(i)%x - N * hx_input(i)%x) + abs(hx(i)%y - N * hx_input(i)%y) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_c2c_1d_z diff --git a/bindings/fortran/rocfft/test/rocfft_c2c_1d_z.f08 b/bindings/fortran/rocfft/test/rocfft_c2c_1d_z.f08 new file mode 100644 index 000000000000..ee9b6123cece --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_c2c_1d_z.f08 @@ -0,0 +1,108 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_c2c_1d_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + + complex(8), allocatable, target, dimension(:) :: hx, hx_input + complex(8), pointer, dimension(:) :: dx => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i) = cmplx(dble(i), dble(N - i), kind=8) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, source=hx)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 1D double (z)' (Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! Forward transform (in-place). + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse transform (in-place). rocFFT is unnormalized, so this yields N*input. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! After forward+inverse the data should equal N times the original input. + do i = 1, N + error = abs(hx(i) - N * hx_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), " expected=", N * hx_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_c2c_1d_z diff --git a/bindings/fortran/rocfft/test/rocfft_c2c_2d_z.f03 b/bindings/fortran/rocfft/test/rocfft_c2c_2d_z.f03 new file mode 100644 index 000000000000..6b963f9683e4 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_c2c_2d_z.f03 @@ -0,0 +1,119 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_c2c_2d_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: Nx = 4, Ny = 8 + integer(c_size_t), parameter :: Ntot = Nx * Ny + integer(c_size_t), parameter :: Nbytes = Ntot*8*2 + + type double2 + double precision :: x + double precision :: y + end type double2 + + type(double2), allocatable, target, dimension(:) :: hx, hx_input + type(c_ptr) :: dx = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: two = 2 + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 2D double (z)' (Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(2)) + lengths(1) = Nx + lengths(2) = Ny + + allocate(hx(Ntot)) + allocate(hx_input(Ntot)) + do i = 1, Ntot + hx(i)%x = dble(i) + hx(i)%y = dble(Ntot - i) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, Nbytes)) + call hipCheck(hipMemcpy(dx, c_loc(hx(1)), Nbytes, hipMemcpyHostToDevice)) + + ! Forward 2D transform (in-place). + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + two,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse 2D transform (in-place). rocFFT unnormalized -> Ntot*input. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + two,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(c_loc(hx(1)), dx, Nbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! After forward+inverse the data should equal Ntot times the original input. + do i = 1, Ntot + error = abs(hx(i)%x - Ntot * hx_input(i)%x) + abs(hx(i)%y - Ntot * hx_input(i)%y) + if (.not. (error <= error_max * Ntot)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_c2c_2d_z diff --git a/bindings/fortran/rocfft/test/rocfft_c2c_2d_z.f08 b/bindings/fortran/rocfft/test/rocfft_c2c_2d_z.f08 new file mode 100644 index 000000000000..11e35bd61182 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_c2c_2d_z.f08 @@ -0,0 +1,111 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_c2c_2d_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: Nx = 4, Ny = 8 + integer(c_size_t), parameter :: Ntot = Nx * Ny + + complex(8), allocatable, target, dimension(:) :: hx, hx_input + complex(8), pointer, dimension(:) :: dx => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: two = 2 + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + allocate(lengths(2)) + lengths(1) = Nx + lengths(2) = Ny + + allocate(hx(Ntot)) + allocate(hx_input(Ntot)) + do i = 1, Ntot + hx(i) = cmplx(dble(i), dble(Ntot - i), kind=8) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, source=hx)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 2D double (z)' (Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! Forward 2D transform (in-place). + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + two,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse 2D transform (in-place). rocFFT unnormalized -> Ntot*input. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + two,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! After forward+inverse the data should equal Ntot times the original input. + do i = 1, Ntot + error = abs(hx(i) - Ntot * hx_input(i)) + if (.not. (error <= error_max * Ntot)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), " expected=", Ntot * hx_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_c2c_2d_z diff --git a/bindings/fortran/rocfft/test/rocfft_c2c_3d_z.f03 b/bindings/fortran/rocfft/test/rocfft_c2c_3d_z.f03 new file mode 100644 index 000000000000..475a098b2484 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_c2c_3d_z.f03 @@ -0,0 +1,120 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_c2c_3d_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: Nx = 4, Ny = 4, Nz = 8 + integer(c_size_t), parameter :: Ntot = Nx * Ny * Nz + integer(c_size_t), parameter :: Nbytes = Ntot*8*2 + + type double2 + double precision :: x + double precision :: y + end type double2 + + type(double2), allocatable, target, dimension(:) :: hx, hx_input + type(c_ptr) :: dx = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: three = 3 + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 3D double (z)' (Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(3)) + lengths(1) = Nx + lengths(2) = Ny + lengths(3) = Nz + + allocate(hx(Ntot)) + allocate(hx_input(Ntot)) + do i = 1, Ntot + hx(i)%x = dble(i) + hx(i)%y = dble(Ntot - i) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, Nbytes)) + call hipCheck(hipMemcpy(dx, c_loc(hx(1)), Nbytes, hipMemcpyHostToDevice)) + + ! Forward 3D transform (in-place). + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + three,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse 3D transform (in-place). rocFFT unnormalized -> Ntot*input. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + three,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(c_loc(hx(1)), dx, Nbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! After forward+inverse the data should equal Ntot times the original input. + do i = 1, Ntot + error = abs(hx(i)%x - Ntot * hx_input(i)%x) + abs(hx(i)%y - Ntot * hx_input(i)%y) + if (.not. (error <= error_max * Ntot)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_c2c_3d_z diff --git a/bindings/fortran/rocfft/test/rocfft_c2c_3d_z.f08 b/bindings/fortran/rocfft/test/rocfft_c2c_3d_z.f08 new file mode 100644 index 000000000000..8ca6063b4328 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_c2c_3d_z.f08 @@ -0,0 +1,112 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_c2c_3d_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: Nx = 4, Ny = 4, Nz = 8 + integer(c_size_t), parameter :: Ntot = Nx * Ny * Nz + + complex(8), allocatable, target, dimension(:) :: hx, hx_input + complex(8), pointer, dimension(:) :: dx => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: three = 3 + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + allocate(lengths(3)) + lengths(1) = Nx + lengths(2) = Ny + lengths(3) = Nz + + allocate(hx(Ntot)) + allocate(hx_input(Ntot)) + do i = 1, Ntot + hx(i) = cmplx(dble(i), dble(Ntot - i), kind=8) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, source=hx)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 3D double (z)' (Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! Forward 3D transform (in-place). + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + three,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse 3D transform (in-place). rocFFT unnormalized -> Ntot*input. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + three,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! After forward+inverse the data should equal Ntot times the original input. + do i = 1, Ntot + error = abs(hx(i) - Ntot * hx_input(i)) + if (.not. (error <= error_max * Ntot)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), " expected=", Ntot * hx_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_c2c_3d_z diff --git a/bindings/fortran/rocfft/test/rocfft_cache_z.f03 b/bindings/fortran/rocfft/test/rocfft_cache_z.f03 new file mode 100644 index 000000000000..90c765e2e308 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_cache_z.f03 @@ -0,0 +1,137 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Serializes the runtime-compiled kernel cache and loads it back. +program rocfft_cache_z + use iso_c_binding + use hip + use rocfft + + implicit none + + ! A length that is not covered by the kernels built into rocFFT, so that plan + ! creation has to compile kernels at runtime and the cache is not empty. + integer(c_size_t), parameter :: N = 10007 + integer(c_size_t), parameter :: Nbytes = N*8*2 + + type double2 + double precision :: x + double precision :: y + end type double2 + + type(double2), allocatable, target, dimension(:) :: hx, hx_input + type(c_ptr) :: dx = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: cache = c_null_ptr + integer(c_size_t) :: cache_bytes = 0 + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT kernel cache double (z)' & + &(Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i)%x = dble(mod(i,11)) - 5.0d0 + hx(i)%y = dble(mod(i,7)) - 3.0d0 + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, Nbytes)) + call hipCheck(hipMemcpy(dx, c_loc(hx(1)), Nbytes, hipMemcpyHostToDevice)) + + ! Creating the plan populates the in-memory kernel cache. + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + + ! rocFFT allocates the buffer; its size comes back by reference. + call rocfftCheck(rocfft_cache_serialize(cache, cache_bytes)) + if (.not. c_associated(cache) .or. cache_bytes == 0) then + write(*,*) "FAILED! empty kernel cache" + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + + call rocfftCheck(rocfft_cache_deserialize(cache, cache_bytes)) + + ! The buffer belongs to the caller once it has been handed out. + call rocfftCheck(rocfft_cache_buffer_free(cache)) + cache = c_null_ptr + + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + + call rocfftCheck(rocfft_execute(plan_fwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_execute(plan_bwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(c_loc(hx(1)), dx, Nbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! rocFFT is unnormalized, so the round trip yields N*input. + do i = 1, N + error = abs(hx(i)%x - N * hx_input(i)%x) + abs(hx(i)%y - N * hx_input(i)%y) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED! kernel cache size: ", cache_bytes, " bytes" + +end program rocfft_cache_z diff --git a/bindings/fortran/rocfft/test/rocfft_cache_z.f08 b/bindings/fortran/rocfft/test/rocfft_cache_z.f08 new file mode 100644 index 000000000000..0fa1cf56c1f5 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_cache_z.f08 @@ -0,0 +1,137 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Demonstrates the compiled-kernel cache described in the rocFFT +! "Runtime compilation" guide: rocFFT compiles the kernels a plan needs when the +! plan is created, and the resulting cache can be serialized, moved and loaded +! back into another process. +program rocfft_cache_z + use iso_c_binding + use hip + use rocfft + + implicit none + + ! A length that is not covered by the kernels built into rocFFT, so that plan + ! creation has to compile kernels at runtime and the cache is not empty. + integer(c_size_t), parameter :: N = 10007 + + complex(8), allocatable, target, dimension(:) :: hx, hx_input + complex(8), pointer, dimension(:) :: dx => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: cache = c_null_ptr + integer(c_size_t), target :: cache_bytes = 0 + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i) = cmplx(dble(mod(i,11)) - 5.0d0, dble(mod(i,7)) - 3.0d0, kind=8) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, source=hx)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT kernel cache double (z)' & + &(Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! Creating the plan populates the in-memory kernel cache. + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + + ! Copy the cache into a buffer that rocFFT allocates. The buffer address is + ! written to the first argument, its size in bytes to the second. + call rocfftCheck(rocfft_cache_serialize(cache, cache_bytes)) + if (.not. c_associated(cache) .or. cache_bytes == 0) then + write(*,*) "FAILED! empty kernel cache" + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + + ! The buffer can be written to a file and loaded in another process. Loading + ! it back replaces matching kernels and leaves the rest of the cache alone, + ! so plans created afterwards reuse the kernels instead of recompiling them. + call rocfftCheck(rocfft_cache_deserialize(cache, cache_bytes)) + + ! The buffer belongs to the caller once it has been handed out. + call rocfftCheck(rocfft_cache_buffer_free(cache)) + cache = c_null_ptr + + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! rocFFT is unnormalized, so the round trip yields N*input. + do i = 1, N + error = abs(hx(i) - N * hx_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), & + " expected=", N * hx_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED! kernel cache size: ", cache_bytes, " bytes" + +end program rocfft_cache_z diff --git a/bindings/fortran/rocfft/test/rocfft_callback_z.f03 b/bindings/fortran/rocfft/test/rocfft_callback_z.f03 new file mode 100644 index 000000000000..a422ed612ee5 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_callback_z.f03 @@ -0,0 +1,148 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Fortran counterpart of clients/tests/default_callbacks_test.cpp: register the +! load and store callbacks of an execution info handle and transform with them. +! Fortran cannot take the address of a device function, so the callbacks are the +! null (default) ones, which rocFFT documents as clearing any registration. +program rocfft_callback_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 1024 + integer(c_size_t), parameter :: Nbytes = N*8*2 + + type double2 + double precision :: x + double precision :: y + end type double2 + + type(double2), allocatable, target, dimension(:) :: hx, hx_input + type(c_ptr) :: dx = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: info = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer(kind(rocfft_status_success)) :: stat + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT callbacks double (z)' & + &(Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i)%x = dble(mod(i,7)) - 3.0d0 + hx(i)%y = dble(mod(i,5)) - 2.0d0 + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, Nbytes)) + call hipCheck(hipMemcpy(dx, c_loc(hx(1)), Nbytes, hipMemcpyHostToDevice)) + + call rocfftCheck(rocfft_execution_info_create(info)) + + call rocfftCheck(rocfft_execution_info_set_load_callback(info, c_null_ptr, c_null_ptr, & + 0_c_size_t)) + call rocfftCheck(rocfft_execution_info_set_store_callback(info, c_null_ptr, c_null_ptr, & + 0_c_size_t)) + + ! rocFFT allocates no LDS for callbacks, so a shared memory request is rejected. + stat = rocfft_execution_info_set_load_callback(info, c_null_ptr, c_null_ptr, 16_c_size_t) + if (stat /= rocfft_status_invalid_arg_value) then + write(*,*) "FAILED! load callback with shared memory returned", stat, & + " expected", int(rocfft_status_invalid_arg_value) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + + stat = rocfft_execution_info_set_store_callback(info, c_null_ptr, c_null_ptr, 16_c_size_t) + if (stat /= rocfft_status_invalid_arg_value) then + write(*,*) "FAILED! store callback with shared memory returned", stat, & + " expected", int(rocfft_status_invalid_arg_value) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, dx, c_null_ptr, info)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, dx, c_null_ptr, info)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call rocfftCheck(rocfft_execution_info_destroy(info)) + + call hipCheck(hipMemcpy(c_loc(hx(1)), dx, Nbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! The default callbacks are plain loads and stores, so the unnormalized round + ! trip still yields N*input. + do i = 1, N + error = abs(hx(i)%x - N * hx_input(i)%x) + abs(hx(i)%y - N * hx_input(i)%y) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_callback_z diff --git a/bindings/fortran/rocfft/test/rocfft_callback_z.f08 b/bindings/fortran/rocfft/test/rocfft_callback_z.f08 new file mode 100644 index 000000000000..f35963663fcd --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_callback_z.f08 @@ -0,0 +1,139 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Fortran counterpart of clients/tests/default_callbacks_test.cpp: register the +! load and store callbacks of an execution info handle and transform with them. +! Fortran cannot take the address of a device function, so the callbacks are the +! null (default) ones, which rocFFT documents as clearing any registration. +program rocfft_callback_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 1024 + + complex(8), allocatable, target, dimension(:) :: hx, hx_input + complex(8), pointer, dimension(:) :: dx => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: info = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer(kind(rocfft_status_success)) :: stat + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT callbacks double (z)' & + &(Fortran 2008 interfaces) - " + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i) = cmplx(dble(mod(i,7)) - 3.0d0, dble(mod(i,5)) - 2.0d0, kind=8) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, source=hx)) + + call rocfftCheck(rocfft_setup()) + call rocfftCheck(rocfft_execution_info_create(info)) + + call rocfftCheck(rocfft_execution_info_set_load_callback(info, c_null_ptr, c_null_ptr, & + 0_c_size_t)) + call rocfftCheck(rocfft_execution_info_set_store_callback(info, c_null_ptr, c_null_ptr, & + 0_c_size_t)) + + ! rocFFT allocates no LDS for callbacks, so a shared memory request is rejected. + stat = rocfft_execution_info_set_load_callback(info, c_null_ptr, c_null_ptr, 16_c_size_t) + if (stat /= rocfft_status_invalid_arg_value) then + write(*,*) "FAILED! load callback with shared memory returned", stat, & + " expected", int(rocfft_status_invalid_arg_value) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + + stat = rocfft_execution_info_set_store_callback(info, c_null_ptr, c_null_ptr, 16_c_size_t) + if (stat /= rocfft_status_invalid_arg_value) then + write(*,*) "FAILED! store callback with shared memory returned", stat, & + " expected", int(rocfft_status_invalid_arg_value) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, info)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, info)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call rocfftCheck(rocfft_execution_info_destroy(info)) + + call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! The default callbacks are plain loads and stores, so the unnormalized round + ! trip still yields N*input. + do i = 1, N + error = abs(hx(i) - N * hx_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), " expected=", N * hx_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_callback_z diff --git a/bindings/fortran/rocfft/test/rocfft_field_brick_z.f03 b/bindings/fortran/rocfft/test/rocfft_field_brick_z.f03 new file mode 100644 index 000000000000..a0dd287070a4 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_field_brick_z.f03 @@ -0,0 +1,245 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_field_brick_z + use iso_c_binding + use hip + use rocfft + + implicit none + + ! 2D transform: N0 x N1, decomposed over 2 devices along dim-1 + integer(c_size_t), parameter :: N0 = 8, N1 = 8 + integer(c_size_t), parameter :: Ntot = N0 * N1 + integer(c_size_t), parameter :: nbytes_elem = 16 ! complex(8) = 16 bytes + integer(c_size_t), parameter :: dim_wb = 3 ! 2 transform dims + 1 batch + integer(c_size_t), parameter :: two = 2, one = 1 + + ! rocFFT handles + type(c_ptr) :: description = c_null_ptr + type(c_ptr) :: infield = c_null_ptr, outfield = c_null_ptr + type(c_ptr) :: abrick = c_null_ptr + type(c_ptr) :: gpu_plan = c_null_ptr + type(c_ptr) :: planinfo = c_null_ptr + + ! Per-brick GPU buffers and work buffers (2 bricks / 2 devices) + type(c_ptr) :: gpu_in(2), gpu_out(2), work_buf(2) + integer(c_size_t) :: wbuf_size + + ! Brick geometry (computed per brick in the loop) + integer(c_size_t) :: brick_len1, brick_lower1, brick_upper1 + integer(c_size_t), target :: blower(3), bupper(3), bstride(3) + + ! Plan lengths + integer(c_size_t), allocatable, target :: lengths(:) + + ! Host staging buffer + complex(kind=8), allocatable, target :: hbuf(:) + + ! Verification + double precision, parameter :: tol = 1.0d-8 + double precision :: pi2, err + integer(c_int) :: ndev + integer :: idx, ix, iy, pos + + gpu_in(1) = c_null_ptr ; gpu_in(2) = c_null_ptr + gpu_out(1) = c_null_ptr ; gpu_out(2) = c_null_ptr + work_buf(1) = c_null_ptr; work_buf(2) = c_null_ptr + + write(*,"(a)",advance="no") & + "-- Running test 'rocFFT field/brick 2D z' (Fortran 2003 interfaces) - " + + call hipCheck(hipGetDeviceCount(ndev)) + if (ndev < 2_c_int) then + write(*,*) "SKIPPED (requires 2 GPUs)" + stop 77 + end if + + ! The field/brick API is marked experimental in the rocFFT headers. + call rocfftCheck(rocfft_setup()) + + pi2 = 8.0d0 * atan(1.0d0) + + ! Plan description: null strides/offsets because the field defines the layout + call rocfftCheck(rocfft_plan_description_create(description)) + call rocfftCheck(rocfft_plan_description_set_data_layout( & + description, & + rocfft_array_type_complex_interleaved, & + rocfft_array_type_complex_interleaved, & + c_null_ptr, c_null_ptr, 0_c_size_t, c_null_ptr, 0_c_size_t, & + 0_c_size_t, c_null_ptr, 0_c_size_t)) + + ! Build input field: split N1 across 2 devices (C++ sample formula) + call rocfftCheck(rocfft_field_create(infield)) + do idx = 0, 1 + brick_len1 = N1 / 2 + & + merge(1_c_size_t, 0_c_size_t, int(idx,c_size_t) < mod(N1, 2_c_size_t)) + brick_lower1 = int(idx,c_size_t) * (N1/2) + & + min(int(idx,c_size_t), mod(N1, 2_c_size_t)) + brick_upper1 = brick_lower1 + brick_len1 + + blower = [0_c_size_t, brick_lower1, 0_c_size_t] + bupper = [N0, brick_upper1, 1_c_size_t] + bstride = [1_c_size_t, N0, brick_len1] + + ! rocfft_brick_create takes c_ptr,value for lower/upper/stride; no f2008 + ! rank-N wrapper exists, so c_loc is used in both dialect files. + call rocfftCheck(rocfft_brick_create(abrick, & + c_loc(blower(1)), c_loc(bupper(1)), c_loc(bstride(1)), & + dim_wb, int(idx, c_int))) + call rocfftCheck(rocfft_field_add_brick(infield, abrick)) + call rocfftCheck(rocfft_brick_destroy(abrick)) + abrick = c_null_ptr + + call hipCheck(hipSetDevice(int(idx, c_int))) + call hipCheck(hipMalloc(gpu_in(idx+1), N0 * brick_len1 * nbytes_elem)) + + ! Fill with known harmonic exp(2*pi*i*(1*ix/N0 + 2*iy/N1)) + allocate(hbuf(N0 * brick_len1)) + do iy = int(brick_lower1), int(brick_upper1) - 1 + do ix = 0, int(N0) - 1 + pos = ix + (iy - int(brick_lower1)) * int(N0) + 1 + hbuf(pos) = exp(cmplx(0.0d0, & + pi2*(dble(ix)/dble(N0) + 2.0d0*dble(iy)/dble(N1)), kind=8)) + end do + end do + call hipCheck(hipMemcpy(gpu_in(idx+1), c_loc(hbuf(1)), & + N0 * brick_len1 * nbytes_elem, hipMemcpyHostToDevice)) + deallocate(hbuf) + end do + + call rocfftCheck(rocfft_plan_description_add_infield(description, infield)) + call rocfftCheck(rocfft_field_destroy(infield)) + infield = c_null_ptr + + ! Build output field (same decomposition as input) + call rocfftCheck(rocfft_field_create(outfield)) + do idx = 0, 1 + brick_len1 = N1 / 2 + & + merge(1_c_size_t, 0_c_size_t, int(idx,c_size_t) < mod(N1, 2_c_size_t)) + brick_lower1 = int(idx,c_size_t) * (N1/2) + & + min(int(idx,c_size_t), mod(N1, 2_c_size_t)) + brick_upper1 = brick_lower1 + brick_len1 + + blower = [0_c_size_t, brick_lower1, 0_c_size_t] + bupper = [N0, brick_upper1, 1_c_size_t] + bstride = [1_c_size_t, N0, brick_len1] + + call rocfftCheck(rocfft_brick_create(abrick, & + c_loc(blower(1)), c_loc(bupper(1)), c_loc(bstride(1)), & + dim_wb, int(idx, c_int))) + call rocfftCheck(rocfft_field_add_brick(outfield, abrick)) + call rocfftCheck(rocfft_brick_destroy(abrick)) + abrick = c_null_ptr + + call hipCheck(hipSetDevice(int(idx, c_int))) + call hipCheck(hipMalloc(gpu_out(idx+1), N0 * brick_len1 * nbytes_elem)) + end do + + call rocfftCheck(rocfft_plan_description_add_outfield(description, outfield)) + call rocfftCheck(rocfft_field_destroy(outfield)) + outfield = c_null_ptr + + ! Create not-in-place forward double plan on device 0 + call hipCheck(hipSetDevice(0_c_int)) + allocate(lengths(2)) + lengths(1) = N0 + lengths(2) = N1 + call rocfftCheck(rocfft_plan_create(gpu_plan, & + rocfft_placement_notinplace, & + rocfft_transform_type_complex_forward, & + rocfft_precision_double, & + two, c_loc(lengths(1)), one, description)) + + ! Per-device work buffer allocation (mirrors C++ sample loop) + do idx = 0, 1 + call hipCheck(hipSetDevice(int(idx, c_int))) + wbuf_size = 0_c_size_t + call rocfftCheck(rocfft_plan_get_work_buffer_size(gpu_plan, wbuf_size)) + if (wbuf_size > 0_c_size_t) then + if (.not. c_associated(planinfo)) & + call rocfftCheck(rocfft_execution_info_create(planinfo)) + call hipCheck(hipMalloc(work_buf(idx+1), wbuf_size)) + call rocfftCheck(rocfft_execution_info_set_work_buffer( & + planinfo, work_buf(idx+1), wbuf_size)) + end if + end do + + ! Execute: in_buffer / out_buffer are arrays of per-brick device pointers + call hipCheck(hipSetDevice(0_c_int)) + call rocfftCheck(rocfft_execute(gpu_plan, gpu_in(1), gpu_out(1), planinfo)) + call hipCheck(hipDeviceSynchronize()) + + ! Verify: forward FFT of exp(2pi*i*(1*ix/N0 + 2*iy/N1)) must give Ntot at + ! bin (kx=1, ky=2) and zero at every other bin. + do idx = 0, 1 + brick_len1 = N1 / 2 + & + merge(1_c_size_t, 0_c_size_t, int(idx,c_size_t) < mod(N1, 2_c_size_t)) + brick_lower1 = int(idx,c_size_t) * (N1/2) + & + min(int(idx,c_size_t), mod(N1, 2_c_size_t)) + brick_upper1 = brick_lower1 + brick_len1 + + allocate(hbuf(N0 * brick_len1)) + call hipCheck(hipSetDevice(int(idx, c_int))) + call hipCheck(hipMemcpy(c_loc(hbuf(1)), gpu_out(idx+1), & + N0 * brick_len1 * nbytes_elem, hipMemcpyDeviceToHost)) + do iy = int(brick_lower1), int(brick_upper1) - 1 + do ix = 0, int(N0) - 1 + pos = ix + (iy - int(brick_lower1)) * int(N0) + 1 + if (ix == 1 .and. iy == 2) then + err = abs(hbuf(pos) - cmplx(dble(Ntot), 0.0d0, kind=8)) + else + err = abs(hbuf(pos)) + end if + if (.not. (err <= tol * dble(Ntot))) then + write(*,*) "FAILED! bin=(", ix, ",", iy, & + ") got=", hbuf(pos), " err=", err + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + end do + deallocate(hbuf) + end do + + if (c_associated(planinfo)) & + call rocfftCheck(rocfft_execution_info_destroy(planinfo)) + call rocfftCheck(rocfft_plan_description_destroy(description)) + call rocfftCheck(rocfft_plan_destroy(gpu_plan)) + call rocfftCheck(rocfft_cleanup()) + + do idx = 0, 1 + call hipCheck(hipSetDevice(int(idx, c_int))) + if (c_associated(work_buf(idx+1))) & + call hipCheck(hipFree(work_buf(idx+1))) + call hipCheck(hipFree(gpu_in(idx+1))) + call hipCheck(hipFree(gpu_out(idx+1))) + end do + + deallocate(lengths) + + write(*,*) "PASSED!" + +end program rocfft_field_brick_z diff --git a/bindings/fortran/rocfft/test/rocfft_field_brick_z.f08 b/bindings/fortran/rocfft/test/rocfft_field_brick_z.f08 new file mode 100644 index 000000000000..a1aa1081f5b9 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_field_brick_z.f08 @@ -0,0 +1,249 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_field_brick_z + use iso_c_binding + use hip + use rocfft + + implicit none + + ! 2D transform: N0 x N1, decomposed over 2 devices along dim-1 + integer(c_size_t), parameter :: N0 = 8, N1 = 8 + integer(c_size_t), parameter :: Ntot = N0 * N1 + integer(c_size_t), parameter :: nbytes_elem = 16 ! complex(8) = 16 bytes + integer(c_size_t), parameter :: dim_wb = 3 ! 2 transform dims + 1 batch + integer(c_size_t), parameter :: two = 2, one = 1 + + ! rocFFT handles + type(c_ptr) :: description = c_null_ptr + type(c_ptr) :: infield = c_null_ptr, outfield = c_null_ptr + type(c_ptr) :: abrick = c_null_ptr + type(c_ptr) :: gpu_plan = c_null_ptr + type(c_ptr) :: planinfo = c_null_ptr + + ! Per-brick GPU buffers and work buffers (2 bricks / 2 devices) + type(c_ptr) :: gpu_in(2), gpu_out(2), work_buf(2) + integer(c_size_t) :: wbuf_size + + ! Brick geometry (computed per brick in the loop) + integer(c_size_t) :: brick_len1, brick_lower1, brick_upper1 + integer(c_size_t), target :: blower(3), bupper(3), bstride(3) + + ! Plan lengths + integer(c_size_t), allocatable, target :: lengths(:) + + ! Host staging buffer + complex(8), allocatable, target :: hbuf(:) + + ! Verification + double precision, parameter :: tol = 1.0d-8 + double precision :: pi2, err + integer(c_int) :: ndev + integer :: idx, ix, iy, pos + + gpu_in(1) = c_null_ptr ; gpu_in(2) = c_null_ptr + gpu_out(1) = c_null_ptr ; gpu_out(2) = c_null_ptr + work_buf(1) = c_null_ptr; work_buf(2) = c_null_ptr + + write(*,"(a)",advance="no") & + "-- Running test 'rocFFT field/brick 2D z' (Fortran 2008 interfaces) - " + + ! Mandatory guard: need at least 2 GPUs + call hipCheck(hipGetDeviceCount(ndev)) + ! Exit 77, which CTest reports as skipped: nothing was checked. + if (ndev < 2_c_int) then + write(*,*) "SKIPPED (requires 2 GPUs)" + stop 77 + end if + + ! The field/brick API is marked experimental in the rocFFT headers. + call rocfftCheck(rocfft_setup()) + + pi2 = 8.0d0 * atan(1.0d0) + + ! Plan description: null strides/offsets because the field defines the layout + call rocfftCheck(rocfft_plan_description_create(description)) + call rocfftCheck(rocfft_plan_description_set_data_layout( & + description, & + rocfft_array_type_complex_interleaved, & + rocfft_array_type_complex_interleaved, & + c_null_ptr, c_null_ptr, 0_c_size_t, c_null_ptr, 0_c_size_t, & + 0_c_size_t, c_null_ptr, 0_c_size_t)) + + ! Build input field: split N1 across 2 devices (C++ sample formula) + call rocfftCheck(rocfft_field_create(infield)) + do idx = 0, 1 + brick_len1 = N1 / 2 + & + merge(1_c_size_t, 0_c_size_t, int(idx,c_size_t) < mod(N1, 2_c_size_t)) + brick_lower1 = int(idx,c_size_t) * (N1/2) + & + min(int(idx,c_size_t), mod(N1, 2_c_size_t)) + brick_upper1 = brick_lower1 + brick_len1 + + blower = [0_c_size_t, brick_lower1, 0_c_size_t] + bupper = [N0, brick_upper1, 1_c_size_t] + bstride = [1_c_size_t, N0, brick_len1] + + ! rocfft_brick_create has no rank-N array wrapper; + ! c_loc is used for the coordinate arrays in both dialect files. + call rocfftCheck(rocfft_brick_create(abrick, & + c_loc(blower(1)), c_loc(bupper(1)), c_loc(bstride(1)), & + dim_wb, int(idx, c_int))) + call rocfftCheck(rocfft_field_add_brick(infield, abrick)) + call rocfftCheck(rocfft_brick_destroy(abrick)) + abrick = c_null_ptr + + call hipCheck(hipSetDevice(int(idx, c_int))) + call hipCheck(hipMalloc(gpu_in(idx+1), N0 * brick_len1 * nbytes_elem)) + + ! Fill with known harmonic exp(2*pi*i*(1*ix/N0 + 2*iy/N1)) + allocate(hbuf(N0 * brick_len1)) + do iy = int(brick_lower1), int(brick_upper1) - 1 + do ix = 0, int(N0) - 1 + pos = ix + (iy - int(brick_lower1)) * int(N0) + 1 + hbuf(pos) = exp(cmplx(0.0d0, & + pi2*(dble(ix)/dble(N0) + 2.0d0*dble(iy)/dble(N1)), kind=8)) + end do + end do + call hipCheck(hipMemcpy(gpu_in(idx+1), c_loc(hbuf(1)), & + N0 * brick_len1 * nbytes_elem, hipMemcpyHostToDevice)) + deallocate(hbuf) + end do + + call rocfftCheck(rocfft_plan_description_add_infield(description, infield)) + call rocfftCheck(rocfft_field_destroy(infield)) + infield = c_null_ptr + + ! Build output field (same decomposition as input) + call rocfftCheck(rocfft_field_create(outfield)) + do idx = 0, 1 + brick_len1 = N1 / 2 + & + merge(1_c_size_t, 0_c_size_t, int(idx,c_size_t) < mod(N1, 2_c_size_t)) + brick_lower1 = int(idx,c_size_t) * (N1/2) + & + min(int(idx,c_size_t), mod(N1, 2_c_size_t)) + brick_upper1 = brick_lower1 + brick_len1 + + blower = [0_c_size_t, brick_lower1, 0_c_size_t] + bupper = [N0, brick_upper1, 1_c_size_t] + bstride = [1_c_size_t, N0, brick_len1] + + call rocfftCheck(rocfft_brick_create(abrick, & + c_loc(blower(1)), c_loc(bupper(1)), c_loc(bstride(1)), & + dim_wb, int(idx, c_int))) + call rocfftCheck(rocfft_field_add_brick(outfield, abrick)) + call rocfftCheck(rocfft_brick_destroy(abrick)) + abrick = c_null_ptr + + call hipCheck(hipSetDevice(int(idx, c_int))) + call hipCheck(hipMalloc(gpu_out(idx+1), N0 * brick_len1 * nbytes_elem)) + end do + + call rocfftCheck(rocfft_plan_description_add_outfield(description, outfield)) + call rocfftCheck(rocfft_field_destroy(outfield)) + outfield = c_null_ptr + + ! Create not-in-place forward double plan on device 0 + call hipCheck(hipSetDevice(0_c_int)) + allocate(lengths(2)) + lengths(1) = N0 + lengths(2) = N1 + ! f2008: c_loc on the whole rank-1 array (rank_1 wrapper uses c_loc(lengths)) + call rocfftCheck(rocfft_plan_create(gpu_plan, & + rocfft_placement_notinplace, & + rocfft_transform_type_complex_forward, & + rocfft_precision_double, & + two, c_loc(lengths), one, description)) + + ! Per-device work buffer allocation (mirrors C++ sample loop) + do idx = 0, 1 + call hipCheck(hipSetDevice(int(idx, c_int))) + wbuf_size = 0_c_size_t + call rocfftCheck(rocfft_plan_get_work_buffer_size(gpu_plan, wbuf_size)) + if (wbuf_size > 0_c_size_t) then + if (.not. c_associated(planinfo)) & + call rocfftCheck(rocfft_execution_info_create(planinfo)) + call hipCheck(hipMalloc(work_buf(idx+1), wbuf_size)) + call rocfftCheck(rocfft_execution_info_set_work_buffer( & + planinfo, work_buf(idx+1), wbuf_size)) + end if + end do + + ! Execute: in_buffer / out_buffer are arrays of per-brick device pointers + call hipCheck(hipSetDevice(0_c_int)) + call rocfftCheck(rocfft_execute(gpu_plan, gpu_in(1), gpu_out(1), planinfo)) + call hipCheck(hipDeviceSynchronize()) + + ! Verify: forward FFT of exp(2pi*i*(1*ix/N0 + 2*iy/N1)) must give Ntot at + ! bin (kx=1, ky=2) and zero at every other bin. + do idx = 0, 1 + brick_len1 = N1 / 2 + & + merge(1_c_size_t, 0_c_size_t, int(idx,c_size_t) < mod(N1, 2_c_size_t)) + brick_lower1 = int(idx,c_size_t) * (N1/2) + & + min(int(idx,c_size_t), mod(N1, 2_c_size_t)) + brick_upper1 = brick_lower1 + brick_len1 + + allocate(hbuf(N0 * brick_len1)) + call hipCheck(hipSetDevice(int(idx, c_int))) + call hipCheck(hipMemcpy(c_loc(hbuf(1)), gpu_out(idx+1), & + N0 * brick_len1 * nbytes_elem, hipMemcpyDeviceToHost)) + do iy = int(brick_lower1), int(brick_upper1) - 1 + do ix = 0, int(N0) - 1 + pos = ix + (iy - int(brick_lower1)) * int(N0) + 1 + if (ix == 1 .and. iy == 2) then + err = abs(hbuf(pos) - cmplx(dble(Ntot), 0.0d0, kind=8)) + else + err = abs(hbuf(pos)) + end if + if (.not. (err <= tol * dble(Ntot))) then + write(*,*) "FAILED! bin=(", ix, ",", iy, & + ") got=", hbuf(pos), " err=", err + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + end do + deallocate(hbuf) + end do + + ! Teardown in sample order: info, description, plan, cleanup, buffers + if (c_associated(planinfo)) & + call rocfftCheck(rocfft_execution_info_destroy(planinfo)) + call rocfftCheck(rocfft_plan_description_destroy(description)) + call rocfftCheck(rocfft_plan_destroy(gpu_plan)) + call rocfftCheck(rocfft_cleanup()) + + do idx = 0, 1 + call hipCheck(hipSetDevice(int(idx, c_int))) + if (c_associated(work_buf(idx+1))) & + call hipCheck(hipFree(work_buf(idx+1))) + call hipCheck(hipFree(gpu_in(idx+1))) + call hipCheck(hipFree(gpu_out(idx+1))) + end do + + deallocate(lengths) + + write(*,*) "PASSED!" + +end program rocfft_field_brick_z diff --git a/bindings/fortran/rocfft/test/rocfft_plan_print_z.f03 b/bindings/fortran/rocfft/test/rocfft_plan_print_z.f03 new file mode 100644 index 000000000000..23a4a0e8c656 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_plan_print_z.f03 @@ -0,0 +1,136 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! rocfft_plan_get_print writes a summary of the plan to stdout. +program rocfft_plan_print_z + use iso_c_binding + use iso_fortran_env, only: output_unit + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: Nx = 4 + integer(c_size_t), parameter :: Ny = 6 + integer(c_size_t), parameter :: N = Nx*Ny + integer(c_size_t), parameter :: Nbytes = N*8*2 + + type double2 + double precision :: x + double precision :: y + end type double2 + + type(double2), allocatable, target, dimension(:,:) :: hx, hx_input + type(c_ptr) :: dx = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer(c_size_t), parameter :: two = 2 + integer :: i, j + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + call rocfftCheck(rocfft_setup()) + + ! lengths starts with the fastest-varying dimension, matching Fortran's + ! column-major storage. + allocate(lengths(2)) + lengths(1) = Nx + lengths(2) = Ny + + allocate(hx(Nx,Ny)) + allocate(hx_input(Nx,Ny)) + do j = 1, Ny + do i = 1, Nx + hx(i,j)%x = dble(i) + hx(i,j)%y = dble(j) + end do + end do + hx_input(:,:) = hx(:,:) + + call hipCheck(hipMalloc(dx, Nbytes)) + call hipCheck(hipMemcpy(dx, c_loc(hx(1,1)), Nbytes, hipMemcpyHostToDevice)) + + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + two,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + + ! rocfft_plan_get_print writes to stdout from C, so flush the Fortran unit + ! first to keep the two output streams in order. + write(*,"(a)") "-- rocfft_plan_get_print output for a 2D complex plan:" + flush(output_unit) + call rocfftCheck(rocfft_plan_get_print(plan_fwd)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT plan print double (z)' & + &(Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + two,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + + call rocfftCheck(rocfft_execute(plan_fwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_execute(plan_bwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(c_loc(hx(1,1)), dx, Nbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! rocFFT is unnormalized, so the round trip yields Nx*Ny*input. + do j = 1, Ny + do i = 1, Nx + error = abs(hx(i,j)%x - N * hx_input(i,j)%x) & + + abs(hx(i,j)%y - N * hx_input(i,j)%y) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " j=", j, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_plan_print_z diff --git a/bindings/fortran/rocfft/test/rocfft_plan_print_z.f08 b/bindings/fortran/rocfft/test/rocfft_plan_print_z.f08 new file mode 100644 index 000000000000..d07574a20ac1 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_plan_print_z.f08 @@ -0,0 +1,131 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Demonstrates rocfft_plan_get_print, the plan introspection helper described in +! the rocFFT "Working with rocFFT" guide. It writes a human-readable summary of +! a plan to stdout, which is the quickest way to check how rocFFT interpreted +! the lengths, strides and distances of a plan description. +program rocfft_plan_print_z + use iso_c_binding + use iso_fortran_env, only: output_unit + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: Nx = 4 + integer(c_size_t), parameter :: Ny = 6 + integer(c_size_t), parameter :: N = Nx*Ny + + complex(8), allocatable, target, dimension(:,:) :: hx, hx_input + complex(8), pointer, dimension(:,:) :: dx => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer(c_size_t), parameter :: two = 2 + integer :: i, j + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + ! lengths starts with the fastest-varying dimension, matching Fortran's + ! column-major storage. + allocate(lengths(2)) + lengths(1) = Nx + lengths(2) = Ny + + allocate(hx(Nx,Ny)) + allocate(hx_input(Nx,Ny)) + do j = 1, Ny + do i = 1, Nx + hx(i,j) = cmplx(dble(i), dble(j), kind=8) + end do + end do + hx_input(:,:) = hx(:,:) + + call hipCheck(hipMalloc(dx, source=hx)) + + call rocfftCheck(rocfft_setup()) + + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + two,& + c_loc(lengths),& + one,& + c_null_ptr)) + + ! rocfft_plan_get_print writes to stdout from C, so flush the Fortran unit + ! first to keep the two output streams in order. + write(*,"(a)") "-- rocfft_plan_get_print output for a 2D complex plan:" + flush(output_unit) + call rocfftCheck(rocfft_plan_get_print(plan_fwd)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT plan print double (z)' & + &(Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + two,& + c_loc(lengths),& + one,& + c_null_ptr)) + + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! rocFFT is unnormalized, so the round trip yields Nx*Ny*input. + do j = 1, Ny + do i = 1, Nx + error = abs(hx(i,j) - N * hx_input(i,j)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " j=", j, " error=", error, & + " got=", hx(i,j), " expected=", N * hx_input(i,j) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_plan_print_z diff --git a/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_d.f03 b/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_d.f03 new file mode 100644 index 000000000000..93d144c1213c --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_d.f03 @@ -0,0 +1,115 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_r2c_c2r_1d_d + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Ncomplex = N/2 + 1 + integer(c_size_t), parameter :: Rbytes = N*8 + integer(c_size_t), parameter :: Cbytes = Ncomplex*8*2 + + real(8), allocatable, target, dimension(:) :: hr, hr_input + type(c_ptr) :: dr = c_null_ptr + type(c_ptr) :: dc = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT R2C/C2R 1D double (d)' (Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(1)) + lengths(1) = N + + allocate(hr(N)) + allocate(hr_input(N)) + do i = 1, N + hr(i) = dble(i) + dble(mod(i,3)) - dble(mod(i,7)) + end do + hr_input(:) = hr(:) + + call hipCheck(hipMalloc(dr, Rbytes)) + call hipCheck(hipMalloc(dc, Cbytes)) + call hipCheck(hipMemcpy(dr, c_loc(hr(1)), Rbytes, hipMemcpyHostToDevice)) + + ! Forward real-to-complex (out-of-place): dr -> dc. + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_notinplace,& + rocfft_transform_type_real_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, dr, dc, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse complex-to-real (out-of-place): dc -> dr. rocFFT unnormalized -> N*input. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_notinplace,& + rocfft_transform_type_real_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, dc, dr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(c_loc(hr(1)), dr, Rbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dr)) + call hipCheck(hipFree(dc)) + + ! After forward+inverse the real data should equal N times the original input. + do i = 1, N + error = abs(hr(i) - N * hr_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hr) + deallocate(hr_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_r2c_c2r_1d_d diff --git a/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_d.f08 b/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_d.f08 new file mode 100644 index 000000000000..0d6419ee6b08 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_d.f08 @@ -0,0 +1,113 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_r2c_c2r_1d_d + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Ncomplex = N/2 + 1 + + real(8), allocatable, target, dimension(:) :: hr, hr_input + complex(8), pointer, dimension(:) :: dc => null() + real(8), pointer, dimension(:) :: dr => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + allocate(lengths(1)) + lengths(1) = N + + allocate(hr(N)) + allocate(hr_input(N)) + do i = 1, N + hr(i) = dble(i) + dble(mod(i,3)) - dble(mod(i,7)) + end do + hr_input(:) = hr(:) + + ! Device buffers: real input of length N, complex output of length N/2+1. + call hipCheck(hipMalloc(dr, source=hr)) + call hipCheck(hipMalloc(dc, Ncomplex)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT R2C/C2R 1D double (d)' (Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! Forward real-to-complex (out-of-place): dr -> dc. + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_notinplace,& + rocfft_transform_type_real_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dr), c_loc(dc), c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse complex-to-real (out-of-place): dc -> dr. rocFFT unnormalized -> N*input. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_notinplace,& + rocfft_transform_type_real_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dc), c_loc(dr), c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(hr, dr, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dr)) + call hipCheck(hipFree(dc)) + + ! After forward+inverse the real data should equal N times the original input. + do i = 1, N + error = abs(hr(i) - N * hr_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hr(i), " expected=", N * hr_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hr) + deallocate(hr_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_r2c_c2r_1d_d diff --git a/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_inplace_d.f03 b/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_inplace_d.f03 new file mode 100644 index 000000000000..f037858466aa --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_inplace_d.f03 @@ -0,0 +1,142 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Fortran counterpart of the in-place path of the rocFFT sample +! clients/samples/rocfft/rocfft_example_realcomplex.cpp: an in-place real +! transform, where the real buffer is padded so that it can also hold the +! Hermitian-symmetric spectrum. +program rocfft_r2c_c2r_1d_inplace_d + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Ncomplex = N/2 + 1 + ! An in-place real transform needs room for N/2+1 complex values, that is + ! two extra reals in the contiguous dimension. + integer(c_size_t), parameter :: Npad = 2 * Ncomplex + integer(c_size_t), parameter :: Nbytes = Npad*8 + + real(8), allocatable, target, dimension(:) :: hr, hr_input + type(c_ptr) :: dr = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: desc_fwd = c_null_ptr + type(c_ptr) :: desc_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), target :: strides(1) + integer(c_size_t), parameter :: one = 1 + integer(c_size_t), parameter :: zero = 0 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT R2C/C2R 1D in-place double (d)' & + &(Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(1)) + lengths(1) = N + strides(1) = 1 + + allocate(hr(Npad)) + allocate(hr_input(Npad)) + hr(:) = 0.0d0 + do i = 1, N + hr(i) = dble(i) + dble(mod(i,3)) - dble(mod(i,7)) + end do + hr_input(:) = hr(:) + + call hipCheck(hipMalloc(dr, Nbytes)) + call hipCheck(hipMemcpy(dr, c_loc(hr(1)), Nbytes, hipMemcpyHostToDevice)) + + ! Forward: real input, Hermitian-interleaved output, both in the same buffer. + call rocfftCheck(rocfft_plan_description_create(desc_fwd)) + call rocfftCheck(rocfft_plan_description_set_data_layout(desc_fwd,& + rocfft_array_type_real,& + rocfft_array_type_hermitian_interleaved,& + c_null_ptr, c_null_ptr,& + one, c_loc(strides(1)), zero,& + one, c_loc(strides(1)), zero)) + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_real_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + desc_fwd)) + call rocfftCheck(rocfft_execute(plan_fwd, dr, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + call rocfftCheck(rocfft_plan_description_destroy(desc_fwd)) + + ! Inverse: the array types are swapped. rocFFT is unnormalized, so the round + ! trip yields N*input. + call rocfftCheck(rocfft_plan_description_create(desc_bwd)) + call rocfftCheck(rocfft_plan_description_set_data_layout(desc_bwd,& + rocfft_array_type_hermitian_interleaved,& + rocfft_array_type_real,& + c_null_ptr, c_null_ptr,& + one, c_loc(strides(1)), zero,& + one, c_loc(strides(1)), zero)) + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_real_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + desc_bwd)) + call rocfftCheck(rocfft_execute(plan_bwd, dr, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + call rocfftCheck(rocfft_plan_description_destroy(desc_bwd)) + + call hipCheck(hipMemcpy(c_loc(hr(1)), dr, Nbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dr)) + + ! Only the first N reals hold the signal; the padding is scratch space. + do i = 1, N + error = abs(hr(i) - N * hr_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hr) + deallocate(hr_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_r2c_c2r_1d_inplace_d diff --git a/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_inplace_d.f08 b/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_inplace_d.f08 new file mode 100644 index 000000000000..3193f62f0c2f --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_inplace_d.f08 @@ -0,0 +1,141 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Fortran counterpart of the in-place path of the rocFFT sample +! clients/samples/rocfft/rocfft_example_realcomplex.cpp: an in-place real +! transform, where the real buffer is padded so that it can also hold the +! Hermitian-symmetric spectrum. +program rocfft_r2c_c2r_1d_inplace_d + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Ncomplex = N/2 + 1 + ! An in-place real transform needs room for N/2+1 complex values, that is + ! two extra reals in the contiguous dimension. + integer(c_size_t), parameter :: Npad = 2 * Ncomplex + + real(8), allocatable, target, dimension(:) :: hr, hr_input + real(8), pointer, dimension(:) :: dr => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: desc_fwd = c_null_ptr + type(c_ptr) :: desc_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), target :: strides(1) + integer(c_size_t), parameter :: one = 1 + integer(c_size_t), parameter :: zero = 0 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + allocate(lengths(1)) + lengths(1) = N + strides(1) = 1 + + allocate(hr(Npad)) + allocate(hr_input(Npad)) + hr(:) = 0.0d0 + do i = 1, N + hr(i) = dble(i) + dble(mod(i,3)) - dble(mod(i,7)) + end do + hr_input(:) = hr(:) + + call hipCheck(hipMalloc(dr, source=hr)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT R2C/C2R 1D in-place double (d)' & + &(Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! Forward: real input, Hermitian-interleaved output, both in the same buffer. + call rocfftCheck(rocfft_plan_description_create(desc_fwd)) + call rocfftCheck(rocfft_plan_description_set_data_layout(desc_fwd,& + rocfft_array_type_real,& + rocfft_array_type_hermitian_interleaved,& + c_null_ptr, c_null_ptr,& + one, c_loc(strides), zero,& + one, c_loc(strides), zero)) + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_real_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + desc_fwd)) + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dr), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + call rocfftCheck(rocfft_plan_description_destroy(desc_fwd)) + + ! Inverse: the array types are swapped. rocFFT is unnormalized, so the round + ! trip yields N*input. + call rocfftCheck(rocfft_plan_description_create(desc_bwd)) + call rocfftCheck(rocfft_plan_description_set_data_layout(desc_bwd,& + rocfft_array_type_hermitian_interleaved,& + rocfft_array_type_real,& + c_null_ptr, c_null_ptr,& + one, c_loc(strides), zero,& + one, c_loc(strides), zero)) + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_real_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + desc_bwd)) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dr), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + call rocfftCheck(rocfft_plan_description_destroy(desc_bwd)) + + call hipCheck(hipMemcpy(hr, dr, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dr)) + + ! Only the first N reals hold the signal; the padding is scratch space. + do i = 1, N + error = abs(hr(i) - N * hr_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hr(i), & + " expected=", N * hr_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hr) + deallocate(hr_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_r2c_c2r_1d_inplace_d diff --git a/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_s.f03 b/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_s.f03 new file mode 100644 index 000000000000..db3c5db136db --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_s.f03 @@ -0,0 +1,115 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_r2c_c2r_1d_s + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Ncomplex = N/2 + 1 + integer(c_size_t), parameter :: Rbytes = N*4 + integer(c_size_t), parameter :: Cbytes = Ncomplex*4*2 + + real(4), allocatable, target, dimension(:) :: hr, hr_input + type(c_ptr) :: dr = c_null_ptr + type(c_ptr) :: dc = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + real :: error + real, parameter :: error_max = 1.0e-4 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT R2C/C2R 1D single (s)' (Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(1)) + lengths(1) = N + + allocate(hr(N)) + allocate(hr_input(N)) + do i = 1, N + hr(i) = real(i) + real(mod(i,3)) - real(mod(i,7)) + end do + hr_input(:) = hr(:) + + call hipCheck(hipMalloc(dr, Rbytes)) + call hipCheck(hipMalloc(dc, Cbytes)) + call hipCheck(hipMemcpy(dr, c_loc(hr(1)), Rbytes, hipMemcpyHostToDevice)) + + ! Forward real-to-complex (out-of-place): dr -> dc. + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_notinplace,& + rocfft_transform_type_real_forward,& + rocfft_precision_single,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, dr, dc, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse complex-to-real (out-of-place): dc -> dr. rocFFT unnormalized -> N*input. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_notinplace,& + rocfft_transform_type_real_inverse,& + rocfft_precision_single,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, dc, dr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(c_loc(hr(1)), dr, Rbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dr)) + call hipCheck(hipFree(dc)) + + ! After forward+inverse the real data should equal N times the original input. + do i = 1, N + error = abs(hr(i) - N * hr_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hr) + deallocate(hr_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_r2c_c2r_1d_s diff --git a/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_s.f08 b/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_s.f08 new file mode 100644 index 000000000000..3531911bbd3b --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_r2c_c2r_1d_s.f08 @@ -0,0 +1,113 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +program rocfft_r2c_c2r_1d_s + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Ncomplex = N/2 + 1 + + real(4), allocatable, target, dimension(:) :: hr, hr_input + complex(4), pointer, dimension(:) :: dc => null() + real(4), pointer, dimension(:) :: dr => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + real :: error + real, parameter :: error_max = 1.0e-4 + + allocate(lengths(1)) + lengths(1) = N + + allocate(hr(N)) + allocate(hr_input(N)) + do i = 1, N + hr(i) = real(i) + real(mod(i,3)) - real(mod(i,7)) + end do + hr_input(:) = hr(:) + + ! Device buffers: real input of length N, complex output of length N/2+1. + call hipCheck(hipMalloc(dr, source=hr)) + call hipCheck(hipMalloc(dc, Ncomplex)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT R2C/C2R 1D single (s)' (Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! Forward real-to-complex (out-of-place): dr -> dc. + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_notinplace,& + rocfft_transform_type_real_forward,& + rocfft_precision_single,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dr), c_loc(dc), c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! Inverse complex-to-real (out-of-place): dc -> dr. rocFFT unnormalized -> N*input. + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_notinplace,& + rocfft_transform_type_real_inverse,& + rocfft_precision_single,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dc), c_loc(dr), c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(hr, dr, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dr)) + call hipCheck(hipFree(dc)) + + ! After forward+inverse the real data should equal N times the original input. + do i = 1, N + error = abs(hr(i) - N * hr_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hr(i), " expected=", N * hr_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hr) + deallocate(hr_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_r2c_c2r_1d_s diff --git a/bindings/fortran/rocfft/test/rocfft_scale_factor_z.f03 b/bindings/fortran/rocfft/test/rocfft_scale_factor_z.f03 new file mode 100644 index 000000000000..0b71249b0e08 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_scale_factor_z.f03 @@ -0,0 +1,126 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. +! +! SPDX-License-Identifier: MIT +! +! Permission is hereby granted, free of charge, to any person obtaining a copy +! of this software and associated documentation files (the "Software"), to deal +! in the Software without restriction, including without limitation the rights +! to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +! copies of the Software, and to permit persons to whom the Software is +! furnished to do so, subject to the following conditions: +! +! The above copyright notice and this permission notice shall be included in +! all copies or substantial portions of the Software. +! +! THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +! IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +! FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +! AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +! LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +! OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +! THE SOFTWARE. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! The 1/N normalization is folded into the transform with +! rocfft_plan_description_set_scale_factor. +program rocfft_scale_factor_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Nbytes = N*8*2 + + type double2 + double precision :: x + double precision :: y + end type double2 + + type(double2), allocatable, target, dimension(:) :: hx, hx_input + type(c_ptr) :: dx = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: desc = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT scale factor double (z)' & + &(Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i)%x = dble(i) + hx(i)%y = dble(N - i) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, Nbytes)) + call hipCheck(hipMemcpy(dx, c_loc(hx(1)), Nbytes, hipMemcpyHostToDevice)) + + ! Forward transform (in-place), unscaled. + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! rocFFT transforms are unnormalized. Attaching a scale factor of 1/N to the + ! inverse plan folds the normalization into the transform itself, so the + ! round trip reproduces the input instead of N*input. + call rocfftCheck(rocfft_plan_description_create(desc)) + call rocfftCheck(rocfft_plan_description_set_scale_factor(desc, 1.0d0/dble(N))) + + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + desc)) + call rocfftCheck(rocfft_execute(plan_bwd, dx, c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call rocfftCheck(rocfft_plan_description_destroy(desc)) + + call hipCheck(hipMemcpy(c_loc(hx(1)), dx, Nbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + do i = 1, N + error = abs(hx(i)%x - hx_input(i)%x) + abs(hx(i)%y - hx_input(i)%y) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_scale_factor_z diff --git a/bindings/fortran/rocfft/test/rocfft_scale_factor_z.f08 b/bindings/fortran/rocfft/test/rocfft_scale_factor_z.f08 new file mode 100644 index 000000000000..58781decce6b --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_scale_factor_z.f08 @@ -0,0 +1,131 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Demonstrates rocfft_plan_description_set_scale_factor, described in the +! rocFFT "Working with rocFFT" guide: fold the 1/N normalization into the +! transform instead of scaling the result with a separate kernel. +program rocfft_scale_factor_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + + complex(8), allocatable, target, dimension(:) :: hx, hx_input + complex(8), pointer, dimension(:) :: dx => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: desc = c_null_ptr + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i) = cmplx(dble(i), dble(N - i), kind=8) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, source=hx)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT scale factor double (z)' & + &(Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! Forward transform (in-place), unscaled. + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + + ! The round trip alone would pass even if no transform ran, so check that the + ! forward transform did run: its DC bin must equal the sum of the input. + call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost)) + error = abs(hx(1) - sum(hx_input)) + if (.not. (error <= error_max * N * abs(sum(hx_input)))) then + write(*,*) "FAILED! forward DC bin error=", error, " got=", hx(1), & + " expected=", sum(hx_input) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + + ! rocFFT transforms are unnormalized. Attaching a scale factor of 1/N to the + ! inverse plan folds the normalization into the transform itself, so the + ! round trip reproduces the input instead of N*input. + call rocfftCheck(rocfft_plan_description_create(desc)) + call rocfftCheck(rocfft_plan_description_set_scale_factor(desc, 1.0d0/dble(N))) + + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + desc)) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call rocfftCheck(rocfft_plan_description_destroy(desc)) + + call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + do i = 1, N + error = abs(hx(i) - hx_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), & + " expected=", hx_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_scale_factor_z diff --git a/bindings/fortran/rocfft/test/rocfft_stream_z.f03 b/bindings/fortran/rocfft/test/rocfft_stream_z.f03 new file mode 100644 index 000000000000..979ee18eb09f --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_stream_z.f03 @@ -0,0 +1,170 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Fortran counterpart of the rocFFT sample +! clients/samples/rocfft/rocfft_example_set_stream.cpp: two independent in-place +! transforms, each queued on its own HIP stream. +program rocfft_stream_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer(c_size_t), parameter :: Nbytes = N*8*2 + integer, parameter :: Nfft = 2 + + type double2 + double precision :: x + double precision :: y + end type double2 + + ! One bundle of handles per transform, mirroring the fft_fixture_t struct + ! used by the rocFFT sample. + type fft_t + type(c_ptr) :: buf = c_null_ptr + type(c_ptr) :: stream = c_null_ptr + type(c_ptr) :: info = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: work = c_null_ptr + integer(c_size_t) :: work_size = 0 + end type fft_t + + type(fft_t) :: ffts(Nfft) + type(double2), allocatable, target, dimension(:,:) :: hx, hx_input + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer(c_size_t) :: work_fwd, work_bwd + integer :: i, k + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT stream double (z)' & + &(Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N,Nfft)) + allocate(hx_input(N,Nfft)) + do k = 1, Nfft + do i = 1, N + hx(i,k)%x = dble(i*k) + hx(i,k)%y = dble(N - i) + end do + end do + hx_input(:,:) = hx(:,:) + + ! Per transform: one buffer, stream, execution info (carrying the stream) + ! and two plans. + do k = 1, Nfft + call hipCheck(hipMalloc(ffts(k)%buf, Nbytes)) + call hipCheck(hipMemcpy(ffts(k)%buf, c_loc(hx(1,k)), Nbytes, hipMemcpyHostToDevice)) + + call hipCheck(hipStreamCreate(ffts(k)%stream)) + call rocfftCheck(rocfft_execution_info_create(ffts(k)%info)) + ! The stream must be a hipStream_t value, not the address of one. + call rocfftCheck(rocfft_execution_info_set_stream(ffts(k)%info, ffts(k)%stream)) + + call rocfftCheck(rocfft_plan_create(ffts(k)%plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_plan_create(ffts(k)%plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + + ! Each transform runs on its own stream, so each gets its own work buffer, + ! sized from what its two plans report. + call rocfftCheck(rocfft_plan_get_work_buffer_size(ffts(k)%plan_fwd, work_fwd)) + call rocfftCheck(rocfft_plan_get_work_buffer_size(ffts(k)%plan_bwd, work_bwd)) + ffts(k)%work_size = max(work_fwd, work_bwd) + if (ffts(k)%work_size > 0) then + call hipCheck(hipMalloc(ffts(k)%work, ffts(k)%work_size)) + call rocfftCheck(rocfft_execution_info_set_work_buffer(ffts(k)%info, & + ffts(k)%work, & + ffts(k)%work_size)) + end if + end do + + ! Execution: the calls return as soon as the work is queued on the streams, + ! so the two transforms can overlap. + do k = 1, Nfft + call rocfftCheck(rocfft_execute(ffts(k)%plan_fwd, ffts(k)%buf, & + c_null_ptr, ffts(k)%info)) + call rocfftCheck(rocfft_execute(ffts(k)%plan_bwd, ffts(k)%buf, & + c_null_ptr, ffts(k)%info)) + end do + + ! Wait for each stream before reading its results back. + do k = 1, Nfft + call hipCheck(hipStreamSynchronize(ffts(k)%stream)) + call hipCheck(hipMemcpy(c_loc(hx(1,k)), ffts(k)%buf, Nbytes, hipMemcpyDeviceToHost)) + end do + + do k = 1, Nfft + call rocfftCheck(rocfft_plan_destroy(ffts(k)%plan_fwd)) + call rocfftCheck(rocfft_plan_destroy(ffts(k)%plan_bwd)) + call rocfftCheck(rocfft_execution_info_destroy(ffts(k)%info)) + call hipCheck(hipStreamDestroy(ffts(k)%stream)) + call hipCheck(hipFree(ffts(k)%buf)) + if (c_associated(ffts(k)%work)) call hipCheck(hipFree(ffts(k)%work)) + end do + + ! rocFFT is unnormalized, so each round trip yields N*input. + do k = 1, Nfft + do i = 1, N + error = abs(hx(i,k)%x - N * hx_input(i,k)%x) & + + abs(hx(i,k)%y - N * hx_input(i,k)%y) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! k=", k, " i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_stream_z diff --git a/bindings/fortran/rocfft/test/rocfft_stream_z.f08 b/bindings/fortran/rocfft/test/rocfft_stream_z.f08 new file mode 100644 index 000000000000..db9c2b02cc90 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_stream_z.f08 @@ -0,0 +1,164 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Fortran counterpart of the rocFFT sample +! clients/samples/rocfft/rocfft_example_set_stream.cpp: two independent in-place +! transforms, each queued on its own HIP stream. +program rocfft_stream_z + use iso_c_binding + use hip + use rocfft + + implicit none + + integer(c_size_t), parameter :: N = 16 + integer, parameter :: Nfft = 2 + + ! One bundle of handles per transform, mirroring the fft_fixture_t struct + ! used by the rocFFT sample. + type fft_t + complex(8), pointer :: buf(:) => null() + type(c_ptr) :: stream = c_null_ptr + type(c_ptr) :: info = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: work = c_null_ptr + integer(c_size_t) :: work_size = 0 + end type fft_t + + type(fft_t) :: ffts(Nfft) + complex(8), allocatable, target, dimension(:,:) :: hx, hx_input + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer(c_size_t) :: work_fwd, work_bwd + integer :: i, k + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N,Nfft)) + allocate(hx_input(N,Nfft)) + do k = 1, Nfft + do i = 1, N + hx(i,k) = cmplx(dble(i*k), dble(N - i), kind=8) + end do + end do + hx_input(:,:) = hx(:,:) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT stream double (z)' & + &(Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! Preparation: one buffer, one stream, one execution info and two plans per + ! transform. The stream is carried by the execution info handle, which is + ! then passed to every rocfft_execute call that should use that stream. + do k = 1, Nfft + call hipCheck(hipMalloc(ffts(k)%buf, N)) + call hipCheck(hipMemcpy(ffts(k)%buf, hx(:,k), hipMemcpyHostToDevice)) + + call hipCheck(hipStreamCreate(ffts(k)%stream)) + call rocfftCheck(rocfft_execution_info_create(ffts(k)%info)) + ! The stream must be a hipStream_t value, not the address of one. + call rocfftCheck(rocfft_execution_info_set_stream(ffts(k)%info, ffts(k)%stream)) + + call rocfftCheck(rocfft_plan_create(ffts(k)%plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_plan_create(ffts(k)%plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + + ! Each transform runs on its own stream, so each gets its own work buffer, + ! sized from what its two plans report. + call rocfftCheck(rocfft_plan_get_work_buffer_size(ffts(k)%plan_fwd, work_fwd)) + call rocfftCheck(rocfft_plan_get_work_buffer_size(ffts(k)%plan_bwd, work_bwd)) + ffts(k)%work_size = max(work_fwd, work_bwd) + if (ffts(k)%work_size > 0) then + call hipCheck(hipMalloc(ffts(k)%work, ffts(k)%work_size)) + call rocfftCheck(rocfft_execution_info_set_work_buffer(ffts(k)%info, & + ffts(k)%work, & + ffts(k)%work_size)) + end if + end do + + ! Execution: the calls return as soon as the work is queued on the streams, + ! so the two transforms can overlap. + do k = 1, Nfft + call rocfftCheck(rocfft_execute(ffts(k)%plan_fwd, c_loc(ffts(k)%buf), & + c_null_ptr, ffts(k)%info)) + call rocfftCheck(rocfft_execute(ffts(k)%plan_bwd, c_loc(ffts(k)%buf), & + c_null_ptr, ffts(k)%info)) + end do + + ! Wait for each stream before reading its results back. + do k = 1, Nfft + call hipCheck(hipStreamSynchronize(ffts(k)%stream)) + call hipCheck(hipMemcpy(hx(:,k), ffts(k)%buf, hipMemcpyDeviceToHost)) + end do + + do k = 1, Nfft + call rocfftCheck(rocfft_plan_destroy(ffts(k)%plan_fwd)) + call rocfftCheck(rocfft_plan_destroy(ffts(k)%plan_bwd)) + call rocfftCheck(rocfft_execution_info_destroy(ffts(k)%info)) + call hipCheck(hipStreamDestroy(ffts(k)%stream)) + call hipCheck(hipFree(ffts(k)%buf)) + if (c_associated(ffts(k)%work)) call hipCheck(hipFree(ffts(k)%work)) + end do + + ! rocFFT is unnormalized, so each round trip yields N*input. + do k = 1, Nfft + do i = 1, N + error = abs(hx(i,k) - N * hx_input(i,k)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! k=", k, " i=", i, " error=", error, & + " got=", hx(i,k), " expected=", N * hx_input(i,k) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_stream_z diff --git a/bindings/fortran/rocfft/test/rocfft_version.f03 b/bindings/fortran/rocfft/test/rocfft_version.f03 new file mode 100644 index 000000000000..80e1a65b2fe2 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_version.f03 @@ -0,0 +1,67 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Demonstrates rocfft_get_version_string, the rocFFT library version query. +program rocfft_version + use iso_c_binding + use hip + use rocfft + + implicit none + + ! rocfft_get_version_string requires a buffer of at least 30 characters. + integer(c_size_t), parameter :: buflen = 64 + + character(kind=c_char), target :: buf(buflen) + character(kind=c_char, len=buflen) :: version + integer :: i + + write(*,"(a)",advance="no") "-- Running test 'rocFFT version string' & + &(Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! The C API writes a NUL-terminated string into the buffer, so pass the + ! address of the first element and the buffer capacity. + buf(:) = c_null_char + call rocfftCheck(rocfft_get_version_string(c_loc(buf(1)), buflen)) + + call rocfftCheck(rocfft_cleanup()) + + ! Copy the C string into a Fortran character variable, stopping at the NUL. + version = c_char_'' + do i = 1, buflen + if (buf(i) == c_null_char) exit + version(i:i) = buf(i) + end do + + if (len_trim(version) == 0) then + write(*,*) "FAILED! rocfft_get_version_string returned an empty string" + STOP 1 + end if + + write(*,*) "PASSED! rocFFT version: ", trim(version) + +end program rocfft_version diff --git a/bindings/fortran/rocfft/test/rocfft_version.f08 b/bindings/fortran/rocfft/test/rocfft_version.f08 new file mode 100644 index 000000000000..55725f592630 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_version.f08 @@ -0,0 +1,67 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Demonstrates rocfft_get_version_string, the rocFFT library version query. +program rocfft_version + use iso_c_binding + use hip + use rocfft + + implicit none + + ! rocfft_get_version_string requires a buffer of at least 30 characters. + integer(c_size_t), parameter :: buflen = 64 + + character(kind=c_char), target :: buf(buflen) + character(kind=c_char, len=buflen) :: version + integer :: i + + write(*,"(a)",advance="no") "-- Running test 'rocFFT version string' & + &(Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + ! The C API writes a NUL-terminated string into the buffer, so pass the + ! address of the first element and the buffer capacity. + buf(:) = c_null_char + call rocfftCheck(rocfft_get_version_string(c_loc(buf(1)), buflen)) + + call rocfftCheck(rocfft_cleanup()) + + ! Copy the C string into a Fortran character variable, stopping at the NUL. + version = c_char_'' + do i = 1, buflen + if (buf(i) == c_null_char) exit + version(i:i) = buf(i) + end do + + if (len_trim(version) == 0) then + write(*,*) "FAILED! rocfft_get_version_string returned an empty string" + STOP 1 + end if + + write(*,*) "PASSED! rocFFT version: ", trim(version) + +end program rocfft_version diff --git a/bindings/fortran/rocfft/test/rocfft_work_buffer_z.f03 b/bindings/fortran/rocfft/test/rocfft_work_buffer_z.f03 new file mode 100644 index 000000000000..e5d8cbce537d --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_work_buffer_z.f03 @@ -0,0 +1,140 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Fortran counterpart of the work-buffer handling in the rocFFT sample +! clients/samples/rocfft/rocfft_example_complexcomplex.cpp: query the scratch +! memory a plan needs and supply it through an execution info handle. +program rocfft_work_buffer_z + use iso_c_binding + use hip + use rocfft + + implicit none + + ! Large transforms are split into several kernels and need scratch memory. + integer(c_size_t), parameter :: N = 262144 + integer(c_size_t), parameter :: Nbytes = N*8*2 + + type double2 + double precision :: x + double precision :: y + end type double2 + + type(double2), allocatable, target, dimension(:) :: hx, hx_input + type(c_ptr) :: dx = c_null_ptr + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: info = c_null_ptr + type(c_ptr) :: dwork = c_null_ptr + integer(c_size_t) :: work_fwd = 0 + integer(c_size_t) :: work_bwd = 0 + integer(c_size_t) :: work_bytes = 0 + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + write(*,"(a)",advance="no") "-- Running test 'rocFFT work buffer double (z)' & + &(Fortran 2003 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i)%x = dble(mod(i,7)) - 3.0d0 + hx(i)%y = dble(mod(i,5)) - 2.0d0 + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, Nbytes)) + call hipCheck(hipMemcpy(dx, c_loc(hx(1)), Nbytes, hipMemcpyHostToDevice)) + + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths(1)),& + one,& + c_null_ptr)) + + ! One explicit work buffer, shared by both plans. + call rocfftCheck(rocfft_plan_get_work_buffer_size(plan_fwd, work_fwd)) + call rocfftCheck(rocfft_plan_get_work_buffer_size(plan_bwd, work_bwd)) + work_bytes = max(work_fwd, work_bwd) + + call rocfftCheck(rocfft_execution_info_create(info)) + if (work_bytes > 0) then + call hipCheck(hipMalloc(dwork, work_bytes)) + call rocfftCheck(rocfft_execution_info_set_work_buffer(info, dwork, work_bytes)) + end if + + call rocfftCheck(rocfft_execute(plan_fwd, dx, c_null_ptr, info)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_execute(plan_bwd, dx, c_null_ptr, info)) + call hipCheck(hipDeviceSynchronize()) + + call rocfftCheck(rocfft_execution_info_destroy(info)) + if (work_bytes > 0) then + call hipCheck(hipFree(dwork)) + end if + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(c_loc(hx(1)), dx, Nbytes, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! rocFFT is unnormalized, so the round trip yields N*input. + do i = 1, N + error = abs(hx(i)%x - N * hx_input(i)%x) + abs(hx(i)%y - N * hx_input(i)%y) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_work_buffer_z diff --git a/bindings/fortran/rocfft/test/rocfft_work_buffer_z.f08 b/bindings/fortran/rocfft/test/rocfft_work_buffer_z.f08 new file mode 100644 index 000000000000..9668e8c49898 --- /dev/null +++ b/bindings/fortran/rocfft/test/rocfft_work_buffer_z.f08 @@ -0,0 +1,135 @@ +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! +! 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. +! +!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + +! Fortran counterpart of the work-buffer handling in the rocFFT sample +! clients/samples/rocfft/rocfft_example_complexcomplex.cpp: query the scratch +! memory a plan needs and supply it through an execution info handle. +program rocfft_work_buffer_z + use iso_c_binding + use hip + use rocfft + + implicit none + + ! Large transforms are split into several kernels and need scratch memory. + integer(c_size_t), parameter :: N = 262144 + + complex(8), allocatable, target, dimension(:) :: hx, hx_input + complex(8), pointer, dimension(:) :: dx => null() + type(c_ptr) :: plan_fwd = c_null_ptr + type(c_ptr) :: plan_bwd = c_null_ptr + type(c_ptr) :: info = c_null_ptr + type(c_ptr) :: dwork = c_null_ptr + integer(c_size_t) :: work_fwd = 0 + integer(c_size_t) :: work_bwd = 0 + integer(c_size_t) :: work_bytes = 0 + integer(c_size_t), allocatable, target, dimension(:) :: lengths + integer(c_size_t), parameter :: one = 1 + integer :: i + double precision :: error + double precision, parameter :: error_max = 1.0d-8 + + allocate(lengths(1)) + lengths(1) = N + + allocate(hx(N)) + allocate(hx_input(N)) + do i = 1, N + hx(i) = cmplx(dble(mod(i,7)) - 3.0d0, dble(mod(i,5)) - 2.0d0, kind=8) + end do + hx_input(:) = hx(:) + + call hipCheck(hipMalloc(dx, source=hx)) + + write(*,"(a)",advance="no") "-- Running test 'rocFFT work buffer double (z)' & + &(Fortran 2008 interfaces) - " + + call rocfftCheck(rocfft_setup()) + + call rocfftCheck(rocfft_plan_create(plan_fwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_forward,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + call rocfftCheck(rocfft_plan_create(plan_bwd,& + rocfft_placement_inplace,& + rocfft_transform_type_complex_inverse,& + rocfft_precision_double,& + one,& + c_loc(lengths),& + one,& + c_null_ptr)) + + ! Ask each plan how much scratch memory it needs. rocFFT allocates the work + ! buffer itself when none is supplied; providing one explicitly lets the + ! application control its lifetime and share a single buffer between plans. + call rocfftCheck(rocfft_plan_get_work_buffer_size(plan_fwd, work_fwd)) + call rocfftCheck(rocfft_plan_get_work_buffer_size(plan_bwd, work_bwd)) + work_bytes = max(work_fwd, work_bwd) + + call rocfftCheck(rocfft_execution_info_create(info)) + if (work_bytes > 0) then + call hipCheck(hipMalloc(dwork, work_bytes)) + call rocfftCheck(rocfft_execution_info_set_work_buffer(info, dwork, work_bytes)) + end if + + call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, info)) + call hipCheck(hipDeviceSynchronize()) + call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, info)) + call hipCheck(hipDeviceSynchronize()) + + call rocfftCheck(rocfft_execution_info_destroy(info)) + if (work_bytes > 0) then + call hipCheck(hipFree(dwork)) + end if + call rocfftCheck(rocfft_plan_destroy(plan_fwd)) + call rocfftCheck(rocfft_plan_destroy(plan_bwd)) + + call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost)) + call hipCheck(hipFree(dx)) + + ! rocFFT is unnormalized, so the round trip yields N*input. + do i = 1, N + error = abs(hx(i) - N * hx_input(i)) + if (.not. (error <= error_max * N)) then + write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), & + " expected=", N * hx_input(i) + call rocfftCheck(rocfft_cleanup()) + STOP 1 + end if + end do + + deallocate(hx) + deallocate(hx_input) + deallocate(lengths) + + call rocfftCheck(rocfft_cleanup()) + + write(*,*) "PASSED!" + +end program rocfft_work_buffer_z