Merge ~bullwinkle-team/ubuntu/+source/hipfft:bullwinkle/llvm-21/ubuntu/devel into ubuntu/+source/hipfft:ubuntu/devel

[Bojan Aleksovski]

Update to new upstream version 7.1.0

[Bojan Aleksovski (b0b0a)]

Uploaded package to this PPA: https://launchpad.net/~b0b0a/+archive/ubuntu/hipfft-2140314

(-proposed and amd64, amd64v3, arm64 archs enabled)

Will trigger autopkgtest when it is built and published

[Bojan Aleksovski (b0b0a)]

`reverse-depends --arch ppc64el src:hipfft -x` returns No reverse dependencies found.

[Andreas Hasenack (ahasenack)]

--- a/debian/control
+++ b/debian/control
+Package: libhipfftw0
+Section: libs
+Architecture: amd64 arm64
+XB-X-ROCm-GPU-Architecture: ${rocm:GPU-Architecture}
+Depends: ${misc:Depends}, ${shlibs:Depends},
+Description: portable interface for Fast Fourier Transforms on the GPU - library
+ hipFFTW is a GPU-aware library for fast Fourier transforms using rocFFT as the
+ backend. It exports an interface borrowing the most commonly-used symbols of
+ FFTW. hipFFTW does not require its computational input and output to be
+ directly accessible by the GPU.
+ .
+ This package provides the AMD ROCm hipFFTW library.

Do you know if debian is going to add this same binary package, and also have libhipfft-dev pull in both libhipfft0 and this new libhipfftw0, like below?

 Package: libhipfft-dev
 Section: libdevel
-Architecture: amd64 arm64 ppc64el
-Depends: libhipfft0 (= ${binary:Version}),${misc:Depends}, ${shlibs:Depends},
- libamdhip64-dev,
+Architecture: amd64 arm64
+Depends: libhipfft0 (= ${binary:Version}),libhipfftw0 (= ${binary:Version}),
+ ${misc:Depends}, ${shlibs:Depends}, libamdhip64-dev,

Dealing with differences in packaging is a very annoying delta do carry in Ubuntu specially if Debian decides on a different split.

review: Needs Information

[Andreas Hasenack (ahasenack)]

--- /dev/null
+++ b/debian/patches/add-so-version-for-libhipfftw0.patch
@@ -0,0 +1,45 @@
+From: Bojan Aleksovski <email address hidden>
+Date: Wed, 23 Jan 2026 16:05:35 +0100
+Subject: Add versioning for hipfftw.so
+
+Summary: This patch is fully based on upstream commit
+dfa828c45abcbeb804876b973314cf39d2b3b2e8 and only has tiny
+additional fixes in the checking that hipfftw is not an alias
+
+Origin: upstream, https://github.com/ROCm/rocm-libraries/commit/dfa828c45abcbeb804876b973314cf39d2b3b2e8.patch
+Bug: https://github.com/ROCm/rocm-libraries/issues/2400
+Forwarded: not-needed

I see the difference:
-+ if (aliased_hipfftw_name STREQUAL "")
++ if (NOT aliased_hipfftw_name)

Could you please elaborate in the patch header why this was necessary? It's better than just to say "tiny change" and leave us guessing where it was and why :)

And why forwarding to upstream is not needed? Maybe it will become clear after the reasoning, but not yet.

review: Needs Information

[Andreas Hasenack (ahasenack)]

--- a/debian/patches/doxygen-refs.patch
+++ b/debian/patches/doxygen-refs.patch
@@ -1,28 +1,19 @@
-From: Kari Pahula <email address hidden>
-Date: Fri, 28 Jun 2024 16:07:05 +0300
-Subject: doxygen-refs
+From: Bojan Aleksovski <email address hidden>
+Date: Wed, 20 Jan 2026 20:55:15 +0100
+Subject: fix doxygen refs
...

What's the nature of these changes? This diff looks like the whole patch is being rewritten, including authorship. If it's indeed necessary, I would prefer if the original patch is dropped, and a new one is added with a different name.

Alternatively, to keep the existing name and authorship, you could add something about the change below the description, and add the Last-Updated DEP3 header with an updated date indicating when this change was made. But at a glance it looks like the first option (delete old, add new with new name) would be clearer?

And is this submittable to upstream?

review: Needs Information

[Andreas Hasenack (ahasenack)]

--- a/debian/rules
+++ b/debian/rules
@@ -17,8 +24,7 @@ CMAKE_FLAGS = \

 # The hipfft test client requires the rocfft sources to build.
 ifeq (,$(filter nocheck,$(DEB_BUILD_PROFILES)))
-CMAKE_FLAGS += -DBUILD_CLIENTS_TESTS=ON \
- -DUSE_HIPRAND=OFF
+CMAKE_FLAGS += -DBUILD_CLIENTS_TESTS=ON
 endif

 %:

I don't think this change is described in d/changelog, and why it's needed (remove USE_HIPRAND=OFF).

review: Needs Fixing

[Andreas Hasenack (ahasenack)]

d/changelog:
  * debian: new binary package libhipfftw0

Nit: that should be "* d/control: new binary package libhipfftw0"

While at it, you should probably also add a line about adding libhipfftw0 to the list of dependencies of libhipfft-dev.

Also in d/changelog:
  * d/control: fix hipfft-doc build

What in d/control fixed the hipfft-doc build? I think you meant d/rules?

review: Needs Fixing

[Bojan Aleksovski (b0b0a)]

Hello Andreas,

Regarding how the packages are organized, we took this direction after discussing with Cory in order to not have diff with Debian.

libhipfftw0 will be introduced as a new binary package.

Regarding the dev package, initially on the table was having separate hipfft and hipfftw dev packages but as they shared one header file we needed to introduce common-dev package that both would depend on. This option seemed far-fetched and it was agreed to leave the dev package as it is.

Regarding the tests package, on the table was renaming it to libhipfft-tests (dropping the 0 as the tests are for both bin pkgs), but as libhipfftw is basically a shim this seemed as not needed complexity as well.
Btw, this made me realize that I left out libhipfftw0 as a dependency of libhipfft0-tests, will fix that.

[Bojan Aleksovski (b0b0a)]

While I was writing the reasoning about how the packages are organized, I saw that you left the other notes. Thank you, I will provide reasoning and fixes soon. Have a great weekend!

[Bojan Aleksovski (b0b0a)]

Hello Andreas,

Please find the effort to have the comments fixed:
- Regarding add-so-version-for-libhipfftw0.patch, the "tiny change" to the fix has been elaborated in patch header and has been forwarded upstream
- Regarding doxygen-refs.patch, it has been dropped. A new fix-doxygen-refs.patch has been added which has been forwarded upstream
- Regarding the change that does `remove USE_HIPRAND=OFF`, it has been forgotten in the previous changelog entry 7.1.0-0ubuntu1. I've added it. That change refers to commit https://git.launchpad.net/~bullwinkle-team/ubuntu/+source/hipfft/commit/?h=bullwinkle/llvm-21/ubuntu/devel&id=88ca75c121a36f50d2e4e686a45058ff3548cab2
- Regarding the changelog entry for the new binary package libhipfftw0, it has been updated and a new line about having libhipfftw0 as a dependency for libhipfft-dev and libhipfft0-tests has been added
- The typo in the changelog about fix hipfft-doc build has been resolved

The easiest place to view the mentioned fixes is https://launchpadlibrarian.net/846723009/hipfft_7.1.0-0ubuntu2~ppa1_7.1.0-0ubuntu2~ppa2.diff.gz

Thank you!

[Andreas Hasenack (ahasenack)]

I noticed there is no symbols file for arm64, just amd64, but that is not a bug being introduced in this PR. At some point it should be fixed, though.

+1

review: Approve

[Andreas Hasenack (ahasenack)]

Sponsored:

Uploading hipfft_7.1.0-0ubuntu2.dsc
Uploading hipfft_7.1.0.orig.tar.gz
Uploading hipfft_7.1.0-0ubuntu2.debian.tar.xz
Uploading hipfft_7.1.0-0ubuntu2_source.buildinfo
Uploading hipfft_7.1.0-0ubuntu2_source.changes

This has a binary NEW package, so it will require an archive admin review before it lands in proposed.

[Bojan Aleksovski (b0b0a)]

Thank you Andreas. Bug has been created regarding the arm64 symbols (using rocblas bug as reference), so it gets fixed soon: https://bugs.launchpad.net/ubuntu/+source/hipfft/+bug/2141347

diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS
old mode 100755
new mode 100644
index eff670f..056a93f
--- a/.github/CODEOWNERS
+++ b/.github/CODEOWNERS
@@ -1,8 +1,8 @@
-* @af-ayala @eng-flavio-teixeira @evetsso @feizheng10 @malcolmroberts
+* @af-ayala @eng-flavio-teixeira @evetsso @malcolmroberts @regan-amd
 # Documentation files
 docs/ @ROCm/rocm-documentation
 *.md @ROCm/rocm-documentation
 *.rst @ROCm/rocm-documentation
 .readthedocs.yaml @ROCm/rocm-documentation
 # Header directory for Doxygen documentation
-library/include/ @ROCm/rocm-documentation @af-ayala @eng-flavio-teixeira @evetsso @feizheng10 @malcolmroberts
+library/include/ @ROCm/rocm-documentation @af-ayala @eng-flavio-teixeira @evetsso @malcolmroberts @regan-amd
diff --git a/.jenkins/common.groovy b/.jenkins/common.groovy
deleted file mode 100644
index e0aea53..0000000
--- a/.jenkins/common.groovy
+++ /dev/null
@@ -1,92 +0,0 @@
-import static groovy.io.FileType.FILES
-
-def runCompileCommand(platform, project, jobName, boolean sameOrg = false)
-{
-    project.paths.construct_build_prefix()
-
-    def getDependenciesCommand = ""
-    if (project.installLibraryDependenciesFromCI)
-    {
-        project.libraryDependencies.each
-        { libraryName ->
-            getDependenciesCommand += auxiliary.getLibrary(libraryName, platform.jenkinsLabel, null, sameOrg)
-        }
-    }
-
-    String cmake = platform.jenkinsLabel.contains('centos') ? "cmake3" : "cmake" 
-    String warningArgs = platform.jenkinsLabel.contains('cuda') ? '':'-DWERROR=ON'
-    String hipClang = platform.jenkinsLabel.contains('hipClang') ? "HIP_COMPILER=clang" : ""
-    String path = platform.jenkinsLabel.contains('centos7') ? "export PATH=/opt/rh/devtoolset-7/root/usr/bin:$PATH" : ":"
-    String dir = jobName.contains('Debug') ? "debug" : "release"
-
-    // hipcc with CUDA backend needs HIP_PLATFORM set accordingly in the environment
-    String hipPlatformCommand = platform.jenkinsLabel.contains("cuda") ? "export HIP_PLATFORM=nvidia" : ""
-
-    def command = """#!/usr/bin/env bash
-                set -x
-                
-                ls /fftw/lib
-                export FFTW_ROOT=/fftw
-                export FFTW_INCLUDE_PATH=\${FFTW_ROOT}/include
-                export FFTW_LIB_PATH=\${FFTW_ROOT}/lib
-                export LD_LIBRARY_PATH=\${FFTW_LIB_PATH}:/opt/rocm/lib:/opt/rocm/hip/lib
-                export CPLUS_INCLUDE_PATH=\${FFTW_INCLUDE_PATH}:\${CPLUS_INCLUDE_PATH}
-                export CMAKE_PREFIX_PATH=\${FFTW_LIB_PATH}/cmake/fftw3:\${CMAKE_PREFIX_PATH}
-                export CMAKE_PREFIX_PATH=\${FFTW_LIB_PATH}/cmake/fftw3f:\${CMAKE_PREFIX_PATH}
-		# default container flags cause problems for CUDA backend, and aren't useful for ROCm
-		unset HIPCC_COMPILE_FLAGS_APPEND
-		unset HIPCC_LINK_FLAGS_APPEND
-		${hipPlatformCommand}
-                
-                cd ${project.paths.project_build_prefix}
-                mkdir -p build/${dir} && cd build/${dir}
-                ${getDependenciesCommand}
-                ${path}
-                ${hipClang} ${cmake} ${warningArgs} ${project.paths.build_command}
-                make -j\$(nproc)
-            """
-
-    platform.runCommand(this, command)
-}
-
-def runTestCommand (platform, project, gfilter)
-{
-    String cudaArgs = platform.jenkinsLabel.contains('cuda') ? '--double_epsilon=5e-11'
-                        : '--precompile=rocfft-test-precompile.db'
-
-    def command = """#!/usr/bin/env bash
-                    set -x
-                    cd ${project.paths.project_build_prefix}/build/release/clients/staging
-                    GTEST_LISTENER=NO_PASS_LINE_IN_LOG ./hipfft-test ${cudaArgs} --gtest_output=xml --gtest_color=yes --gtest_filter=${gfilter}
-                """
-
-    platform.runCommand(this, command)
-    //junit "${project.paths.project_build_prefix}/build/release/clients/staging/*.xml"
-}
-
-def runPackageCommand(platform, project, jobName, label='')
-{
-    def command
-
-    label = label != '' ? '-' + label.toLowerCase() : ''
-    String ext = platform.jenkinsLabel.contains('ubuntu') ? "deb" : "rpm"
-    String dir = jobName.contains('Debug') ? "debug" : "release"
-
-    command = """
-            set -x
-            cd ${project.paths.project_build_prefix}/build/${dir}
-            make package
-            mkdir -p package
-            for f in hipfft*.$ext
-            do 
-                mv "\$f" "hipfft${label}-\${f#*-}"
-            done
-            mv *.${ext} package/
-        """
-
-    platform.runCommand(this, command)
-    platform.archiveArtifacts(this, """${project.paths.project_build_prefix}/build/${dir}/package/*.${ext}""")
-}
-
-
-return this
diff --git a/.jenkins/debug.groovy b/.jenkins/debug.groovy
deleted file mode 100644
index abac808..0000000
--- a/.jenkins/debug.groovy
+++ /dev/null
@@ -1,83 +0,0 @@
-#!/usr/bin/env groovy
-@Library('rocJenkins@pong') _
-
-import com.amd.project.*
-import com.amd.docker.*
-import java.nio.file.Path
-
-def runCI =
-{
-    nodeDetails, jobName, buildCommand ->
-
-    def prj = new rocProject('hipFFT', 'Debug')
-    // customize for project
-    prj.paths.build_command = buildCommand
-    prj.libraryDependencies = ['rocRAND', 'rocFFT', 'hipRAND']
-
-    // Define test architectures, optional rocm version argument is available
-    def nodes = new dockerNodes(nodeDetails, jobName, prj)
-
-    def commonGroovy
-
-    boolean formatCheck = false
-
-    def compileCommand =
-    {
-        platform, project->
-
-        project.paths.construct_build_prefix()
-
-        commonGroovy = load "${project.paths.project_src_prefix}/.jenkins/common.groovy"
-        commonGroovy.runCompileCommand(platform, project, jobName)
-    }
-
-    buildProject(prj, formatCheck, nodes.dockerArray, compileCommand, null, null)
-}
-
-def setupCI(urlJobName, jobNameList, buildCommand, runCI, label)
-{
-    jobNameList = auxiliary.appendJobNameList(jobNameList)
-
-    jobNameList.each
-    {
-        jobName, nodeDetails->
-        if (urlJobName == jobName)
-            stage(label + ' ' + jobName) {
-                runCI(nodeDetails, jobName, buildCommand, label)
-            }
-    }
-
-    // For url job names that are not listed by the jobNameList i.e. compute-rocm-dkms-no-npi-1901
-    if(!jobNameList.keySet().contains(urlJobName))
-    {
-        properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 1 * * *')])]))
-        stage(label + ' ' + urlJobName) {
-            runCI([ubuntu18:['gfx906']], urlJobName, buildCommand, label)
-        }
-    }
-}
-
-ci: {
-    String urlJobName = auxiliary.getTopJobName(env.BUILD_URL)
-
-    def propertyList = ["compute-rocm-dkms-no-npi":[pipelineTriggers([cron('0 1 * * 0')])]]
-
-    propertyList = auxiliary.appendPropertyList(propertyList)
-
-    propertyList.each
-    {
-        jobName, property->
-        if (urlJobName == jobName)
-            properties(auxiliary.addCommonProperties(property))
-    }
-
-    def hostJobNameList = ["compute-rocm-dkms-no-npi":([ubuntu18:['gfx900'],centos7:['gfx906'],sles15sp1:['gfx906']])]
-
-    def hipClangJobNameList = ["compute-rocm-dkms-no-npi":([ubuntu18:['gfx900'],centos7:['gfx906'],sles15sp1:['gfx906']])]
-
-    String hostBuildCommand = '-DCMAKE_CXX_COMPILER=g++ -DCMAKE_BUILD_TYPE=Debug -L ../..'
-    String hipClangBuildCommand = '-DCMAKE_CXX_COMPILER=/opt/rocm/bin/amdclang++ -DCMAKE_BUILD_TYPE=Debug -DBUILD_CLIENTS_TESTS=ON -DBUILD_CLIENTS_SAMPLES=ON -DBUILD_CLIENTS_SAMPLES=ON -L ../..'
-
-    setupCI(urlJobName, hostJobNameList, hostBuildCommand, runCI, 'g++')
-    setupCI(urlJobName, hipClangJobNameList, hipClangBuildCommand, runCI, 'hip-clang')
-}
diff --git a/.jenkins/multigpu.groovy b/.jenkins/multigpu.groovy
deleted file mode 100644
index 9c315df..0000000
--- a/.jenkins/multigpu.groovy
+++ /dev/null
@@ -1,96 +0,0 @@
-#!/usr/bin/env groovy
-@Library('rocJenkins@pong') _
-
-import com.amd.project.*
-import com.amd.docker.*
-import java.nio.file.Path
-
-def runCI =
-{
-    nodeDetails, jobName, buildCommand, label, runTest ->
-
-    def prj = new rocProject('hipFFT', 'multigpu')
-    // customize for project
-    prj.paths.build_command = buildCommand
-    prj.libraryDependencies = ['rocRAND', 'rocFFT', 'hipRAND']
-    prj.timeout.test = 360
-
-    // Define test architectures, optional rocm version argument is available
-    def nodes = new dockerNodes(nodeDetails, jobName, prj)
-
-    def commonGroovy
-
-    boolean formatCheck = false
-
-    def compileCommand =
-    {
-        platform, project->
-
-        project.paths.construct_build_prefix()
-
-        commonGroovy = load "${project.paths.project_src_prefix}/.jenkins/common.groovy"
-        commonGroovy.runCompileCommand(platform, project,jobName)
-    }
-
-    def testCommand =
-    {
-        platform, project->
-
-        def gfilter = "*multi_gpu*"
-        commonGroovy.runTestCommand(platform, project, gfilter)
-    }
-
-    def packageCommand =
-    {
-        platform, project->
-
-        commonGroovy.runPackageCommand(platform, project, jobName, label)
-    }
-
-    buildProject(prj, formatCheck, nodes.dockerArray, compileCommand, runTest ? testCommand : null, packageCommand)
-}
-
-def setupCI(urlJobName, jobNameList, buildCommand, runCI, label, runTest)
-{
-    jobNameList = auxiliary.appendJobNameList(jobNameList)
-
-    jobNameList.each
-    {
-        jobName, nodeDetails->
-        if (urlJobName == jobName)
-            stage(label + ' ' + jobName) {
-                runCI(nodeDetails, jobName, buildCommand, label, runTest)
-            }
-    }
-
-    // For url job names that are not listed by the jobNameList i.e. compute-rocm-dkms-no-npi-1901
-    if(!jobNameList.keySet().contains(urlJobName))
-    {
-        properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 1 * * *')])]))
-        stage(label + ' ' + urlJobName) {
-            runCI([ubuntu18:['gfx906']], urlJobName, buildCommand, label)
-        }
-    }
-}
-
-ci: {
-    String urlJobName = auxiliary.getTopJobName(env.BUILD_URL)
-
-    def propertyList = ["main":[pipelineTriggers([cron('0 6 * * 0')])]]
-
-    propertyList = auxiliary.appendPropertyList(propertyList)
-
-    def jobNameList = ["main":([ubuntu20:['8gfx90a']])]
-    jobNameList = auxiliary.appendJobNameList(jobNameList)
-    
-    propertyList.each
-    {
-        jobName, property->
-        if (urlJobName == jobName)
-            properties(auxiliary.addCommonProperties(property))
-    }
-    
-    String hipClangBuildCommand = '-DCMAKE_CXX_COMPILER=/opt/rocm/bin/amdclang++ -DCMAKE_BUILD_TYPE=RelWithDebInfo -DBUILD_CLIENTS_TESTS=ON -DBUILD_CLIENTS_SAMPLES=ON -L ../..'
-
-    setupCI(urlJobName, jobNameList, hipClangBuildCommand, runCI, 'hip-clang', true)
-}
diff --git a/.jenkins/staticanalysis.groovy b/.jenkins/staticanalysis.groovy
deleted file mode 100644
index dd54dde..0000000
--- a/.jenkins/staticanalysis.groovy
+++ /dev/null
@@ -1,32 +0,0 @@
-#!/usr/bin/env groovy
-@Library('rocJenkins@pong') _
-
-import com.amd.project.*
-import com.amd.docker.*
-import java.nio.file.Path
-
-def runCI =
-{
-    nodeDetails, jobName ->
-
-    def prj = new rocProject('hipFFT-internal', 'PreCheckin')
-    // customize for project
-    prj.libraryDependencies = ['rocRAND','rocFFT']
-
-    // Define test architectures, optional rocm version argument is available
-    def nodes = new dockerNodes(nodeDetails, jobName, prj)
-
-    boolean formatCheck = true
-    boolean staticAnalysis = true
-
-    buildProject(prj, formatCheck, nodes.dockerArray, null, null, null, staticAnalysis)
-}
-
-ci: { 
-    String urlJobName = auxiliary.getTopJobName(env.BUILD_URL)
-
-    properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 1 * * 7')])]))
-    stage(urlJobName) {
-        runCI([ubuntu22:['any']], urlJobName)
-    }
-}
diff --git a/.jenkins/staticlibrary.groovy b/.jenkins/staticlibrary.groovy
deleted file mode 100644
index 1e99ac3..0000000
--- a/.jenkins/staticlibrary.groovy
+++ /dev/null
@@ -1,98 +0,0 @@
-#!/usr/bin/env groovy
-@Library('rocJenkins@pong') _
-
-import com.amd.project.*
-import com.amd.docker.*
-import java.nio.file.Path
-
-def runCI =
-{
-    nodeDetails, jobName, buildCommand, label ->
-
-    def prj = new rocProject('hipFFT-internal', 'StaticLibrary')
-    // customize for project
-    prj.paths.build_command = buildCommand
-    prj.libraryDependencies = ['rocRAND','rocFFT']
-
-    // Define test architectures, optional rocm version argument is available
-    def nodes = new dockerNodes(nodeDetails, jobName, prj)
-
-    def commonGroovy
-
-    boolean formatCheck = false
-
-    def compileCommand =
-    {
-        platform, project->
-
-        project.paths.construct_build_prefix()
-
-        commonGroovy = load "${project.paths.project_src_prefix}/.jenkins/common.groovy"
-        commonGroovy.runCompileCommand(platform, project, jobName, true)
-    }
-
-    def testCommand =
-    {
-        platform, project->
-
-        def gfilter = "*"
-        commonGroovy.runTestCommand(platform, project, gfilter)
-    }
-
-    def packageCommand =
-    {
-        platform, project->
-
-        commonGroovy.runPackageCommand(platform, project, jobName, label)
-    }
-
-    buildProject(prj, formatCheck, nodes.dockerArray, compileCommand, testCommand, packageCommand)
-}
-
-def setupCI(urlJobName, jobNameList, buildCommand, runCI, label)
-{
-    jobNameList = auxiliary.appendJobNameList(jobNameList)
-
-    jobNameList.each
-    {
-        jobName, nodeDetails->
-        if (urlJobName == jobName)
-            stage(label + ' ' + jobName) {
-                runCI(nodeDetails, jobName, buildCommand, label)
-            }
-    }
-
-    // For url job names that are not listed by the jobNameList i.e. compute-rocm-dkms-no-npi-1901
-    if(!jobNameList.keySet().contains(urlJobName))
-    {
-        properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 1 * * *')])]))
-        stage(label + ' ' + urlJobName) {
-            runCI([ubuntu16:['gfx906']], urlJobName, buildCommand, label)
-        }
-    }
-}
-
-ci: {
-    String urlJobName = auxiliary.getTopJobName(env.BUILD_URL)
-
-    def propertyList = ["compute-rocm-dkms-no-npi":[pipelineTriggers([cron('0 1 * * 0')])]]
-
-    propertyList = auxiliary.appendPropertyList(propertyList)
-
-    propertyList.each
-    {
-        jobName, property->
-        if (urlJobName == jobName)
-            properties(auxiliary.addCommonProperties(property))
-    }
-
-    def hostJobNameList = ["compute-rocm-dkms-no-npi-hipclang":([ubuntu18:['gfx900']])]
-
-    def hipClangJobNameList = ["compute-rocm-dkms-no-npi-hipclang":([ubuntu18:['gfx900']])]
-
-    String hostBuildCommand = '-DCMAKE_CXX_COMPILER=g++ -DCMAKE_BUILD_TYPE=RelWithDebInfo -DBUILD_SHARED_LIBS=OFF -L ../..'
-    String hipClangBuildCommand = '-DCMAKE_CXX_COMPILER=/opt/rocm/bin/amdclang++ -DCMAKE_BUILD_TYPE=RelWithDebInfo -DBUILD_CLIENTS_TESTS=ON -DBUILD_CLIENTS_SAMPLES=ON -DBUILD_SHARED_LIBS=OFF -L ../..'
-
-    setupCI(urlJobName, hostJobNameList, hostBuildCommand, runCI, 'g++')
-    setupCI(urlJobName, hipClangJobNameList, hipClangBuildCommand, runCI, 'hip-clang')
-}
diff --git a/CHANGELOG.md b/CHANGELOG.md
index d9c3dd8..257c7a5 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -3,6 +3,26 @@
 Documentation for hipFFT is available at
 [https://rocm.docs.amd.com/projects/hipFFT/en/latest/](https://rocm.docs.amd.com/projects/hipFFT/en/latest/).
 
+## hipFFT 1.0.21 for ROCm 7.1.0
+
+### Added
+
+* Improved test coverage of multi-stream plans.
+* Improved test coverage of user-specified work areas.
+* Improved test coverage of default stride calculation.
+* **[Experimental]** Introduced the hipFFTW library, interfacing rocFFT on AMD platforms using the same symbols as FFTW3 (with partial support).
+
+## hipFFT 1.0.20 for ROCm 7.0.0
+
+### Added
+
+* Added gfx950 support.
+
+### Removed
+
+* Removed hipfft-rider legacy compatibility from clients
+* Remove support for the gfx940 and gfx941 targets from the client programs.
+
 ## hipFFT 1.0.18 for ROCm 6.4.0
 
 ### Added
@@ -10,8 +30,7 @@ Documentation for hipFFT is available at
 * Implemented the `hipfftMpAttachComm`, `hipfftXtSetDistribution`, and `hipfftXtSetSubformatDefault` APIs to allow
   computing FFTs that are distributed between multiple MPI (Message Passing Interface) processes.  These APIs can be enabled
   with the `HIPFFT_MPI_ENABLE` CMake option, which defaults to `OFF`.
-
-  The backend FFT library called by hipFFT must support MPI for these APIs to work.
+* The backend FFT library called by hipFFT must support MPI for these APIs to work.
 
 ### Changed
 
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 7f81897..8f802ac 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -1,5 +1,5 @@
 # #############################################################################
-# Copyright (C) 2020 - 2022 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (C) 2020 - 2025 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -37,7 +37,6 @@ else()
     "Install path prefix, prepended onto install directories" )
 endif()
 
-
 # Workarounds..
 list( APPEND CMAKE_PREFIX_PATH /opt/rocm/llvm /opt/rocm )
 list( APPEND CMAKE_MODULE_PATH ${ROCM_PATH}/lib/cmake/hip /opt/rocm/lib/cmake/hip /opt/rocm/hip/cmake )
@@ -83,9 +82,8 @@ set(DEFAULT_GPUS
   gfx906
   gfx908
   gfx90a
-  gfx940
-  gfx941
   gfx942
+  gfx950
   gfx1030
   gfx1100
   gfx1101
@@ -101,8 +99,6 @@ if(BUILD_ADDRESS_SANITIZER)
   SET(DEFAULT_GPUS
     gfx908:xnack+
     gfx90a:xnack+
-    gfx940:xnack+
-    gfx941:xnack+
     gfx942:xnack+)
   add_link_options(-fuse-ld=lld)
   add_compile_definitions(ADDRESS_SANITIZER)
@@ -167,14 +163,8 @@ if (BUILD_WITH_COMPILER STREQUAL "HIP-NVCC" )
   set( WARNING_FLAGS ${NVCC_WARNING_FLAGS} )
 
 else()
-  # Define GPU targets
-  if(AMDGPU_TARGETS AND NOT GPU_TARGETS)
-    message( DEPRECATION "AMDGPU_TARGETS use is deprecated. Use GPU_TARGETS." )
-  endif()
-  set(AMDGPU_TARGETS "${DEFAULT_GPUS}" CACHE STRING "Target default GPUs if AMDGPU_TARGETS is not defined. (Deprecated, prefer GPU_TARGETS)")
-  rocm_check_target_ids(AMDGPU_TARGETS TARGETS "${AMDGPU_TARGETS}")
-  # Don't force, users should be able to override GPU_TARGETS at the command line if desired
-  set(GPU_TARGETS "${AMDGPU_TARGETS}" CACHE STRING "GPU architectures to build for")
+  # Check support for the GPU target(s)
+  rocm_check_target_ids(GPU_TARGETS TARGETS "${GPU_TARGETS}")
   if( BUILD_WITH_COMPILER STREQUAL "HIP-CLANG" )
     set( HIP_PLATFORM "amd" )
     set( HIP_COMPILER "clang" )
@@ -184,22 +174,11 @@ endif()
 # Show the actual compiler(internal option)
 message(STATUS "BUILD_WITH_COMPILER = " ${BUILD_WITH_COMPILER})
 
-# FOR HANDLING ENABLE/DISABLE OPTIONAL BACKWARD COMPATIBILITY for FILE/FOLDER REORG
-option(BUILD_FILE_REORG_BACKWARD_COMPATIBILITY "Build with file/folder reorg with backward compatibility enabled" OFF)
-if(BUILD_FILE_REORG_BACKWARD_COMPATIBILITY AND NOT WIN32)
-  rocm_wrap_header_dir(
-    ${CMAKE_SOURCE_DIR}/library/include
-    PATTERNS "*.h"
-    GUARDS SYMLINK WRAPPER
-    WRAPPER_LOCATIONS ${CMAKE_INSTALL_INCLUDEDIR}
-  )
-endif()
-
 # Version
-set( VERSION_STRING "1.0.18" )
+set( VERSION_STRING "1.0.21" )
 set( hipfft_SOVERSION 0.1 )
 
-if( ROCM_FOUND )
+if( ROCmCMakeBuildTools_FOUND )
   rocm_setup_version( VERSION ${VERSION_STRING} )
 endif()
 
@@ -262,7 +241,9 @@ endif()
 if(WIN32)
   set(CPACK_SOURCE_GENERATOR "ZIP")
   set(CPACK_GENERATOR "ZIP")
-  set(CMAKE_INSTALL_PREFIX "C:/hipSDK" CACHE PATH "Install path" FORCE)
+  if(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT)
+    set(CMAKE_INSTALL_PREFIX "C:/hipSDK" CACHE PATH "Install path" FORCE)
+  endif()
   set(INSTALL_PREFIX "C:/hipSDK")
   set(CPACK_SET_DESTDIR OFF)
   set(CPACK_PACKAGE_INSTALL_DIRECTORY "C:/hipSDK")
@@ -270,7 +251,7 @@ if(WIN32)
   set(CPACK_INCLUDE_TOPLEVEL_DIRECTORY OFF)
 endif()
 
-if( ROCM_FOUND )
+if( ROCmCMakeBuildTools_FOUND )
   # Package specific CPACK vars
   if( NOT BUILD_WITH_LIB STREQUAL "CUDA" )
     rocm_package_add_dependencies(DEPENDS "rocfft >= 1.0.21")
diff --git a/LICENSE.md b/LICENSE.md
index f081475..ee33c80 100644
--- a/LICENSE.md
+++ b/LICENSE.md
@@ -1,6 +1,6 @@
 MIT License
 
-Copyright (C) 2016 - 2025 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (C) Advanced Micro Devices, Inc.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -20,6 +20,8 @@ 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.
 
+---
+
 This product includes software from copyright holders as shown below, and distributed under their license terms as specified.
 
 CLI11 2.2 Copyright (c) 2017-2024 University of Cincinnati, developed by Henry
diff --git a/README.md b/README.md
index d47fd12..5961742 100644
--- a/README.md
+++ b/README.md
@@ -1,7 +1,7 @@
 # hipFFT
 
 hipFFT is an FFT marshalling library that supports
-[rocFFT](https://github.com/ROCmSoftwarePlatform/rocFFT) and
+[rocFFT](https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocfft) and
 [cuFFT](https://developer.nvidia.com/cufft) backends.
 
 hipFFT exports an interface that doesn't require the client to change, regardless of the chosen backend.
@@ -11,12 +11,12 @@ and marshals results back to your application.
 ## Documentation
 
 > [!NOTE]
-> The published hipFFT documentation is available at [hipFFT](https://rocm.docs.amd.com/projects/hipFFT/en/latest/index.html) in an organized, easy-to-read format, with search and a table of contents. The documentation source files reside in the hipFFT/docs folder of this repository. As with all ROCm projects, the documentation is open source. For more information, see [Contribute to ROCm documentation](https://rocm.docs.amd.com/en/latest/contribute/contributing.html).
+> The published hipFFT documentation is available at [hipFFT](https://rocm.docs.amd.com/projects/hipFFT/en/latest/index.html) in an organized, easy-to-read format, with search and a table of contents. The documentation source files reside in the projects/hipfft/docs folder of the rocm-libraries repository. As with all ROCm projects, the documentation is open source. For more information, see [Contribute to ROCm documentation](https://rocm.docs.amd.com/en/latest/contribute/contributing.html).
 
 To build our documentation locally, run the following code:
 
 ```bash
-cd docs
+cd projects/hipfft/docs
 
 pip3 install -r sphinx/requirements.txt
 
@@ -36,14 +36,13 @@ To build hipFFT from source, follow these steps:
 
 1. Install the library build dependencies:
 
-   * On AMD platforms, you must install [rocFFT](https://github.com/ROCmSoftwarePlatform/rocFFT).
-   * On NVIDIA platforms, you must install [cuFFT](https://developer.nvidia.com/cufft).
+   * On AMD platforms, you must install [rocFFT](https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocfft).
 
 2. Install the client build dependencies:
 
    * The clients (samples, tests, etc) included with the hipFFT source depend on hipRAND, FFTW and GoogleTest.
 
-3. Build hipFFT:
+3. Build hipFFT. Run these commands from the `rocm-libraries/projects/hipfft` directory:
 
     To show all build options:
 
@@ -52,40 +51,42 @@ To build hipFFT from source, follow these steps:
       cmake -LH ..
     ```
 
-Here are some CMake build examples:
+Here are some CMake build examples for an AMD GPU:
 
-* AMD GPU
-  * Case: Build a project using HIP language APIs + hipFFT with standard host compiler
-    * Code: `cmake -DCMAKE_CXX_COMPILER=g++ -DCMAKE_BUILD_TYPE=Release -L ..`
-  * Case: Build a project using HIP language APIs + hipFFT + device kernels with HIP-Clang
-    * Code: `cmake -DCMAKE_CXX_COMPILER=amdclang++ -DCMAKE_BUILD_TYPE=Release -DBUILD_CLIENTS=ON -L ..`
-* NVIDIA GPU
-  * Case: Build a project using HIP language APIs + hipFFT with standard host compiler
-    * Code: `cmake -DCMAKE_CXX_COMPILER=g++ -DCMAKE_BUILD_TYPE=Release -DBUILD_WITH_LIB=CUDA -L ..`
-  * Case: Build a project using HIP language APIs + hipFFT + device kernels with HIP-NVCC
-    * Code: `HIP_PLATFORM=nvidia cmake -DCMAKE_CXX_COMPILER=hipcc -DCMAKE_BUILD_TYPE=Release -DBUILD_CLIENTS=ON -L ..`
+* Case: Build a project using HIP language APIs + hipFFT with standard host compiler
+   * Code: `cmake -DCMAKE_CXX_COMPILER=g++ -DCMAKE_BUILD_TYPE=Release -L ..`
+* Case: Build a project using HIP language APIs + hipFFT + device kernels with HIP-Clang
+   * Code: `cmake -DCMAKE_CXX_COMPILER=amdclang++ -DCMAKE_BUILD_TYPE=Release -DBUILD_CLIENTS=ON -L ..`
 
 ```note
 The `-DBUILD_CLIENTS=ON` option is only allowed with the amdclang++ or HIPCC compilers.
 ```
 
+## Code Coverage
+You can generate a test coverage report with the following:
+
+```bash
+cmake -DCMAKE_CXX_COMPILER=amdclang++ -DBUILD_CLIENTS_SAMPLES=ON -DBUILD_CLIENTS_TESTS=ON -DBUILD_CODE_COVERAGE=ON <optional: -DCOVERAGE_TEST_OPTIONS="cmdline args to pass to hipfft-test (default: --smoketest)"> ..
+make -j coverage
+```
+The commands above will output the coverage report to the terminal and save an html coverage report to `$PWD/coverage-report`.  Note that hipFFT uses llvm for code coverage, which only works with clang compilers.
+
 ## Porting from CUDA
 
 If you have existing CUDA code and want to transition to HIP, follow these steps:
 
 1. [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY) your code and fix all unsupported CUDA
    features and user-defined macros
-2. Build with HIP-NVCC to run on an NVIDIA device
-3. Build with HIP-Clang to run on an AMD device
+2. Build with HIP-Clang to run on an AMD device
 
 More information about porting to HIP is available in the
 [HIP porting guide](https://rocm.docs.amd.com/projects/HIP/en/develop/user_guide/hip_porting_guide.html).
 
 ## Support
 
-You can report bugs and feature requests through the GitHub
-[issue tracker](https://github.com/ROCm/hipFFT/issues).
+You can report bugs and feature requests through the rocm-libraries GitHub
+[issue tracker](https://github.com/ROCm/rocm-libraries/issues).
 
 ## Contribute
 
-If you want to contribute to hipFFT, you must follow our [contribution guidelines](https://github.com/ROCm/hipFFT/blob/develop/.github/CONTRIBUTING.md).
+If you want to contribute to hipFFT, you must follow our [contribution guidelines](https://github.com/ROCm/rocm-libraries/blob/develop/projects/hipfft/.github/CONTRIBUTING.md).
diff --git a/clients/CMakeLists.txt b/clients/CMakeLists.txt
index 739ef03..67de6e8 100644
--- a/clients/CMakeLists.txt
+++ b/clients/CMakeLists.txt
@@ -88,9 +88,9 @@ if( BUILD_CLIENTS_SAMPLES )
 endif( )
 
 if( BUILD_CLIENTS_TESTS )
-  find_package( GTest 1.11.0 )
+  find_package( GTest QUIET )
   include( ExternalProject )
-  if( NOT GTEST_FOUND )
+  if( NOT GTest_FOUND )
     set( GTEST_INCLUDE_DIRS
       ${CMAKE_CURRENT_BINARY_DIR}/src/gtest/googletest/include
       )
diff --git a/clients/bench/CMakeLists.txt b/clients/bench/CMakeLists.txt
index c6da335..4cece34 100644
--- a/clients/bench/CMakeLists.txt
+++ b/clients/bench/CMakeLists.txt
@@ -93,19 +93,3 @@ set_target_properties( hipfft-bench
   ${BENCH_OUT_DIR} )
 
 rocm_install(TARGETS hipfft-bench COMPONENT benchmarks)
-
-# install compatibility for old name of bench program - symlink on
-# unix, hardlink on windows (since privilege is required to create
-# symlinks there)
-if( WIN32 )
-  set( BENCH_LINK_COMMAND create_hardlink )
-  set( BENCH_NEW_NAME ${CMAKE_INSTALL_PREFIX}/${CMAKE_INSTALL_BINDIR}/hipfft-bench${CMAKE_EXECUTABLE_SUFFIX} )
-  set( BENCH_OLD_NAME ${CMAKE_INSTALL_PREFIX}/${CMAKE_INSTALL_BINDIR}/hipfft-rider${CMAKE_EXECUTABLE_SUFFIX} )
-else()
-  set( BENCH_LINK_COMMAND create_symlink )
-  set( BENCH_NEW_NAME hipfft-bench )
-  set( BENCH_OLD_NAME ${CMAKE_INSTALL_PREFIX}/${CMAKE_INSTALL_BINDIR}/hipfft-rider )
-endif()
-install(
-  CODE "execute_process( COMMAND \"${CMAKE_COMMAND}\" -E ${BENCH_LINK_COMMAND} \"${BENCH_NEW_NAME}\" \"${BENCH_OLD_NAME}\" )"
-)
diff --git a/clients/bench/bench.cpp b/clients/bench/bench.cpp
index c9cfed9..52cf888 100644
--- a/clients/bench/bench.cpp
+++ b/clients/bench/bench.cpp
@@ -31,6 +31,9 @@
 #include "../../shared/client_except.h"
 #include "../../shared/gpubuf.h"
 
+// initialize static class member of hipfft_params
+std::vector<gpubuf> hipfft_params::externally_managed_workareas = std::vector<gpubuf>();
+
 int main(int argc, char* argv[])
 {
     // This helps with mixing output of both wide and narrow characters to the screen
@@ -79,6 +82,11 @@ int main(int argc, char* argv[])
                      "forward\n3) real inverse")
         ->default_val(fft_transform_type_complex_forward);
     non_token
+        ->add_option("--auto_allocation",
+                     params.auto_allocate,
+                     "hipFFT's auto-allocation behavior: \"on\", \"off\", or \"default\"")
+        ->default_val("default");
+    non_token
         ->add_option(
             "--precision", params.precision, "Transform precision: single (default), double, half")
         ->excludes("--double");
diff --git a/clients/hipfft_params.h b/clients/hipfft_params.h
index e0fc576..f5bcdb4 100644
--- a/clients/hipfft_params.h
+++ b/clients/hipfft_params.h
@@ -32,12 +32,28 @@
 #include "../shared/hipfft_brick.h"
 #include "hipfft/hipfft.h"
 #include "hipfft/hipfftXt.h"
+#include <random>
 
 #ifdef HIPFFT_MPI_ENABLE
 #include "hipfft/hipfftMp.h"
 #include <mpi.h>
 #endif
 
+template <typename T,
+          typename... Args,
+          std::enable_if_t<std::is_integral_v<T> && (std::is_same_v<T, Args> && ...), bool> = true>
+static void set_with_random_nonnegative_values(const std::string& token, T& val, Args&... args)
+{
+    // using a hash of the token as random seed to avoid
+    // dependencies on externally-defined variables
+    std::hash<std::string>           hasher;
+    std::ranlux24_base               gen(hasher(token));
+    std::uniform_int_distribution<T> dis(static_cast<T>(0), std::numeric_limits<T>::max());
+    val = dis(gen);
+    ((args = dis(gen)), ...);
+    return;
+}
+
 inline fft_status fft_status_from_hipfftparams(const hipfftResult_t val)
 {
     switch(val)
@@ -141,6 +157,13 @@ public:
     std::vector<long long int> ll_inembed;
     std::vector<long long int> ll_onembed;
 
+    template <typename T>
+    struct many_api_layout_args
+    {
+        T *input_embed, *output_embed;
+        T  input_stride, output_stride, input_distance, output_distance;
+    };
+
     struct hipLibXtDesc_deleter
     {
         void operator()(hipLibXtDesc* d)
@@ -160,22 +183,46 @@ public:
 
     // backend library can write N worksize values for N GPUs, so
     // allocate a vector for that if necessary
-    std::vector<size_t> xt_worksize;
+    std::vector<size_t> auto_allocated_worksizes;
+    // if auto_allocate == fft_auto_allocation_off, the hipFFT plan(s)
+    // will be provided with externally-managed work area(s):
+    static std::vector<gpubuf> externally_managed_workareas;
 
-    // pointer we pass to the backend library.  By default point to the
-    // single-GPU workbuffer size.
-    size_t* workbuffersize_ptr;
+    size_t auto_allocated_extra_vram_footprint() const
+    {
+        return std::accumulate(auto_allocated_worksizes.begin(),
+                               auto_allocated_worksizes.end(),
+                               static_cast<size_t>(0));
+    }
+
+    static size_t externally_managed_extra_vram_footprint()
+    {
+        return std::accumulate(externally_managed_workareas.begin(),
+                               externally_managed_workareas.end(),
+                               static_cast<size_t>(0),
+                               [](size_t total, const gpubuf& buf) { return total + buf.size(); });
+    }
 
-    hipfft_params()
+    bool is_preventing_auto_allocation_at_generation() const
     {
-        workbuffersize_ptr = &workbuffersize;
+        if(auto_allocate != fft_auto_allocation_off)
+            return false;
+        // Let hipFFT sometimes auto-allocate nonetheless so that tests cover its
+        // ability to free resources (allocated at generation) when/if some
+        // externally-managed workarea(s) are provided after plan generation
+        // Note: this member function must return the same result even if called
+        // more than once by a given instance, it must be stable for any instance
+        return std::hash<std::string>()(token()) % 2 == 1;
     }
 
+    hipfft_params() = default;
+
     hipfft_params(const fft_params& p)
         : fft_params(p)
     {
-        workbuffersize_ptr = &workbuffersize;
     }
+    hipfft_params(hipfft_params&& p) = default;
+    hipfft_params& operator=(hipfft_params&& other) = default;
 
     ~hipfft_params()
     {
@@ -206,12 +253,14 @@ public:
         }
         catch(fft_params::work_buffer_alloc_failure& e)
         {
-            val += workbuffersize;
+            val += auto_allocated_extra_vram_footprint();
+            val += externally_managed_extra_vram_footprint();
             std::stringstream msg;
             msg << "Plan work buffer size (" << val << " bytes raw data) too large for device";
             throw ROCFFT_SKIP{msg.str()};
         }
-        val += workbuffersize;
+        val += auto_allocated_extra_vram_footprint();
+        val += externally_managed_extra_vram_footprint();
         return val;
     }
 
@@ -336,6 +385,11 @@ public:
             int_inembed[i] = ll_inembed[i];
             int_onembed[i] = ll_onembed[i];
         }
+        // reset auto_allocated_worksizes
+        auto_allocated_worksizes.resize(get_num_used_gpus());
+        std::for_each(auto_allocated_worksizes.begin(),
+                      auto_allocated_worksizes.end(),
+                      [](decltype(auto_allocated_worksizes)::value_type& val) { val = 0; });
 
         hipfftResult ret = HIPFFT_SUCCESS;
         return fft_status_from_hipfftparams(ret);
@@ -402,19 +456,47 @@ public:
         }
         }
 
+        if(ret == HIPFFT_SUCCESS && auto_allocate == fft_auto_allocation_off)
+        {
+            ret = set_externally_managed_work_areas();
+        }
+
         // hipFFT can fail plan creation due to allocation failure -
         // tests are expecting a specific exception in that case,
         // because the test was unable to run.  Doesn't mean the test
         // case failed.
         if(ret == HIPFFT_ALLOC_FAILED)
-            throw fft_params::work_buffer_alloc_failure(
-                "plan create failed due to allocation failure");
+        {
+            if(!final_attempt_at_plan_creation && externally_managed_extra_vram_footprint() > 0)
+            {
+                final_attempt_at_plan_creation = true;
+                // device allocation(s) in externally_managed_workareas might be
+                // larger than needed or even unnecessary for the instance of interest.
+                // Free them up and try again before concluding.
+                externally_managed_workareas.clear();
+                return create_plan();
+            }
+            else
+            {
+                throw fft_params::work_buffer_alloc_failure(
+                    "plan create failed due to allocation failure",
+                    externally_managed_extra_vram_footprint()
+                        + auto_allocated_extra_vram_footprint());
+            }
+        }
 
         // store token to check if plan was already made
         current_token = token();
         return fft_status_from_hipfftparams(ret);
     }
 
+    hipfftResult_t set_stream(hipStream_t stream)
+    {
+        if(plan == INVALID_PLAN_HANDLE)
+            throw std::runtime_error("Plan must be created before setting a desired stream");
+        return hipfftSetStream(plan, stream);
+    }
+
     void validate_fields() const override
     {
         validate_brick_volume();
@@ -454,10 +536,12 @@ public:
         }
     }
 
-    fft_status set_callbacks(void* load_cb_host,
-                             void* load_cb_data,
-                             void* store_cb_host,
-                             void* store_cb_data) override
+    fft_status set_callbacks(void*  load_cb_host,
+                             void*  load_cb_data,
+                             void*  store_cb_host,
+                             void*  store_cb_data,
+                             size_t load_cb_shared_mem_bytes  = 0,
+                             size_t store_cb_shared_mem_bytes = 0) override
     {
         if(run_callbacks)
         {
@@ -476,6 +560,16 @@ public:
                     plan, &store_cb_host, HIPFFT_CB_ST_COMPLEX, &store_cb_data);
                 if(ret != HIPFFT_SUCCESS)
                     return fft_status_from_hipfftparams(ret);
+
+                ret = hipfftXtSetCallbackSharedSize(
+                    plan, HIPFFT_CB_LD_REAL, load_cb_shared_mem_bytes);
+                if(ret != HIPFFT_SUCCESS)
+                    return fft_status_from_hipfftparams(ret);
+
+                ret = hipfftXtSetCallbackSharedSize(
+                    plan, HIPFFT_CB_ST_COMPLEX, store_cb_shared_mem_bytes);
+                if(ret != HIPFFT_SUCCESS)
+                    return fft_status_from_hipfftparams(ret);
                 break;
             case HIPFFT_D2Z:
                 ret = hipfftXtSetCallback(
@@ -487,6 +581,16 @@ public:
                     plan, &store_cb_host, HIPFFT_CB_ST_COMPLEX_DOUBLE, &store_cb_data);
                 if(ret != HIPFFT_SUCCESS)
                     return fft_status_from_hipfftparams(ret);
+
+                ret = hipfftXtSetCallbackSharedSize(
+                    plan, HIPFFT_CB_LD_REAL_DOUBLE, load_cb_shared_mem_bytes);
+                if(ret != HIPFFT_SUCCESS)
+                    return fft_status_from_hipfftparams(ret);
+
+                ret = hipfftXtSetCallbackSharedSize(
+                    plan, HIPFFT_CB_ST_COMPLEX_DOUBLE, store_cb_shared_mem_bytes);
+                if(ret != HIPFFT_SUCCESS)
+                    return fft_status_from_hipfftparams(ret);
                 break;
             case HIPFFT_C2R:
                 ret = hipfftXtSetCallback(plan, &load_cb_host, HIPFFT_CB_LD_COMPLEX, &load_cb_data);
@@ -496,6 +600,16 @@ public:
                 ret = hipfftXtSetCallback(plan, &store_cb_host, HIPFFT_CB_ST_REAL, &store_cb_data);
                 if(ret != HIPFFT_SUCCESS)
                     return fft_status_from_hipfftparams(ret);
+
+                ret = hipfftXtSetCallbackSharedSize(
+                    plan, HIPFFT_CB_LD_COMPLEX, load_cb_shared_mem_bytes);
+                if(ret != HIPFFT_SUCCESS)
+                    return fft_status_from_hipfftparams(ret);
+
+                ret = hipfftXtSetCallbackSharedSize(
+                    plan, HIPFFT_CB_ST_REAL, store_cb_shared_mem_bytes);
+                if(ret != HIPFFT_SUCCESS)
+                    return fft_status_from_hipfftparams(ret);
                 break;
             case HIPFFT_Z2D:
                 ret = hipfftXtSetCallback(
@@ -507,6 +621,16 @@ public:
                     plan, &store_cb_host, HIPFFT_CB_ST_REAL_DOUBLE, &store_cb_data);
                 if(ret != HIPFFT_SUCCESS)
                     return fft_status_from_hipfftparams(ret);
+
+                ret = hipfftXtSetCallbackSharedSize(
+                    plan, HIPFFT_CB_LD_COMPLEX_DOUBLE, load_cb_shared_mem_bytes);
+                if(ret != HIPFFT_SUCCESS)
+                    return fft_status_from_hipfftparams(ret);
+
+                ret = hipfftXtSetCallbackSharedSize(
+                    plan, HIPFFT_CB_ST_REAL_DOUBLE, store_cb_shared_mem_bytes);
+                if(ret != HIPFFT_SUCCESS)
+                    return fft_status_from_hipfftparams(ret);
                 break;
             case HIPFFT_C2C:
                 ret = hipfftXtSetCallback(plan, &load_cb_host, HIPFFT_CB_LD_COMPLEX, &load_cb_data);
@@ -517,6 +641,16 @@ public:
                     plan, &store_cb_host, HIPFFT_CB_ST_COMPLEX, &store_cb_data);
                 if(ret != HIPFFT_SUCCESS)
                     return fft_status_from_hipfftparams(ret);
+
+                ret = hipfftXtSetCallbackSharedSize(
+                    plan, HIPFFT_CB_LD_COMPLEX, load_cb_shared_mem_bytes);
+                if(ret != HIPFFT_SUCCESS)
+                    return fft_status_from_hipfftparams(ret);
+
+                ret = hipfftXtSetCallbackSharedSize(
+                    plan, HIPFFT_CB_ST_COMPLEX, store_cb_shared_mem_bytes);
+                if(ret != HIPFFT_SUCCESS)
+                    return fft_status_from_hipfftparams(ret);
                 break;
             case HIPFFT_Z2Z:
                 ret = hipfftXtSetCallback(
@@ -528,6 +662,16 @@ public:
                     plan, &store_cb_host, HIPFFT_CB_ST_COMPLEX_DOUBLE, &store_cb_data);
                 if(ret != HIPFFT_SUCCESS)
                     return fft_status_from_hipfftparams(ret);
+
+                ret = hipfftXtSetCallbackSharedSize(
+                    plan, HIPFFT_CB_LD_COMPLEX_DOUBLE, load_cb_shared_mem_bytes);
+                if(ret != HIPFFT_SUCCESS)
+                    return fft_status_from_hipfftparams(ret);
+
+                ret = hipfftXtSetCallbackSharedSize(
+                    plan, HIPFFT_CB_ST_COMPLEX_DOUBLE, store_cb_shared_mem_bytes);
+                if(ret != HIPFFT_SUCCESS)
+                    return fft_status_from_hipfftparams(ret);
                 break;
             default:
                 throw std::runtime_error("Invalid execution type");
@@ -921,17 +1065,301 @@ private:
         CREATE_XT_MAKE_PLAN_MANY,
     };
 
+    // check that worksize estimates can be successfully queried with or without a valid plan
+    hipfftResult_t check_worksize_estimate()
+    {
+        hipfftResult_t ret{HIPFFT_INTERNAL_ERROR};
+        if(!hipfft_transform_type)
+        {
+            throw std::runtime_error("Estimating worksize requires a valid type of transform");
+        }
+        std::vector<size_t> worksize_estimate(get_num_used_gpus(), absurd_init_worksize_estimate);
+        switch(get_create_type())
+        {
+        case CREATE_MAKE_PLAN_Nd:
+        {
+            switch(dim())
+            {
+            case 1:
+                if(plan == INVALID_PLAN_HANDLE)
+                    ret = hipfftEstimate1d(
+                        int_length[0], *hipfft_transform_type, nbatch, worksize_estimate.data());
+                else
+                    ret = hipfftGetSize1d(plan,
+                                          int_length[0],
+                                          *hipfft_transform_type,
+                                          nbatch,
+                                          worksize_estimate.data());
+                break;
+            case 2:
+                if(plan == INVALID_PLAN_HANDLE)
+                    ret = hipfftEstimate2d(int_length[0],
+                                           int_length[1],
+                                           *hipfft_transform_type,
+                                           worksize_estimate.data());
+                else
+                    ret = hipfftGetSize2d(plan,
+                                          int_length[0],
+                                          int_length[1],
+                                          *hipfft_transform_type,
+                                          worksize_estimate.data());
+                break;
+            case 3:
+                if(plan == INVALID_PLAN_HANDLE)
+                    ret = hipfftEstimate3d(int_length[0],
+                                           int_length[1],
+                                           int_length[2],
+                                           *hipfft_transform_type,
+                                           worksize_estimate.data());
+                else
+                    ret = hipfftGetSize3d(plan,
+                                          int_length[0],
+                                          int_length[1],
+                                          int_length[2],
+                                          *hipfft_transform_type,
+                                          worksize_estimate.data());
+                break;
+            default:
+                throw std::runtime_error("invalid dim");
+            }
+            break;
+        }
+        case CREATE_MAKE_PLAN_MANY:
+        {
+            auto layout_args = make_valid_layout_args_for_plan_many<int>();
+            if(plan == INVALID_PLAN_HANDLE)
+                ret = hipfftEstimateMany(dim(),
+                                         int_length.data(),
+                                         layout_args.input_embed,
+                                         layout_args.input_stride,
+                                         layout_args.input_distance,
+                                         layout_args.output_embed,
+                                         layout_args.output_stride,
+                                         layout_args.output_distance,
+                                         *hipfft_transform_type,
+                                         nbatch,
+                                         worksize_estimate.data());
+            else
+                ret = hipfftGetSizeMany(plan,
+                                        dim(),
+                                        int_length.data(),
+                                        layout_args.input_embed,
+                                        layout_args.input_stride,
+                                        layout_args.input_distance,
+                                        layout_args.output_embed,
+                                        layout_args.output_stride,
+                                        layout_args.output_distance,
+                                        *hipfft_transform_type,
+                                        nbatch,
+                                        worksize_estimate.data());
+            break;
+        }
+        case CREATE_MAKE_PLAN_MANY64:
+        {
+            if(plan == INVALID_PLAN_HANDLE)
+            {
+                // no direct equivalent in estimate-fetching APIs
+                std::for_each(worksize_estimate.begin(),
+                              worksize_estimate.end(),
+                              [](decltype(worksize_estimate)::value_type& val) { val = 0; });
+                ret = HIPFFT_SUCCESS;
+            }
+            else
+            {
+                auto layout_args = make_valid_layout_args_for_plan_many<long long>();
+                ret              = hipfftGetSizeMany64(plan,
+                                          dim(),
+                                          ll_length.data(),
+                                          layout_args.input_embed,
+                                          layout_args.input_stride,
+                                          layout_args.input_distance,
+                                          layout_args.output_embed,
+                                          layout_args.output_stride,
+                                          layout_args.output_distance,
+                                          *hipfft_transform_type,
+                                          nbatch,
+                                          worksize_estimate.data());
+            }
+            break;
+        }
+        case CREATE_XT_MAKE_PLAN_MANY:
+        {
+            if(plan == INVALID_PLAN_HANDLE)
+            {
+                // no direct equivalent in estimate-fetching APIs
+                std::for_each(worksize_estimate.begin(),
+                              worksize_estimate.end(),
+                              [](decltype(worksize_estimate)::value_type& val) { val = 0; });
+                ret = HIPFFT_SUCCESS;
+            }
+            else
+            {
+                auto executionType = get_xt_api_execution_type();
+                auto layout_args   = make_valid_layout_args_for_plan_many<long long>();
+                ret                = hipfftXtGetSizeMany(plan,
+                                          dim(),
+                                          ll_length.data(),
+                                          layout_args.input_embed,
+                                          layout_args.input_stride,
+                                          layout_args.input_distance,
+                                          inputType,
+                                          layout_args.output_embed,
+                                          layout_args.output_stride,
+                                          layout_args.output_distance,
+                                          outputType,
+                                          nbatch,
+                                          worksize_estimate.data(),
+                                          executionType);
+            }
+            break;
+        }
+        case PLAN_Nd:
+        case PLAN_MANY:
+        default:
+        {
+            // should be indirectly disabled via get_create_type()
+            return HIPFFT_INTERNAL_ERROR;
+        }
+        }
+        // check that the value(s) of worksize_estimate were actually set, assuming that
+        // setting a worksize_estimate equal to absurd_init_worksize_estimate by hipFFT
+        // cannot be considered "correct".
+        // Note: worksize_estimate value(s) are *not* guaranteed to be greater than or equal
+        // to the actual value(s) of the work area(s), queriable after plan generation via
+        // hipfftGetSize.
+        if(ret == HIPFFT_SUCCESS)
+        {
+            // the estimate can't have any knowledge about the number of GPUs being used if
+            // the plan wasn't created first
+            const size_t num_values_to_check
+                = plan == INVALID_PLAN_HANDLE ? 1 : worksize_estimate.size();
+            for(size_t idx = 0; ret == HIPFFT_SUCCESS && idx < num_values_to_check; idx++)
+            {
+                ret = worksize_estimate[idx] != absurd_init_worksize_estimate
+                          ? HIPFFT_SUCCESS
+                          : HIPFFT_INTERNAL_ERROR;
+            }
+        }
+        return ret;
+    }
+
+    // provide a work area to a successfully generated plan
+    hipfftResult_t set_externally_managed_work_areas()
+    {
+        std::vector<size_t> req_workarea_sizes(get_num_used_gpus(), absurd_init_worksize_estimate);
+        hipfftResult_t      ret = hipfftGetSize(plan, req_workarea_sizes.data());
+        if(ret != HIPFFT_SUCCESS)
+        {
+            return ret;
+        }
+        else if(std::any_of(req_workarea_sizes.begin(),
+                            req_workarea_sizes.end(),
+                            [](const decltype(req_workarea_sizes)::value_type& val) {
+                                return val == absurd_init_worksize_estimate;
+                            }))
+        {
+            return HIPFFT_INTERNAL_ERROR;
+        }
+        // req_workarea_sizes are known and validated
+        // check if the current externally_managed_workareas can be used as is or not
+        if(externally_managed_workareas.size() < get_num_used_gpus())
+            externally_managed_workareas.resize(get_num_used_gpus());
+        std::vector<void*> workareas(get_num_used_gpus(), nullptr);
+        for(auto workarea_idx = 0; workarea_idx < get_num_used_gpus(); workarea_idx++)
+        {
+            const auto req_size = req_workarea_sizes[workarea_idx];
+            auto&      buf      = externally_managed_workareas[workarea_idx];
+            if(buf.size() < req_size)
+            {
+                // too small, free and reallocate to meet current needs
+                buf.free();
+                if(buf.alloc(req_size) != hipSuccess)
+                {
+                    return HIPFFT_ALLOC_FAILED;
+                }
+            }
+            workareas[workarea_idx] = buf.data();
+        }
+        if(get_num_used_gpus() > 1)
+        {
+            // TODO: enable below once hipfftXtSetWorkArea is enabled
+#if(0)
+            ret = hipfftXtSetWorkArea(plan, workareas.data);
+#else
+            throw unimplemented_exception(
+                "No implementation support for externally-managed work areas with multi-gpu usage");
+#endif
+        }
+        else
+        {
+            ret = hipfftSetWorkArea(plan, workareas[0]);
+        }
+        if(ret == HIPFFT_SUCCESS)
+        {
+            // the above "SetWorkArea" frees auto_allocated worksizes (if any)
+            auto_allocated_worksizes.clear();
+        }
+        return ret;
+    }
+
     // return true if we need to use hipFFT APIs that separate plan
     // allocation and plan init
     bool need_separate_create_make() const
     {
-        // scale factor and multi-GPU need API calls between create +
-        // init
-        if(scale_factor != 1.0 || multiGPU > 1 || mp_lib != fft_mp_lib_none)
+        // scale factor and multi-GPU and disabled auto-allocation need API
+        // calls between create + init
+        if(scale_factor != 1.0 || multiGPU > 1 || mp_lib != fft_mp_lib_none
+           || auto_allocate == fft_auto_allocation_off)
             return true;
         return false;
     }
 
+    template <
+        typename T,
+        std::enable_if_t<std::is_same_v<T, int> || std::is_same_v<T, long long int>, bool> = true>
+    many_api_layout_args<T> make_valid_layout_args_for_plan_many()
+    {
+        many_api_layout_args<T> ret;
+        if constexpr(std::is_same_v<T, int>)
+        {
+            ret.input_embed  = int_inembed.data();
+            ret.output_embed = int_onembed.data();
+        }
+        else
+        {
+            ret.input_embed  = ll_inembed.data();
+            ret.output_embed = ll_onembed.data();
+        }
+        ret.input_stride    = static_cast<T>(istride.back());
+        ret.output_stride   = static_cast<T>(ostride.back());
+        ret.input_distance  = static_cast<T>(idist);
+        ret.output_distance = static_cast<T>(odist);
+        if(is_using_default_layout())
+        {
+            // If using a default layout, users can
+            // (A) either set explicitly inembed, onembed, strides, and distances (like above);
+            // (B) or use nullptr as arguments for inembed and onembed. Strides and
+            //     distances are supposed to be ignored in that case.
+            // --> choose randomly between either valid usage when a default layout is
+            //     used, so that all possible valid use case scenarios are considered.
+            const std::string test_token = token();
+            int               randomizer;
+            set_with_random_nonnegative_values(test_token, randomizer);
+            if(randomizer % 2 == 0)
+            {
+                ret.input_embed  = nullptr;
+                ret.output_embed = nullptr;
+                // FIXME: negative values are not truly ignored for now.
+                set_with_random_nonnegative_values(test_token,
+                                                   ret.input_stride,
+                                                   ret.output_stride,
+                                                   ret.input_distance,
+                                                   ret.output_distance);
+            }
+        }
+        return ret;
+    }
+
     // Not all plan options work with all creation types.  Return a
     // suitable plan creation type for the current FFT parameters.
     int get_create_type()
@@ -996,15 +1424,16 @@ private:
     }
     hipfftResult_t create_plan_many()
     {
-        auto ret = hipfftPlanMany(&plan,
+        auto layout_args = make_valid_layout_args_for_plan_many<int>();
+        auto ret         = hipfftPlanMany(&plan,
                                   dim(),
                                   int_length.data(),
-                                  int_inembed.data(),
-                                  istride.back(),
-                                  idist,
-                                  int_onembed.data(),
-                                  ostride.back(),
-                                  odist,
+                                  layout_args.input_embed,
+                                  layout_args.input_stride,
+                                  layout_args.input_distance,
+                                  layout_args.output_embed,
+                                  layout_args.output_stride,
+                                  layout_args.output_distance,
                                   *hipfft_transform_type,
                                   nbatch);
         return ret;
@@ -1014,7 +1443,14 @@ private:
     // relevant pre-Make APIs (scale factor, XtSetGPUs)
     hipfftResult_t create_with_pre_make()
     {
-        auto ret = hipfftCreate(&plan);
+        hipfftResult_t ret{HIPFFT_INVALID_PLAN};
+        if(auto_allocate == fft_auto_allocation_off)
+        {
+            ret = check_worksize_estimate(); // read worksize estimate before plan creation
+            if(ret != HIPFFT_SUCCESS)
+                return ret;
+        }
+        ret = hipfftCreate(&plan);
         if(ret != HIPFFT_SUCCESS)
             return ret;
         if(scale_factor != 1.0)
@@ -1026,7 +1462,8 @@ private:
         if(multiGPU > 1)
         {
             int deviceCount = 0;
-            (void)hipGetDeviceCount(&deviceCount);
+            if(hipGetDeviceCount(&deviceCount) != hipSuccess)
+                throw std::runtime_error("hipGetDeviceCount failed");
 
             // ensure that users request less than or equal to the total number of devices
             if(static_cast<int>(multiGPU) > deviceCount)
@@ -1035,9 +1472,8 @@ private:
             std::vector<int> GPUs(multiGPU);
             std::iota(GPUs.begin(), GPUs.end(), 0);
             ret = hipfftXtSetGPUs(plan, static_cast<int>(multiGPU), GPUs.data());
-
-            xt_worksize.resize(GPUs.size());
-            workbuffersize_ptr = xt_worksize.data();
+            if(ret != HIPFFT_SUCCESS)
+                return ret;
         }
         if(mp_lib == fft_mp_lib_mpi)
         {
@@ -1090,34 +1526,52 @@ private:
                                               input_stride.data(),
                                               output_stride.data());
             }
+            if(ret != HIPFFT_SUCCESS)
+                return ret;
 #else
             throw std::runtime_error("MPI is not enabled");
 #endif
         }
+        if(auto_allocate == fft_auto_allocation_off)
+        {
+            ret = check_worksize_estimate(); // read worksize estimate again after plan creation
+            if(ret != HIPFFT_SUCCESS)
+                return ret;
+        }
+        if(is_preventing_auto_allocation_at_generation())
+        {
+            ret = hipfftSetAutoAllocation(plan, 0);
+        }
+
         return ret;
     }
 
     hipfftResult_t create_make_plan_Nd()
     {
         auto ret = create_with_pre_make();
+
         if(ret != HIPFFT_SUCCESS)
             return ret;
-
+        // do not register plan's worksizes as "auto-allocated" if auto-allocation was explicitly prevented
+        std::vector<size_t> tmp_worksize(get_num_used_gpus());
+        size_t*             worksize_ptr = is_preventing_auto_allocation_at_generation()
+                                               ? tmp_worksize.data()
+                                               : auto_allocated_worksizes.data();
         switch(dim())
         {
         case 1:
             return hipfftMakePlan1d(
-                plan, int_length[0], *hipfft_transform_type, nbatch, workbuffersize_ptr);
+                plan, int_length[0], *hipfft_transform_type, nbatch, worksize_ptr);
         case 2:
             return hipfftMakePlan2d(
-                plan, int_length[0], int_length[1], *hipfft_transform_type, workbuffersize_ptr);
+                plan, int_length[0], int_length[1], *hipfft_transform_type, worksize_ptr);
         case 3:
             return hipfftMakePlan3d(plan,
                                     int_length[0],
                                     int_length[1],
                                     int_length[2],
                                     *hipfft_transform_type,
-                                    workbuffersize_ptr);
+                                    worksize_ptr);
         default:
             throw std::runtime_error("invalid dim");
         }
@@ -1128,18 +1582,24 @@ private:
         auto ret = create_with_pre_make();
         if(ret != HIPFFT_SUCCESS)
             return ret;
+        // do not register plan's worksizes as "auto-allocated" if auto-allocation was explicitly prevented
+        std::vector<size_t> tmp_worksize(get_num_used_gpus());
+        size_t*             worksize_ptr = is_preventing_auto_allocation_at_generation()
+                                               ? tmp_worksize.data()
+                                               : auto_allocated_worksizes.data();
+        auto                layout_args  = make_valid_layout_args_for_plan_many<int>();
         return hipfftMakePlanMany(plan,
                                   dim(),
                                   int_length.data(),
-                                  int_inembed.data(),
-                                  istride.back(),
-                                  idist,
-                                  int_onembed.data(),
-                                  ostride.back(),
-                                  odist,
+                                  layout_args.input_embed,
+                                  layout_args.input_stride,
+                                  layout_args.input_distance,
+                                  layout_args.output_embed,
+                                  layout_args.output_stride,
+                                  layout_args.output_distance,
                                   *hipfft_transform_type,
                                   nbatch,
-                                  workbuffersize_ptr);
+                                  worksize_ptr);
     }
 
     hipfftResult_t create_make_plan_many64()
@@ -1147,58 +1607,85 @@ private:
         auto ret = create_with_pre_make();
         if(ret != HIPFFT_SUCCESS)
             return ret;
+
+        // do not register plan's worksizes as "auto-allocated" if auto-allocation was explicitly prevented
+        std::vector<size_t> tmp_worksize(get_num_used_gpus());
+        size_t*             worksize_ptr = is_preventing_auto_allocation_at_generation()
+                                               ? tmp_worksize.data()
+                                               : auto_allocated_worksizes.data();
+        auto                layout_args  = make_valid_layout_args_for_plan_many<long long int>();
         return hipfftMakePlanMany64(plan,
                                     dim(),
                                     ll_length.data(),
-                                    ll_inembed.data(),
-                                    istride.back(),
-                                    idist,
-                                    ll_onembed.data(),
-                                    ostride.back(),
-                                    odist,
+                                    layout_args.input_embed,
+                                    layout_args.input_stride,
+                                    layout_args.input_distance,
+                                    layout_args.output_embed,
+                                    layout_args.output_stride,
+                                    layout_args.output_distance,
                                     *hipfft_transform_type,
                                     nbatch,
-                                    workbuffersize_ptr);
+                                    worksize_ptr);
     }
 
-    hipfftResult_t create_xt_make_plan_many()
+    hipDataType get_xt_api_execution_type() const
     {
-        auto ret = create_with_pre_make();
-        if(ret != HIPFFT_SUCCESS)
-            return ret;
-
         // execution type is always complex, matching the precision
         // of the transform
         // Initializing as double by default
-        hipDataType executionType = HIP_C_64F;
+        hipDataType ret = HIP_C_64F;
         switch(precision)
         {
         case fft_precision_half:
-            executionType = HIP_C_16F;
+            ret = HIP_C_16F;
             break;
         case fft_precision_single:
-            executionType = HIP_C_32F;
+            ret = HIP_C_32F;
             break;
         case fft_precision_double:
-            executionType = HIP_C_64F;
+            ret = HIP_C_64F;
             break;
+        default:
+            throw std::runtime_error("Invalid precision");
         }
+        return ret;
+    }
 
+    hipfftResult_t create_xt_make_plan_many()
+    {
+        auto ret = create_with_pre_make();
+        if(ret != HIPFFT_SUCCESS)
+            return ret;
+
+        // do not register plan's worksizes as "auto-allocated" if auto-allocation was explicitly prevented
+        std::vector<size_t> tmp_worksize(get_num_used_gpus());
+        size_t*             worksize_ptr  = is_preventing_auto_allocation_at_generation()
+                                                ? tmp_worksize.data()
+                                                : auto_allocated_worksizes.data();
+        auto                executionType = get_xt_api_execution_type();
+        auto                layout_args   = make_valid_layout_args_for_plan_many<long long int>();
         return hipfftXtMakePlanMany(plan,
                                     dim(),
                                     ll_length.data(),
-                                    ll_inembed.data(),
-                                    istride.back(),
-                                    idist,
+                                    layout_args.input_embed,
+                                    layout_args.input_stride,
+                                    layout_args.input_distance,
                                     inputType,
-                                    ll_onembed.data(),
-                                    ostride.back(),
-                                    odist,
+                                    layout_args.output_embed,
+                                    layout_args.output_stride,
+                                    layout_args.output_distance,
                                     outputType,
                                     nbatch,
-                                    workbuffersize_ptr,
+                                    worksize_ptr,
                                     executionType);
     }
+    static constexpr size_t absurd_init_worksize_estimate  = std::numeric_limits<size_t>::max();
+    bool                    final_attempt_at_plan_creation = false;
+
+    size_t get_num_used_gpus() const
+    {
+        return multiGPU > 1 ? multiGPU : 1;
+    };
 };
 
 #endif
diff --git a/clients/hipfftw_helper.h b/clients/hipfftw_helper.h
new file mode 100644
index 0000000..3e86134
--- /dev/null
+++ b/clients/hipfftw_helper.h
@@ -0,0 +1,1238 @@
+// Copyright (C) 2025 Advanced Micro Devices, Inc. All rights reserved.
+//
+// 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.
+
+#ifndef HIPFFTW_HELPER_H
+#define HIPFFTW_HELPER_H
+
+#include "../shared/environment.h"
+#include "../shared/fft_params.h"
+#include <algorithm>
+#include <fftw3.h>
+#include <memory>
+#include <sstream>
+#include <stdexcept>
+#include <string>
+#include <type_traits>
+
+#ifdef WIN32
+#include <windows.h>
+// psapi.h requires windows.h to be included first
+#include <psapi.h>
+typedef HMODULE LIB_HANDLE_T;
+#else
+#include <dlfcn.h>
+#include <link.h>
+typedef void* LIB_HANDLE_T;
+#endif
+
+template <fft_precision prec>
+struct hipfftw_trait;
+template <>
+struct hipfftw_trait<fft_precision_single>
+{
+    using plan_t    = fftwf_plan;
+    using complex_t = fftwf_complex;
+    using real_t    = float;
+};
+template <>
+struct hipfftw_trait<fft_precision_double>
+{
+    using plan_t    = fftw_plan;
+    using complex_t = fftw_complex;
+    using real_t    = double;
+};
+
+template <fft_precision prec>
+using hipfftw_real_t = typename hipfftw_trait<prec>::real_t;
+template <fft_precision prec>
+using hipfftw_complex_t = typename hipfftw_trait<prec>::complex_t;
+template <fft_precision prec>
+using hipfftw_plan_t = typename hipfftw_trait<prec>::plan_t;
+
+// singleton class encapsulating the dynamically-loaded hipfftw library
+class dynamically_loaded_hipfftw
+{
+private:
+    LIB_HANDLE_T       lib_handle;
+    std::ostringstream load_error_info;
+
+    dynamically_loaded_hipfftw()
+    {
+#ifdef __HIP_PLATFORM_AMD__
+        const std::string lib_basename = "hipfftw";
+#else
+        const std::string lib_basename = "cufftw";
+#endif
+#ifdef WIN32
+        const std::string lib_fullame = lib_basename + ".dll";
+        lib_handle                    = LoadLibraryA(lib_fullame.c_str());
+#else
+        const std::string lib_fullame  = "lib" + lib_basename + ".so";
+        lib_handle                     = dlopen(lib_fullame.c_str(), RTLD_LAZY);
+#endif
+        load_error_info.clear();
+        if(!lib_handle)
+        {
+            load_error_info << "failed to open library " << lib_fullame;
+#ifdef WIN32
+            load_error_info << ". System's error code = " << GetLastError();
+#else
+            load_error_info << ". System's error message = " << dlerror();
+#endif
+            // do not throw from here to ease exception handling
+        }
+    }
+    /* disable copies and moves */
+    dynamically_loaded_hipfftw(const dynamically_loaded_hipfftw&) = delete;
+    dynamically_loaded_hipfftw(dynamically_loaded_hipfftw&&)      = delete;
+    dynamically_loaded_hipfftw& operator=(const dynamically_loaded_hipfftw&) = delete;
+    dynamically_loaded_hipfftw& operator=(dynamically_loaded_hipfftw&&) = delete;
+
+    static const dynamically_loaded_hipfftw& get_instance()
+    {
+        static dynamically_loaded_hipfftw singleton_instance;
+        return singleton_instance;
+    }
+
+public:
+    static LIB_HANDLE_T get_lib()
+    {
+        return get_instance().lib_handle;
+    }
+    static std::string get_load_error_info()
+    {
+        return get_instance().load_error_info.str();
+    }
+    ~dynamically_loaded_hipfftw()
+    {
+        if(lib_handle)
+        {
+#ifdef WIN32
+            (void)FreeLibrary(lib_handle);
+#else
+            (void)dlclose(lib_handle);
+#endif
+        }
+        lib_handle = nullptr;
+    }
+};
+
+// exception specific to issues when loading hipfftw and/or when fetching
+// the address of the supposedly-available functions therefrom
+struct hipfftw_undefined_function_ptr : std::runtime_error
+{
+    using std::runtime_error::runtime_error;
+};
+
+// helper struct for retrieving a function's return type
+template <class T>
+struct func_ret;
+template <typename R, class... Args>
+struct func_ret<R(Args...)>
+{
+    using type = R;
+};
+template <class T>
+using func_ret_t = typename func_ret<T>::type;
+
+template <typename func_type, std::enable_if_t<std::is_function_v<func_type>, bool> = true>
+struct dynamically_loaded_function_t
+{
+private:
+    // address of the desired function, to be fetched from a dynamically loaded shared library
+    func_type* func_ptr;
+    // address of the reference function (from linked fftw3)
+    func_type* const reference_func_ptr;
+    // symbol of said function
+    std::string func_symbol;
+
+public:
+    dynamically_loaded_function_t(const char* symbol, func_type* ref_func_address)
+        : func_ptr(nullptr)
+        , reference_func_ptr(ref_func_address)
+        , func_symbol(symbol)
+    {
+    }
+
+    // forwarding functional calls
+    template <typename... Args>
+    func_ret_t<func_type> operator()(Args... args) const
+    {
+        if(!may_be_used())
+            throw hipfftw_undefined_function_ptr(dynamically_loaded_hipfftw::get_load_error_info());
+        return func_ptr(args...);
+    }
+    template <bool call_reference, typename... Args>
+    func_ret_t<func_type> call(Args... args) const
+    {
+        if constexpr(!call_reference)
+        {
+            return this->operator()(args...);
+        }
+        else
+        {
+            if(!reference_func_ptr)
+                throw hipfftw_undefined_function_ptr(
+                    "Ill-defined reference function pointer for symbol " + func_symbol);
+            return reference_func_ptr(args...);
+        }
+        // unreachable
+    }
+    void load_implementation()
+    {
+        const auto hipfftw_lib = dynamically_loaded_hipfftw::get_lib();
+        if(!hipfftw_lib)
+        {
+            // make func_ptr unambiguously unset to force the dedicated exception
+            // to be thrown at forwarded functional call(s)
+            func_ptr = nullptr;
+            return;
+        }
+#ifdef WIN32
+        func_ptr = reinterpret_cast<func_type*>(GetProcAddress(hipfftw_lib, func_symbol.c_str()));
+#else
+        func_ptr = reinterpret_cast<func_type*>(dlsym(hipfftw_lib, func_symbol.c_str()));
+#endif
+    }
+    bool may_be_used() const
+    {
+        return func_ptr != nullptr;
+    }
+    std::string get_symbol() const
+    {
+        return func_symbol;
+    }
+};
+
+template <typename T, typename... Args>
+static void load_implementations(dynamically_loaded_function_t<T>& first, Args&... others)
+{
+    first.load_implementation();
+    if constexpr(sizeof...(others) > 0)
+        load_implementations(others...);
+}
+
+// define singleton structures encapsulating all the hipfftw function
+// pointers (one specialization per supported precision)
+template <fft_precision prec>
+struct hipfftw_funcs;
+
+#define HIPFFTW_STRINGIFY(x) #x
+#define HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, func)                        \
+    dynamically_loaded_function_t<decltype(prefix##func)> func                                   \
+        = dynamically_loaded_function_t<decltype(prefix##func)>(HIPFFTW_STRINGIFY(prefix##func), \
+                                                                &(prefix##func));
+
+#define HIPFFTW_FUNCS_SPECIALIZATION(prefix, specialization)                         \
+    template <>                                                                      \
+    struct hipfftw_funcs<specialization>                                             \
+    {                                                                                \
+    private:                                                                         \
+        hipfftw_funcs()                                                              \
+        {                                                                            \
+            load_implementations(malloc,                                             \
+                                 alloc_real,                                         \
+                                 alloc_complex,                                      \
+                                 free,                                               \
+                                 destroy_plan,                                       \
+                                 cleanup,                                            \
+                                 execute,                                            \
+                                 plan_dft_1d,                                        \
+                                 plan_dft_2d,                                        \
+                                 plan_dft_3d,                                        \
+                                 plan_dft,                                           \
+                                 plan_dft_r2c_1d,                                    \
+                                 plan_dft_r2c_2d,                                    \
+                                 plan_dft_r2c_3d,                                    \
+                                 plan_dft_r2c,                                       \
+                                 plan_dft_c2r_1d,                                    \
+                                 plan_dft_c2r_2d,                                    \
+                                 plan_dft_c2r_3d,                                    \
+                                 plan_dft_c2r,                                       \
+                                 print_plan,                                         \
+                                 set_timelimit,                                      \
+                                 cost,                                               \
+                                 flops);                                             \
+        }                                                                            \
+        /* disable copies and moves */                                               \
+        hipfftw_funcs(const hipfftw_funcs&) = delete;                                \
+        hipfftw_funcs& operator=(const hipfftw_funcs&) = delete;                     \
+        hipfftw_funcs(hipfftw_funcs&&)                 = delete;                     \
+        hipfftw_funcs& operator=(hipfftw_funcs&&) = delete;                          \
+                                                                                     \
+    public:                                                                          \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, malloc)          \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, alloc_real)      \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, alloc_complex)   \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, free)            \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, destroy_plan)    \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, cleanup)         \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, execute)         \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, plan_dft_1d)     \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, plan_dft_2d)     \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, plan_dft_3d)     \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, plan_dft)        \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, plan_dft_r2c_1d) \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, plan_dft_r2c_2d) \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, plan_dft_r2c_3d) \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, plan_dft_r2c)    \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, plan_dft_c2r_1d) \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, plan_dft_c2r_2d) \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, plan_dft_c2r_3d) \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, plan_dft_c2r)    \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, print_plan)      \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, set_timelimit)   \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, cost)            \
+        HIPFFTW_DECLARE_DYNAMICALLY_LOADED_FUNCTION_POINTER(prefix, flops)           \
+        static const hipfftw_funcs& get_instance()                                   \
+        {                                                                            \
+            static const hipfftw_funcs instance;                                     \
+            return instance;                                                         \
+        }                                                                            \
+    }
+
+HIPFFTW_FUNCS_SPECIALIZATION(fftwf_, fft_precision_single);
+HIPFFTW_FUNCS_SPECIALIZATION(fftw_, fft_precision_double);
+
+// structure enabling verbosity for hipfftw's exception handler and redirecting std::cerr
+// to a runtime buffer throughout its lifetime (unless it was already enabled prior/externally)
+struct hipfftw_exception_logger
+{
+    bool                  active;
+    std::stringstream     buffer;
+    std::streambuf* const original_cerr_rdbuf = nullptr;
+
+    std::unique_ptr<EnvironmentSetTemp> hipfftw_temp_logger_env;
+
+public:
+    hipfftw_exception_logger()
+        : active(false)
+        , original_cerr_rdbuf(std::cerr.rdbuf())
+    {
+#ifdef __HIP_PLATFORM_AMD__
+        const auto env_val = rocfft_getenv("HIPFFTW_LOG_EXCEPTIONS");
+        // activate temporary redirection only if not already used otherwise
+        // (e.g., in test user's environment )
+        if(env_val.empty() || std::stoull(env_val) == 0)
+        {
+            hipfftw_temp_logger_env
+                = std::make_unique<EnvironmentSetTemp>("HIPFFTW_LOG_EXCEPTIONS", "1");
+            const auto temp_env_val = rocfft_getenv("HIPFFTW_LOG_EXCEPTIONS");
+            active                  = !temp_env_val.empty() && std::stoull(temp_env_val) != 0;
+        }
+#endif
+        if(active)
+            std::cerr.rdbuf(buffer.rdbuf());
+    }
+    hipfftw_exception_logger(const hipfftw_exception_logger&) = delete;
+    hipfftw_exception_logger(hipfftw_exception_logger&&)      = delete;
+    hipfftw_exception_logger& operator=(const hipfftw_exception_logger&) = delete;
+    hipfftw_exception_logger& operator=(hipfftw_exception_logger&&) = delete;
+    ~hipfftw_exception_logger()
+    {
+        if(active)
+        {
+            // restore cerr to its original state
+            std::cerr.rdbuf(original_cerr_rdbuf);
+        }
+    }
+    bool is_active() const
+    {
+        return active;
+    }
+    std::string get_log() const
+    {
+        return buffer.str();
+    }
+};
+
+// bit-flagging enum used for representing (combinations of) plan creation
+// function(s) to consider
+enum hipfftw_plan_creation_func : unsigned
+{
+    NONE        = 0x0, // not to be used (exceptfor validating values)
+    PLAN_DFT_ND = 0x1 << 0,
+    PLAN_DFT    = 0x1 << 1,
+    PLAN_MANY   = 0x1 << 2,
+    PLAN_GURU   = 0x1 << 3,
+    PLAN_GURU64 = 0x1 << 4,
+    ANY         = PLAN_DFT_ND | PLAN_DFT | PLAN_MANY | PLAN_GURU | PLAN_GURU64
+};
+static const std::vector<hipfftw_plan_creation_func> hipfftw_plan_creation_func_candidates
+    = {hipfftw_plan_creation_func::PLAN_DFT_ND,
+       hipfftw_plan_creation_func::PLAN_DFT,
+       hipfftw_plan_creation_func::PLAN_MANY,
+       hipfftw_plan_creation_func::PLAN_GURU,
+       hipfftw_plan_creation_func::PLAN_GURU64};
+
+static bool hipfftw_creation_options_are_well_defined(hipfftw_plan_creation_func creation_options)
+{
+    return creation_options == (creation_options & hipfftw_plan_creation_func::ANY);
+}
+
+static std::string hipfftw_creation_options_to_string(hipfftw_plan_creation_func creation_options,
+                                                      fft_transform_type         dft_type,
+                                                      int                        intended_rank)
+{
+    if(!hipfftw_creation_options_are_well_defined(creation_options))
+        throw std::invalid_argument(
+            "invalid creation_options for hipfftw_creation_options_to_string");
+    if(creation_options == hipfftw_plan_creation_func::NONE)
+        return "none";
+    if(creation_options == hipfftw_plan_creation_func::ANY)
+        return "any";
+    if(std::find(hipfftw_plan_creation_func_candidates.begin(),
+                 hipfftw_plan_creation_func_candidates.end(),
+                 creation_options)
+       == hipfftw_plan_creation_func_candidates.end())
+    {
+        // 2 or more qualifying candidates flagged in creation_options
+        std::string ret;
+        for(auto candidate : hipfftw_plan_creation_func_candidates)
+        {
+            if(creation_options & candidate)
+            {
+                if(!ret.empty())
+                    ret += "_or_";
+                ret += hipfftw_creation_options_to_string(candidate, dft_type, intended_rank);
+            }
+        }
+        return ret;
+    }
+    // creation_options is one unique qualifying candidate
+    std::ostringstream ret;
+    const std::string  real_or_empty_qualifier
+        = is_real(dft_type) ? (is_fwd(dft_type) ? "_r2c" : "_c2r") : "";
+    switch(creation_options)
+    {
+    case hipfftw_plan_creation_func::PLAN_DFT_ND:
+        ret << "plan_dft" << real_or_empty_qualifier << "_" << (intended_rank < 0 ? "negative" : "")
+            << std::abs(intended_rank) << "d";
+        break;
+    case hipfftw_plan_creation_func::PLAN_DFT:
+        ret << "plan_dft" << real_or_empty_qualifier;
+        break;
+    case hipfftw_plan_creation_func::PLAN_MANY:
+        ret << "plan_many_dft" << real_or_empty_qualifier;
+        break;
+    case hipfftw_plan_creation_func::PLAN_GURU:
+        ret << "plan_guru_dft" << real_or_empty_qualifier;
+        break;
+    case hipfftw_plan_creation_func::PLAN_GURU64:
+        ret << "plan_guru64_dft" << real_or_empty_qualifier;
+        break;
+    default:
+        throw std::runtime_error("hipfftw_creation_options_to_string: internal error encountered "
+                                 "(unexpected value for creation_options)");
+        break;
+    }
+    return ret.str();
+}
+
+template <
+    fft_precision prec,
+    std::enable_if_t<prec == fft_precision_single || prec == fft_precision_double, bool> = true>
+struct hipfftw_plan_bundle_t
+{
+private:
+    const decltype(hipfftw_funcs<prec>::destroy_plan)& plan_destructor;
+
+public:
+    hipfftw_plan_t<prec>       plan;
+    std::pair<void*, void*>    creation_io; // not owned
+    hipfftw_plan_creation_func creation_func;
+    std::string                plan_token; // <-- plan details, except for creation io data pointers
+    hipfftw_plan_bundle_t(decltype(plan_destructor) plan_destructor_func)
+        : plan_destructor(plan_destructor_func)
+        , plan(nullptr)
+        , creation_io({nullptr, nullptr})
+        , creation_func(hipfftw_plan_creation_func::NONE)
+        , plan_token("")
+    {
+    }
+    ~hipfftw_plan_bundle_t()
+    {
+        // make sure the plan destructor may be used to avoid
+        // throwing from the hipfftw_plan_bundle_t destructor
+        if(plan_destructor.may_be_used())
+        {
+            // should be stable even if plan == nullptr;
+            plan_destructor(plan);
+        }
+        else if(plan)
+        {
+            std::cerr << "WARNING: A " << (prec == fft_precision_single ? "single" : "double")
+                      << "-precision plan was seemingly created but its destructor cannot be used "
+                      << std::endl;
+        }
+    }
+    // disable copies and moves
+    hipfftw_plan_bundle_t(const hipfftw_plan_bundle_t&) = delete;
+    hipfftw_plan_bundle_t& operator=(const hipfftw_plan_bundle_t&) = delete;
+    hipfftw_plan_bundle_t(hipfftw_plan_bundle_t&&)                 = delete;
+    hipfftw_plan_bundle_t& operator=(hipfftw_plan_bundle_t&&) = delete;
+};
+
+static bool rank_is_valid_for_hipfftw(int r)
+{
+    return r > 0;
+}
+template <typename T, std::enable_if_t<std::is_integral_v<T>, bool> = true>
+static bool lengths_are_valid_for_hipfftw_as(const std::vector<ptrdiff_t> len, int intended_rank)
+{
+    if(!rank_is_valid_for_hipfftw(intended_rank))
+        return false; // impossible to validate lengths for an invalid rank
+    // check that lengths are all strictly positive and representable with
+    // type T without data loss
+    return len.size() == intended_rank
+           && std::all_of(len.begin(), len.end(), [](const decltype(len)::value_type& val) {
+                  return val > 0 && val <= std::numeric_limits<T>::max();
+              });
+}
+static bool sign_is_valid_for_hipfftw(int s, const fft_transform_type& dft_kind)
+{
+    if(is_real(dft_kind))
+        return true; // sign is irrelevant for real transforms
+    return s == (is_fwd(dft_kind) ? FFTW_FORWARD : FFTW_BACKWARD);
+}
+static constexpr unsigned hipfftw_valid_flags_mask
+    = FFTW_WISDOM_ONLY | FFTW_MEASURE | FFTW_DESTROY_INPUT | FFTW_UNALIGNED | FFTW_CONSERVE_MEMORY
+      | FFTW_EXHAUSTIVE | FFTW_PRESERVE_INPUT | FFTW_PATIENT | FFTW_ESTIMATE;
+static bool flags_are_valid_for_hipfftw(unsigned f)
+{
+    return (f & hipfftw_valid_flags_mask) == f;
+}
+
+template <
+    fft_precision prec,
+    std::enable_if_t<prec == fft_precision_single || prec == fft_precision_double, bool> = true>
+struct hipfftw_helper
+{
+private:
+    // plan_bundle stores information about the latest plan possibly created by this
+    // object. A shard_ptr is used to make hipfftw_helper safe w.r.t. shallow
+    // copies (as required by gtest for parameterized tests).
+    // This member is also made mutable so we can release/create it even from a
+    // const-qualified objects (e.g., to release owned resources upon test completion,
+    // or to re-create the plan at execution if needed or found necessary)
+    mutable std::shared_ptr<hipfftw_plan_bundle_t<prec>> plan_bundle;
+
+    fft_transform_type     dft_kind;
+    int                    rank = 0;
+    std::vector<ptrdiff_t> lengths;
+    fft_result_placement   plan_placement;
+    int                    sign  = 0;
+    unsigned               flags = std::numeric_limits<unsigned>::max();
+
+    template <typename T>
+    void reset_member_value(T& member, const T& new_value)
+    {
+        if(new_value != member)
+        {
+            member = new_value;
+            plan_bundle.reset();
+        }
+    }
+
+    hipfftw_plan_creation_func get_creation_func(hipfftw_plan_creation_func creation_options) const
+    {
+        if(!hipfftw_creation_options_are_well_defined(creation_options))
+            throw std::invalid_argument("invalid creation_options for get_creation_func");
+        if(!can_use_creation_options(creation_options))
+        {
+            // e.g., rank < 0 with creation_options == hipfftw_plan_creation_func::PLAN_DFT_ND
+            throw std::invalid_argument(
+                "The plan creation options "
+                + hipfftw_creation_options_to_string(creation_options, dft_kind, rank)
+                + " cannot be used with this object");
+        }
+        std::vector<hipfftw_plan_creation_func> valid_candidates;
+        for(auto candidate : hipfftw_plan_creation_func_candidates)
+        {
+            if(!(creation_options & candidate))
+                continue; // candidate is not in given creation_options
+            if(can_use_creation_options(candidate))
+            {
+                // If creation_options != candidate for all candidates, creation_optionsactually
+                // combines 2 or more candidates --> only the candidates actually supporting plan
+                // creation will be considered "valid". If there exists one (usable) candidate s.t.
+                // creation_options == candidate however, this choice is considered "enforced"
+                // (e.g. for function-specific argument validation testing purposes)
+                if(creation_options == candidate || can_create_plan_with(candidate))
+                    valid_candidates.push_back(candidate);
+            }
+        }
+        if(valid_candidates.empty())
+            return hipfftw_plan_creation_func::NONE;
+        // "randomly" (yet reproducibly) choose
+        return valid_candidates[std::hash<std::string>()(token()) % valid_candidates.size()];
+    }
+
+    template <bool make_reference_plan = false>
+    hipfftw_plan_t<prec>
+        make_plan(void* in, void* out, hipfftw_plan_creation_func chosen_creation) const
+    {
+        if(std::find(hipfftw_plan_creation_func_candidates.begin(),
+                     hipfftw_plan_creation_func_candidates.end(),
+                     chosen_creation)
+           == hipfftw_plan_creation_func_candidates.end())
+        {
+            throw std::invalid_argument("Invalid chosen_creation for hipfftw_helper::make_plan");
+        }
+
+        // fetch/infer plan creation function arguments
+        const auto& hipfftw_impl = hipfftw_funcs<prec>::get_instance();
+        const auto  int_len      = get_length_as<int>();
+        const int*  int_len_ptr  = int_len.empty() ? nullptr : int_len.data();
+
+        switch(chosen_creation)
+        {
+        case hipfftw_plan_creation_func::PLAN_DFT_ND:
+        {
+            if(!can_use_creation_options(hipfftw_plan_creation_func::PLAN_DFT_ND))
+                throw std::runtime_error("hipfftw_plan_creation_func::PLAN_DFT_ND cannot be used.");
+            if(rank == 1)
+            {
+                if(dft_kind == fft_transform_type_real_forward)
+                {
+                    return hipfftw_impl.plan_dft_r2c_1d.template call<make_reference_plan>(
+                        int_len_ptr[0],
+                        static_cast<hipfftw_real_t<prec>*>(in),
+                        static_cast<hipfftw_complex_t<prec>*>(out),
+                        flags);
+                }
+                else if(dft_kind == fft_transform_type_real_inverse)
+                {
+                    return hipfftw_impl.plan_dft_c2r_1d.template call<make_reference_plan>(
+                        int_len_ptr[0],
+                        static_cast<hipfftw_complex_t<prec>*>(in),
+                        static_cast<hipfftw_real_t<prec>*>(out),
+                        flags);
+                }
+                else
+                {
+
+                    return hipfftw_impl.plan_dft_1d.template call<make_reference_plan>(
+                        int_len_ptr[0],
+                        static_cast<hipfftw_complex_t<prec>*>(in),
+                        static_cast<hipfftw_complex_t<prec>*>(out),
+                        sign,
+                        flags);
+                }
+            }
+            else if(rank == 2)
+            {
+                if(dft_kind == fft_transform_type_real_forward)
+                {
+                    return hipfftw_impl.plan_dft_r2c_2d.template call<make_reference_plan>(
+                        int_len_ptr[0],
+                        int_len_ptr[1],
+                        static_cast<hipfftw_real_t<prec>*>(in),
+                        static_cast<hipfftw_complex_t<prec>*>(out),
+                        flags);
+                }
+                else if(dft_kind == fft_transform_type_real_inverse)
+                {
+
+                    return hipfftw_impl.plan_dft_c2r_2d.template call<make_reference_plan>(
+                        int_len_ptr[0],
+                        int_len_ptr[1],
+                        static_cast<hipfftw_complex_t<prec>*>(in),
+                        static_cast<hipfftw_real_t<prec>*>(out),
+                        flags);
+                }
+                else
+                {
+                    return hipfftw_impl.plan_dft_2d.template call<make_reference_plan>(
+                        int_len_ptr[0],
+                        int_len_ptr[1],
+                        static_cast<hipfftw_complex_t<prec>*>(in),
+                        static_cast<hipfftw_complex_t<prec>*>(out),
+                        sign,
+                        flags);
+                }
+            }
+            else
+            {
+                if(dft_kind == fft_transform_type_real_forward)
+                {
+                    return hipfftw_impl.plan_dft_r2c_3d.template call<make_reference_plan>(
+                        int_len_ptr[0],
+                        int_len_ptr[1],
+                        int_len_ptr[2],
+                        static_cast<hipfftw_real_t<prec>*>(in),
+                        static_cast<hipfftw_complex_t<prec>*>(out),
+                        flags);
+                }
+                else if(dft_kind == fft_transform_type_real_inverse)
+                {
+                    return hipfftw_impl.plan_dft_c2r_3d.template call<make_reference_plan>(
+                        int_len_ptr[0],
+                        int_len_ptr[1],
+                        int_len_ptr[2],
+                        static_cast<hipfftw_complex_t<prec>*>(in),
+                        static_cast<hipfftw_real_t<prec>*>(out),
+                        flags);
+                }
+                else
+                {
+                    return hipfftw_impl.plan_dft_3d.template call<make_reference_plan>(
+                        int_len_ptr[0],
+                        int_len_ptr[1],
+                        int_len_ptr[2],
+                        static_cast<hipfftw_complex_t<prec>*>(in),
+                        static_cast<hipfftw_complex_t<prec>*>(out),
+                        sign,
+                        flags);
+                }
+            }
+        }
+        break;
+        case hipfftw_plan_creation_func::PLAN_DFT:
+        {
+            if(!can_use_creation_options(hipfftw_plan_creation_func::PLAN_DFT))
+                throw std::runtime_error("hipfftw_plan_creation_func::PLAN_DFT cannot be used.");
+
+            if(dft_kind == fft_transform_type_real_forward)
+            {
+                return hipfftw_impl.plan_dft_r2c.template call<make_reference_plan>(
+                    rank,
+                    int_len_ptr,
+                    static_cast<hipfftw_real_t<prec>*>(in),
+                    static_cast<hipfftw_complex_t<prec>*>(out),
+                    flags);
+            }
+            else if(dft_kind == fft_transform_type_real_inverse)
+            {
+                return hipfftw_impl.plan_dft_c2r.template call<make_reference_plan>(
+                    rank,
+                    int_len_ptr,
+                    static_cast<hipfftw_complex_t<prec>*>(in),
+                    static_cast<hipfftw_real_t<prec>*>(out),
+                    flags);
+            }
+            else
+            {
+                return hipfftw_impl.plan_dft.template call<make_reference_plan>(
+                    rank,
+                    int_len_ptr,
+                    static_cast<hipfftw_complex_t<prec>*>(in),
+                    static_cast<hipfftw_complex_t<prec>*>(out),
+                    sign,
+                    flags);
+            }
+        }
+        break;
+        case hipfftw_plan_creation_func::PLAN_MANY:
+            [[fallthrough]];
+        case hipfftw_plan_creation_func::PLAN_GURU:
+            [[fallthrough]];
+        case hipfftw_plan_creation_func::PLAN_GURU64:
+            throw std::runtime_error("Enforced plan creation is not implemented yet");
+            break;
+        default:
+            throw std::runtime_error("Unknown kind of plan creation");
+            break;
+        }
+        // unreachable
+    }
+
+public:
+    hipfftw_helper()                       = default;
+    ~hipfftw_helper()                      = default;
+    hipfftw_helper(hipfftw_helper&& other) = default;
+    hipfftw_helper& operator=(hipfftw_helper&& other) = default;
+    hipfftw_helper(const hipfftw_helper& other)       = default;
+    hipfftw_helper& operator=(const hipfftw_helper& rhs) = default;
+
+    void set_creation_args(fft_transform_type            dft_kind_to_set,
+                           int                           rank_to_set,
+                           const std::vector<ptrdiff_t>& lengths_to_set,
+                           fft_result_placement          placement_to_set,
+                           int                           sign_to_set,
+                           unsigned                      flags_to_set)
+    {
+        reset_member_value(dft_kind, dft_kind_to_set);
+        reset_member_value(rank, rank_to_set);
+        reset_member_value(lengths, lengths_to_set);
+        reset_member_value(plan_placement, placement_to_set);
+        reset_member_value(sign, sign_to_set);
+        reset_member_value(flags, flags_to_set);
+    }
+    // getters
+    fft_transform_type get_dft_kind() const
+    {
+        return dft_kind;
+    }
+    int get_rank() const
+    {
+        return rank;
+    }
+    // returns the lengths as an std::vector<T> if they may all be safely converted to T
+    // (the returned vector is empty otherwise)
+    template <typename T, std::enable_if_t<std::is_integral_v<T>, bool> = true>
+    std::vector<T> get_length_as() const
+    {
+        if constexpr(std::is_same_v<T, typename decltype(lengths)::value_type>)
+            return lengths;
+        std::vector<T> ret;
+        if(std::any_of(lengths.begin(),
+                       lengths.end(),
+                       [](const typename decltype(lengths)::value_type& val) {
+                           return val < std::numeric_limits<T>::lowest()
+                                  || val > std::numeric_limits<T>::max();
+                       }))
+        {
+            // not a safe conversion, return empty lengths
+            return ret;
+        }
+        ret.assign(lengths.begin(), lengths.end());
+        return ret;
+    }
+    fft_result_placement get_placement() const
+    {
+        return plan_placement;
+    }
+    int get_sign() const
+    {
+        return sign;
+    }
+    unsigned get_flags() const
+    {
+        return flags;
+    }
+    std::shared_ptr<hipfftw_plan_bundle_t<prec>> get_plan_bundle() const
+    {
+        return plan_bundle;
+    }
+    template <typename T, std::enable_if_t<std::is_integral_v<T>, bool> = true>
+    std::vector<T> get_strides_as(fft_io io) const
+    {
+        if(!rank_is_valid_for_hipfftw(rank) || !has_valid_lengths())
+            throw std::runtime_error(
+                "cannot calculate default strides with invalid rank or invalid lengths");
+        // only default strides for now
+        std::vector<ptrdiff_t> strides(rank, 1);
+        if(rank > 1)
+        {
+            if(is_complex(dft_kind))
+                strides[rank - 2] = lengths.back();
+            else
+            {
+                if(is_fwd(dft_kind) == (io == fft_io::fft_io_out))
+                    strides[rank - 2] = lengths.back() / 2 + 1;
+                else
+                {
+                    if(plan_placement == fft_placement_inplace)
+                        strides[rank - 2] = 2 * (lengths.back() / 2 + 1);
+                    else
+                        strides[rank - 2] = lengths.back();
+                }
+            }
+        }
+        for(auto dim = rank - 3; dim >= 0; dim--)
+            strides[dim] = strides[dim + 1] * lengths[dim + 1];
+
+        std::vector<T> ret;
+        if(std::any_of(strides.begin(),
+                       strides.end(),
+                       [](const typename decltype(strides)::value_type& val) {
+                           return val < std::numeric_limits<T>::lowest()
+                                  || val > std::numeric_limits<T>::max();
+                       }))
+        {
+            // not a safe conversion, return empty lengths
+            return ret;
+        }
+        ret.assign(strides.begin(), strides.end());
+        return ret;
+    }
+    template <typename T, std::enable_if_t<std::is_integral_v<T>, bool> = true>
+    T get_dist_as(fft_io io) const
+    {
+        if(!rank_is_valid_for_hipfftw(rank) || !has_valid_lengths())
+            throw std::runtime_error(
+                "cannot calculate default distance(s) with invalid rank or invalid lengths");
+        // only default distances for now
+        ptrdiff_t dist = 0;
+        if(rank == 1)
+        {
+            if(is_complex(dft_kind))
+                dist = lengths.back();
+            else
+            {
+                if(is_fwd(dft_kind) == (io == fft_io::fft_io_out))
+                    dist = lengths.back() / 2 + 1;
+                else
+                {
+                    if(plan_placement == fft_placement_inplace)
+                        dist = 2 * (lengths.back() / 2 + 1);
+                    else
+                        dist = lengths.back();
+                }
+            }
+        }
+        else
+        {
+            const auto strides = get_strides_as<ptrdiff_t>(io);
+            dist               = strides.front() * lengths.front();
+        }
+        if(dist < std::numeric_limits<T>::lowest() || dist > std::numeric_limits<T>::max())
+            throw std::runtime_error("distance cannot be safely converted to the desired type");
+        return static_cast<T>(dist);
+    }
+    template <typename T, std::enable_if_t<std::is_integral_v<T>, bool> = true>
+    T get_nbatch_as(fft_io io) const
+    {
+        // only unbatched for now
+        T ret = 1;
+        return ret;
+    }
+
+    // validity checks
+    bool has_valid_rank() const
+    {
+        return rank_is_valid_for_hipfftw(rank);
+    }
+    bool has_valid_lengths() const
+    {
+        return lengths_are_valid_for_hipfftw_as<ptrdiff_t>(lengths, rank);
+    }
+    bool has_valid_sign() const
+    {
+        return sign_is_valid_for_hipfftw(sign, dft_kind);
+    }
+    bool has_valid_flags() const
+    {
+        return flags_are_valid_for_hipfftw(flags);
+    }
+    // checks if the current parameters can be used with (any of) the given option(s) of
+    // plan creation (NOT whether they're valid or not). For instance, one cannot possibly
+    // communicate rank > 3 with hipfftw_plan_creation_func::PLAN_DFT_ND, or communicate
+    // non-default strides with hipfftw_plan_creation_func::PLAN_DFT_ND or
+    // hipfftw_plan_creation_func::PLAN_DFT...
+    // TODO: expand logic when extra configuration parameters are added (e.g. batch sizes,
+    // strides, etc.)
+    bool can_use_creation_options(hipfftw_plan_creation_func creation_options) const
+    {
+        if(!hipfftw_creation_options_are_well_defined(creation_options))
+            throw std::invalid_argument(
+                "ill-defined creation_options used in can_use_creation_options");
+        if(creation_options == hipfftw_plan_creation_func::NONE)
+            return false;
+        if(std::find(hipfftw_plan_creation_func_candidates.begin(),
+                     hipfftw_plan_creation_func_candidates.end(),
+                     creation_options)
+           == hipfftw_plan_creation_func_candidates.end())
+        {
+            // creation_options combines several candidates in hipfftw_plan_creation_func_candidates
+            // --> parse them individually and find out if any applicable can be used
+            return std::any_of(hipfftw_plan_creation_func_candidates.begin(),
+                               hipfftw_plan_creation_func_candidates.end(),
+                               [=](const hipfftw_plan_creation_func& candidate) {
+                                   return (creation_options & candidate)
+                                          && can_use_creation_options(candidate);
+                               });
+        }
+        // "creation_options" actually is an individual value in hipfftw_plan_creation_func_candidates
+        switch(creation_options)
+        {
+        case hipfftw_plan_creation_func::PLAN_DFT_ND:
+            // rank is not passed as an argument but dictated by the called function,
+            // (must be 1, 2, or 3), and as many lengths must be passed as individual
+            // integer values
+            return (rank == 1 || rank == 2 || rank == 3) && get_length_as<int>().size() == rank;
+            break;
+        case hipfftw_plan_creation_func::PLAN_DFT:
+            // the lengths must be representable as integers, if not empty (supposedly
+            // intentionally, e.g., for input validation testing purposes)
+            return lengths.empty() || get_length_as<int>().size() == rank;
+            break;
+        case hipfftw_plan_creation_func::PLAN_MANY:
+            [[fallthrough]];
+        case hipfftw_plan_creation_func::PLAN_GURU:
+            [[fallthrough]];
+        case hipfftw_plan_creation_func::PLAN_GURU64:
+            return false;
+            break;
+        default:
+            throw std::runtime_error("hipfftw_helper: internal error encountered (unexpected value "
+                                     "for creation_options)");
+            break;
+        }
+        // unreachable
+    }
+
+    // checks validity of configuration parameters and whether creation can be
+    // attempted via (any of) the given option(s)
+    bool is_valid_for_creation_with(hipfftw_plan_creation_func creation_options) const
+    {
+        if(!hipfftw_creation_options_are_well_defined(creation_options))
+            throw std::invalid_argument("invalid creation_options for is_valid_for_creation_with");
+
+        // TODO: expand the global validity checks below when this struct is
+        // expanded to cover more configurations (e.g., batching, srides, etc.)
+        return has_valid_rank() && has_valid_lengths() && has_valid_sign() && has_valid_flags()
+               && can_use_creation_options(creation_options);
+    }
+    bool is_valid_for_creation() const
+    {
+        return is_valid_for_creation_with(hipfftw_plan_creation_func::ANY);
+    }
+    // check expected support by (any of) the given option(s)
+    bool has_unsupported_args_for(hipfftw_plan_creation_func creation_options) const
+    {
+        // extra conditions for configurations supported by hipfftw:
+        if(rank > 3)
+            return true;
+        if(flags & FFTW_WISDOM_ONLY)
+            return true;
+        if(dft_kind == fft_transform_type_real_inverse && rank > 1 && (flags & FFTW_PRESERVE_INPUT))
+            return true;
+        if(!(creation_options & hipfftw_plan_creation_func::PLAN_GURU64) && has_valid_rank()
+           && has_valid_lengths())
+        {
+            // cannot handle data sizes involving more elements than the
+            // largest representable int value
+            if(get_num_elements_in(fft_io_in) > std::numeric_limits<int>::max()
+               || get_num_elements_in(fft_io_out) > std::numeric_limits<int>::max())
+                return true;
+        }
+        return false;
+    }
+    bool can_create_plan_with(hipfftw_plan_creation_func creation_options) const
+    {
+        if(!hipfftw_creation_options_are_well_defined(creation_options))
+            throw std::invalid_argument("invalid creation_option for can_create_plan_with");
+
+        if(!is_valid_for_creation_with(creation_options))
+            return false;
+        if(has_unsupported_args_for(creation_options))
+            return false;
+        return true;
+    }
+    bool can_create_plan() const
+    {
+        return can_create_plan_with(hipfftw_plan_creation_func::ANY);
+    }
+    // create a token consistent with other tests to enable kernel precompilation
+    // for valid cases, and/or capturing all required details about members otherwise
+    std::string token() const
+    {
+        std::ostringstream ret;
+        switch(dft_kind)
+        {
+        case fft_transform_type_complex_forward:
+            ret << "complex_forward";
+            break;
+        case fft_transform_type_complex_inverse:
+            ret << "complex_inverse";
+            break;
+        case fft_transform_type_real_forward:
+            ret << "real_forward";
+            break;
+        case fft_transform_type_real_inverse:
+            ret << "real_inverse";
+            break;
+        default:
+            throw std::runtime_error("unknown type of transform");
+        }
+
+        // report rank if invalid
+        if(!has_valid_rank() || lengths.empty())
+            ret << "_invalid_rank" << (rank < 0 ? "_negative_" : "_") << std::abs(rank);
+        ret << "_len";
+        if(lengths.empty())
+            ret << "_none";
+        else
+        {
+            for(const auto& len : lengths)
+                ret << (len < 0 ? "_negative_" : "_") << std::abs(len);
+        }
+        if constexpr(prec == fft_precision_single)
+            ret << "_single";
+        else
+            ret << "_double";
+        ret << (plan_placement == fft_placement_inplace ? "_ip" : "_op");
+        // only supporting unbatched cases as of now
+        ret << "_batch_1";
+        if(has_valid_rank() && has_valid_lengths())
+        {
+            ret << "_istride";
+            for(const auto& stride : get_strides_as<size_t>(fft_io::fft_io_in))
+                ret << "_" << stride;
+            if(!is_real(dft_kind))
+                ret << "_CI";
+            else if(dft_kind == fft_transform_type_real_forward)
+                ret << "_R";
+            else
+                ret << "_HI";
+            ret << "_ostride";
+            for(const auto& stride : get_strides_as<size_t>(fft_io::fft_io_out))
+                ret << "_" << stride;
+            if(!is_real(dft_kind))
+                ret << "_CI";
+            else if(dft_kind == fft_transform_type_real_forward)
+                ret << "_HI";
+            else
+                ret << "_R";
+            ret << "_idist_" << get_dist_as<size_t>(fft_io::fft_io_in);
+            ret << "_odist_" << get_dist_as<size_t>(fft_io::fft_io_out);
+            ret << "_ioffset_0_0_ooffset_0_0";
+        }
+        if(!has_valid_sign())
+            ret << "_invalid_sign" << (sign < 0 ? "_negative_" : "_") << std::abs(sign);
+        ret << "_flags_" << flags;
+        return ret.str();
+    }
+    // create_plan invokes an hipfftw plan creation function for the object's configuration
+    // parameters, the corresponding plan pointer returned by hipfftw is stored internally.
+    // IMPORTANT NOTE: if one wants to target a specific creation function (as represented
+    // by any value in hipfftw_plan_creation_func_candidates), setting the creation_options
+    // argument to that specific value effectively bypasses the verification that the
+    // object's configuration is actually (expected to be) supported and attempts the plan
+    // creation anyways (unless it simply cannot be done, e.g., attempting
+    // creation_options = hipfftw_plan_creation_func::PLAN_DFT_ND herein on an object
+    // holding a value for rank > 3 simply cannot be done)
+    void create_plan(void*                      in,
+                     void*                      out,
+                     hipfftw_plan_creation_func creation_options
+                     = hipfftw_plan_creation_func::ANY) const
+    {
+        const auto&                      hipfftw_impl  = hipfftw_funcs<prec>::get_instance();
+        const hipfftw_plan_creation_func chosen_option = get_creation_func(creation_options);
+        if(chosen_option == hipfftw_plan_creation_func::NONE)
+        {
+            plan_bundle = std::make_shared<hipfftw_plan_bundle_t<prec>>(hipfftw_impl.destroy_plan);
+            plan_bundle->creation_io   = {in, out};
+            plan_bundle->plan          = nullptr;
+            plan_bundle->creation_func = chosen_option;
+            plan_bundle->plan_token    = "";
+            return;
+        }
+        // early return if there is no need to (re)build
+        if(plan_bundle && plan_bundle->plan_token == token() && plan_bundle->creation_io.first == in
+           && plan_bundle->creation_io.second == out && plan_bundle->creation_func == chosen_option)
+            return;
+
+        // create the desired plan
+        plan_bundle = std::make_shared<hipfftw_plan_bundle_t<prec>>(hipfftw_impl.destroy_plan);
+        plan_bundle->plan          = make_plan(in, out, chosen_option);
+        plan_bundle->creation_io   = {in, out};
+        plan_bundle->creation_func = chosen_option;
+        plan_bundle->plan_token    = token();
+    }
+
+    // returns a reference FFTW plan for the current configuration
+    // The returned plan is NOT owned by this object!
+    hipfftw_plan_t<prec> get_reference_plan(void*                      in,
+                                            void*                      out,
+                                            hipfftw_plan_creation_func creation_options
+                                            = hipfftw_plan_creation_func::ANY) const
+    {
+        const hipfftw_plan_creation_func chosen_option = get_creation_func(creation_options);
+        if(chosen_option == hipfftw_plan_creation_func::NONE)
+        {
+            return nullptr;
+        }
+        constexpr bool make_reference_plan = true;
+        return make_plan<make_reference_plan>(in, out, chosen_option);
+    }
+
+    void execute(void* execute_in, void* execute_out) const
+    {
+        if(!plan_bundle || plan_bundle->plan_token != token())
+        {
+            // plan is not created or possibly not up-to-date
+            create_plan(execute_in, execute_out);
+        }
+
+        const auto& hipfftw_impl = hipfftw_funcs<prec>::get_instance();
+        if(execute_in == plan_bundle->creation_io.first
+           && execute_out == plan_bundle->creation_io.second)
+        {
+            hipfftw_impl.execute(plan_bundle->plan);
+        }
+        else
+        {
+            throw std::runtime_error("New-array execution functions not implemented yet.");
+        }
+    }
+
+    // TODO: revise/expand logic below when the structure is expanded for more cases (batches,
+    // non-default strides, etc.)
+    size_t get_num_elements_in(fft_io in_or_out) const
+    {
+        if(in_or_out != fft_io_in && in_or_out != fft_io_out)
+            throw std::invalid_argument("invalid in_or_out for get_num_elements_in");
+        if(!has_valid_rank() || !has_valid_lengths())
+            throw std::runtime_error("get_num_elements_in requires valid rank and lengths");
+        const auto tmp = get_length_as<size_t>();
+        if(tmp.empty() || tmp.size() != rank)
+        {
+            throw std::runtime_error(
+                "get_num_elements_in failed to correctly convert lengths to size_t values");
+        }
+        size_t num_elems = 1;
+        if(is_complex(dft_kind))
+        {
+            num_elems *= tmp[rank - 1];
+        }
+        else
+        {
+            const size_t cmplx_len = tmp[rank - 1] / 2 + 1;
+            if(is_fwd(dft_kind) == (in_or_out == fft_io_out))
+                num_elems *= cmplx_len;
+            else
+                num_elems
+                    *= plan_placement == fft_placement_inplace ? 2 * cmplx_len : tmp[rank - 1];
+        }
+        num_elems *= product(tmp.begin(), tmp.begin() + rank - 1);
+        return num_elems;
+    }
+
+    size_t get_data_byte_size(fft_io in_or_out) const
+    {
+        if(in_or_out != fft_io_in && in_or_out != fft_io_out)
+            throw std::invalid_argument("invalid in_or_out for get_data_byte_size");
+        // for in-place, input and output data sizes are enforced equal
+        std::vector<fft_io> io_range_to_consider = {in_or_out};
+        if(plan_placement == fft_placement_inplace)
+            io_range_to_consider.push_back(in_or_out == fft_io::fft_io_in ? fft_io::fft_io_out
+                                                                          : fft_io::fft_io_in);
+
+        size_t ret = 0;
+        for(auto io : io_range_to_consider)
+        {
+            const size_t num_elems = get_num_elements_in(io);
+            if(is_complex(dft_kind) || (is_fwd(dft_kind) == (io == fft_io_out)))
+                ret = std::max(ret, num_elems * sizeof(hipfftw_complex_t<prec>));
+            else
+                ret = std::max(ret, num_elems * sizeof(hipfftw_real_t<prec>));
+        }
+        return ret;
+    }
+    void release_plan() const
+    {
+        plan_bundle.reset();
+    }
+};
+
+#endif
diff --git a/clients/samples/CMakeLists.txt b/clients/samples/CMakeLists.txt
index 488e09d..37cc924 100644
--- a/clients/samples/CMakeLists.txt
+++ b/clients/samples/CMakeLists.txt
@@ -124,8 +124,15 @@ foreach( sample ${sample_list} )
   
 endforeach()
 
-# cuFFT callback code must be compiled with -dc to enable relocatable
-# device code
-if( BUILD_WITH_LIB STREQUAL "CUDA" AND hipfft_callback IN_LIST sample_list )
-  target_compile_options( hipfft_callback PRIVATE -dc )
+# callback code must be compiled as relocatable device code
+if( hipfft_callback IN_LIST sample_list )
+  if( BUILD_WITH_LIB STREQUAL "CUDA" )
+    target_compile_options( hipfft_callback PRIVATE -dc )
+  else()
+    # -fgpu-rdc causes failure at link stage on Windows
+    if (NOT WIN32)
+      target_compile_options( hipfft_callback PRIVATE -fgpu-rdc )
+      target_link_options( hipfft_callback PRIVATE -fgpu-rdc )
+    endif()
+  endif()
 endif()
diff --git a/clients/samples/hipfft_callback.cpp b/clients/samples/hipfft_callback.cpp
index b365e87..8a77266 100644
--- a/clients/samples/hipfft_callback.cpp
+++ b/clients/samples/hipfft_callback.cpp
@@ -20,6 +20,7 @@
 // THE SOFTWARE.
 
 #include <iostream>
+#ifndef WIN32
 #include <vector>
 
 #include <hip/hip_runtime.h>
@@ -48,9 +49,15 @@ __device__ hipfftDoubleComplex load_callback(hipfftDoubleComplex* input,
 }
 
 __device__ auto load_callback_dev = load_callback;
+#endif
 
 int main()
 {
+#ifdef WIN32
+    std::cout << "This sample is temporarily disabled on Windows" << std::endl;
+    return EXIT_SUCCESS;
+#else
+
     std::cout << "hipfft 1D double-precision complex-to-complex transform with callback\n";
 
     const int Nx        = 8;
@@ -156,4 +163,5 @@ int main()
         throw std::runtime_error("hipFree failed");
 
     return 0;
+#endif
 }
diff --git a/clients/tests/CMakeLists.txt b/clients/tests/CMakeLists.txt
index 28c6594..91024cf 100644
--- a/clients/tests/CMakeLists.txt
+++ b/clients/tests/CMakeLists.txt
@@ -32,6 +32,13 @@ else( )
     "Install path prefix, prepended onto install directories" )
 endif( )
 
+# Dependencies
+
+find_package( ROCmCMakeBuildTools REQUIRED CONFIG PATHS /opt/rocm )
+include(ROCMInstallTargets)
+list( APPEND CMAKE_MODULE_PATH ${CMAKE_CURRENT_SOURCE_DIR}/../cmake )
+
+
 # This has to be initialized before the project() command appears
 # Set the default of CMAKE_BUILD_TYPE to be release, unless user
 # specifies with -D.  MSVC_IDE does not use CMAKE_BUILD_TYPE
@@ -42,13 +49,19 @@ endif()
 
 project( hipfft-clients-tests LANGUAGES CXX )
 
+if( NOT HIPFFT_BUILD_SCOPE )
+  find_package( hipfft REQUIRED CONFIG PATHS )
+endif()
 
 find_package( Boost REQUIRED)
 
 set( Boost_USE_STATIC_LIBS OFF )
 
+
 find_package( FFTW 3.0 REQUIRED MODULE COMPONENTS FLOAT DOUBLE )
 
+set( BUILD_WITH_LIB "ROCM" CACHE STRING "Build ${PROJECT_NAME} with ROCM or CUDA libraries" )
+
 set( THREADS_PREFER_PTHREAD_FLAG ON )
 find_package( Threads REQUIRED )
 
@@ -60,7 +73,9 @@ set( hipfft-test_source
   accuracy_test_2D.cpp
   accuracy_test_3D.cpp
   accuracy_test_callback.cpp
+  hipfftw_test.cpp
   multi_device_test.cpp
+  multi_stream_test.cpp
   ../../shared/array_validator.cpp
   )
 
@@ -84,7 +99,7 @@ if( NOT BUILD_WITH_LIB STREQUAL "CUDA" )
   if( WIN32 )
     find_package( HIP CONFIG REQUIRED )
   else()
-    find_package( HIP MODULE REQUIRED )
+    find_package( hip REQUIRED CONFIG PATHS /opt/rocm/lib/cmake/hip/ )
   endif()
 endif()
 
@@ -98,6 +113,20 @@ endif()
 string( CONCAT TESTS_OUT_DIR "${PROJECT_BINARY_DIR}" ${TESTS_OUT_DIR} )
 
 option( BUILD_CLIENTS_TESTS_OPENMP "Build tests with OpenMP" ON )
+if( BUILD_CLIENTS_TESTS_OPENMP AND NOT BUILD_WITH_LIB STREQUAL "CUDA" )
+  # Attempt to find a config version, which provides openmp_LIB_DIR.
+  find_package( OpenMP CONFIG PATHS "${HIP_CLANG_ROOT}/lib/cmake" )
+  if( NOT OPENMP_FOUND OR NOT DEFINED ${openmp_LIB_DIR} )
+     # Fall-back to module mode.
+     find_package( OpenMP REQUIRED )
+     set( BUILD_RPATH "${HIP_CLANG_ROOT}/lib" )
+     set( INSTALL_RPATH "$ORIGIN/../llvm/lib" )
+  else()
+     set( BUILD_RPATH "${HIP_CLANG_ROOT}/${openmp_LIB_DIR}" )
+     set( INSTALL_RPATH "$ORIGIN/../llvm/${openmp_LIB_DIR}" )
+
+  endif()
+endif()
 
 foreach( target ${TEST_TARGETS} )
   set_target_properties( ${target} PROPERTIES
@@ -105,6 +134,15 @@ foreach( target ${TEST_TARGETS} )
     CXX_STANDARD_REQUIRED ON
     )
 
+    if( BUILD_CLIENTS_TESTS_OPENMP )
+      set_target_properties( ${TEST_TARGETS} PROPERTIES
+         BUILD_RPATH "${BUILD_RPATH}"
+      )
+      set_target_properties( ${TEST_TARGETS} PROPERTIES
+        INSTALL_RPATH "${INSTALL_RPATH}"
+      )
+    endif()
+
   if( BUILD_WITH_LIB STREQUAL "ROCM" )
     target_compile_options( ${target} PRIVATE ${WARNING_FLAGS} )
     target_link_libraries( ${target}
@@ -112,7 +150,7 @@ foreach( target ${TEST_TARGETS} )
       hip::host
       hip::device
     )
-    foreach( gpu_target ${AMDGPU_TARGETS} )
+    foreach( gpu_target ${GPU_TARGETS} )
       target_compile_options( ${target} PRIVATE --offload-arch=${gpu_target} )
     endforeach()
 
@@ -140,30 +178,12 @@ foreach( target ${TEST_TARGETS} )
     target_compile_definitions( ${target} PUBLIC _CUFFT_BACKEND )
   endif()
 
-  if( BUILD_CLIENTS_TESTS_OPENMP )
-    if( BUILD_WITH_LIB STREQUAL "CUDA" )
-      message( STATUS "OpenMP is not supported on CUDA, building tests without it" )
-    else()
-      target_compile_options( ${target} PRIVATE -DBUILD_CLIENTS_TESTS_OPENMP )
-      if(NOT (CMAKE_CXX_COMPILER MATCHES ".*hipcc$" OR CMAKE_CXX_COMPILER MATCHES ".*clang\\+\\+"))
-        target_compile_options( ${target} PRIVATE -fopenmp )
-        target_link_libraries( ${target} PRIVATE -fopenmp -L${HIP_CLANG_ROOT}/lib -Wl,-rpath=${HIP_CLANG_ROOT}/lib )
-        target_include_directories( ${target} PRIVATE ${HIP_CLANG_ROOT}/include )
-      else()
-        if(CMAKE_CXX_COMPILER_ID STREQUAL "Clang")
-          target_compile_options( ${target} PRIVATE -fopenmp=libomp )
-          target_link_options( ${target} PRIVATE -fopenmp=libomp )
-        endif()
-      endif()
-    endif()
-  endif()
 
   target_include_directories( ${target}
     PRIVATE
     $<BUILD_INTERFACE:${Boost_INCLUDE_DIRS}>
     $<BUILD_INTERFACE:${FFTW_INCLUDE_DIRS}>
     $<BUILD_INTERFACE:${hip_INCLUDE_DIRS}>
-    $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/../../library/include>
   )
 
   target_link_libraries( ${target}
@@ -171,6 +191,18 @@ foreach( target ${TEST_TARGETS} )
     hip::hipfft
     ${FFTW_LIBRARIES}
   )
+  
+  if( BUILD_CLIENTS_TESTS_OPENMP )
+    if( BUILD_WITH_LIB STREQUAL "CUDA" )
+      message( STATUS "OpenMP is not supported on CUDA, building tests without it" )
+    else()
+      if( DEFINED ${openmp_LIB_DIR} )
+        set_target_properties( ${target} PROPERTIES BUILD_RPATH "${HIP_CLANG_ROOT}/${openmp_LIB_DIR}" )
+        set_target_properties( ${target} PROPERTIES INSTALL_RPATH "${HIP_CLANG_ROOT}/${openmp_LIB_DIR}" )
+      endif()
+      target_link_libraries( ${target} PRIVATE OpenMP::OpenMP_CXX )
+    endif()
+  endif()
 
   if( HIPFFT_MPI_ENABLE )
     target_link_libraries( ${target}
@@ -187,15 +219,29 @@ foreach( target ${TEST_TARGETS} )
   rocm_install(TARGETS ${target} COMPONENT tests)
 endforeach()
 
-if( GTEST_FOUND )
+find_package( GTest 1.11.0 )
+
+if( GTest_FOUND )
   target_include_directories( hipfft-test PRIVATE $<BUILD_INTERFACE:${GTEST_INCLUDE_DIRS}> )
-  target_link_libraries( hipfft-test PRIVATE ${GTEST_LIBRARIES} )
+  target_link_libraries( hipfft-test PRIVATE GTest::gtest )
 else()
+  # gtest build by the hipFFT
   add_dependencies( hipfft-test gtest )
   target_include_directories( hipfft-test PRIVATE hipfft-test_include_dirs ${GTEST_INCLUDE_DIRS} )
   target_link_libraries( hipfft-test PRIVATE ${GTEST_LIBRARIES} )
 endif()
 
+# tests have callback functions, which need to be built as relocatable device code
+if( BUILD_WITH_LIB STREQUAL "CUDA" )
+  target_compile_options( hipfft-test PRIVATE -dc )
+else()
+  # -fgpu-rdc causes failure at link stage on Windows
+  if (NOT WIN32)
+    target_compile_options( hipfft-test PRIVATE -fgpu-rdc )
+    target_link_options( hipfft-test PRIVATE -fgpu-rdc )
+  endif()
+endif()
+
 if(FFTW_MULTITHREAD)
   target_compile_options( hipfft-test PRIVATE -DFFTW_MULTITHREAD )
 endif( )
@@ -203,8 +249,14 @@ endif( )
 target_link_libraries( hipfft-test
   PRIVATE
   Threads::Threads
+  ${CMAKE_DL_LIBS}
   )
 
+# hipfft-test will opens the hipfftw library but does not link to it
+if( TARGET hipfftw )
+  add_dependencies( hipfft-test hipfftw )
+endif()
+
 if (WIN32)
 
   # Ensure tests run with HIP DLLs and not anything the driver owns
@@ -222,3 +274,42 @@ if (WIN32)
     add_custom_command( TARGET hipfft-test POST_BUILD COMMAND ${CMAKE_COMMAND} ARGS -E copy ${file_i} $<TARGET_FILE_DIR:hipfft-test> )
   endforeach( file_i )
 endif()
+
+option(BUILD_CODE_COVERAGE "Build with code coverage flags (clang only)" OFF)
+set(COVERAGE_TEST_OPTIONS "--smoketest;--gtest_filter=-*call*" CACHE STRING "Command line arguments for hipfft-test when generating a code coverage report (Note: an additional run of hipfft-test targeting multi_gpu* and callback* tests is always executed and coverage results are aggregated)")
+if (BUILD_CODE_COVERAGE)
+  add_custom_target(
+    code_cov_tests
+    DEPENDS hipfft-test
+    COMMAND ${CMAKE_COMMAND} -E rm -rf ./coverage-report
+    COMMAND ${CMAKE_COMMAND} -E make_directory ./coverage-report/profraw
+    COMMAND ${CMAKE_COMMAND} -E env LLVM_PROFILE_FILE="./coverage-report/profraw/hipfft-coverage_%p.profraw" GTEST_LISTENER=NO_PASS_LINE_IN_LOG $<TARGET_FILE:hipfft-test> --precompile=./clients/staging/hipfft-test-precompile.db ${COVERAGE_TEST_OPTIONS}
+    COMMAND ${CMAKE_COMMAND} -E env LLVM_PROFILE_FILE="./coverage-report/profraw/hipfft-coverage_%p.profraw" GTEST_LISTENER=NO_PASS_LINE_IN_LOG $<TARGET_FILE:hipfft-test> --precompile=./clients/staging/hipfft-test-precompile-multi_gpu-plus-callback.db --gtest_filter=multi_gpu*:callback*
+    WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+  )
+
+  find_program(
+    LLVM_PROFDATA
+    llvm-profdata
+    REQUIRED
+    HINTS ${ROCM_PATH}/llvm/bin
+    PATHS /opt/rocm/llvm/bin
+  )
+
+  find_program(
+    LLVM_COV
+    llvm-cov
+    REQUIRED
+    HINTS ${ROCM_PATH}/llvm/bin
+    PATHS /opt/rocm/llvm/bin
+  )
+
+  add_custom_target(
+    coverage
+    DEPENDS code_cov_tests
+    COMMAND ${LLVM_PROFDATA} merge -sparse ./coverage-report/profraw/hipfft-coverage_*.profraw -o ./coverage-report/hipfft.profdata
+    COMMAND ${LLVM_COV} report -object ./library/libhipfftw.so -object ./library/libhipfft.so -instr-profile=./coverage-report/hipfft.profdata
+    COMMAND ${LLVM_COV} show -object ./library/libhipfftw.so -object ./library/libhipfft.so -instr-profile=./coverage-report/hipfft.profdata -format=html -output-dir=coverage-report
+    WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+  )
+endif()
diff --git a/clients/tests/accuracy_test_1D.cpp b/clients/tests/accuracy_test_1D.cpp
index 91e238e..b2aa382 100644
--- a/clients/tests/accuracy_test_1D.cpp
+++ b/clients/tests/accuracy_test_1D.cpp
@@ -81,14 +81,8 @@ static std::vector<size_t> small_1D_sizes()
     static const size_t SMALL_1D_MAX = 8192;
 
     // generate a list of sizes from 2 and up, skipping any sizes that are already covered
-    std::vector<size_t> covered_sizes;
-    std::copy(pow2_range.begin(), pow2_range.end(), std::back_inserter(covered_sizes));
-    std::copy(pow3_range.begin(), pow3_range.end(), std::back_inserter(covered_sizes));
-    std::copy(pow5_range.begin(), pow5_range.end(), std::back_inserter(covered_sizes));
-    std::copy(radX_range.begin(), radX_range.end(), std::back_inserter(covered_sizes));
-    std::copy(mix_range.begin(), mix_range.end(), std::back_inserter(covered_sizes));
-    std::copy(prime_range.begin(), prime_range.end(), std::back_inserter(covered_sizes));
-    std::sort(covered_sizes.begin(), covered_sizes.end());
+    std::vector<size_t> covered_sizes = merge_and_sort_values<size_t>(
+        {pow2_range, pow3_range, pow5_range, radX_range, mix_range, prime_range});
 
     std::vector<size_t> output;
     for(size_t i = 2; i < SMALL_1D_MAX; ++i)
@@ -319,29 +313,24 @@ INSTANTIATE_TEST_SUITE_P(DISABLED_offset_mix_1D,
 
 // small 1D sizes just need to make sure our factorization isn't
 // completely broken, so we just check simple C2C outplace interleaved
-INSTANTIATE_TEST_SUITE_P(small_1D,
-                         accuracy_test,
-                         ::testing::ValuesIn(param_generator_base(
-                             test_prob,
-                             {fft_transform_type_complex_forward},
-                             generate_lengths({small_1D_sizes()}),
-                             {fft_precision_single},
-                             {1},
-                             [](fft_transform_type                       t,
-                                const std::vector<fft_result_placement>& place_range,
-                                const bool                               planar) {
-                                 return std::vector<type_place_io_t>{
-                                     std::make_tuple(t,
-                                                     place_range[0],
-                                                     fft_array_type_complex_interleaved,
-                                                     fft_array_type_complex_interleaved)};
-                             },
-                             stride_range,
-                             stride_range,
-                             ioffset_range_zero,
-                             ooffset_range_zero,
-                             {fft_placement_notinplace})),
-                         accuracy_test::TestName);
+const static std::vector<size_t> small_1D_lengths = small_1D_sizes();
+
+INSTANTIATE_TEST_SUITE_P(
+    small_1D,
+    accuracy_test,
+    ::testing::ValuesIn(param_generator_base(test_prob,
+                                             {fft_transform_type_complex_forward},
+                                             generate_lengths({small_1D_lengths}),
+                                             {fft_precision_single},
+                                             {1},
+                                             generate_types,
+                                             stride_range,
+                                             stride_range,
+                                             ioffset_range_zero,
+                                             ooffset_range_zero,
+                                             {fft_placement_notinplace},
+                                             false /* planar */)),
+    accuracy_test::TestName);
 
 // NB:
 // We have known non-unit strides issues for 1D:
@@ -352,7 +341,7 @@ INSTANTIATE_TEST_SUITE_P(small_1D,
 // main tests.
 //
 // The below test covers non-unit strides, pow of 2, middle sizes, which has SBCC/SBRC kernels
-// invloved.
+// involved.
 const static std::vector<size_t>              pow2_range_for_stride      = {4096, 8192, 524288};
 const static std::vector<size_t>              pow2_range_for_stride_half = {4096, 8192};
 const static std::vector<std::vector<size_t>> stride_range_for_pow2      = {{2}, {3}};
@@ -364,7 +353,7 @@ INSTANTIATE_TEST_SUITE_P(
     ::testing::ValuesIn(param_generator_complex(test_prob,
                                                 generate_lengths({pow2_range_for_stride}),
                                                 precision_range_sp_dp,
-                                                batch_range_1D,
+                                                batch_range_for_stride,
                                                 stride_range_for_pow2,
                                                 stride_range_for_pow2,
                                                 ioffset_range_zero,
@@ -380,7 +369,7 @@ INSTANTIATE_TEST_SUITE_P(
     ::testing::ValuesIn(param_generator_real(test_prob,
                                              generate_lengths({pow2_range_for_stride}),
                                              precision_range_sp_dp,
-                                             batch_range_1D,
+                                             batch_range_for_stride,
                                              stride_range_for_pow2,
                                              stride_range_for_pow2,
                                              ioffset_range_zero,
@@ -396,7 +385,7 @@ INSTANTIATE_TEST_SUITE_P(
     ::testing::ValuesIn(param_generator_real(test_prob,
                                              generate_lengths({pow2_range_for_stride_half}),
                                              {fft_precision_half},
-                                             batch_range_1D,
+                                             batch_range_for_stride,
                                              stride_range_for_pow2,
                                              stride_range_for_pow2,
                                              ioffset_range_zero,
@@ -537,3 +526,24 @@ INSTANTIATE_TEST_SUITE_P(
                                                               ooffset_range_zero,
                                                               place_range)),
     accuracy_test::TestName);
+
+const static std::vector<size_t> lengths_for_disabled_autoalloc = merge_and_sort_values<size_t>(
+    {pow2_range, pow3_range, pow5_range, radX_range, mix_range, small_1D_lengths, prime_range},
+    128);
+
+INSTANTIATE_TEST_SUITE_P(
+    various_1D,
+    accuracy_test,
+    ::testing::ValuesIn(param_generator(test_prob,
+                                        generate_lengths({lengths_for_disabled_autoalloc}),
+                                        precision_range_sp_dp,
+                                        batch_range_1D,
+                                        stride_range,
+                                        stride_range,
+                                        ioffset_range_zero,
+                                        ooffset_range_zero,
+                                        place_range,
+                                        false,
+                                        false,
+                                        fft_auto_allocation_off)),
+    accuracy_test::TestName);
diff --git a/clients/tests/accuracy_test_2D.cpp b/clients/tests/accuracy_test_2D.cpp
index 5f1b8a9..6f45abe 100644
--- a/clients/tests/accuracy_test_2D.cpp
+++ b/clients/tests/accuracy_test_2D.cpp
@@ -278,3 +278,24 @@ INSTANTIATE_TEST_SUITE_P(len1_swap_2D,
                              false,
                              false)),
                          accuracy_test::TestName);
+
+const static std::vector<size_t> lengths_for_disabled_autoalloc
+    = merge_and_sort_values<size_t>({pow2_range, pow3_range, prime_range, mix_range}, 12);
+
+INSTANTIATE_TEST_SUITE_P(
+    various_2D,
+    accuracy_test,
+    ::testing::ValuesIn(param_generator(test_prob,
+                                        generate_lengths({lengths_for_disabled_autoalloc,
+                                                          lengths_for_disabled_autoalloc}),
+                                        precision_range_sp_dp,
+                                        batch_range,
+                                        stride_range,
+                                        stride_range,
+                                        ioffset_range_zero,
+                                        ooffset_range_zero,
+                                        place_range,
+                                        false,
+                                        false,
+                                        fft_auto_allocation_off)),
+    accuracy_test::TestName);
diff --git a/clients/tests/accuracy_test_3D.cpp b/clients/tests/accuracy_test_3D.cpp
index f338695..f3f780d 100644
--- a/clients/tests/accuracy_test_3D.cpp
+++ b/clients/tests/accuracy_test_3D.cpp
@@ -284,3 +284,25 @@ INSTANTIATE_TEST_SUITE_P(
         false,
         false)),
     accuracy_test::TestName);
+
+const static std::vector<size_t> lengths_for_disabled_autoalloc = merge_and_sort_values<size_t>(
+    {pow2_range, pow3_range, pow5_range, prime_range, sbrc_range}, 5);
+
+INSTANTIATE_TEST_SUITE_P(
+    various_3D,
+    accuracy_test,
+    ::testing::ValuesIn(param_generator(test_prob,
+                                        generate_lengths({lengths_for_disabled_autoalloc,
+                                                          lengths_for_disabled_autoalloc,
+                                                          lengths_for_disabled_autoalloc}),
+                                        precision_range_sp_dp,
+                                        batch_range,
+                                        stride_range,
+                                        stride_range,
+                                        ioffset_range_zero,
+                                        ooffset_range_zero,
+                                        place_range,
+                                        false,
+                                        false,
+                                        fft_auto_allocation_off)),
+    accuracy_test::TestName);
diff --git a/clients/tests/accuracy_test_callback.cpp b/clients/tests/accuracy_test_callback.cpp
index 00e6026..0bae646 100644
--- a/clients/tests/accuracy_test_callback.cpp
+++ b/clients/tests/accuracy_test_callback.cpp
@@ -104,22 +104,27 @@ const static std::vector<std::vector<size_t>> ooffset_range = {{0, 0}, {1, 1}};
 auto transform_types = {fft_transform_type_complex_forward, fft_transform_type_real_forward};
 
 #ifdef __HIP__
-INSTANTIATE_TEST_SUITE_P(callback,
-                         accuracy_test,
-                         ::testing::ValuesIn(param_generator_base(test_prob,
-                                                                  transform_types,
-                                                                  callback_sizes,
-                                                                  precision_range_sp_dp,
-                                                                  batch_range,
-                                                                  generate_types,
-                                                                  stride_range,
-                                                                  stride_range,
-                                                                  ioffset_range_zero,
-                                                                  ooffset_range_zero,
-                                                                  place_range,
-                                                                  false,
-                                                                  true)),
-                         accuracy_test::TestName);
+INSTANTIATE_TEST_SUITE_P(
+#ifdef WIN32
+    DISABLED_callback_no_offset,
+#else
+    callback,
+#endif
+    accuracy_test,
+    ::testing::ValuesIn(param_generator_base(test_prob,
+                                             transform_types,
+                                             callback_sizes,
+                                             precision_range_sp_dp,
+                                             batch_range,
+                                             generate_types,
+                                             stride_range,
+                                             stride_range,
+                                             ioffset_range_zero,
+                                             ooffset_range_zero,
+                                             place_range,
+                                             false,
+                                             true)),
+    accuracy_test::TestName);
 
 INSTANTIATE_TEST_SUITE_P(DISABLED_callback,
                          accuracy_test,
diff --git a/clients/tests/gtest_main.cpp b/clients/tests/gtest_main.cpp
index b574106..ee2f5ec 100644
--- a/clients/tests/gtest_main.cpp
+++ b/clients/tests/gtest_main.cpp
@@ -41,20 +41,18 @@
 #include "../hipfft_params.h"
 #include "hipfft/hipfft.h"
 #include "hipfft_accuracy_test.h"
-#include "hipfft_test_params.h"
+
+// initialize static class member of hipfft_params
+std::vector<gpubuf> hipfft_params::externally_managed_workareas = std::vector<gpubuf>();
 
 // Control output verbosity:
 int verbose;
 
-// Run a short (~5 min) test suite by setting test_prob to an appropriate value
-bool smoketest = false;
-
 // User-defined random seed
-size_t random_seed;
+size_t             random_seed;
+std::random_device default_seed_dev;
 // Overall probability of running conventional tests
 double test_prob;
-// Probability of running tests from the emulation suite
-double emulation_prob;
 // Modifier for probability of running tests with complex interleaved data
 double complex_interleaved_prob_factor;
 // Modifier for probability of running tests with real data
@@ -63,9 +61,12 @@ double real_prob_factor;
 double complex_planar_prob_factor;
 // Modifier for probability of running tests with callbacks
 double callback_prob_factor;
+// Constraints for the hipfftw tests
+size_t max_length_for_hipfftw_test;
+size_t max_io_gb_for_hipfftw_test;
 
 // Transform parameters for manual test:
-fft_params manual_params;
+hipfft_params manual_params;
 
 // Host memory limitation for tests (GiB):
 size_t ramgb;
@@ -104,7 +105,7 @@ last_cpu_fft_cache last_cpu_fft_data;
 fft_params::fft_mp_lib mp_lib = fft_params::fft_mp_lib_none;
 // Number of multi-process ranks to launch
 int mp_ranks = 1;
-// Multi-process launch command (e.g. mpirun --np 4 /path/to/rocfft_mpi_worker)
+// Multi-process launch command (e.g. mpirun --np 4 /path/to/hipfft_mpi_worker)
 std::string mp_launch;
 
 void init_gtest_flags()
@@ -174,9 +175,10 @@ void precompile_test_kernels(const std::string& precompile_file)
                 continue;
 
             // only care about accuracy tests
-            if(name.find("vs_fftw/") != std::string::npos)
+            const auto pos = name.find("vs_fftw/");
+            if(pos != std::string::npos)
             {
-                name.erase(0, 8);
+                name.erase(0, pos + 8);
 
                 // change batch to 1, so we don't waste time creating
                 // multiple plans that differ only by batch
@@ -199,7 +201,7 @@ void precompile_test_kernels(const std::string& precompile_file)
     std::mt19937       dist(dev());
     std::shuffle(tokens.begin(), tokens.end(), dist);
     auto precompile_begin = std::chrono::steady_clock::now();
-    std::cout << "precompiling " << tokens.size() << " FFT plans...\n";
+    std::cout << "precompiling kernels for " << tokens.size() << " tokens...\n";
 
     for(auto&& t : tokens)
         tokenQueue.push(std::move(t));
@@ -222,6 +224,13 @@ void precompile_test_kernels(const std::string& precompile_file)
                     params.from_token(token);
                     params.validate();
                     params.create_plan();
+                    if(params.is_forward())
+                    {
+                        hipfft_params inverse_params;
+                        inverse_params.inverse_from_forward(params);
+                        inverse_params.validate();
+                        inverse_params.create_plan();
+                    }
                 }
                 catch(fft_params::work_buffer_alloc_failure&)
                 {
@@ -271,26 +280,45 @@ int main(int argc, char* argv[])
         "      HP - hermitian planar\n"
         "\n"
         "Usage"};
-
-    // Override CLI11 help to print after later CLI11 options that are defined, and allow gtest's help
+    // Override CLI11 help to print it along gtest's help
     app.set_help_flag("");
-    CLI::Option* opt_help = app.add_flag("-h, --help", "Produces this help message");
+    const auto opt_help = app.add_flag("-h, --help", "Produces this help message");
     app.add_option("-v, --verbose", verbose, "Print out detailed information for the tests")
         ->default_val(0);
     app.add_option("--test_prob", test_prob, "Probability of running individual tests")
         ->default_val(1.0)
         ->check(CLI::Range(0.0, 1.0));
+    app.add_option("--real_prob",
+                   real_prob_factor,
+                   "Probability multiplier for running individual real/complex transforms")
+        ->default_val(1.0)
+        ->check(CLI::PositiveNumber);
+    app.add_option("--planar_prob",
+                   complex_planar_prob_factor,
+                   "Probability multiplier for running individual planar transforms")
+        ->default_val(0.1)
+        ->check(CLI::PositiveNumber);
     app.add_option(
            "--complex_interleaved_prob_factor",
            complex_interleaved_prob_factor,
            "Probability multiplier for running individual transforms with complex interleaved data")
         ->default_val(1)
-        ->check(CLI::NonNegativeNumber);
+        ->check(CLI::PositiveNumber);
     app.add_option("--callback_prob",
                    callback_prob_factor,
                    "Probability multiplier for running individual callback transforms")
         ->default_val(0.1)
         ->check(CLI::NonNegativeNumber);
+    app.add_option("--max_hipfftw_test_len",
+                   max_length_for_hipfftw_test,
+                   "Maximum length to be considered in hipfftw tests")
+        ->default_val(8192)
+        ->check(CLI::PositiveNumber);
+    app.add_option("--max_io_gb_for_hipfftw_test",
+                   max_io_gb_for_hipfftw_test,
+                   "Maximum size of I/O to be considered in hipfftw tests in GiB")
+        ->default_val(1) /* 1 GiB */
+        ->check(CLI::PositiveNumber);
 
     app.add_option("--fftw_compare", fftw_compare, "Compare to FFTW in accuracy tests")
         ->default_val(true);
@@ -302,71 +330,43 @@ int main(int argc, char* argv[])
     app.add_option("--mp_launch",
                    mp_launch,
                    "Command line prefix to launch multi-process transforms, e.g. \"mpirun --np 4 "
-                   "/path/to/rocfft_mpi_worker\"")
+                   "/path/to/hipfft_mpi_worker\"")
         ->default_val("")
         ->each([&](const std::string&) {
             if(mp_lib == fft_params::fft_mp_lib_none)
             {
-                std::cout << "--mp_launch requires an mp library (see mp_lib in --help).\n";
-                std::exit(-1);
+                throw CLI::ValidationError(
+                    "--mp_launch requires an mp library (see mp_lib in --help)");
             }
         })
         ->needs("--mp_lib");
-    // FIXME: Seed has no use currently
-    // CLI::Option* opt_seed =
-    app.add_option("--seed", random_seed, "Random seed; if unset, use an actual random seed");
+    app.add_option("--seed", random_seed, "Random seed; if unset, use an actual random seed")
+        ->default_val(default_seed_dev());
     app.add_flag("--smoketest", "Run a short (approx 5 minute) randomized selection of tests")
         ->each([&](const std::string&) {
-            // The objective is to have an test that takes about 5 minutes, so just set the probability
-            // per test to a small value to achieve this result.
+            // The objective is to have an test that takes about 5 minutes, so just set the
+            // probability per test to a small value to achieve this result.
             test_prob = 0.002;
         });
-
-    // Try parsing initial args that will be used to configure tests
-    // Allow extras to pass on gtest and hipFFT arguments without error
-    app.allow_extras();
-    try
-    {
-        app.parse(argc, argv);
-    }
-    catch(const CLI::ParseError& e)
-    {
-        return app.exit(e);
-    }
-
-    // NB: If we initialize gtest first, then it removes all of its own command-line
-    // arguments and sets argc and argv correctly;
-    ::testing::InitGoogleTest(&argc, argv);
-
-    // Filename for fftw and fftwf wisdom.
-    std::string fftw_wisdom_filename;
-
     // Token string to fully specify fft params for the manual test.
     std::string test_token;
-
-    // Filename for precompiled kernels to be written to
-    std::string precompile_file;
-
-    // Declare the supported options. Some option pointers are declared to track passed opts.
-    app.add_flag("--callback", "Inject load/store callbacks")->each([&](const std::string&) {
-        manual_params.run_callbacks = true;
-    });
-    // app.add_flag("--version", "Print queryable version information from the rocfft library")
-    //     ->each([](const std::string&) {
-    //         rocfft_setup();
-    //         char v[256];
-    //         rocfft_get_version_string(v, 256);
-    //         std::cout << "rocFFT version: " << v << std::endl;
-    //         return EXIT_SUCCESS;
-    //     });
-    CLI::Option* opt_token
+    auto*       opt_token
         = app.add_option("--token", test_token, "Test token name for manual test")->default_val("");
     // Group together options that conflict with --token
     auto* non_token = app.add_option_group("Token Conflict", "Options excluded by --token");
+    non_token->excludes(opt_token);
+    // Declare the supported options. Some option pointers are declared to track passed opts.
+    non_token->add_flag("--callback", "Inject load/store callbacks")->each([&](const std::string&) {
+        manual_params.run_callbacks = true;
+    });
+    non_token
+        ->add_option("--auto_allocation",
+                     manual_params.auto_allocate,
+                     "hipFFT's auto-allocation behavior: \"on\", \"off\", or \"default\"")
+        ->default_val("default");
     non_token
         ->add_flag("--double", "Double precision transform (deprecated: use --precision double)")
         ->each([&](const std::string&) { manual_params.precision = fft_precision_double; });
-    non_token->excludes(opt_token);
     non_token
         ->add_option("-t, --transformType",
                      manual_params.transform_type,
@@ -406,8 +406,20 @@ int main(int argc, char* argv[])
         ->default_val(0);
     non_token->add_option("--ioffset", manual_params.ioffset, "Input offset");
     non_token->add_option("--ooffset", manual_params.ooffset, "Output offset");
-    app.add_option("--isize", manual_params.isize, "Logical size of input buffer");
-    app.add_option("--osize", manual_params.osize, "Logical size of output buffer");
+    non_token->add_option("--isize", manual_params.isize, "Logical size of input buffer");
+    non_token->add_option("--osize", manual_params.osize, "Logical size of output buffer");
+    non_token->add_option(
+        "--scalefactor", manual_params.scale_factor, "Scale factor to apply to output");
+    // Default value is set in fft_params.h based on if device-side PRNG was enabled.
+    non_token->add_option("-g, --inputGen",
+                          manual_params.igen,
+                          "Input data generation:\n0) PRNG sequence (device)\n"
+                          "1) PRNG sequence (host)\n"
+                          "2) linearly-spaced sequence (device)\n"
+                          "3) linearly-spaced sequence (host)");
+    const auto* opt_version = app.add_flag(
+        "--version",
+        "Print queryable version information from the hipfft library's backend (and return)");
     app.add_option("--R", ramgb, "RAM limit in GiB for tests")
         ->default_val(host_memory::singleton().get_total_gbytes());
     app.add_option("--V", vramgb, "VRAM limit in GiB for tests")->default_val(0);
@@ -419,22 +431,19 @@ int main(int argc, char* argv[])
                    "Skip the test if there is a runtime failure")
         ->default_val(true);
     app.add_option("-w, --wise", use_fftw_wisdom, "Use FFTW wisdom");
+    // Filename for fftw and fftwf wisdom.
+    std::string fftw_wisdom_filename;
     app.add_option("-W, --wisdomfile", fftw_wisdom_filename, "FFTW3 wisdom filename")
         ->default_val("wisdom3.txt");
-    app.add_option("--scalefactor", manual_params.scale_factor, "Scale factor to apply to output");
+    // Filename for precompiled kernels to be written to
+    std::string precompile_file;
     app.add_option("--precompile",
                    precompile_file,
                    "Precompile kernels to a file for all test cases before running tests")
         ->default_val("");
-    // Default value is set in fft_params.h based on if device-side PRNG was enabled.
-    app.add_option("-g, --inputGen",
-                   manual_params.igen,
-                   "Input data generation:\n0) PRNG sequence (device)\n"
-                   "1) PRNG sequence (host)\n"
-                   "2) linearly-spaced sequence (device)\n"
-                   "3) linearly-spaced sequence (host)");
-
-    // Parse rest of args and catch any errors here
+    // Try parsing initial args that will be used to configure tests
+    // Allow extras to pass on gtest arguments without error
+    app.allow_extras();
     try
     {
         app.parse(argc, argv);
@@ -444,26 +453,21 @@ int main(int argc, char* argv[])
         return app.exit(e);
     }
 
-    if(*opt_help)
+    if(!test_token.empty())
     {
-        std::cout << app.help() << "\n";
-        return EXIT_SUCCESS;
-    }
+        std::cout << "Reading fft params from token:\n" << test_token << std::endl;
 
-    // Ensure there are no leftover options used by neither gtest nor CLI11
-    std::vector<std::string> remaining_args = app.remaining();
-    if(!remaining_args.empty())
-    {
-        std::cout << "Unrecognised option(s) found:\n  ";
-        for(auto i : app.remaining())
-            std::cout << i << " ";
-        std::cout << "\nRun with --help for more information.\n";
-        return EXIT_FAILURE;
+        try
+        {
+            manual_params.from_token(test_token);
+            std::cout << "manual_params.token() = " << manual_params.token() << std::endl;
+        }
+        catch(...)
+        {
+            std::cout << "Unable to parse token." << std::endl;
+            return 1;
+        }
     }
-
-    std::cout << "half epsilon: " << half_epsilon << "\tsingle epsilon: " << single_epsilon
-              << "\tdouble epsilon: " << double_epsilon << std::endl;
-
     if(manual_params.length.empty())
     {
         manual_params.length.push_back(8);
@@ -482,6 +486,66 @@ int main(int argc, char* argv[])
         // TODO: add random size?
     }
 
+    // User-settable options defining the values of all the actual test parameters
+    // (e.g., probability factors and value of manual_params) must be handled
+    // before invoking ::testing::InitGoogleTest as it triggers evaluation of said
+    // parameters (e.g., args of "::testing::Values{In}" in instantiations of test
+    // suites).
+    // set any "unset" parameters of manual_params before initiating gtests
+    // (makes the token reported by gtest less ambiguous)
+    manual_params.validate();
+
+    // extract remaining arguments for subsequent gtest initialization
+    std::vector<std::string> remaining_args = app.remaining();
+    std::string              gtest_help_opt = "--help";
+    // NB: If we initialize gtest first, then it removes all of its own command-line
+    // arguments and sets argc and argv correctly;
+    std::vector<char*> gtest_argv;
+    gtest_argv.insert(gtest_argv.begin(), argv[0]);
+    for(std::string& s : remaining_args)
+    {
+        gtest_argv.push_back(&s[0]);
+    }
+    if(*opt_help)
+    {
+        // make sure gtest prints its help as well
+        gtest_argv.push_back(&gtest_help_opt[0]);
+    }
+    gtest_argv.push_back(NULL);
+    decltype(argc) gtest_argc = gtest_argv.size() - 1;
+    ::testing::InitGoogleTest(&gtest_argc, gtest_argv.data()); // gtest-relevant args are removed
+
+    if(*opt_help)
+    {
+        std::cout << app.help() << "\n";
+        return EXIT_SUCCESS;
+    }
+    // no help was used, gtest_argc is expected to be 1 at this point. If not, some of the
+    // used options were not recognized at all
+    if(gtest_argc > 1)
+    {
+        std::cout << "Unrecognised option(s) found:\n  ";
+        for(auto i = 1; i < gtest_argc; i++)
+            std::cout << gtest_argv[i] << " ";
+        std::cout << "\nRun with --help for more information.\n";
+        return EXIT_FAILURE;
+    }
+
+    if(*opt_version || verbose > 0)
+    {
+        int hipfft_version;
+        hipfftGetVersion(&hipfft_version);
+        std::cout << "hipFFT version: " << hipfft_version << std::endl;
+        if(*opt_version)
+        {
+            return EXIT_SUCCESS;
+        }
+    }
+
+    std::cout << "Using random_seed = " << random_seed << std::endl;
+    std::cout << "half epsilon: " << half_epsilon << "\tsingle epsilon: " << single_epsilon
+              << "\tdouble epsilon: " << double_epsilon << std::endl;
+
     // if precompiling, tell rocFFT to use the specified cache file
     // to write kernels to
     //
@@ -494,11 +558,6 @@ int main(int argc, char* argv[])
                                                               precompile_file.c_str());
     }
 
-    // rocfft_setup();
-    // char v[256];
-    // rocfft_get_version_string(v, 256);
-    // std::cout << "rocFFT version: " << v << std::endl;
-
 #ifdef FFTW_MULTITHREAD
     fftw_init_threads();
     fftwf_init_threads();
@@ -508,6 +567,7 @@ int main(int argc, char* argv[])
 
     // Set host memory limit from command-line options
     host_memory::singleton().set_limit_gbytes(ramgb);
+    std::cout << "Host memory limit: " << ramgb << " GiB" << std::endl;
 
     if(use_fftw_wisdom)
     {
@@ -555,21 +615,6 @@ int main(int argc, char* argv[])
         fftwf_import_wisdom_from_string(fftwf_wisdom.c_str());
     }
 
-    if(!test_token.empty())
-    {
-        std::cout << "Reading fft params from token:\n" << test_token << std::endl;
-
-        try
-        {
-            manual_params.from_token(test_token);
-        }
-        catch(...)
-        {
-            std::cout << "Unable to parse token." << std::endl;
-            return 1;
-        }
-    }
-
     if(!precompile_file.empty())
         precompile_test_kernels(precompile_file);
 
@@ -593,38 +638,14 @@ int main(int argc, char* argv[])
     std::cout << "double precision max l-inf epsilon: " << max_linf_eps_double << std::endl;
     std::cout << "double precision max l2 epsilon:     " << max_l2_eps_double << std::endl;
 
-    // rocfft_cleanup();
+    hipfft_params::externally_managed_workareas.clear();
+
     return retval;
 }
 
-TEST(manual, vs_fftw)
-{
-    // Run an individual test using the provided command-line parameters.
-
-    std::cout << "Manual test:" << std::endl;
-
-    manual_params.validate();
-
-    std::cout << "Token: " << manual_params.token() << std::endl;
-
-    hipfft_params params(manual_params);
-
-    try
-    {
-        fft_vs_reference(params, false);
-    }
-    catch(HOSTBUF_MEM_USAGE& e)
-    {
-        // explicitly clear test cache
-        last_cpu_fft_data = last_cpu_fft_cache();
-        GTEST_SKIP() << e.msg;
-    }
-    catch(ROCFFT_SKIP& e)
-    {
-        GTEST_SKIP() << e.msg;
-    }
-    catch(ROCFFT_FAIL& e)
-    {
-        GTEST_FAIL() << e.msg;
-    }
-}
+// instantiation of the paramameterized accuracy_test for the
+// configuration set manually:
+INSTANTIATE_TEST_SUITE_P(manual,
+                         accuracy_test,
+                         ::testing::Values(static_cast<const fft_params&>(manual_params)),
+                         accuracy_test::TestName);
diff --git a/clients/tests/hipfft_accuracy_test.cpp b/clients/tests/hipfft_accuracy_test.cpp
index e228c5f..a5c763e 100644
--- a/clients/tests/hipfft_accuracy_test.cpp
+++ b/clients/tests/hipfft_accuracy_test.cpp
@@ -34,6 +34,7 @@
 #include "../../shared/accuracy_test.h"
 #include "../../shared/fftw_transform.h"
 #include "../../shared/gpubuf.h"
+#include "../../shared/params_gen.h"
 #include "../../shared/rocfft_against_fftw.h"
 #include "../../shared/rocfft_complex.h"
 #include "../../shared/subprocess.h"
@@ -42,6 +43,21 @@ extern std::string mp_launch;
 
 extern last_cpu_fft_cache last_cpu_fft_data;
 
+// clang-format off
+// tokens of tests found to be symptomatic
+static const std::vector<std::string> symptomatic_tokens = {
+#ifndef _CUFFT_BACKEND
+// cases specific to ROCM backend
+#else
+    // cases specific to CUFFT backend
+    "real_forward_len_16384_half_ip_batch_4_istride_1_R_ostride_1_HI_idist_16386_odist_8193_ioffset_0_0_ooffset_0_0",
+    "real_forward_len_32768_half_ip_batch_4_istride_1_R_ostride_1_HI_idist_32770_odist_16385_ioffset_0_0_ooffset_0_0",
+    "real_forward_len_65536_half_ip_batch_2_istride_1_R_ostride_1_HI_idist_65538_odist_32769_ioffset_0_0_ooffset_0_0",
+#endif
+    // common  to both backends
+};
+// clang-format on
+
 void fft_vs_reference(hipfft_params& params, bool round_trip)
 {
     switch(params.precision)
@@ -78,29 +94,43 @@ TEST_P(accuracy_test, vs_fftw)
     {
     case fft_params::fft_mp_lib_none:
     {
-        // only do round trip for single-GPU FFTs
-        bool round_trip = params.multiGPU <= 1;
+        // skipping symptomatic case(s), unless forcefully/knowingly executing normally-disabled
+        // test tokens (e.g., by using --gtest_also_run_disabled_tests)
+        const char* test_suite_name
+            = ::testing::UnitTest::GetInstance()->current_test_info()->test_suite_name();
+        if(!symptomatic_tokens.empty() && std::strstr(test_suite_name, "DISABLED") == nullptr
+           && std::find(symptomatic_tokens.begin(), symptomatic_tokens.end(), params.token())
+                  != symptomatic_tokens.end())
+        {
+            GTEST_SKIP()
+                << "Symptomatic test that's currently disabled by default (force-skipping). Use "
+                   "CLI arguments '--gtest_also_run_disabled_tests' to force the test execution "
+                   "(via another test suite).";
+        }
+        // only do round trip for forward FFTs
+        const bool do_round_trip = params.is_forward();
 
-        if(!params.run_callbacks)
+        try
         {
-            try
-            {
-                fft_vs_reference(params, round_trip);
-            }
-            catch(HOSTBUF_MEM_USAGE& e)
-            {
-                // explicitly clear cache
-                last_cpu_fft_data = last_cpu_fft_cache();
-                GTEST_SKIP() << e.msg;
-            }
-            catch(ROCFFT_SKIP& e)
-            {
-                GTEST_SKIP() << e.msg;
-            }
-            catch(ROCFFT_FAIL& e)
-            {
-                GTEST_FAIL() << e.msg;
-            }
+            fft_vs_reference(params, do_round_trip);
+        }
+        catch(HOSTBUF_MEM_USAGE& e)
+        {
+            // explicitly clear cache
+            last_cpu_fft_data = last_cpu_fft_cache();
+            GTEST_SKIP() << e.msg;
+        }
+        catch(ROCFFT_SKIP& e)
+        {
+            GTEST_SKIP() << e.msg;
+        }
+        catch(const fft_params::unimplemented_exception& e)
+        {
+            GTEST_SKIP() << "Unimplemented exception: " << e.what();
+        }
+        catch(ROCFFT_FAIL& e)
+        {
+            GTEST_FAIL() << e.msg;
         }
         break;
     }
@@ -140,6 +170,11 @@ TEST_P(accuracy_test, vs_fftw)
     SUCCEED();
 }
 
+INSTANTIATE_TEST_SUITE_P(DISABLED_symptomatic_tokens,
+                         accuracy_test,
+                         ::testing::ValuesIn(param_generator_token(test_prob, symptomatic_tokens)),
+                         accuracy_test::TestName);
+
 #ifdef __HIP__
 
 // load/store callbacks - cbdata in each is actually a scalar double
@@ -166,6 +201,37 @@ __device__ auto load_callback_dev_complex_float  = load_callback<rocfft_complex<
 __device__ auto load_callback_dev_double         = load_callback<double>;
 __device__ auto load_callback_dev_complex_double = load_callback<rocfft_complex<double>>;
 
+// load/store callbacks - cbdata in each is actually a scalar double
+// with a number to apply to each element
+template <typename Tdata>
+__host__ __device__ Tdata
+    load_callback_round_trip_inverse(Tdata* input, size_t offset, void* cbdata, void* sharedMem)
+{
+    auto testdata = static_cast<const callback_test_data*>(cbdata);
+    // subtract each element by scalar
+    if(input == testdata->base)
+        return input[offset] - testdata->scalar;
+    // wrong base address passed, return something obviously wrong
+    else
+    {
+        // wrong base address passed, return something obviously wrong
+        return input[0];
+    }
+}
+
+__device__ auto load_callback_round_trip_inverse_dev_half
+    = load_callback_round_trip_inverse<rocfft_fp16>;
+__device__ auto load_callback_round_trip_inverse_dev_complex_half
+    = load_callback_round_trip_inverse<rocfft_complex<rocfft_fp16>>;
+__device__ auto load_callback_round_trip_inverse_dev_float
+    = load_callback_round_trip_inverse<float>;
+__device__ auto load_callback_round_trip_inverse_dev_complex_float
+    = load_callback_round_trip_inverse<rocfft_complex<float>>;
+__device__ auto load_callback_round_trip_inverse_dev_double
+    = load_callback_round_trip_inverse<double>;
+__device__ auto load_callback_round_trip_inverse_dev_complex_double
+    = load_callback_round_trip_inverse<rocfft_complex<double>>;
+
 void* get_load_callback_host(fft_array_type itype,
                              fft_precision  precision,
                              bool           round_trip_inverse = false)
@@ -179,22 +245,55 @@ void* get_load_callback_host(fft_array_type itype,
         switch(precision)
         {
         case fft_precision_half:
-            EXPECT_EQ(hipMemcpyFromSymbol(&load_callback_host,
-                                          HIP_SYMBOL(load_callback_dev_complex_half),
-                                          sizeof(void*)),
-                      hipSuccess);
+            if(round_trip_inverse)
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(
+                              &load_callback_host,
+                              HIP_SYMBOL(load_callback_round_trip_inverse_dev_complex_half),
+                              sizeof(void*)),
+                          hipSuccess);
+            }
+            else
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&load_callback_host,
+                                              HIP_SYMBOL(load_callback_dev_complex_half),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
             return load_callback_host;
         case fft_precision_single:
-            EXPECT_EQ(hipMemcpyFromSymbol(&load_callback_host,
-                                          HIP_SYMBOL(load_callback_dev_complex_float),
-                                          sizeof(void*)),
-                      hipSuccess);
+            if(round_trip_inverse)
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(
+                              &load_callback_host,
+                              HIP_SYMBOL(load_callback_round_trip_inverse_dev_complex_float),
+                              sizeof(void*)),
+                          hipSuccess);
+            }
+            else
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&load_callback_host,
+                                              HIP_SYMBOL(load_callback_dev_complex_float),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
             return load_callback_host;
         case fft_precision_double:
-            EXPECT_EQ(hipMemcpyFromSymbol(&load_callback_host,
-                                          HIP_SYMBOL(load_callback_dev_complex_double),
-                                          sizeof(void*)),
-                      hipSuccess);
+            if(round_trip_inverse)
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(
+                              &load_callback_host,
+                              HIP_SYMBOL(load_callback_round_trip_inverse_dev_complex_double),
+                              sizeof(void*)),
+                          hipSuccess);
+            }
+            else
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&load_callback_host,
+                                              HIP_SYMBOL(load_callback_dev_complex_double),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
             return load_callback_host;
         }
     }
@@ -203,19 +302,55 @@ void* get_load_callback_host(fft_array_type itype,
         switch(precision)
         {
         case fft_precision_half:
-            EXPECT_EQ(hipMemcpyFromSymbol(
-                          &load_callback_host, HIP_SYMBOL(load_callback_dev_half), sizeof(void*)),
-                      hipSuccess);
+            if(round_trip_inverse)
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&load_callback_host,
+                                              HIP_SYMBOL(load_callback_round_trip_inverse_dev_half),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
+            else
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&load_callback_host,
+                                              HIP_SYMBOL(load_callback_dev_half),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
             return load_callback_host;
         case fft_precision_single:
-            EXPECT_EQ(hipMemcpyFromSymbol(
-                          &load_callback_host, HIP_SYMBOL(load_callback_dev_float), sizeof(void*)),
-                      hipSuccess);
+            if(round_trip_inverse)
+            {
+                EXPECT_EQ(
+                    hipMemcpyFromSymbol(&load_callback_host,
+                                        HIP_SYMBOL(load_callback_round_trip_inverse_dev_float),
+                                        sizeof(void*)),
+                    hipSuccess);
+            }
+            else
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&load_callback_host,
+                                              HIP_SYMBOL(load_callback_dev_float),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
             return load_callback_host;
         case fft_precision_double:
-            EXPECT_EQ(hipMemcpyFromSymbol(
-                          &load_callback_host, HIP_SYMBOL(load_callback_dev_double), sizeof(void*)),
-                      hipSuccess);
+
+            if(round_trip_inverse)
+            {
+                EXPECT_EQ(
+                    hipMemcpyFromSymbol(&load_callback_host,
+                                        HIP_SYMBOL(load_callback_round_trip_inverse_dev_double),
+                                        sizeof(void*)),
+                    hipSuccess);
+            }
+            else
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&load_callback_host,
+                                              HIP_SYMBOL(load_callback_dev_double),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
             return load_callback_host;
         }
     }
@@ -245,6 +380,31 @@ __device__ auto store_callback_dev_complex_float  = store_callback<rocfft_comple
 __device__ auto store_callback_dev_double         = store_callback<double>;
 __device__ auto store_callback_dev_complex_double = store_callback<rocfft_complex<double>>;
 
+template <typename Tdata>
+__host__ __device__ static void store_callback_round_trip_inverse(
+    Tdata* output, size_t offset, Tdata element, void* cbdata, void* sharedMem)
+{
+    auto testdata = static_cast<callback_test_data*>(cbdata);
+    // divide each element by scalar
+    if(output == testdata->base)
+    {
+        output[offset] = element / testdata->scalar;
+    }
+    // otherwise, wrong base address passed, just don't write
+}
+__device__ auto store_callback_round_trip_inverse_dev_half
+    = store_callback_round_trip_inverse<rocfft_fp16>;
+__device__ auto store_callback_round_trip_inverse_dev_complex_half
+    = store_callback_round_trip_inverse<rocfft_complex<rocfft_fp16>>;
+__device__ auto store_callback_round_trip_inverse_dev_float
+    = store_callback_round_trip_inverse<float>;
+__device__ auto store_callback_round_trip_inverse_dev_complex_float
+    = store_callback_round_trip_inverse<rocfft_complex<float>>;
+__device__ auto store_callback_round_trip_inverse_dev_double
+    = store_callback_round_trip_inverse<double>;
+__device__ auto store_callback_round_trip_inverse_dev_complex_double
+    = store_callback_round_trip_inverse<rocfft_complex<double>>;
+
 void* get_store_callback_host(fft_array_type otype,
                               fft_precision  precision,
                               bool           round_trip_inverse = false)
@@ -258,22 +418,55 @@ void* get_store_callback_host(fft_array_type otype,
         switch(precision)
         {
         case fft_precision_half:
-            EXPECT_EQ(hipMemcpyFromSymbol(&store_callback_host,
-                                          HIP_SYMBOL(store_callback_dev_complex_half),
-                                          sizeof(void*)),
-                      hipSuccess);
+            if(round_trip_inverse)
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(
+                              &store_callback_host,
+                              HIP_SYMBOL(store_callback_round_trip_inverse_dev_complex_half),
+                              sizeof(void*)),
+                          hipSuccess);
+            }
+            else
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&store_callback_host,
+                                              HIP_SYMBOL(store_callback_dev_complex_half),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
             return store_callback_host;
         case fft_precision_single:
-            EXPECT_EQ(hipMemcpyFromSymbol(&store_callback_host,
-                                          HIP_SYMBOL(store_callback_dev_complex_float),
-                                          sizeof(void*)),
-                      hipSuccess);
+            if(round_trip_inverse)
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(
+                              &store_callback_host,
+                              HIP_SYMBOL(store_callback_round_trip_inverse_dev_complex_float),
+                              sizeof(void*)),
+                          hipSuccess);
+            }
+            else
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&store_callback_host,
+                                              HIP_SYMBOL(store_callback_dev_complex_float),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
             return store_callback_host;
         case fft_precision_double:
-            EXPECT_EQ(hipMemcpyFromSymbol(&store_callback_host,
-                                          HIP_SYMBOL(store_callback_dev_complex_double),
-                                          sizeof(void*)),
-                      hipSuccess);
+            if(round_trip_inverse)
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(
+                              &store_callback_host,
+                              HIP_SYMBOL(store_callback_round_trip_inverse_dev_complex_double),
+                              sizeof(void*)),
+                          hipSuccess);
+            }
+            else
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&store_callback_host,
+                                              HIP_SYMBOL(store_callback_dev_complex_double),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
             return store_callback_host;
         }
     }
@@ -282,21 +475,55 @@ void* get_store_callback_host(fft_array_type otype,
         switch(precision)
         {
         case fft_precision_half:
-            EXPECT_EQ(hipMemcpyFromSymbol(
-                          &store_callback_host, HIP_SYMBOL(store_callback_dev_half), sizeof(void*)),
-                      hipSuccess);
+            if(round_trip_inverse)
+            {
+                EXPECT_EQ(
+                    hipMemcpyFromSymbol(&store_callback_host,
+                                        HIP_SYMBOL(store_callback_round_trip_inverse_dev_half),
+                                        sizeof(void*)),
+                    hipSuccess);
+            }
+            else
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&store_callback_host,
+                                              HIP_SYMBOL(store_callback_dev_half),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
             return store_callback_host;
         case fft_precision_single:
-            EXPECT_EQ(hipMemcpyFromSymbol(&store_callback_host,
-                                          HIP_SYMBOL(store_callback_dev_float),
-                                          sizeof(void*)),
-                      hipSuccess);
+            if(round_trip_inverse)
+            {
+                EXPECT_EQ(
+                    hipMemcpyFromSymbol(&store_callback_host,
+                                        HIP_SYMBOL(store_callback_round_trip_inverse_dev_float),
+                                        sizeof(void*)),
+                    hipSuccess);
+            }
+            else
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&store_callback_host,
+                                              HIP_SYMBOL(store_callback_dev_float),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
             return store_callback_host;
         case fft_precision_double:
-            EXPECT_EQ(hipMemcpyFromSymbol(&store_callback_host,
-                                          HIP_SYMBOL(store_callback_dev_double),
-                                          sizeof(void*)),
-                      hipSuccess);
+            if(round_trip_inverse)
+            {
+                EXPECT_EQ(
+                    hipMemcpyFromSymbol(&store_callback_host,
+                                        HIP_SYMBOL(store_callback_round_trip_inverse_dev_double),
+                                        sizeof(void*)),
+                    hipSuccess);
+            }
+            else
+            {
+                EXPECT_EQ(hipMemcpyFromSymbol(&store_callback_host,
+                                              HIP_SYMBOL(store_callback_dev_double),
+                                              sizeof(void*)),
+                          hipSuccess);
+            }
             return store_callback_host;
         }
     }
diff --git a/clients/tests/hipfft_mpi_worker.cpp b/clients/tests/hipfft_mpi_worker.cpp
index 9a1e219..a4b85c6 100644
--- a/clients/tests/hipfft_mpi_worker.cpp
+++ b/clients/tests/hipfft_mpi_worker.cpp
@@ -23,6 +23,9 @@
 #include "../../shared/mpi_worker.h"
 #include "../hipfft_params.h"
 
+// initialize static class member of hipfft_params
+std::vector<gpubuf> hipfft_params::externally_managed_workareas = std::vector<gpubuf>();
+
 int main(int argc, char* argv[])
 {
     return mpi_worker_main<std::array<hipfft_params, 1>, false>(
diff --git a/clients/tests/hipfft_test_params.h b/clients/tests/hipfft_test_params.h
deleted file mode 100644
index 4d22fbb..0000000
--- a/clients/tests/hipfft_test_params.h
+++ /dev/null
@@ -1,32 +0,0 @@
-// Copyright (C) 2022 - 2022 Advanced Micro Devices, Inc. All rights reserved.
-//
-// 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.
-
-#pragma once
-#ifndef TESTCONSTANTS_H
-#define TESTCONSTANTS_H
-
-#include "hipfft/hipfft.h"
-
-#include <stdexcept>
-
-extern int    verbose;
-extern size_t ramgb;
-
-#endif
diff --git a/clients/tests/hipfftw_test.cpp b/clients/tests/hipfftw_test.cpp
new file mode 100644
index 0000000..b0ad0b0
--- /dev/null
+++ b/clients/tests/hipfftw_test.cpp
@@ -0,0 +1,2214 @@
+// Copyright (C) 2025 Advanced Micro Devices, Inc. All rights reserved.
+//
+// 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.
+
+#include "../hipfftw_helper.h"
+
+#include "../../shared/environment.h"
+#include "../../shared/gpubuf.h"
+#include "../../shared/hostbuf.h"
+#include "../../shared/params_gen.h"
+#include "../../shared/test_params.h"
+
+#include <cstdint>
+#include <cstring>
+#include <fftw3.h>
+#include <future>
+#include <gtest/gtest.h>
+#include <memory>
+#include <numeric>
+
+#ifdef WIN32
+#include <windows.h>
+#else
+#include <cstdlib>
+#endif
+
+#ifdef _OPENMP
+#include <omp.h>
+#endif
+
+// test details
+namespace
+{
+    //
+    //---------------------------------------------------------------------------------------------
+    //                                  COMMONS AND HELPERS
+    //---------------------------------------------------------------------------------------------
+    //
+
+    size_t max_byte_size_for_hipfftw_tests()
+    {
+        auto get_io_byte_size_limit = []() {
+            size_t tmp = vramgb * ONE_GiB;
+            if(tmp == 0)
+            {
+                size_t free = 0, total = 0;
+                if(hipMemGetInfo(&free, &total) == hipSuccess)
+                    tmp = total / 8;
+            }
+            tmp = std::min(tmp, ramgb * ONE_GiB);
+            tmp = std::min(tmp, max_io_gb_for_hipfftw_test * ONE_GiB);
+            if(verbose > 0)
+            {
+                std::cout << "Limit for the size of I/O data used in hipfftw tests: ";
+                if(tmp >= ONE_GiB)
+                    std::cout << float(tmp) / ONE_GiB << " GiB." << std ::endl;
+                else if(tmp >= ONE_MiB)
+                    std::cout << float(tmp) / ONE_MiB << " MiB." << std ::endl;
+                else
+                    std::cout << float(tmp) / ONE_KiB << " KiB." << std ::endl;
+            }
+            return tmp;
+        };
+        static const size_t io_byte_size_limit = get_io_byte_size_limit();
+        return io_byte_size_limit;
+    }
+
+    std::ranlux24_base& get_pseudo_rng()
+    {
+        static std::ranlux24_base gen(random_seed);
+        return gen;
+    }
+
+    // NOTE: this function makes use of comparison operator < and != which must be defined for the
+    // specialization type T
+    template <typename T>
+    void insert_into_unique_sorted_params(std::vector<T>& unique_sorted_params,
+                                          const T&        param_to_insert)
+    {
+        auto it = std::lower_bound(
+            unique_sorted_params.begin(), unique_sorted_params.end(), param_to_insert);
+        if(it != unique_sorted_params.end() && *it == param_to_insert)
+            return; // it's already in there generated
+        unique_sorted_params.insert(it, param_to_insert);
+    }
+
+    enum class hipfftw_internal_exception
+    {
+        none,
+        flow_redirection,
+        invalid_args,
+        unsupported_args,
+        ill_defined
+    };
+
+    // for well-defined internal exceptions, we may expect a specific report thereof
+    // in the hipfftw exception log (if logging is activated)
+    template <hipfftw_internal_exception>
+    constexpr std::string_view hipfftw_expected_log_instance;
+    template <>
+    constexpr std::string_view hipfftw_expected_log_instance<
+        hipfftw_internal_exception::invalid_args> = R"(Invalid argument reported)";
+    template <>
+    constexpr std::string_view hipfftw_expected_log_instance<
+        hipfftw_internal_exception::unsupported_args> = R"(Unsupported usage reported)";
+    template <>
+    constexpr std::string_view hipfftw_expected_log_instance<
+        hipfftw_internal_exception::flow_redirection> = R"(Redirecting execution flow)";
+
+    // randomizers
+    // Note: albeit not supported, ranks > 3 are "valid" rank argument
+    // --> limiting rank value to max of 10 by default to avoid ridiculously long
+    // lengths possibly created in automated parameter generations;
+    template <bool validity_flag,
+              int  min_rank          = validity_flag ? 1 : std::numeric_limits<int>::lowest(),
+              int  max_rank          = validity_flag ? 10 : 0,
+              std::enable_if_t<(min_rank <= max_rank) && (!validity_flag || min_rank > 0)
+                                   && (validity_flag || max_rank <= 0),
+                               bool> = true>
+    int get_random_rank()
+    {
+        static std::uniform_int_distribution<int> rank_rng(min_rank, max_rank);
+
+        auto ret = rank_rng(get_pseudo_rng());
+        if(rank_is_valid_for_hipfftw(ret) != validity_flag)
+        {
+            throw std::runtime_error(
+                "failed to generate a rank value of desired validity randomly");
+        }
+        return ret;
+    }
+    template <bool validity_flag, typename type_to_consider_for_validity>
+    std::vector<ptrdiff_t>
+        get_random_lengths(int       desired_rank,
+                           ptrdiff_t max_abs_len
+                           = std::numeric_limits<type_to_consider_for_validity>::max(),
+                           ptrdiff_t min_abs_len = 0)
+    {
+        std::vector<ptrdiff_t> ret;
+        // cannot generate lengths for invalid ranks --> return empty lengths in that case
+        if(!rank_is_valid_for_hipfftw(desired_rank))
+            return ret;
+        if(min_abs_len < 0 || min_abs_len > max_abs_len)
+            throw std::invalid_argument("invalid bounds used for get_random_lengths");
+        // generate values that are all representable as integers
+        auto&                                    pseudo_rng = get_pseudo_rng();
+        std::uniform_int_distribution<ptrdiff_t> length_rng(min_abs_len, max_abs_len);
+        // setter lambda
+        auto set_random_len = [&]() {
+            for(auto& l : ret)
+            {
+                const ptrdiff_t val = length_rng(pseudo_rng);
+                if constexpr(validity_flag)
+                    l = val;
+                else
+                {
+                    if(pseudo_rng() % 2)
+                        l = -val;
+                    else
+                        l = val;
+                }
+            }
+        };
+
+        ret.resize(desired_rank);
+        set_random_len();
+        while(lengths_are_valid_for_hipfftw_as<type_to_consider_for_validity>(ret, desired_rank)
+              != validity_flag)
+            set_random_len();
+
+        return ret;
+    }
+    template <bool validity_flag>
+    int get_random_sign(fft_transform_type intended_dft_kind)
+    {
+        if(!validity_flag && is_real(intended_dft_kind))
+            throw std::invalid_argument("An invalid sign cannot be generated for real transforms "
+                                        "(sign is ignored in that case)");
+        std::uniform_int_distribution<int> sign_rng(std::numeric_limits<int>::lowest(),
+                                                    std::numeric_limits<int>::max());
+
+        int tmp = validity_flag && is_complex(intended_dft_kind)
+                      ? (is_fwd(intended_dft_kind) ? FFTW_FORWARD : FFTW_BACKWARD)
+                      : sign_rng(get_pseudo_rng());
+
+        while(sign_is_valid_for_hipfftw(tmp, intended_dft_kind) != validity_flag)
+            tmp = sign_rng(get_pseudo_rng());
+        return tmp;
+    }
+    template <bool validity_flag>
+    unsigned get_random_flags()
+    {
+        std::uniform_int_distribution<unsigned> flags_rng(std::numeric_limits<unsigned>::lowest(),
+                                                          std::numeric_limits<unsigned>::max());
+
+        auto tmp = flags_rng(get_pseudo_rng());
+        if constexpr(validity_flag)
+        {
+            tmp &= hipfftw_valid_flags_mask;
+            if(!flags_are_valid_for_hipfftw(tmp))
+                throw std::runtime_error("failed to create random valid flags");
+            return tmp;
+        }
+        while(flags_are_valid_for_hipfftw(tmp))
+            tmp = flags_rng(get_pseudo_rng());
+        return tmp;
+    }
+    size_t get_random_idx(size_t upper_bound)
+    {
+        if(upper_bound == 0)
+            throw std::invalid_argument("upper_bound must be strictly positive for get_random_idx");
+        std::uniform_int_distribution<size_t> idx_rng(0, upper_bound - 1);
+        return idx_rng(get_pseudo_rng());
+    }
+
+    // calculates the threshold value X such that max_data_idx is no greater
+    // (resp. larger) than num_elems, if all elements of lengths are all no greater
+    // (resp. larger) than X, and lengths.size() == rank [using bisection]
+    ptrdiff_t get_len_threshold(size_t num_elems, int rank, bool is_real_inplace)
+    {
+        if(rank < 1)
+            throw std::invalid_argument("invalid rank used in get_len_threshold");
+        if(num_elems == 0)
+            return 1;
+        constexpr ptrdiff_t X_max = std::numeric_limits<ptrdiff_t>::max();
+        // we need to find X in [0, X_max] s.t.
+        // largest_idx(X) <= num_elems && largest_idx(X + 1) > num_elems
+        auto largest_idx = [&](ptrdiff_t X) {
+            size_t ret = rank > 1 && is_real_inplace ? 2 * (X / 2 + 1) : X;
+            for(auto i = 1; i < rank; i++)
+                ret *= X;
+            return ret;
+        };
+        // initialization
+        ptrdiff_t X_down
+            = rank == 1 ? static_cast<ptrdiff_t>(num_elems)
+                        : static_cast<ptrdiff_t>(std::floor(std::pow(num_elems, 1.0 / rank)));
+        ptrdiff_t diff = 1;
+        ptrdiff_t X_up = X_down;
+        while(largest_idx(X_up) <= num_elems && X_up < X_max)
+        {
+            X_down = X_up;
+            X_up   = X_up <= X_max - diff ? X_up + diff : X_max;
+            diff *= 2;
+        }
+        diff = 1;
+        while(largest_idx(X_down) > num_elems && X_down > 0)
+        {
+            X_up   = X_down;
+            X_down = X_down >= diff ? X_down - diff : 0;
+            diff *= 2;
+        }
+        // bisection
+        while(X_up - X_down > 1)
+        {
+            const auto tmp = (X_up + X_down) / 2;
+            if(largest_idx(tmp) <= num_elems)
+                X_down = tmp;
+            else
+                X_up = tmp;
+        }
+        return X_down;
+    }
+
+    template <fft_precision prec>
+    size_t max_num_elems_for_data_size(size_t data_byte_size, fft_transform_type dft_kind)
+    {
+        return data_byte_size
+               / (is_real(dft_kind) ? sizeof(hipfftw_real_t<prec>)
+                                    : sizeof(hipfftw_complex_t<prec>));
+    }
+
+    // exception for hip runtime error(s) specifically
+    struct hip_runtime_error : public std::runtime_error
+    {
+        const hipError_t hip_error;
+        hip_runtime_error(const std::string& info, hipError_t hip_status)
+            : std::runtime_error::runtime_error(info)
+            , hip_error(hip_status)
+
+        {
+        }
+    };
+    int get_current_device_id()
+    {
+        int        ret        = hipInvalidDeviceId;
+        const auto hip_status = hipGetDevice(&ret);
+        if(hip_status != hipSuccess)
+            throw hip_runtime_error("hipGetDevice failed.", hip_status);
+        return ret;
+    }
+
+    //
+    //---------------------------------------------------------------------------------------------
+    //                             EXISTENCE OF UTILITY FUNCTIONS
+    //---------------------------------------------------------------------------------------------
+    //
+
+    template <fft_precision prec>
+    void test_existence_of_utility_functions()
+    {
+        try
+        {
+            // call utility functions - they need to exist but don't need to work
+            const auto& hipfftw_ = hipfftw_funcs<prec>::get_instance();
+            hipfftw_.print_plan(nullptr);
+            hipfftw_.set_timelimit(0.0);
+            hipfftw_.cost(nullptr);
+            hipfftw_.flops(nullptr, nullptr, nullptr, nullptr);
+            hipfftw_.cleanup();
+        }
+        catch(const hipfftw_undefined_function_ptr& e)
+        {
+            GTEST_FAIL() << "Undefined function pointers detected. Error info: " << e.what();
+        }
+        catch(...)
+        {
+            GTEST_FAIL() << "Unexpected failure";
+        }
+    }
+
+    //
+    //---------------------------------------------------------------------------------------------
+    //                                   ALLOCATION AND FREE
+    //---------------------------------------------------------------------------------------------
+    //
+
+    enum class hipfftw_alloc_func_type
+    {
+        unspecified,
+        real,
+        complex
+    };
+    bool hipfftw_alloc_func_type_is_valid(hipfftw_alloc_func_type func)
+    {
+        return func == hipfftw_alloc_func_type::unspecified || func == hipfftw_alloc_func_type::real
+               || func == hipfftw_alloc_func_type::complex;
+    }
+    // bit mask to prevent allocation kind(s), by increasing "rank" to enable meaningful
+    // comparison operators (implicitly defined based on the underlying type)
+    enum hipfftw_alloc_memkind : unsigned
+    {
+        none          = 0x0,
+        pageable_host = 0x1 << 0,
+        pinned_host   = 0x1 << 1,
+        any           = pageable_host | pinned_host
+    };
+
+    const std::vector<hipfftw_alloc_memkind> hipfftw_possible_memkinds
+        = {pinned_host, pageable_host};
+
+    bool hipfftw_alloc_kind_is_valid(hipfftw_alloc_memkind kind)
+    {
+        return kind == (kind & hipfftw_alloc_memkind::any);
+    }
+    std::string hipfftw_alloc_kind_to_string(hipfftw_alloc_memkind kind)
+    {
+        if(!hipfftw_alloc_kind_is_valid(kind))
+            throw std::invalid_argument("alloc_kind_to_string: invalid kind");
+        if(kind == hipfftw_alloc_memkind::none)
+            return "none";
+        if(std::find(hipfftw_possible_memkinds.begin(), hipfftw_possible_memkinds.end(), kind)
+           == hipfftw_possible_memkinds.end())
+        {
+            // several values enabled
+            std::string ret;
+            for(auto to_consider : hipfftw_possible_memkinds)
+            {
+                if(!(kind & to_consider))
+                    continue;
+                if(!ret.empty())
+                    ret += "_or_";
+                ret += hipfftw_alloc_kind_to_string(to_consider);
+            }
+            return ret;
+        }
+        // kind is a well-defined value in hipfftw_possible_memkinds
+        switch(kind)
+        {
+        case hipfftw_alloc_memkind::pinned_host:
+            return "pinned_host";
+            break;
+        case hipfftw_alloc_memkind::pageable_host:
+            return "pageable_host";
+            break;
+        default:
+            throw std::runtime_error("alloc_kind_to_string: internal error encountered "
+                                     "(unexpected value for kind)");
+            break;
+        }
+        // unreachable
+    }
+
+    template <fft_precision prec>
+    struct hipfftw_malloc_params
+    {
+        size_t                  alloc_arg;
+        hipfftw_alloc_func_type alloc_func;
+        hipfftw_alloc_memkind   alloc_kind;
+
+        size_t get_byte_size() const
+        {
+            return alloc_arg
+                   * (alloc_func == hipfftw_alloc_func_type::unspecified
+                          ? sizeof(char)
+                          : (alloc_func == hipfftw_alloc_func_type::real
+                                 ? sizeof(hipfftw_real_t<prec>)
+                                 : sizeof(hipfftw_complex_t<prec>)));
+        }
+
+        std::string to_string() const
+        {
+            if(!hipfftw_alloc_func_type_is_valid(alloc_func))
+                throw std::runtime_error("invalid type of allocation function");
+            if(!hipfftw_alloc_kind_is_valid(alloc_kind))
+                throw std::runtime_error("invalid allocation kind(s)");
+
+            std::string ret;
+            if constexpr(prec == fft_precision_single)
+                ret += "fftwf_";
+            else
+                ret += "fftw_";
+
+            if(alloc_func == hipfftw_alloc_func_type::unspecified)
+                ret += "malloc_";
+            else if(alloc_func == hipfftw_alloc_func_type::real)
+                ret += "alloc_real_";
+            else
+                ret += "alloc_complex_";
+            ret += std::to_string(alloc_arg);
+            ret += "_alloc_kind_" + hipfftw_alloc_kind_to_string(alloc_kind);
+            return ret;
+        }
+        // for using with insert_into_unique_sorted_params
+        bool operator<(const hipfftw_malloc_params& other) const
+        {
+            return to_string() < other.to_string();
+        }
+        bool operator==(const hipfftw_malloc_params& other) const
+        {
+            return to_string() == other.to_string();
+        }
+    };
+
+    template <fft_precision prec>
+    std::vector<hipfftw_malloc_params<prec>> params_for_testing_hipfftw_malloc()
+    {
+        std::vector<hipfftw_malloc_params<prec>> ret;
+        // testing argument value 0 and a randomly chosen one (max 64MiB in byte size, arbitrarily chosen)
+        constexpr size_t max_test_alloc_size = 1ULL << 26;
+
+        const std::vector<hipfftw_alloc_func_type> func_range
+            = {hipfftw_alloc_func_type::unspecified,
+               hipfftw_alloc_func_type::real,
+               hipfftw_alloc_func_type::complex};
+        std::vector<hipfftw_alloc_memkind> memkind_range;
+        // add all possible combinations of memory kinds:
+        for(auto kind = static_cast<std::underlying_type<hipfftw_alloc_memkind>::type>(
+                hipfftw_alloc_memkind::none);
+            kind <= static_cast<std::underlying_type<hipfftw_alloc_memkind>::type>(
+                hipfftw_alloc_memkind::any);
+            kind++)
+        {
+            memkind_range.push_back(static_cast<hipfftw_alloc_memkind>(kind));
+        }
+
+        hipfftw_malloc_params<prec> to_add;
+        for(auto func : func_range)
+        {
+            size_t max_arg = max_test_alloc_size;
+            if(func == hipfftw_alloc_func_type::real)
+                max_arg /= sizeof(hipfftw_real_t<prec>);
+            else if(func == hipfftw_alloc_func_type::complex)
+                max_arg /= sizeof(hipfftw_complex_t<prec>);
+            std::uniform_int_distribution<size_t> arg_rng(1, max_arg);
+            for(auto kind : memkind_range)
+            {
+                for(auto arg : {size_t(0), arg_rng(get_pseudo_rng())})
+                {
+                    to_add.alloc_arg  = arg;
+                    to_add.alloc_func = func;
+                    to_add.alloc_kind = kind;
+                    insert_into_unique_sorted_params(ret, to_add);
+                }
+            }
+        }
+        return ret;
+    }
+
+    template <fft_precision prec>
+    class hipfftw_allocation_test : public ::testing::TestWithParam<hipfftw_malloc_params<prec>>
+    {
+    protected:
+        void* test_allocation = nullptr;
+        bool  expect_no_allocation;
+        std::map<hipfftw_alloc_memkind, std::unique_ptr<EnvironmentSetTemp>> temp_alloc_limit_env;
+
+        void SetUp() override
+        {
+            if(test_allocation)
+                GTEST_FAIL() << "Starting from an unclean slate (test_allocation is not nullptr)";
+            const hipfftw_malloc_params<prec>& params = this->GetParam();
+            // check validity of params
+            if(!hipfftw_alloc_kind_is_valid(params.alloc_kind))
+                GTEST_FAIL() << "invalid value for allocation kind";
+            if(!hipfftw_alloc_func_type_is_valid(params.alloc_func))
+                GTEST_FAIL() << "unknown allocation function";
+
+            size_t limit_for_alloc_kind = 0;
+            for(auto alloc_kind_candidate : hipfftw_possible_memkinds)
+            {
+                if(alloc_kind_candidate != hipfftw_alloc_memkind::pinned_host
+                   && alloc_kind_candidate != hipfftw_alloc_memkind::pageable_host)
+                {
+                    throw std::runtime_error("unexpected memory allocation kind "
+                                             + hipfftw_alloc_kind_to_string(alloc_kind_candidate));
+                }
+                const std::string control_env_var
+                    = alloc_kind_candidate == hipfftw_alloc_memkind::pinned_host
+                          ? "HIPFFTW_BYTE_SIZE_LIMIT_PINNED_HOST_ALLOC"
+                          : "HIPFFTW_BYTE_SIZE_LIMIT_PAGEABLE_HOST_ALLOC";
+
+                if(alloc_kind_candidate & params.alloc_kind)
+                {
+                    const auto test_user_limit = rocfft_getenv(control_env_var.c_str());
+                    limit_for_alloc_kind
+                        = std::max(limit_for_alloc_kind,
+                                   test_user_limit.empty() ? std::numeric_limits<size_t>::max()
+                                                           : size_t(std::stoull(test_user_limit)));
+                }
+                else
+                {
+                    // disable the other possible allocation kind(s) by temporarily
+                    // setting the corresponding byte size limit to 0
+                    temp_alloc_limit_env[alloc_kind_candidate]
+                        = std::make_unique<EnvironmentSetTemp>(control_env_var.c_str(), "0");
+                    // skip if temporary limit(s) was(were) not successfully set
+                    const auto tmp_limit = rocfft_getenv(control_env_var.c_str());
+                    if(tmp_limit.empty() || std::stoull(tmp_limit) != 0)
+                    {
+                        GTEST_SKIP() << "failed to set environment variable disabling "
+                                     << hipfftw_alloc_kind_to_string(alloc_kind_candidate)
+                                     << " allocation(s) by hipFFTW";
+                    }
+                }
+            }
+            const size_t req_byte_size = params.get_byte_size();
+            expect_no_allocation       = params.alloc_kind == hipfftw_alloc_memkind::none
+                                   || req_byte_size == 0 || req_byte_size > limit_for_alloc_kind;
+        }
+        void TearDown() override
+        {
+            temp_alloc_limit_env.clear();
+            const hipfftw_funcs<prec>& hipfftw_impl = hipfftw_funcs<prec>::get_instance();
+            if(test_allocation && !hipfftw_impl.free.may_be_used())
+                GTEST_FAIL() << "An allocation was created but it can't be freed";
+            // note: free should be stable even with nullptr
+            if(hipfftw_impl.free.may_be_used())
+                hipfftw_impl.free(test_allocation);
+        }
+
+        void test_malloc_write_and_read()
+        {
+            const hipfftw_malloc_params<prec>& params       = this->GetParam();
+            const hipfftw_funcs<prec>&         hipfftw_impl = hipfftw_funcs<prec>::get_instance();
+            hipfftw_exception_logger           exception_logger;
+
+            struct allocation_test_to_be_skipped : std::runtime_error
+            {
+                using std::runtime_error::runtime_error;
+            };
+            struct allocation_test_failed : std::runtime_error
+            {