Skip to content

Prefer IVF-PQ compute_similarity blockDim=512 on AMD (gfx1100) #7

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

Open
konkolchin wants to merge 1 commit into
base: release/rocmds-25.10
Choose a base branch
from
Open

Prefer IVF-PQ compute_similarity blockDim=512 on AMD (gfx1100) #7

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
35 changes: 35 additions & 0 deletions cpp/src/neighbors/ivf_pq/ivf_pq_compute_similarity_impl.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,8 @@
*/
#pragma once

#include <cstdlib> // getenv (AMD IVF-PQ blockDim prefer)

#include "../ivf_common.cuh" // dummy_block_sort_t
#include "../sample_filter.cuh" // none_sample_filter
#include <cuvs/distance/distance.hpp> // cuvs::distance::DistanceType
Expand Down Expand Up @@ -890,6 +892,9 @@ auto compute_similarity_select(const cudaDeviceProp& dev_props,
continue;
}

// Cap before auto-shrink (used by AMD prefer-512 below).
const uint32_t n_threads_cap = n_threads;

{
// Try to reduce the number of threads to increase occupancy and data locality
auto n_threads_tmp = n_threads_min;
Expand Down Expand Up @@ -933,6 +938,36 @@ auto compute_similarity_select(const cudaDeviceProp& dev_props,
}
}

#if defined(__HIP_PLATFORM_AMD__)
{
// DXC gfx1100 PoC: prefer blockDim=512 for IVF-PQ compute_similarity.
// Measured +28–36% search QPS vs stock auto-shrink on RX 7900 XTX
// (ann-harness results/lib_bench/LAUNCH_KNOBS.md, 2026-08-14); recall flat.
// Env HIPVS_IVF_PQ_BLOCK_THREADS (patch 0005) still overrides when set.
const char* env_bt = std::getenv("HIPVS_IVF_PQ_BLOCK_THREADS");
const bool env_overrides = (env_bt != nullptr && env_bt[0] != '\0');
constexpr uint32_t kAmdIvfPqPreferThreads = 512u;
if (!env_overrides && kAmdIvfPqPreferThreads >= n_threads_gty &&
kAmdIvfPqPreferThreads <= n_threads_cap) {
auto smem_pref = smem_size_f(kAmdIvfPqPreferThreads);
cudaError_t st_pref =
cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_pref);
if (st_pref == cudaSuccess) {
occupancy_t<OutT, LutT, IvfSampleFilterT> pref(
smem_pref, kAmdIvfPqPreferThreads, kernel, dev_props);
if (pref.blocks_per_sm > 0) {
n_threads = kAmdIvfPqPreferThreads;
smem_size = smem_pref;
cur = pref;
}
} else {
RAFT_EXPECTS(st_pref == cudaGetLastError(),
"Tried to reset the expected cuda error code, but it didn't match the expectation");
}
}
}
#endif

@kevinjosephamd kevinjosephamd Aug 24, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for this contribution and for sharing the performance results, those improvements look promising. Would you mind removing the HIPVS_IVF_PQ_BLOCK_THREADS environment-variable handling? We’d prefer not to expose an environment variable for configuring this launch parameter. Instead, the preferred block size should be derived from the device’s reported wavefront size.
Could you also update the comments to describe the device-derived policy and remove the reference to the “patch 0005” override? Thanks again for working on this optimization.


{
if (selected_perf.occupancy <= 0.0 // no candidate yet
|| (selected_perf.occupancy < cur.occupancy * kTargetOccupancy &&
Expand Down