Support pre-Volta (Maxwell/Pascal) GPU architectures - #2529
Conversation
Part 1 of 2. On its own this commit changes nothing observable: it is a
prerequisite for building cuVS with a CUDA toolkit older than 12.8, and that
becomes possible only in combination with
rapidsai/librtcx#17 - "Support CUDA toolkits older than 12.8 via the
driver API"
Neither is useful without the other, and either may merge first. The split
follows the code: since NVIDIA#2311 the JIT-LTO plumbing lives in librtcx, which loads
the linked image (cudaLibraryLoadData / cudaLibraryGetKernel /
cudaLibraryUnload) and launches it (cudaLaunchKernelExC); cuVS keeps only what
it does with the resulting handle. Both halves hit the same wall -- the runtime
library-management API arrived in CUDA 12.8 -- but they are in different
repositories, so they are fixed separately. Without librtcx#17 every cuVS
translation unit that includes the rtcx launcher still fails with
rtcx/algorithm_launcher.hpp:24: error: identifier "cudaLibrary_t" is undefined
and this commit does not, and cannot, change that.
The cuVS half. cuVS queries and sets attributes on the cudaKernel_t handles that
rtcx hands out for run-time-linked kernels. Both things it does with those
handles are CUDA 12.8 features:
* Passing a cudaKernel_t where a "const void* func" entry point is expected.
The runtime documents this from 12.8 onwards -- "If the specified function
does not exist, then it is assumed to be a cudaKernel_t and used as is",
which appears in the 12.8 documentation of cudaFuncGetAttributes,
cudaFuncSetAttribute and cudaOccupancyMaxActiveBlocksPerMultiprocessor and in
no earlier version. On an older toolkit the handle is not a valid entry
address and the call fails with cudaErrorInvalidDeviceFunction.
* cudaKernelSetAttributeForDevice, which does not exist before 12.8 in either
the headers or libcudart.
The driver API has always taken the equivalent CUkernel / CUfunction, and
cudaKernel_t is a typedef for struct CUkern_st*, i.e. exactly CUkernel, so the
handles carry over unchanged. Add src/util/jit_kernel_compat.hpp, which forwards
to the runtime API on 12.8+ and to cuKernelSetAttribute, cuFuncGetAttribute and
cuOccupancyMaxActiveBlocksPerMultiprocessor otherwise, and route the six call
sites through it. On 12.8+ every wrapper compiles down to the original runtime
call and CUDA::cuda_driver is not linked, so nothing changes for the toolkits
cuVS supports today.
Two details worth recording:
* The runtime's primary context is forced to exist before any driver call,
because cuKernelGetFunction resolves against the current context. Skipping
this yields CUDA_ERROR_INVALID_CONTEXT from an otherwise correct sequence.
* Driver statuses are translated to the nearest runtime status so that existing
RAFT_CUDA_TRY diagnostics stay meaningful.
In launchConfigGenerator the choice between the two occupancy APIs is an
if constexpr on the argument type, because that helper is also called with
genuine __global__ function pointers, for which the runtime API is correct.
The IVF-PQ kernel selector needed one further change. It asserted that a failed
cudaFuncSetAttribute was also observable through cudaGetLastError(), which only
holds for the runtime API; the driver path never sets the runtime's sticky
error. The intent -- skip a kernel candidate that cannot get the shared memory
it wants -- is preserved, and the sticky error is cleared explicitly.
Verified with CUDA 12.6.85: with librtcx#17 applied (-DCPM_rtcx_SOURCE=...),
libcuvs.so builds and links for sm_50, and CAGRA build and search run correctly
on a Maxwell device. Without it, the same tree fails only inside the rtcx
headers, which is what "part 1 of 2" means in practice.
Note that dependencies.yaml still lists CUDA 12.2 and 12.5 in its cuda_version
matrix, so the toolkit range the project declares and the one it can actually
build have disagreed since NVIDIA#1405.
cuVS could not be compiled for anything older than sm_70. Building with --gpu-arch="50;61" failed in <cuda/semaphore>, and several further blockers were hiding behind that first error. Introduce a CMake-generated cuvs/detail/arch_config.hpp that exposes the architecture range of the build (CUVS_MIN_CUDA_ARCH, CUVS_CUTLASS_ENABLED, ...). Unlike __CUDA_ARCH__ it is visible in the host pass too, so it can gate includes and template instantiations consistently; unlike a global -D it does not invalidate every object file when the architecture list changes. Compile-time blockers addressed: * CUTLASS. cuda::binary_semaphore hard-errors below sm_70 and the reduced-vec output iterator needs cg::binary_partition. The CUTLASS kernels are only ever dispatched to on sm_80+, so below sm_70 the whole path is compiled out and the pre-existing SIMT/JIT fallbacks are used. This also cuts build time. * __match_any_sync (sm_70) in the sparse SpMV kernel. Replaced with a helper in util/arch_compat.cuh that uses the intrinsic where available and otherwise a ballot-based emulation whose loop trip count is warp-uniform, which is required on pre-Volta hardware. * atomicAdd_block (sm_60). Falls back to the device-scoped atomic. * atomicAdd(double*, bool) in bin_distance. Below sm_60 the native double overload does not exist, leaving only RAFT's atomicAdd(T*, T) template, whose deduction then fails. Fixed with an explicit cast. * cuco's block-scoped atomics (atom.cas.cta.*, sm_60) in the sparse hash SpMV strategy, and the system-scoped atomics dynamic batching uses for its host/device handshake. Neither has a meaningful pre-Pascal emulation, so the device code is compiled out and both paths now fail with an explicit error instead of silently misbehaving. The dense_smem SpMV strategy is unaffected. * __half/__half2 arithmetic (sm_53) in the CAGRA VPQ distance, the faiss-select comparators and the l_inf distance op. Added util/half_compat.cuh; pre-Pascal targets widen to fp32, which is also a small accuracy improvement. * atomicAdd(double*, double) and the atomicAdd_block family are only declared for __CUDA_ARCH__ >= 600, so in an older device pass they do not exist and every header that calls them unqualified fails to compile -- RAFT's strided_reduction.cuh (instantiated with double by cuVS) and CUB's agent_histogram.cuh among them. Added util/atomic_compat.cuh, which supplies atomicAdd(double*) as a 64-bit CAS loop comparing bit patterns rather than values, so a stored NaN cannot make it spin forever, plus atomicAdd_block overloads forwarding to the device-scoped atomics. It expands to nothing on sm_60+ and in the host pass. The calls live in templates but their argument types are fundamental, so ADL contributes nothing and only declarations visible at the definition context are considered: the header must be included before any RAFT/CUB header, which is recorded in a comment at each of the four translation units that need it. * raft::sqrt(__half) lowers to hsqrt, an FP16 ALU instruction requiring sm_53, and its RAFT overload static_asserts below that. Added cuvs::util::sqrt_op, which widens __half to fp32 on such targets and forwards to raft::sqrt otherwise, and used it in the two element-wise passes over possibly-fp16 data. Fixes for RAFT, CUB and cuco headers all live in cuVS: those are CPM-fetched, so patching them in place would be invisible to a reviewer, lost on a version bump and untestable in CI. cuVS instead makes the missing declarations visible before the dependency header is parsed. Runtime blockers addressed: * JIT-LTO fragments were pinned to a 70-real baseline regardless of CMAKE_CUDA_ARCHITECTURES. LTO-IR is forward but not backward compatible, so the baseline now tracks the oldest architecture being built for; otherwise nvJitLink cannot link CAGRA/IVF search kernels for an older device. * NN-descent was disabled entirely below sm_70, so CAGRA had no NN-descent graph build. Only the WMMA kernel needs tensor cores; the SIMT kernel uses nothing newer than sm_50 and its static shared memory stays under the 48 KiB per-block limit for every supported dtype. Lower its floor to sm_50 and fall back to it on pre-Volta devices instead of throwing.
Two IVF-PQ search settings assume hardware that pre-Volta and pre-Pascal GPUs do not have. Both are performance knobs whose fp32 equivalent computes the same quantity, so widen rather than fail, and say so once. * coarse_search_dtype selects the element type of the cluster-probing GEMM. cuBLAS has no fp16 GEMM below sm_53 and no int8 GEMM below sm_61 (that needs dp4a), and cublasLtMatmulAlgoGetHeuristic reports CUBLAS_STATUS_NOT_SUPPORTED rather than picking a different kernel. This is not hypothetical for anyone building CAGRA: cagra::graph_build_params::ivf_pq_params sets the dtype to CUDA_R_16F, so the default IVF-PQ graph build fails on such a device. Widen to CUDA_R_32F, which is the same computation at higher precision, and log the substitution once. Everything downstream reads the adjusted parameters, so there is a single source of truth for the dtype. * cudaFuncAttributePreferredSharedMemoryCarveout is only meaningful from Volta onwards; older hardware has a fixed shared-memory/L1 partition. Skip the call there. It is a hint the driver is free to ignore in any case, so no other behaviour changes. Devices from sm_61 on are unaffected: the dtype is returned unchanged and the carveout hint is still set.
CAGRA fails to compile for sm_50 or sm_60: cuda/__barrier/barrier_block_scope.h(171): error: the global scope has no "__match_any_sync" __match_any_sync is an sm_70 instruction, and CUDA only declares it for __CUDA_ARCH__ >= 700. libcu++ calls it from the body of the block-scoped barrier's __arrive_sm70. That body is only executed under an NV_IF_TARGET(NV_PROVIDES_SM_70, ...), but it is still parsed for every device pass, so an older target fails even though it would never run the code. cuVS reaches it through cuco::bloom_filter -> cuda/annotated_ptr -> cuda/__memcpy_async -> cuda/__barrier, i.e. from the CAGRA sample filter. Add util/warp_intrinsic_compat.cuh, which defines __match_any_sync at global scope for device passes below sm_70, and include it before <cuco/bloom_filter.cuh> in sample_filter_data.cuh, which is the single point through which the offending header is reached. As with atomic_compat.cuh, the fix lives in cuVS rather than in the dependency: CCCL and cuco are CPM-fetched, so patching them there would be invisible to a reviewer, lost on a version bump and untestable in CI. The definition is the warp-uniform ballot emulation already used by arch_compat::match_any_sync, so it is correct if it is ever called, which on a pre-Volta device it should not be. It expands to nothing on sm_70+ and in the host pass. The proper home for this fix is libcu++, which should guard the body of __arrive_sm70 the way it guards its call sites; this header can be dropped once it does.
|
Hi @drzraf, thanks for the contribution, but we very intentionally support Ampere architectures and above. It's very challenging to maintain support beyond that, and for cuda versions prior to the last 4 minor versions. Can you explain why this change is needed? Is this to enable an integration into an existing product? |
|
They are hundred of thousand pre-Volta GPU in service as we write. Some of their users already rely on vector search and more are to come as semantic (local) search is only going to increase. pgvector today, maybe sqlite-vec tomorrow (and all the semantically-enabled agent/MCP/browsers/local search-engines/desktop/addons/... popping up on top of them) Cagra is efficient and provides very noticeable gain (×2 to ×4 on a sm_50 wrt a CPU of its generation). So it's basically about maximizing the usage and the efficiency of existing hardware for use-cases that will only gain traction in the coming months/years. On a personal note, I (like most people) absolutely despise planned obsolescence (when a perfectly working device is made useless or locked-down due to arbitrary limitations). I consider (free) software to be the very solution against this waste (energy/hardware/headaches/frustrations/...) The overall patchset is ~ 1200 lines (#2527 + #2528 + this + rapidsai/librtcx#17) while providing support to hundreds of thousands of devices, across many versions of CUDA 12.x for a hardware-optimized algorithm whose usage will likely increase in the nearby future, bringing 200-400% performance improvement. So when it comes to the trade-off between these additional LoC vs hardware support I'm absolutely convinced it's worth it. (The Linux kernel demonstrates every day how far hardware support can go while keeping a maintainable codebase, and... doing so on a budget if a comparison was to be done) I tried to get (... the LLM to get) the patch-set as logical, self-contained and understandable as possible and these ~1200 lines will not become 2000 or 10000 later on nor do they refrain future evolution of the cuvs library nor do they settle a negative precedent AFAICT. Only ~350 of these LoC actually touch existing code. The rest are simple wrappers, inert at compile-time for sm_70+ cases: zero overhead for existing hardware (unless I've overlooked something) Once the one-time review effort has been paid by someone (close to the codebase), the new LoC could just be forgotten. (If the reviewer(s) are among these persons who experience good conscience feeling when they avoid waste, that would be even better. Merge/release will pay for itself) Taken from the opposite side, aka "rejecting support of two generations of (less than 10 years old) Nvidia GPUs, because it would cost only 1200 LoC" (aka 1.5% of cuvs 80k LoC codebase) would likely (and rightfully?) be considered a bad decision by most and a very unfortunate signal even for sm_70+ owners (not to say CUDA 13.x users stuck to a given minor version due to [notoriously debatable] Nvida CUDA hardware support^Wdeprecation policies) Psychosocioemotional aspects aside, if you can suggest any strategy/improvement to further reduce the impact on the codebase, so that these 320 "intrusive" LoC could be 200 or even 100, I'd would glad to follow your suggestions : older hardware support is a duty but making it as discreet as possible is another legitimate one. |
2 participants
Support pre-Volta (Maxwell/Pascal) GPU architectures
Context: rapidsai/librtcx#17
Includes/applies on top of #2527
To consider: #2528
cuVS could not be compiled for anything older than sm_70. Building with
--gpu-arch="50;61"failed in<cuda/semaphore>, and several further blockers were hiding behind that first error.Introduce a CMake-generated
cuvs/detail/arch_config.hppthat exposes the architecture range of the build (CUVS_MIN_CUDA_ARCH,CUVS_CUTLASS_ENABLED, ...). Unlike__CUDA_ARCH__it is visible in the host pass too, so it can gate includes and template instantiations consistently; unlike a global -D it does not invalidate every object file when the architecture list changes.Compile-time blockers addressed:
CUTLASS.cuda::binary_semaphorehard-errors below sm_70 and the reduced-vec output iterator needscg::binary_partition. TheCUTLASSkernels are only ever dispatched to on sm_80+, so below sm_70 the whole path is compiled out and the pre-existing SIMT/JIT fallbacks are used. This also cuts build time.__match_any_sync(sm_70) in the sparse SpMV kernel. Replaced with a helper inutil/arch_compat.cuhthat uses the intrinsic where available and otherwise a ballot-based emulation whose loop trip count is warp-uniform, which is required on pre-Volta hardware.atomicAdd_block(sm_60). Falls back to the device-scoped atomic.atomicAdd(double*, bool)in bin_distance. Below sm_60 the native double overload does not exist, leaving only RAFT'satomicAdd(T*, T)template, whose deduction then fails. Fixed with an explicit cast.cuco's block-scoped
atomics (atom.cas.cta.*, sm_60)in the sparse hash SpMV strategy, and the system-scoped atomics dynamic batching uses for its host/device handshake. Neither has a meaningful pre-Pascal emulation, so the device code is compiled out and both paths now fail with an explicit error instead of silently misbehaving. Thedense_smemSpMV strategy is unaffected.__half/__half2arithmetic (sm_53) in the CAGRA VPQ distance, the faiss-select comparators and the l_inf distance op. Addedutil/half_compat.cuh; pre-Pascal targets widen to fp32, which is also a small accuracy improvement.atomicAdd(double*, double)and theatomicAdd_blockfamily are only declared for__CUDA_ARCH__ >= 600, so in an older device pass they do not exist and every header that calls them unqualified fails to compile -- RAFT'sstrided_reduction.cuh(instantiated with double by cuVS) and CUB'sagent_histogram.cuhamong them. Addedutil/atomic_compat.cuh, which suppliesatomicAdd(double*)as a 64-bit CAS loop comparing bit patterns rather than values, so a stored NaN cannot make it spin forever, plusatomicAdd_blockoverloads forwarding to the device-scoped atomics. It expands to nothing on sm_60+ and in the host pass. The calls live in templates but their argument types are fundamental, so ADL contributes nothing and only declarations visible at the definition context are considered: the header must be included before any RAFT/CUB header, which is recorded in a comment at each of the four translation units that need it.raft::sqrt(__half)lowers to hsqrt, an FP16 ALU instruction requiring sm_53, and its RAFT overload static_asserts below that. Addedcuvs::util::sqrt_op, which widens __half to fp32 on such targets and forwards toraft::sqrtotherwise, and used it in the two element-wise passes over possibly-fp16 data.Fixes for RAFT, CUB and cuco headers all live in cuVS: those are CPM-fetched, so patching them in place would be invisible to a reviewer, lost on a version bump and untestable in CI. cuVS instead makes the missing declarations visible before the dependency header is parsed.
Runtime blockers addressed:
JIT-LTO fragments were pinned to a 70-real baseline regardless of
CMAKE_CUDA_ARCHITECTURES. LTO-IR is forward but not backward compatible, so the baseline now tracks the oldest architecture being built for; otherwise nvJitLink cannot link CAGRA/IVF search kernels for an older device.NN-descent was disabled entirely below sm_70, so CAGRA had no NN-descent graph build. Only the WMMA kernel needs tensor cores; the SIMT kernel uses nothing newer than sm_50 and its static shared memory stays under the 48 KiB per-block limit for every supported dtype. Lower its floor to sm_50 and fall back to it on pre-Volta devices instead of throwing.
Opus 5
Benchmarked on SIFT with a Maxwell device (940MX) :
=> 2.5–4.1× at recall 0.99