Skip to content

[hipfftw tests] Use reference_fft_data in hipfftw tests + avoid H2H copies for host-residing results #7189

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
regan-amd merged 3 commits into from
Jul 8, 2026
Merged

[hipfftw tests] Use reference_fft_data in hipfftw tests + avoid H2H copies for host-residing results #7189

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
04b6e08
[hipfftw tests] Use reference_fft_data in hipfftw tests + avoid H2H c…
regan-amd May 8, 2026
8821687
Aligning with PR#7892
regan-amd Jun 8, 2026
869f43f
Use a non-owned hostbuffer encapsulating a raw pointer to a managed a…
regan-amd Jun 23, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
31 changes: 31 additions & 0 deletions projects/hipfft/clients/hipfftw_helper.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2111,6 +2111,37 @@ struct hipfftw_helper
== default_distances(
dft_kind, plan_placement, fft_io::fft_io_out, lengths, batches);
}

operator fft_params() const
{
fft_params ret;
ret.length = convert_vector_to<decltype(ret.length)::value_type>(lengths);
ret.precision = prec;
ret.placement = plan_placement;
ret.transform_type = dft_kind;
if(batch_rank != 1 || batches[0] < 0)
throw std::runtime_error(
"Conversion to fft_params impossible for batch_rank != 1 or negative batch sizes");
ret.nbatch = batches[0];
ret.run_callbacks = fft_callback_type_none;
ret.scale_factor = 1.0;
ret.istride = convert_vector_to<decltype(ret.istride)::value_type>(istrides);
ret.ostride = convert_vector_to<decltype(ret.ostride)::value_type>(ostrides);
if(idist[0] < 0 || odist[0] < 0)
throw std::runtime_error("Conversion to fft_params impossible for negative distances");
ret.idist = idist[0];
ret.odist = odist[0];
ret.validate(); // sets itype, otype, isize, osize, etc. from the above
// other ret's members should be irrelevant/fine with default values
return ret;
}

fft_params make_params_for_reference_cpu() const
{
fft_params converted = *this;
return converted.make_params_for_reference_cpu();
}

// public ad-hoc exceptions, specific to hipfftw_helper
struct type_conversion_exception : std::runtime_error
{
Expand Down
213 changes: 110 additions & 103 deletions projects/hipfft/clients/tests/hipfftw_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,6 @@
#include "../hipfftw_helper.h"

#include "../../shared/environment.h"
#include "../../shared/fftw_transform.h"
#include "../../shared/gpubuf.h"
#include "../../shared/hostbuf.h"
#include "../../shared/params_gen.h"
Expand Down Expand Up @@ -2620,20 +2619,16 @@ namespace
{
const hipfftw_functional_validation_params<prec>& params = this->GetParam();

if(!fftw_compare)
GTEST_SKIP() << "skipped due to disabled fftw comparison";

if(!params.can_be_tested())
GTEST_FAIL() << "invalid parameters, cannot be tested";
#ifdef __HIP_PLATFORM_AMD__
if(!params.is_manual_test() && is_known_defect(params.plan_helper))
GTEST_SKIP()
<< "skipped automatically-generated test due to known rocFFT defect";
#endif
if(reference_plan)
GTEST_FAIL()
<< "Starting from an unclean slate (reference plan is not nullptr)";

execution_results_on_host.resize(1);
execution_results_on_host[0].alloc(
params.plan_helper.get_data_byte_size(fft_io::fft_io_out));

std::vector<fft_io> io_range = {fft_io::fft_io_in};
if(params.get_placement() == fft_placement_notinplace)
Expand All @@ -2644,10 +2639,6 @@ namespace
for(auto io : io_range)
{
const size_t data_size = params.plan_helper.get_data_byte_size(io);
auto& io_verification_vec
= io == fft_io::fft_io_in ? verification_input : verification_output;
io_verification_vec.resize(1);
io_verification_vec[0].alloc(data_size);
for(auto step : step_range)
{
const std::pair<hipfftw_step, fft_io> map_key = {step, io};
Expand Down Expand Up @@ -2678,75 +2669,75 @@ namespace
}
}
}
if(verification_input.size() != 1
|| (params.get_placement() != fft_placement_inplace
&& verification_output.size() != 1))
GTEST_FAIL() << "Verification IO buffer incorrectly initialized";
// generate input data
const std::vector<size_t> ioffset = {0};
const std::vector<size_t> field_lower(params.get_rank(), 0);
set_input<hostbuf, hipfftw_real_t<prec>>(verification_input,
fft_input_random_generator_host,
params.get_array_type(fft_io::fft_io_in),
params.get_lengths(),
params.get_ilengths(),
params.get_istride(),
ioffset,
params.get_idist(),
params.get_nbatch(),
get_curr_device_prop(),
field_lower,
0 /* field_lower_batch */,
params.get_contiguous_istrides(),
params.get_contiguous_idist());
// create the reference plan (systematically using the most general guru64 creation)
reference_plan = fftw_trait<hipfftw_real_t<prec>>::make_wrapper(
params.plan_helper.get_reference_plan(verification_input[0].data(),
params.get_placement()
== fft_placement_inplace
? verification_input[0].data()
: verification_output[0].data()));

if(!reference_plan)

execution_results_on_host.resize(1);
const std::pair<hipfftw_step, fft_io> output_key
= {params.use_creation_io_at_execution() ? hipfftw_step::plan_creation
: hipfftw_step::plan_execution,
params.get_placement() == fft_placement_inplace ? fft_io::fft_io_in
: fft_io::fft_io_out};
switch(params.mem_type.at(output_key))
{
GTEST_FAIL() << "could not create a reference plan";
case hipfftw_data_memory_type::device:
// a clean allocation is needed
execution_results_on_host[0].alloc(
params.plan_helper.get_data_byte_size(fft_io::fft_io_out));
break;
#ifndef _WIN32
case hipfftw_data_memory_type::managed:
{
// Actually verify that there is no need to copy data explicitly with
// managed memory.
// NOTE: the raw pointer to a *managed* allocation is encapsulated in
// a non-owned hostbuffer in this case...
execution_results_on_host[0] = hostbuf::make_nonowned(
gpu_io_buffer.at(output_key).data(), gpu_io_buffer.at(output_key).size());
}
break;
#endif
case hipfftw_data_memory_type::pinned_host:
[[fallthrough]];
case hipfftw_data_memory_type::pageable_host:
// no copy needed at all
execution_results_on_host[0] = hostbuf::make_nonowned(
host_io_buffer.at(output_key).data(), host_io_buffer.at(output_key).size());
break;
default:
throw std::runtime_error("Unexpected output memory type");
break;
}

verification_data = std::make_unique<reference_fft_data_t>(
params.plan_helper.make_params_for_reference_cpu());
// generate input data and launch compute
verification_data->initialize_input(fft_input_random_generator_host);
verification_data->launch_async_compute();
}
ROCFFT_CATCH_TEST_EXCEPTIONS;
}
void TearDown() override
{
verification_input.clear();
verification_output.clear();
verification_data.reset();
execution_results_on_host.clear();
host_io_buffer.clear();
gpu_io_buffer.clear();
reference_plan.reset();
this->GetParam().plan_helper.release_plan();
}
// verification buffers (set_input and other common routines require std::vector's of size 1 for these)
std::vector<hostbuf> verification_input;
std::vector<hostbuf> verification_output;
std::vector<hostbuf> execution_results_on_host;
std::unique_ptr<reference_fft_data_t> verification_data;
std::vector<hostbuf> execution_results_on_host;
// possible host buffers (pageable or pinned host allocation)
std::map<std::pair<hipfftw_step, fft_io>, hostbuf> host_io_buffer;
// possible nonhost buffers (may be current/other device or runtime-managed)
std::map<std::pair<hipfftw_step, fft_io>, gpubuf> gpu_io_buffer;
// reference plan
fftw_plan_wrapper_t<hipfftw_real_t<prec>> reference_plan
= fftw_trait<hipfftw_real_t<prec>>::make_wrapper(nullptr);

void functional_test() const
void functional_test()
{
const hipfftw_functional_validation_params<prec>& params = this->GetParam();
try
{
std::ostringstream gtest_info;
if(verification_input.size() != 1
|| (params.get_placement() != fft_placement_inplace
&& verification_output.size() != 1))
GTEST_FAIL() << "The verification I/O buffer(s) were not initialized as "
"needed; host buffer(s) are required";
if(!verification_data || verification_data->needs_computing())
GTEST_FAIL() << "The verification data was not initialized as needed";
if(execution_results_on_host.size() != 1)
GTEST_FAIL() << "Improper test initialization: no host buffer to copy the "
"execution results";
Expand Down Expand Up @@ -2823,25 +2814,35 @@ namespace
= {hipfftw_step::plan_execution, fft_io::fft_io_in};
if(test_io_type.at(exec_in_key) == hipfftw_data_memory_type::device)
{
// an explicit host-to-device copy is needed
const auto hip_status
= hipMemcpyAsync(test_io_ptr.at(exec_in_key),
verification_input[0].data(),
params.plan_helper.get_data_byte_size(fft_io::fft_io_in),
hipMemcpyHostToDevice);
if(hip_status != hipSuccess)
throw hip_runtime_error("hipMemcpyAsync failed.", hip_status);
std::vector<gpubuf> tmp_gpu_buf_vec;
tmp_gpu_buf_vec.emplace_back(gpubuf::make_nonowned(
test_io_ptr.at(exec_in_key),
params.plan_helper.get_data_byte_size(fft_io::fft_io_in)));
verification_data->copy_input_data_in_device_buffers(tmp_gpu_buf_vec,
params.plan_helper);
}
else
{
std::memcpy(test_io_ptr.at(exec_in_key),
verification_input[0].data(),
params.plan_helper.get_data_byte_size(fft_io::fft_io_in));
std::vector<hostbuf> tmp_host_buf_vec;
tmp_host_buf_vec.emplace_back(hostbuf::make_nonowned(
test_io_ptr.at(exec_in_key),
params.plan_helper.get_data_byte_size(fft_io::fft_io_in)));

copy_buffers(verification_data->get_buffers<fft_io::fft_io_in>(),
tmp_host_buf_vec,
verification_data->get_params().ilength(),
verification_data->get_params().nbatch,
verification_data->get_params().precision,
verification_data->get_params().itype,
verification_data->get_params().istride,
verification_data->get_params().idist,
/* destination array type is same as source array type*/
verification_data->get_params().itype,
params.get_istride(),
params.get_idist(),
verification_data->get_params().ioffset,
{0} /* destination offset */);
}

std::shared_future<void> reference_cpu_dft = std::async(std::launch::async, [&]() {
fftw_execute_type<hipfftw_real_t<prec>>(reference_plan);
});
hipfftw_exception_logger exception_logger;

params.plan_helper.create_plan(
Expand Down Expand Up @@ -2874,13 +2875,15 @@ namespace
}
}

// copy hipfftw results back into the execution_results_on_host[0] buffer
// for verification purposes
const std::pair<hipfftw_step, fft_io> exec_out_key
= {hipfftw_step::plan_execution, fft_io::fft_io_out};
if(test_io_type.at(exec_out_key) == hipfftw_data_memory_type::device)
switch(test_io_type.at(exec_out_key))
{
case hipfftw_data_memory_type::device:
{
// making this copy synchronous as the next step is verifying the results
// copy hipfftw results back into the execution_results_on_host[0] buffer
// for verification purposes. Making this copy synchronous as the next step
// is verifying the results
const auto hip_status
= hipMemcpy(execution_results_on_host[0].data(),
test_io_ptr.at(exec_out_key),
Expand All @@ -2890,47 +2893,51 @@ namespace
throw hip_runtime_error("hipMemcpy failed (copying output results).",
hip_status);
}
else
break;
#ifndef _WIN32
case hipfftw_data_memory_type::managed:
[[fallthrough]];
#endif
case hipfftw_data_memory_type::pageable_host:
[[fallthrough]];
case hipfftw_data_memory_type::pinned_host:
{
std::memcpy(execution_results_on_host[0].data(),
test_io_ptr.at(exec_out_key),
params.plan_helper.get_data_byte_size(fft_io::fft_io_out));
// no need to actually copy at all, check that execution_results_on_host
// was correctly set in SetUp
if(execution_results_on_host[0].data() != test_io_ptr.at(exec_out_key))
throw std::logic_error("Logic error detected (definition of "
"execution_results_on_host with "
"host-accessible output data)");
}
break;
default:
throw std::runtime_error("Unexpected output memory type");
break;
}
// compare results
if(reference_cpu_dft.valid())
reference_cpu_dft.get();
const auto test_lengths = params.get_lengths();
const auto total_length = product(test_lengths.begin(), test_lengths.end());
const auto& reference_output = params.get_placement() == fft_placement_inplace
? verification_input
: verification_output;

const auto ref_norm = norm(reference_output,
params.get_olengths(),
params.get_nbatch(),
prec,
params.get_array_type(fft_io::fft_io_out),
params.get_ostride(),
params.get_odist(),
{0} /* offset */);
const auto test_lengths = params.get_lengths();
const auto total_length = product(test_lengths.begin(), test_lengths.end());

const auto ref_norm
= verification_data->get_norm<fft_io::fft_io_out>(params.get_nbatch()).get();
const double test_epsilon
= prec == fft_precision_single ? single_epsilon : double_epsilon;
const double linf_cutoff = test_epsilon * ref_norm.l_inf * log(total_length);
// compare results
const auto diff = distance(reference_output,
const auto diff = distance(verification_data->get_buffers<fft_io::fft_io_out>(),
execution_results_on_host,
params.get_olengths(),
params.get_nbatch(),
prec,
params.get_array_type(fft_io::fft_io_out),
params.get_ostride(),
params.get_odist(),
verification_data->get_params().otype,
verification_data->get_params().ostride,
verification_data->get_params().odist,
params.get_array_type(fft_io::fft_io_out),
params.get_ostride(),
params.get_odist(),
nullptr,
linf_cutoff,
{0} /* offset */,
verification_data->get_params().ooffset,
{0} /* offset */);
EXPECT_LE(diff.l_inf, linf_cutoff);
EXPECT_LE(diff.l_2 / ref_norm.l_2, sqrt(log2(total_length)) * test_epsilon);
Expand Down
Loading