github pytorch/FBGEMM v1.8.0
FBGEMM_GPU v1.8.0
on GitHub

latest release: ciflow/rocm/6040
6 hours ago

FBGEMM_GPU v1.8.0 Release Notes

Highlights

New DRAM+SSD Composite Backend

  • Introduced DRAM_SSD backend type enabling DRAM as L1 cache with SSD backing store for embeddings (#5774, #5802, #5926)
  • Full training pipeline integration with metadata tracking, eviction callbacks, and stats reporting (#5951, #5957, #5958, #5961)
  • DramSsdKVEmbeddingCache composite backend with automatic DRAM-SSD coordination (#5929, #5946)

Sparse Operations Performance

  • Major kernel speedups: binary search for unique_indices (#5766), flat-grid linearize (#5768), and optimized delinearize (#5770)
  • Replaced at::_unique with CUB pipeline for better performance (#5769)
  • Added INT8 support to index_select_dim0 forward pass (#5782)
  • Preallocated output buffers for permute_2D_sparse operations (#5461)

TBE Determinism & Observability

  • Deterministic backward support with GPU coverage for INT-weight nobag aligned-D (#5773, #5801)
  • Silent-failure and queue-depth ODS counters for production monitoring (#5850, #5913)
  • VBE batch size validation and improved bounds checking (#5797, #5839)

Platform & Hardware Support

  • ROCm MI350 optimizations: asynchronous LDS loads (#5858) and subwarp_reduce_add (#5875)
  • ARM SVE-FP16 version of EmbeddingSpMDMNbit (#5728)
  • BF16 AVX2 8-bit and N-bit dequantization kernels (#5815)
  • Extensive ROCm grid overflow fixes across TBE, sparse, and quantization kernels

PyTorch 2.x Compatibility

  • FakeTensor and meta implementations for padded-FP8, 8bit→bf16 quantize, and FusedNBit ops (#5910, #5918)
  • Improved opcheck coverage with data-dependent meta hardening (#5911, #5919)

New Features

  • Adaptive regularization kernel for sparse embeddings (#5949)
  • TBE metadata write/lifecycle methods (#5904, #5915)

Software Requirements

FBGEMM_GPU v1.8.0 has been tested and known to work on the following setups:

  • PyTorch: v2.13.x
  • CUDA: v12.6, 13.0
  • Python: v3.10, 3.11, 3.12, 3.13, 3.14
  • ROCm: v7.0, 7.1

It is recommended to prepare an isolated environment for installing and running FBGEMM_GPU, such as Conda and/or Docker.

Availability

FBGEMM_GPU can be fetched directly from PyPI: 

# FBGEMM_GPU CUDA variant
pip install fbgemm-gpu==1.8.0

# FBGEMM_GPU CPU variant
pip install fbgemm-gpu-cpu==1.8.0

Alternatively, it can be fetched from PyTorch PIP: 

# FBGEMM_GPU CUDA variant
pip install fbgemm-gpu==1.8.0 --index-url https://download.pytorch.org/whl/cu126/
pip install fbgemm-gpu==1.8.0 --index-url https://download.pytorch.org/whl/cu130/

# FBGEMM_GPU CPU variant
pip install fbgemm-gpu-cpu==1.8.0 --index-url https://download.pytorch.org/whl/cpu

Changes

Table Batched Embedding (TBE) Operators

For GPU

  • [New] Add SVE-FP16 version of EmbeddingSpMDMNbit (#5728)
  • [Fix] Fix invocations of deprecated torch api (#5772)
  • [Improvement] Add TBE backward determinism tests (#5773)
  • [Improvement] Option to exclude UVM weights from core dump (#5784)
  • [Fix] Fix backward determinism tests (#5786)
  • [Improvement] Add OOB benchmarking and V2 bounds check mode support to bounds_check_indices benchmark (#5797)
  • [Improvement] Add deterministic GPU coverage for INT-weight nobag aligned-D fast path (#5801)
  • [Improvement] Migrate embedding_bounds_check_v2 to cap_grid_dim_x (#5812)
  • [Fix] Fix V2 bounds check printf: add __syncthreads() after tree reduction (#5818)
  • [Improvement] Migrate TBE forward template to cap_grid_dim_x (#5828)
  • [Fix] Fix ROCm deadlock in get_pruning_lengths kernel (#5831)
  • [Improvement] TBE inference INT-weight nobag: shifted-index store fast path (ROCm) (#5778) (#5833)
  • [Improvement] Wrap .item() args to torch._check* helpers in cast for stricter typing (#5834)
  • [Improvement] Add PROCESS_BLOCK macro for grad accumulation loop unrolling (#5835)
  • [Improvement] Add VBE batch size validation in TBE (#5839)
  • [Improvement] Silent-failure ODS counters (#5850)
  • [Improvement] Migrate TBE backward template to cap_grid_dim_x; clean up legacy get_max_thread_blocks helpers (#5853)
  • [Improvement] Implement asynchronous LDS loads for MI350 (#5858)
  • [Improvement] Make tbe.cache.cache_config canonical for TBE cache types (#5865)
  • [Improvement] Add subwarp_reduce_add for GROUP_REDUCE_ALL_SUM on ROCm (#5875)
  • [Fix] Fix HIP grid overflow in lru_cache_insert_kernel (#5881)
  • [Fix] Fix HIP grid overflow in direct_mapped_lxu_cache_lookup_kernel (#5882)
  • [Fix] Register Meta and Autograd dispatch keys for dense TBE op (#5886)
  • [Improvement] Improve ROCm detection and centralize differing_executors health check fixes (#5889)
  • [Improvement] Assert bad offsets in bounds_check_indices (Part 2) (#5895)
  • [Improvement] Replace deprecated PyTorch APIs in fbgemm_gpu (#5887) (#5896)
  • [Improvement] Remove legacy get_max_thread_blocks helper from cuda_utilities (#5897)
  • [Fix] Fix test_lru_cache_insert_large_grid associativity on ROCm wavefront64 (#5899)
  • [Improvement] Remove unreachable test_direct_mapped_lxu_cache_lookup_large_grid (#5900)
  • [Fix] Fix test_lxu_cache_lookup: only test kWarpSize associativity (#5903)
  • [New] Add metadata write/lifecycle methods (#5904)
  • [Fix] Fix unit test for rowwise adagrad with counter (#5908)
  • [Improvement] Queue-depth ODS counters (#5913)
  • [Fix] Make sparse_type_int_to_dtype import resilient to stale frozen torch.package depots (#5942)
  • [Fix] Fix integer overflow in csr_offsets subscripts (#5950)
  • [Fix] Re-export TBE optimizer config types to restore single class identity (#5968)

For CPU

  • [Improvement] Add compare CPU baseline support to TBE benchmarks (#5749)
  • [Fix] Fix uvm_to_cpu/uvm_to_device/uvm_get_guard_index for host-mapped tensors (#5816)
  • [Improvement] Make V1 codegen conditional on split_function_args_v1 (#5824)
  • [Fix] Fix optional TBE V1 codegen (#5842)
  • [Improvement] Add parametrized NaN-poison zero-fill test for async embedding loads (#5888)
  • [Improvement] Migrate OMP use to Aten/Parallel.h (#5947)
  • [New] Adaptive Regularization for sparse embeddings — kernel + codegen (#5949)

SSD Table Batched Embedding (TBE) Operators

  • [Improvement] make stats logging a dictionary (#5733)
  • [Improvement] add e2e test (#5762)
  • [New] Add BackendType.DRAM_SSD enum value and DRAM_SSD_VIRTUAL_TABLE enum to torchrec and sparsenn_configs (#5774)
  • [Improvement] add overall L1→DRAM hit rate metric (#5777)
  • [Improvement] Migrate internal tbe/* and torchrec callers to canonical type packages (#5780)
  • [New] Add DRAM_SSD backend type to BackendType enum (#5802)
  • [Fix] Fix DramKV race: hold rlock during inplace update writes (#5807)
  • [New] Add dirty bit tracking to FixedBlockPool (#5817)
  • [New] Add eviction writeback callback to FeatureEvict (#5832)
  • [Improvement] Make tbe.config.embedding_config canonical for TBE config types (#5864)
  • [New] Add virtual methods to EmbeddingKVDB interface (#5868)
  • [Fix] Fix int64 key truncation in DramKV inference eviction loop (#5876)
  • [Improvement] Add metadata_dim_ and total_rows_written_ members to EmbeddingRocksDB (#5883)
  • [New] Add metadata column family fields to EmbeddingRocksDB (#5890)
  • [New] Add metadata write/lifecycle and split-read methods (#5915)
  • [New] Wire metadata column family into existing EmbeddingRocksDB methods (#5920)
  • [New] Add DramSsdKVEmbeddingCache composite backend (DRAM forwarding) (#5926)
  • [New] Add SSD backend integration and metadata API to DramKVEmbeddingCache (#5929)
  • [Improvement] Add MetaHeader parity test (#5935)
  • [New] Add is_eviction_enabled() predicate for DramKV eviction guards (#5946)
  • [New] Implement DRAM+SSD composite logic in DramSsdKVEmbeddingCache (#5951)
  • [Fix] Fix use-after-free race between DramKV eviction and inplace updates (#5954)
  • [New] Add DRAM_SSD support to KVTensorWrapper (#5957)
  • [New] Add DRAM_SSD backend to training.py (#5958)
  • [New] Add DRAM_SSD stats reporting to tbe_stats_reporters (#5961)
  • [New] Return zero optimizer state for DRAM_SSD embedding cache (#5965)

GenAI Support and Operators

Triton GEMM Support

  • [Fix] Fix MoE scatter_add_padded_tokens test timeout (#5863)
  • [Improvement] Tma load grouped gemm (#5937) (#5960)

Quantization Operators

  • [Improvement] Switch bf16 quantize ops to at::kBFloat16 (#5735)
  • [Improvement] Implement PackDepthwiseConvMatrix in NEON + deprecate aarch64 compat layers (#5779)
  • [New] Add bf16 avx2 8bit dequant and Nbit dequant (#5815)
  • [Fix] Cpp20 clang tidy gpu (#5827)
  • [Fix] Fix ROCm __syncthreads deadlock in gemv_quantized_fp8_fp8 (#5856)
  • [Fix] Fix ROCm __syncthreads deadlock in quantize_float_to_mx4_kernel (#5857)
  • [Fix] Register missing FakeTensor/meta impls for padded-FP8 and 8bit->bf16 quantize ops (#5910)
  • [Fix] De-flake quantize tests and break pt2_compliant_tag cascade (#5911)
  • [Fix] Add FusedNBit metas + harden data-dependent metas + xfail harness-incompatible opcheck tests (#5918)
  • [Fix] De-flake kmeans/comm_codec/permute_pooled_embedding tests (#5919)
  • [Fix] Fix residual fbgemm_dev OMH test-health issues (#5931)
  • [Fix] Fix deadline/flaky + suppress GPU-only opcheck SKIPPING (OMH round 2) (#5932)
  • [Fix] Fix comm_codec/zipf/preproc OMH residuals (round 3) (#5939)
  • [Improvement] Gate MTIA mx4 dep of quantize_utils behind ovr_config//gpu:mtia (#5956)

Sparse Operators

  • [New] Add permute_2D_sparse_preallocated_out with pre-allocated output buffers (#5461)
  • [Fix] pack_segments backward CUDA: zero-init gradient buffer to fix uninit-memory NaNs (#5754)
  • [Fix] Make lengths tensor contiguous in pack_segments CUDA kernels (#5761)
  • [Fix] Fix HIP grid overflow in permute_1D_sparse_data_cuda (#5763)
  • [Improvement] Speedup unique_indices_length_kernel via binary search (#5766)
  • [Improvement] Speedup linearize_index via flat-grid kernel (#5768)
  • [Improvement] Replace at::_unique with cub pipeline + end_bit trim (#5769)
  • [Improvement] Speedup delinearize via feature-lookup over num_unique (#5770)
  • [Fix] Fix index_select library type-checking test (#5775)
  • [New] Add INT8 (uint8) support to FBGEMM index_select_dim0 forward (#5782)
  • [Fix] Fix jagged_unique_indices out-of-bounds regression (#5799)
  • [Fix] Fix permute_embeddings library type-checking test (#5808)
  • [Fix] Fix SVE no_bag uint8 overflow (#5810)
  • [Improvement] Replace ALWAYS_INLINE with constexpr for specialization helpers (#5821)
  • [Improvement] Migrate sparse_permute_1d to cap_grid_dim_x_from_workload (#5826)
  • [Fix] Fix wavefront64 ballot truncation in populate_bucketized_permute warp kernel (#5830)
  • [Improvement] Replace autovec embedding specialization macros with C++20 optimisations (#5823) (#5838)
  • [Improvement] Add missing headers and use std::min and std::max (#5836) (#5845)
  • [Improvement] Use 256-thread blocks for jagged dense-output kernel (#5848)
  • [Fix] Fix block_bucketize_sparse_features meta optional-output semantics (#5861)
  • [Fix] Add populate_bucketized_permute meta + fix block_bucketize opcheck (#5869)
  • [Fix] Fix block_bucketize 2DWeightsTest opcheck faketensor failures (#5870)
  • [Fix] Implement ragged inputs in group_index_select_2d_bench (#5871)
  • [Fix] Fix int32 overflow in segment_sum_csr and harden grid-size / num_valid narrowings (#5877)
  • [Fix] Guard dense_to_jagged_forward against negative total_L (#5879)
  • [Fix] Guard jagged_index_select against negative output rows (#5880)
  • [Fix] Guard jagged_index_add against negative output rows (#5884)
  • [Fix] Guard repeat_arange against negative output_size (#5885)
  • [Fix] Improve error validation and guard int32 overflow in dense_to_jagged (#5891)
  • [Improvement] Guard jagged_slice against negative row counts (#5893)
  • [Fix] Fix the grid problem with sparse permute 2d (V2: threshold-guarded) (#5898)
  • [Improvement] Normalize group_index_select_2d_bench param names (#5901)
  • [Improvement] Normalize config-column names in benchmark reporting + analysis (#5902)
  • [Fix] Fix HIP grid overflow in zipf_cuda (#5905)
  • [Fix] Fix HIP grid overflow in expand_into_jagged_permute_cuda (#5906)
  • [Fix] Fix HIP grid overflow in pack_segments_cuda{,_v2} (#5907)
  • [Fix] Add grid-stride loop and ROCm cap to _populate_length_to_feature_id_inplace_kernel (2D weights) (#5916)
  • [Fix] Skip opcheck for pack_segments_large_grid stress test (#5922)
  • [Fix] Add grid-stride loops and ROCm cap to _populate_length_to_feature_id_inplace_kernel and _populate_bucketized_permute_warp_parallel_kernel (#5923)
  • [Fix] Add grid-stride loop and ROCm cap to _segment_sum_csr_cuda_kernel (#5925)
  • [Fix] Add grid-stride loop and ROCm cap to _batched_complete_cumsum_kernel (#5927)
  • [Improvement] Skip dense_to_jagged aot_dispatch_dynamic opcheck variants (#5928)
  • [Fix] Add grid-stride loop and ROCm cap to batched_unary_embeddings_{forward,backward}_kernel (#5933)
  • [Fix] Add grid-stride loop and ROCm cap to index_add_2d_with_unique_indices_kernel (#5934)
  • [Fix] Add grid-stride loops and ROCm cap to reorder_batched_ad_{lengths,indices,sequence_embeddings}_kernel (#5938)
  • [Improvement] Round-4 OMH residuals: GPU-only opcheck suppression, timeouts, flaky test fixes (#5948)
  • [Fix] Apply ROCm grid-overflow cap to index_select_2d_kernel (#5959)
  • [Fix] Fix HIP grid overflow in jagged_softmax_forward_cuda (#5973)
  • [Fix] Fix HIP grid overflow in jagged_softmax_backward_cuda (#5974)

Build / CI Improvements and Better Engineering

  • [Improvement] ci: declare workflow-level contents: read on 11 workflows (#5764)
  • [Improvement] Update repository references (#5767)
  • [Improvement] Remove Windows CI from fbgemm_ci.yml to address OSS CI budget overage (#5781)
  • [Improvement] Simplify cuda_calc_xblock_count / cuda_calc_block_count via if constexpr (#5783)
  • [Fix] Upgrade macOS CI runner from macos-13 to macos-14 (#5785)
  • [Improvement] Use Python 3.10+ typing in codegen scripts (#5787) (#5792)
  • [Improvement] Use Python 3.10+ typing in tests (#5788) (#5793)
  • [Improvement] Disable build-macos to confirm cost driver (#5798)
  • [New] Add determine_grid_blocks_from_{items,blocks} helper pair + BlockCapPolicy enum (#5800)
  • [Improvement] Migrate fb-exported to meta-exported where possible (best effort) (#5809)
  • [Fix] Fix aarch64+gcc OSS build break (#5813)
  • [Improvement] Use Python 3.10+ typing in CI scripts (#5819)
  • [Improvement] Add ISA-pinned FP16/FP32 test runs and apply C++17/20 modernization (#5820)
  • [Fix] FindAVX: probe AVX support via compile-only check (#5846)
  • [Improvement] Use Python 3.10+ typing in setup and list_versions (#5847)
  • [Improvement] Add seed_all() deterministic-seed test helper (#5851)
  • [Fix] De-flake gen_ai gqa + hstu attention tests (#5854)
  • [Fix] Fix off-by-one rank in MoE shuffling test (#5855)
  • [Improvement] Remove CUDA 12.9.1 support from FBGEMM, following PyTorch (#5859)
  • [Fix] Exclude free-threaded py3.14t + ROCm wheel builds (#5862)
  • [Improvement] Add FBGEMM 1.6/1.7/1.8 entries to PyTorch compat table (#5941)
  • [Fix] Add visibility("hidden") to KernelLauncher to fix cross-module TLS mismatch (#5963)
  • [Fix] Drop CUDA 13.2 (cu132) from all Linux wheel build matrices (#5966)
  • [Improvement] Drop CUDA 12.8 from release workflows (torch 2.13 ships no cu128) (#5976)

Tests and Benchmarks

  • [Improvement] Use Python 3.10+ typing in benchmarks (#5841) (#5844)
  • [Improvement] Remove usage of aarch64 compatibility layers (#5921)
Check out latest releases or
releases around pytorch/FBGEMM v1.8.0

Don't miss a new FBGEMM release

NewReleases is sending notifications on new releases.

Get notifications