Skip to content

test(rocfft): stable, position-aware hash for bitwise repro buffers #12752

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
Abuudiii merged 7 commits into from
Oct 12, 2026
Merged

test(rocfft): stable, position-aware hash for bitwise repro buffers #12752

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
377f19d
test(rocfft): use a stable hash for bitwise repro buffer hashing (AIF…
Abuudiii Sep 29, 2026
a3b9e61
comment cleanup
Abuudiii Oct 1, 2026
bf9581d
style(rocfft): match CI clang-format in fft_hash.h
Abuudiii Oct 1, 2026
fc47fa0
docs(rocfft): describe the splitmix64 bitwise repro hash in testing.md
Abuudiii Oct 5, 2026
3e957fb
test(rocfft): hash raw bits with logical position in bitwise repro hash
Abuudiii Oct 7, 2026
ec7d018
test(rocfft): use PROB_SKIP_UNITTEST in new buffer hash tests
Abuudiii Oct 7, 2026
b531323
Merge branch 'develop' into users/Abdullah/rocfft-bitwise-repro-stabl…
Abuudiii Oct 9, 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
130 changes: 63 additions & 67 deletions projects/hipfft/shared/fft_hash.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,8 @@
#include "../../../shared/rocfft_complex.h"

#include <algorithm>
#include <bit>
#include <cstdint>
#include <iostream>
#include <memory>
#include <string>
Expand Down Expand Up @@ -82,59 +84,59 @@ struct hash_output
Tint buffer_imag;
};

static inline double get_weight(const size_t counter, const size_t max_counter)
// Fixed, platform-independent 64-bit mix (splitmix64 finalizer).
// Unlike std::hash, the result is defined here and does not depend on compiler,
// standard library, or process.
static inline uint64_t stable_hash(uint64_t z)
{
return (static_cast<double>(counter) / static_cast<double>(max_counter));
z = (z ^ (z >> 30)) * 0xbf58476d1ce4e5b9ULL;
z = (z ^ (z >> 27)) * 0x94d049bb133111ebULL;
return z ^ (z >> 31);
}

template <typename Tint>
static inline void hash_value(Tint& hash_value,
const size_t counter,
const size_t max_counter,
const rocfft_fp16& input_value)
// Raw IEEE-754 bit pattern of a value, widened to 64 bits. Hashing the bits
// (not a rounded product) makes every bit of the value count.
static inline uint64_t raw_bits(rocfft_fp16 v)
{
auto weight = get_weight(counter, max_counter);
hash_value += std::hash<float>{}(weight * input_value);
return std::bit_cast<uint16_t>(v);
}
template <typename Tint>
static inline void hash_value(Tint& hash_value,
const size_t counter,
const size_t max_counter,
const float& input_value)
static inline uint64_t raw_bits(float v)
{
auto weight = get_weight(counter, max_counter);
hash_value += std::hash<float>{}(weight * input_value);
return std::bit_cast<uint32_t>(v);
}
template <typename Tint>
static inline void hash_value(Tint& hash_value,
const size_t counter,
const size_t max_counter,
const double& input_value)
static inline uint64_t raw_bits(double v)
{
auto weight = get_weight(counter, max_counter);
hash_value = std::hash<double>{}(weight * input_value) + hash_value;
return std::bit_cast<uint64_t>(v);
}

template <typename Tint>
static inline size_t get_max_counter(const Tint& whole_length, const size_t nbatch)
// Row-major index of an element within one batch. Uses the logical lengths,
// not the memory strides, so padding and layout do not change the hash.
template <typename T1>
static inline size_t logical_index(const T1& index, [[maybe_unused]] const T1& length)
{
return static_cast<size_t>(count_iters(whole_length) * nbatch);
return index;
}

template <typename T1>
static inline T1 get_unit_value(const T1& val)
static inline size_t logical_index(const std::tuple<T1, T1>& index,
const std::tuple<T1, T1>& length)
{
return static_cast<T1>(1);
return std::get<0>(index) * std::get<1>(length) + std::get<1>(index);
}
template <typename T1>
static inline std::tuple<T1, T1> get_unit_value(const std::tuple<T1, T1>& val)
static inline size_t logical_index(const std::tuple<T1, T1, T1>& index,
const std::tuple<T1, T1, T1>& length)
{
return std::make_tuple(static_cast<T1>(1), static_cast<T1>(1));
return (std::get<0>(index) * std::get<1>(length) + std::get<1>(index)) * std::get<2>(length)
+ std::get<2>(index);
}
template <typename T1>
static inline std::tuple<T1, T1, T1> get_unit_value(const std::tuple<T1, T1, T1>& val)

// Mix a value's bits with its global logical position. Per-element results
// are summed, so the total does not depend on how threads split the buffer,
// while swapped elements (transpose, batch swap) still change it.
template <typename Tint, typename Tfloat>
static inline void hash_value(Tint& hash, const size_t position, const Tfloat& value)
{
return std::make_tuple(static_cast<T1>(1), static_cast<T1>(1), static_cast<T1>(1));
hash += stable_hash(raw_bits(value) ^ stable_hash(static_cast<uint64_t>(position)));
}

template <typename T1>
Expand All @@ -153,29 +155,27 @@ static inline void compute_real_buffer_hash(const std::vector<hostbuf>& ibuffer,
Tint2& hash_real,
Tint2& hash_imag)
{
auto unit_stride = get_unit_value(whole_stride);

size_t max_counter = get_max_counter<Tint1>(whole_length, nbatch);

auto idata = (Tfloat*)ibuffer[0].data();
size_t i_base = 0;
auto partitions = partition_rowmajor(whole_length);
auto idata = (Tfloat*)ibuffer[0].data();
size_t i_base = 0;
auto partitions = partition_rowmajor(whole_length);
const size_t batch_elems = count_iters(whole_length);

std::vector<Tint2> partition_hash_real(partitions.size(), 0);

for(unsigned int b = 0; b < nbatch; b++, i_base += idist)
{
const size_t batch_base = b * batch_elems;
#pragma omp parallel for num_threads(partitions.size())
for(size_t part = 0; part < partitions.size(); ++part)
{
auto index = partitions[part].first;
const auto length = partitions[part].second;
do
{
const auto i = compute_index(index, whole_stride, i_base);
const auto counter = compute_index(index, unit_stride, i_base) + 1;
const auto i = compute_index(index, whole_stride, i_base);
const size_t position = batch_base + logical_index(index, whole_length);

hash_value<Tint2>(partition_hash_real[part], counter, max_counter, idata[i]);
hash_value<Tint2>(partition_hash_real[part], position, idata[i]);
} while(increment_rowmajor(index, length));
}
}
Expand All @@ -193,32 +193,30 @@ static inline void compute_planar_buffer_hash(const std::vector<hostbuf>& ibuffe
Tint2& hash_real,
Tint2& hash_imag)
{
auto unit_stride = get_unit_value(whole_stride);

size_t max_counter = get_max_counter<Tint1>(whole_length, nbatch);

auto ireal = (Tfloat*)ibuffer[0].data();
auto iimag = (Tfloat*)ibuffer[1].data();
size_t i_base = 0;
auto partitions = partition_rowmajor(whole_length);
auto ireal = (Tfloat*)ibuffer[0].data();
auto iimag = (Tfloat*)ibuffer[1].data();
size_t i_base = 0;
auto partitions = partition_rowmajor(whole_length);
const size_t batch_elems = count_iters(whole_length);

std::vector<Tint2> partition_hash_real(partitions.size(), 0);
std::vector<Tint2> partition_hash_imag(partitions.size(), 0);

for(unsigned int b = 0; b < nbatch; b++, i_base += idist)
{
const size_t batch_base = b * batch_elems;
#pragma omp parallel for num_threads(partitions.size())
for(size_t part = 0; part < partitions.size(); ++part)
{
auto index = partitions[part].first;
const auto length = partitions[part].second;
do
{
const auto i = compute_index(index, whole_stride, i_base);
const auto counter = compute_index(index, unit_stride, i_base) + 1;
const auto i = compute_index(index, whole_stride, i_base);
const size_t position = batch_base + logical_index(index, whole_length);

hash_value<Tint2>(partition_hash_real[part], counter, max_counter, ireal[i]);
hash_value<Tint2>(partition_hash_imag[part], counter, max_counter, iimag[i]);
hash_value<Tint2>(partition_hash_real[part], position, ireal[i]);
hash_value<Tint2>(partition_hash_imag[part], position, iimag[i]);
} while(increment_rowmajor(index, length));
}
}
Expand All @@ -236,31 +234,29 @@ static inline void compute_interleaved_buffer_hash(const std::vector<hostbuf>& i
Tint2& hash_real,
Tint2& hash_imag)
{
auto unit_stride = get_unit_value(whole_stride);

size_t max_counter = get_max_counter<Tint1>(whole_length, nbatch);

auto idata = (rocfft_complex<Tfloat>*)ibuffer[0].data();
size_t i_base = 0;
auto partitions = partition_rowmajor(whole_length);
auto idata = (rocfft_complex<Tfloat>*)ibuffer[0].data();
size_t i_base = 0;
auto partitions = partition_rowmajor(whole_length);
const size_t batch_elems = count_iters(whole_length);

std::vector<Tint2> partition_hash_real(partitions.size(), 0);
std::vector<Tint2> partition_hash_imag(partitions.size(), 0);

for(unsigned int b = 0; b < nbatch; b++, i_base += idist)
{
const size_t batch_base = b * batch_elems;
#pragma omp parallel for num_threads(partitions.size())
for(size_t part = 0; part < partitions.size(); ++part)
{
auto index = partitions[part].first;
const auto length = partitions[part].second;
do
{
const auto i = compute_index(index, whole_stride, i_base);
const auto counter = compute_index(index, unit_stride, i_base) + 1;
const auto i = compute_index(index, whole_stride, i_base);
const size_t position = batch_base + logical_index(index, whole_length);

hash_value<Tint2>(partition_hash_real[part], counter, max_counter, idata[i].real());
hash_value<Tint2>(partition_hash_imag[part], counter, max_counter, idata[i].imag());
hash_value<Tint2>(partition_hash_real[part], position, idata[i].real());
hash_value<Tint2>(partition_hash_imag[part], position, idata[i].imag());
} while(increment_rowmajor(index, length));
}
}
Expand Down
133 changes: 133 additions & 0 deletions projects/rocfft/clients/tests/buffer_hash_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,7 @@
#include "../../shared/params_gen.h"
#include "../../shared/rocfft_params.h"
#include <algorithm>
#include <cmath>
#include <gtest/gtest.h>
#include <memory>
#include <random>
Expand Down Expand Up @@ -431,3 +432,135 @@ TEST(rocfft_UnitTest, buffer_hashing_double)
GTEST_SKIP() << e.what();
}
}

// Hash a host buffer laid out exactly as given (strides/dist describe memory).
template <typename Telem>
static hash_output<size_t> hash_host_buffer(rocfft_precision precision,
rocfft_array_type type,
const std::vector<Telem>& mem,
const std::vector<size_t>& length,
const std::vector<size_t>& stride,
size_t dist,
size_t nbatch)
{
std::vector<hostbuf> buf(1);
buf[0].alloc(mem.size() * sizeof(Telem));
std::copy(mem.begin(), mem.end(), static_cast<Telem*>(buf[0].data()));
auto hash_in = hash_input(precision, length, stride, dist, type, nbatch);
auto hash_out = hash_output<size_t>();
compute_hash(buf, hash_in, hash_out);
return hash_out;
}

template <typename Tfloat>
static hash_output<size_t> hash_real_1234(rocfft_precision precision)
{
std::vector<Tfloat> mem;
for(size_t i = 0; i < 4; ++i)
mem.push_back(static_cast<Tfloat>(i + 1)); // {1, 2, 3, 4}
return hash_host_buffer(precision, rocfft_array_type_real, mem, {4}, {1}, 4, 1);
}

// The buffer hash must be a fixed function of the data (no std::hash), so a
// known buffer always produces the same fingerprint on every toolchain.
TEST(rocfft_UnitTest, buffer_hash_golden_values)
{
PROB_SKIP_UNITTEST();
auto h_half = hash_real_1234<rocfft_fp16>(rocfft_precision_half);
Comment thread
Abuudiii marked this conversation as resolved.
Resolved
EXPECT_EQ(h_half.buffer_real, 6010172260665940913ULL);
EXPECT_EQ(h_half.buffer_imag, 0u);

auto h_single = hash_real_1234<float>(rocfft_precision_single);
EXPECT_EQ(h_single.buffer_real, 13029437891923012638ULL);
EXPECT_EQ(h_single.buffer_imag, 0u);

auto h_double = hash_real_1234<double>(rocfft_precision_double);
EXPECT_EQ(h_double.buffer_real, 4918654574207189942ULL);
EXPECT_EQ(h_double.buffer_imag, 0u);
}

template <typename Tfloat>
static void expect_real_transpose_detected(rocfft_precision precision)
{
std::vector<Tfloat> a(16), t(16);
for(size_t r = 0; r < 4; ++r)
for(size_t c = 0; c < 4; ++c)
a[r * 4 + c] = static_cast<Tfloat>(r * 4 + c + 1);
for(size_t r = 0; r < 4; ++r)
for(size_t c = 0; c < 4; ++c)
t[r * 4 + c] = a[c * 4 + r];
auto ha = hash_host_buffer(precision, rocfft_array_type_real, a, {4, 4}, {4, 1}, 16, 1);
auto ht = hash_host_buffer(precision, rocfft_array_type_real, t, {4, 4}, {4, 1}, 16, 1);
EXPECT_NE(ha.buffer_real, ht.buffer_real);
}

// Elements on an anti-diagonal must not share a position, so a transposed
// square output hashes differently.
TEST(rocfft_UnitTest, buffer_hash_detects_transpose)
{
PROB_SKIP_UNITTEST();
expect_real_transpose_detected<float>(rocfft_precision_single);
expect_real_transpose_detected<double>(rocfft_precision_double);

std::vector<rocfft_complex<float>> a(16), t(16);
for(size_t r = 0; r < 4; ++r)
for(size_t c = 0; c < 4; ++c)
{
const float v = static_cast<float>(r * 4 + c + 1);
a[r * 4 + c] = rocfft_complex<float>(v, -v);
}
for(size_t r = 0; r < 4; ++r)
for(size_t c = 0; c < 4; ++c)
t[r * 4 + c] = a[c * 4 + r];
const auto type = rocfft_array_type_complex_interleaved;
auto ha = hash_host_buffer(rocfft_precision_single, type, a, {4, 4}, {4, 1}, 16, 1);
auto ht = hash_host_buffer(rocfft_precision_single, type, t, {4, 4}, {4, 1}, 16, 1);
EXPECT_NE(ha.buffer_real, ht.buffer_real);
EXPECT_NE(ha.buffer_imag, ht.buffer_imag);
}

// A 1-ULP change in any element must change the hash.
TEST(rocfft_UnitTest, buffer_hash_detects_one_ulp)
{
PROB_SKIP_UNITTEST();
std::vector<float> a = {1.0f, 2.0f, 1.7f, 4.0f};
std::vector<float> b = a;
b[2] = std::nextafter(1.7f, 2.0f);
auto ha = hash_host_buffer(rocfft_precision_single, rocfft_array_type_real, a, {4}, {1}, 4, 1);
auto hb = hash_host_buffer(rocfft_precision_single, rocfft_array_type_real, b, {4}, {1}, 4, 1);
EXPECT_NE(ha.buffer_real, hb.buffer_real);
}

// Swapping two batches must change the hash.
TEST(rocfft_UnitTest, buffer_hash_detects_batch_swap)
{
PROB_SKIP_UNITTEST();
std::vector<float> a = {1, 2, 3, 4, 5, 6, 7, 8};
std::vector<float> b = {5, 6, 7, 8, 1, 2, 3, 4};
auto ha = hash_host_buffer(rocfft_precision_single, rocfft_array_type_real, a, {4}, {1}, 4, 2);
auto hb = hash_host_buffer(rocfft_precision_single, rocfft_array_type_real, b, {4}, {1}, 4, 2);
EXPECT_NE(ha.buffer_real, hb.buffer_real);
}

// The hash depends on logical data only: strides, dist and padding contents
// must not change it.
TEST(rocfft_UnitTest, buffer_hash_ignores_padding)
{
PROB_SKIP_UNITTEST();
const size_t nbatch = 2, rows = 3, cols = 4;
std::vector<float> contig(nbatch * rows * cols);
std::vector<float> padded(nbatch * 17, -999.0f);
for(size_t b = 0; b < nbatch; ++b)
for(size_t r = 0; r < rows; ++r)
for(size_t c = 0; c < cols; ++c)
{
const float v = static_cast<float>(b * 100 + r * 10 + c);
contig[b * 12 + r * 4 + c] = v;
padded[b * 17 + r * 5 + c] = v;
}
const auto prec = rocfft_precision_single;
const auto type = rocfft_array_type_real;
auto hc = hash_host_buffer(prec, type, contig, {rows, cols}, {4, 1}, 12, nbatch);
auto hp = hash_host_buffer(prec, type, padded, {rows, cols}, {5, 1}, 17, nbatch);
EXPECT_EQ(hc.buffer_real, hp.buffer_real);
}
Loading
Loading