Skip to content

Releases: pytorch/FBGEMM

FBGEMM_GPU v1.1.0 Release Notes

29 Jan 20:54
@q10 q10
v1.1.0
7243043
Compare
Choose a tag to compare
FBGEMM_GPU v1.1.0 Release Notes Latest
Latest

Highlights

TBE GPU

  • Added support for int64_t table indices and offsets in TBE inference
  • Introducing support for int32_t indices in TBE training
  • Extended TBE support for larger embedding dimensions
  • Made the learning rate a tensor value
  • Improvements on indices bounds checking

TBE CPU

  • Improved ARM support with SVE implementations for matrix multiplication and float matrix transpose
  • Improved the EmbeddingSpMDMAutovec API
  • Migrated FP32 ops to OSS

TBE SSD

  • Enabled VBE in SSD-TBE
  • Async initialization of RockDB SSD tensors and pad before writing to rocksDB
  • Improvements on indices bounds and other constraints checking

Gen AI Ops

  • Added nccl_alltoall function
  • Custom allgather support multiple dtypes, with dtype checking to prevent silent failures

ROCm

  • Add CK FP8 Batched GEMM and Rowwise GEMM kernels along with heuristic tuning
  • Fixed CK FP8 rowwise quantization for some GEMM shapes
  • Introduced HIP-specific optimizations to the TBE forward and backward passes

SLL ops

  • Migrated Sequence Learning Library (SLL) ops to OSS

Better Engineering

  • Restructured the build to produce multipiple smaller shared libraries instead of a single large binary
  • New and improved tests and benchmarks
  • Improved ROCm build variant support
  • Add build support for CUDA 12.6 and Python 3.13

Software Requirements

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

  • PyTorch: v2.6
  • CUDA: v11.8, 12.4, 12.6
  • Python: v3.9, 3.10, 3.11, 3.12, 3.13

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 (only the CUDA 12.4 variant is available)
pip install fbgemm-gpu==1.1.0

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

Alternatively, it can be fetched from PyTorch PIP:

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

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

Changes

Table Batched Embedding (TBE) operators

For GPU

  • [New] Add support for int32_t indices in TBE training (#3377, #3375, #3374, #3372, #3371, #3324, #3267, #3264, #3263 #3257)
  • [New] Add support for int64_t indices and offsets in TBE inference (#3254, #3233)
  • [New] Extend TBE support for larger embedding dimensions (#3462, #3467)
  • [New] Make learning rate tensor (Backend) (#3287, #3310, #3332)
  • [New] Add PTA checks to embedding_bounds_check kernels" (#3318)
  • [Fix] Fix PackedTensorAccessor for batch_index_select (#3281)
  • [Fix] Set cache_precision = weights_precision in TBE if it is not explicitly set (#3370)
  • [Fix] Fix pt2_wrapper registration for unified TBE interface (#3238)
  • [Fix] Fix PT2 compliant opcheck tests (#3404)
  • [Fix] Fix FBGEMM_GPU_MEMCHECK in Split optimizers (#3416)
  • [Fix] Fix learning rate as tensor for PT2 compile (#3407)
  • [New] Add new optimizer state row_counter for Adam [Frontend] (#3558)
  • [New] Add new optimizer state row_counter for Adam [Backend] (#3342)
  • [Fix] Back out ""Add support for int64_t indices and offsets in TBE inference [7C/N]"" (#3258)
  • [Fix] Back out ""Add support for int64_t indices and offsets in TBE inference [8/N]"" (#3255)
  • [Fix] Fix global weight decay Faketensor test (#3341)
  • [Fix] Fix pt2_wrapper registration for unified TBE interface (#3237)
  • [Fix] Fix ""Cannot call numel() on tensor with symbolic sizes/strides"" (#3368)
  • [Fix] Fix grid size overflow in generate_vbe_metadata (#3484)
  • [Fix] Fix an integer overflow in permute_multi_embedding() (#3465)
  • [Fix] Fix the sync point caused by iter_cpu.item() (#3401)
  • [Fix] Fix global weight decay Faketensor test (#3341)
  • [Fix] Hot fix to skip VBE CPU reshaping for MTIA (#3466)
  • [Fix] address mem over used during flushing (#3460)
  • [Improvement] Add iter singular value into TBE optimizer state (#3228)
  • [Improvement] V2 fwd modified warps (#3570)
  • [Improvement] Add enable_async_update into tbe signature and config (#3431, #3461)"
  • [Improvement] Adjust kNumThreads for bounds_check_indices_kernel (#3299)
  • [Improvement] Reduce registers in bounds_check_indices" (#3298)
  • [Improvement] Mark unified autograd function traceable (#3378)
  • [Improvement] Improve bounds_check_indices for VBE (#3388, #3386)
  • [Improvement] Do not call scalar_type (#3394)
  • [Improvement] optimizer 1d -- EMA in place (fbgemm part) (#3402)
  • [Improvement] Clean up nbit_forward tests (#3286)
  • [Improvement] Remove unused-variable in some generated code (#3327)
  • [Improvement] Limit grid size of bounds_check_indices" (#3282)
  • [Improvement] Support config based bound check version via extended modes (#3418)
  • [Improvement] Use int64_t index for SplitOptimizer grad (#3447)
  • [Improvement] Remove unused arg from generate_vbe_metadata frontend (#3453)
  • [Improvement] Add generate_vbe_metadata test (#3483)
  • [Improvement] Support config based bound check version via extended modes (#3454)
  • [Improvement] make iter PT2 compatible (#3253)
  • [Improvement] Add meta function for PT2 wrappers (#3240)
  • [Improvement] Nesterov (#3232)

For CPU

  • [New] Introduce SVE function for matrix multiplication (#3348)
  • [New] Add sve implementation for float matrix transpose (#3421)
  • [New] autovec specialization framework (#3393)
  • [New] Move FP32 kernels to OSS (#3568)
  • [Improvement] Pull in PR for Kleidi-based FP16 kernel (#3507)
  • [Improvement] Use local buffer where possible (#3304)
  • [Improvement] Refactor GenerateEmbeddingXXX functions (#3307)
  • [Improvement] Increase local_storage size to 512 floats (#3357)
  • [Improvement] Adjust EmbeddingSpMDMAutovec API (#3366)
  • [Improvement] Split loops to work around loop vectorizer weakness (#3406)
  • [Improvement] Do an early check that data_size is not negative (#3305)
  • [Improvement] Fix strict aliasing violation, code cleanup (#3306)

SSD TBE Operators

  • [New] Enable VBE in SSD-TBE (#3247)
  • [Improvement] put KVTensorWrapper in its own header (#3575)
  • [Improvement] Moving KVTensorWrapper to a header file to be used in ModelStore checkpointing code (#3276)
  • [Improvement] Async initialization of RockDB SSD tensors (#3520)
  • [Improvement] pad before writing to rocksDB (#3245)
  • [Improvement] use RocksDB iterator to read key range from ssd embedding (#3495)
  • [Improvement] Log total duration spent prefetching (#3487)
  • [Improvement] address mem over used during flushing (#3460)
  • [Improvement] Create move TBE to right device, and set Cache Load in TBE class (#3438)
  • [Improvement] Unit test for new move tbe from device/cache_load method (#3437)
  • [Improvement] make L2/rocksdb update async optional (#3429)
  • [Improvement] Drop RoPE when filling KV cache (#3346)
  • [Improvement] Remove setting total_cache_hash_size as buffer (#3441)
  • [Improvement] Add meta registrations for kv_cache operators (#3442)
  • [Improvement] remove output dtype restriction in SSD TBE (#3524)
  • [Improvement] change pmt require grad to false when detached (#3525)
  • [Improvement] add more attributes to PartiallyMaterializedTensor (#3300)
  • [Improvement] skip broken inference test that uses ssd TBE (#3494)
  • [Improvement] "coro => fut" (#3430)
  • [Improvement] Reland of D65489998 Optimize sharding performance of embeddings (#3549)
  • [Improvement] Remove torch.jit.script (#3562)

GenAI Support and Operators

  • [New] Add nccl_alltoall function (#3551)
  • [New] custom allgather support multiple dtypes (#3498)
  • [Improvement] Make sure fake tensor functions return on proper device (#3258)
  • [Improvement] Add CPU registrations to custom operators (#3262)
  • [Improvement] Check src & dst dtypes in allgather to prevent silent failures (#3523)
  • [Improvement] Better shape function registration (#3237, #3340)
  • [Improvement] Package re-organization improvements (#3546, #3251, #3419, #3268, #3512)

FP8 and other Quantization support

  • [New] New autotune config for M=4 (#3277)
  • [New] MoE FP8 grouped GEMM (#3321)
  • [New] Add shape check on GroupedGEMM kernel (#3449)
  • [New] Tuning for fp8 gemm with emu1.7 shapes (#3436)
  • [Improvement] more fp8 tuning for decode and not need to pad (#3576)
  • [Improvement] llm decode shapes fp8 rowwise gemm tuning (#3565)
  • [Improvement] Split FP8 Grouped Gemm into dynamic and static version (#3543)
  • [Improvement] Warp-specialized FP8 rowsise GEMM kernel (#3532)
  • [Improvement] Add Cutlass FP8 Grouped Gemm to Quantize Bench (#3530)
  • [Improvement] Fixed FBGEMM fp8 rowwise for irregular shapes (#3491)
  • [Improvement] Properly define preallocated output as mutable in fp8 rowwise gemm (#3476)
  • [Improvement] Fix FP8 Rowwise Gemm Compilation with Auto-functionalize V2 (#3457)
  • [Improvement] Support zero-size inputs in FP8 cuda quantize kernel (#3448)
  • [Improvement] update FP8 GEMM tuning for emu1.7 7B shapes (#3391)
  • [Improvement] Customize FP8 grouped GEMM for non-zero calculation for token choice MoE (#3383)
  • [Improvement] Support FP8 grouped GEMM with cudagraph (#3373)
  • [Improvement] Refactor FP8 grouped GEMM to prepare cudagraph support (#3369)
  • [Improvement] Improve FP8 BMM heuristic for large shapes and MoE E2E performance (#3344)
  • [Improvement] retune some of the EMU1.6 7B FP8 GEMM shapes (#3328)
  • [Improvement] Make FP8 BMM output contiguous (#3270)
  • [Improvement] Tune FP8 rowwise bmm tile hueristic (#3256)
  • [Improvement] more FP8 GEMM tuning for LDM shapes (#3414)
  • [Improvement] Split up f8f8bf16_rowwise_batched.cu (#3381)
  • [Improvement] use sym int in quant...
Read more
Assets 2
Loading

FBGEMM_GPU v1.0.0 Release Notes

19 Oct 20:57
@spcyppt spcyppt
v1.0.0
73a64e7
Compare
Choose a tag to compare
FBGEMM_GPU v1.0.0 Release Notes

Stable API

We provide the stable API support starting from FBGEMM_GPU v1.0.0. This includes Table batched embedding (TBE) modules, Pooled embedding operators and modules, Sparse operators, Jagged tensor operators and Quantization operators.

  • API backward compatibility guarantees via thorough testing. We guarantee that our stable APIs will be backward compatible within a major version, meaning that the stable APIs for v1.0.0 will be compatible with every future release unless explicitly announced in advance
    *Enhanced documentation, ensuring that every stable API has comprehensive and up-to-date documentation.
  • Functionality guarantees are only provided through unit testing framework. We do NOT guarantee any functionalities that are NOT explicitly tested and documented in our unit tests.
  • No performance guarantees. However, we are committed to providing support on a best-effort basis.

More details can be found in stable API documentation

Highlights

Table Batched Embedding (TBE)

  • New optimizer support for TBE Training
  • Enhanced Global weight decay support in TBE
  • Improvement and bug fixes for TBE training and inference modules and sparse operators

For SSD

  • New pipeline prefetching enabled
  • New cache and indices related ops
  • Integration of L3 cache to TBE operators
  • Many improvements to kernel and logging

For CPU

  • New type support for CPU Sequence TBE
  • Kernel improvements and bug fixes

Generative AI

  • Gen AI Ops support and improvement
  • Improvements to Triton-based and CUTLASS-based operators
  • New and optimized FP8 GEMM and quantization operators

Others

  • Optimized MX4 quantization operators
  • New dequantization operator
  • Removal of python 3.8 Support

Better engineering

  • Code refactoring and reorganization for faster builds
  • New and improved tests and benchmarks
  • Improved AMD support

Software Requirements

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

  • PyTorch: v2.5
  • CUDA: v11.8, 12.1, 12.4
  • Python: v3.9, 3.10, 3.11, 3.12

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 (only the CUDA 12.4 variant is available)
pip install fbgemm-gpu==1.0.0

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

Alternatively, it can be fetched from PyTorch PIP:

# FBGEMM_GPU CUDA variant
pip install fbgemm-gpu==1.0.0 --index-url https://download.pytorch.org/whl/cu118/
pip install fbgemm-gpu==1.0.0 --index-url https://download.pytorch.org/whl/cu121/

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

Changes

Table batched embedding (TBE) operators

For GPU

  • [New] Ensemble adagrad optimizer (#3197, #2955, #2954, #3161, #3091, #2981, #2889, #3180, #3158)
  • [New] Bounds check in prefetch in TBE training (#3015)
  • [New] Method to update internal hyperparameters for FBGEMM TBE (#3025)
  • [Improvement] Enhanced Global Weight Decay and state tracking (#2904, #2897, #2882, #2896, #2890, #2884, #2883 )
  • [Improvement] masked_index_* values index type fix (#2979)
  • [Improvement] generate_vbe_metadata fixes (#3095, #3087)
  • [Improvement] Fixes on the efficiency of VBE TBE forward due to blocking D2H copy (#2862)
  • [Improvement] Work around on offsets and indices type mismatch int TBE training (#3037)
  • [Improvement] Add a host map option for a UVM tensor alloc (#3073)
  • [Improvement] uvm_to_device expose device as interface (#3030)
  • [Improvement] Add Meta backend/dispatcher for new_unified_tensor (#3005)
  • [Improvement] General TBE enhancements and bug fixes (#2892, #3114, #3022, #2958)
  • [Improvement] Consolidate repeat code in TBE inference (#3028)

For CPU

  • [New] Add int4 to int4 CPU Sequence TBE kernel (#2996, #2994)
  • [New] Use auto-vec kernel in CPU sequential embedding lookup for int8 tables (#2863, #2878)
  • [Improvement] Work around OMP barrier issue with MSVCand unused var error (#2918, #3084)

SSD Table batched embedding (TBE) operators

  • [New] Enable pipeline prefetching (#2963)
  • [New] Enable cache line locking support in SSD kernel (#2949)
  • [New] Add L2 flush (#3110)
  • [New] Added SSD ODS and IO/mem stats (#2906, #2913, #3035)
  • [New] Add SSDScratchPadIndicesQueue (#2911, #2948)
  • [New] Integrate l2 cache to TBE operator (#2959, #3032, #3031 )
  • [New] Add ssd_update_row_addrs (#2953)
  • [New] Add bounds check in SSD-TBE (#3013)
  • [New] Add 32-bit index support in SSD kernels (#3064)
  • [New] Add kv cache related ops (#3001, #2968)
  • [New] Add compact_indices op (#3075 )
  • [New] Create embedding cache interface and impl RocksDB cache (#2858)
  • [New] Reduce prefetch SM usage when using pipeline prefetching (#2991)
  • [New] Add a host map option for a UVM tensor alloc (#3003)
  • [New] Add masked_index_select and refactor masked_index_put (#2910)
  • [Improvement] Add parallelism on cache update (#3062)
  • [Improvement] add parameter server attributes (#2947)
  • [Improvement] Make the scratch pad tensor UVA (#2844)
  • [Improvement] Use less thread blocks for find_uncached kernel (#3101)
  • [Improvement] Fix stream sync for scratch pad eviction (#2843)
  • [Improvement] Make indices related to cache eviction UVA tensors (#3077
  • [Improvement] Split cachelib cache into header and src (#3063)
  • [Improvement] Record more functions and logging in SSD TBE (#2854, #2867, #2975)
  • [Improvement] Attach eviction filling logic to set_cache (#3034)
  • [Improvement] Move set_cache and set_async to background thread (#3033)
  • [Improvement] Refactoring vec copy in masked_index_put_kernel (#2861, #2908)
  • [Improvement] Increase memcpy and compute overlap (#2860)
  • [Improvement] Add set_async in background thread (#3036 )
  • [Improvement] Make evicted_rows a UVA buffer (#3079 )
  • [Improvement] General enhancement and bug fixes (#2937, #2993, #3151, #3089, #2898, #2930)

GenAI Support and Operators

  • [New] Decode and Prefill support (#3009 )
  • [New] Support rope with block tables (#3146)
  • [New] EP support (#3071)
  • [New] Implement SDPA kernel wrapper to use run_kernel flow for perf (#2820)
  • [Improvement] Move mqa code (#3011)
  • [Improvement] BE improvements to init_comms #3103

Triton GEMM support

  • [New] Enable torch.compile compatibility for triton fp8 rowwise gemm (#2978)
  • [New] Add 3D+ input support for fp8 rowwise GEMM (#2845)
  • [New] GEMM custom op enablement (#3046)
  • [New] Add 3D+ input support for fp8 rowwise GEMM (#2845)
  • [Improvement] Add fused bias to Triton FP8 Rowwise Kernels (#2852)
  • [Improvement] Triton dependency ( #3027)
  • [Improvement] Fix triton fp8 handling of non-contiguous inputs (#2919)
  • [Improvement] More autotune configs and bug fixes in TMA kernel (#3078, #3066, #3072)
  • [Improvement] Fp8 gemm tweak for 405B Decoding (#3104 )

FP8 and other Quantization support

  • [New] CK FP8 Optimizations and fixes (#2940, #2912, #2987, #3017, (#2893 )
  • [New] FP8 kernel development and enablement (#2866)
  • [New] GenAI CK Version update and integration (#2865, #2971)
  • [Improvement] Also hipify the fp8 related cuda functions (#2834 )
  • [Improvement] Auto-generation of CUTLASS Extension Kernel Templates (#2932)
  • [Improvement] Marlin Mixed Input Kernel Productionization (#3008)
  • [Improvement] Remove redundant torch.abs (#3020, #2822 )
  • [Improvement] Tuning for 405B/70B Prefill with small seqlen (#3042)
  • [Improvement] Added new instances for 405B decoding (#2936 )

Permute and Pooled Embeddings Ops

  • [New] Implementation of permute_multi_embedding (#2833)
  • [Improvement] Clean up and removal of unused exception (#2832, #2891)
  • [Improvement] Use at::parallel_for in cpu kernel (#2817)
  • [Improvement] Add dispatch_to_cpu for the operators (#2874, #2881)
  • [Improvement] Print the exact variable values triggering the alert in Merge Pooled Embedding (#3038)

Sparse Operators

  • [New] Support original indices for FBGEMM block bucketization flag (#2999, #2925)
  • [Improvement] Fix pack_segments backward when grad is non-contig (#3006)
  • [Improvement] Fix FBGEMM_GPU_MEMCHECK in sparse_ops_cuda (#2943 )
  • [Improvement] Update sparse_ops.py to use generic gpu target fbgemm_gpu:input_combine to support both nvidia and AMD(#2905)
  • [Improvement] Add abstract impl and functions (#2962, #2983, #3000 )
  • [Improvement] Use guard_size_oblivious in tbe_input_combine_abstract fake kernel (#2923)
  • [Improvement] Out variant for asynchronous_exclusive_cumsum_cpu + some more static dispatch kernels (#3090)

Quantize ops

  • [New] Add a CPU nbit to float dequantization op that supports torch.quintMxN type (#2995)

MX4 Ops

  • [New] Optimize FBGEMM Triton MX4 Quantize-Dequantize (#2838, #2837)
  • [New] Rounding Mode Support (#2821, #2816, #2933, #2859 )
  • [New] FBGEMM/TorchRec MX4 padding support (#3055, #3047, #3010 )
  • [New] Add Stochastic downcasting to MX4 Quantization (#2899)
  • [New] Support for other MX4 formats in Triton kernels (#2900)
  • [Improvement] Refactor MX4 Kernel to operate on flat tensors (#2836)
  • [Improvement] Optimize MX4 padding to minimize need for tuning (#3040)

Benchmarks / Tests

  • [New] Add schema compatibility test (#3130)
  • [New] Add SSD/UVM caching in TBE device benchmark (#3076)
  • [New] Add EmbeddingSpMDM8BitBenchmarkOutTypeFloat16 (#2952 )
  • [New] Add benchmark EmbeddingSpMDMNBitBenchmarkOutTypeFloat16 (#2901 )
  • [New] Add unit test for int4 to int4 sequence CPU TBE (#2997)
  • [New] Add rocm support for fp8 benchmarks (#2965)
  • [New] Add rotating buffer feature to quantize_bench #2857)
  • [New] Benchmark of fbge...
Read more
Assets 2
Loading

FBGEMM_GPU v0.8.0 Release Notes

29 Jul 22:40
@spcyppt spcyppt
v0.8.0
8846c27
Compare
Choose a tag to compare
FBGEMM_GPU v0.8.0 Release Notes

Release Note

Highlights

Table Batched Embedding

For GPU

  • New Table Batched Embedding (TBE) operators and momentum type support
  • New Intraining Embedding Pruning (ITEP) operators
  • VBE support for Dense TBE
  • Global weight decay support in TBE
  • New type support and improvement to SSD TBE
  • Improvement and bug fixes for TBE training and inference modules and sparse operators

For MTIA

  • MTIA support for DenseTBE

Generative AI

  • Gen AI Ops integration
  • Support for Triton-based and CUTLASS-based operators (#2552, #2537)
  • New FP8 GEMM and quantization operators
  • New query attention operators
  • New Car and All-To-All (NCCL-based) communication operators
  • AMD Support for FP8

Others

  • New MX4 quantization operators
  • Support for CUDA 12.4

Better engineering

  • Code refactoring and reorganization for faster builds
  • New tests and benchmarks
  • Improved AMD support

Software Requirements

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

  • PyTorch: v2.4
  • CUDA: v11.8, 12.1, 12.4
  • Python: v3.8, 3.9, 3.10, 3.11, 3.12

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 (only the CUDA 12.1 variant is available)
pip install fbgemm-gpu==0.8.0

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

Alternatively, it can be fetched from PyTorch PIP:

# FBGEMM_GPU CUDA variant
pip install fbgemm-gpu==0.8.0 --index-url https://download.pytorch.org/whl/cu118/
pip install fbgemm-gpu==0.8.0 --index-url https://download.pytorch.org/whl/cu121/

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

Changes

Table batched embedding (TBE) operators

For GPU

  • [New] VBE support for Dense TBE (#2628, #2620, #2641)
  • [New] BF16 momentum support in PARTIAL_ROWWISE_ADAM (#2524, #2522, #2518)
  • [New] Global weight decay support (#2516, #2507, #2506)
  • [New] Multi-pass prefetch for memory efficiency (#2566)
  • [Improvement] Work around masked_select for numel > MAX_INT (#2648)
  • [Improvement] Fused optim in backward capability with aot_autograd (#2651)
  • [Improvement] Weights mutations declaration in TBE backward ops schemas (#2698)
  • [Improvement] Helper ops to support cache conflict misses (#2571)
  • [Improvement] Fixed the hang issue in some TBE GPU optimizers (#2509)
  • [Improvement] Misc TBE fixes and refactoring (#2583, #2597, #2529)
  • [Improvement] Cache prefetch and conflict miss improvements (#2596, #2514)

For MTIA

  • [New] Support MTIA in DenseTableBatchedEmbeddingBagsCodegen (#2680)

SSD Table batched embedding (TBE) operators

  • [New] Add FP16 weight and output support to SSD TBE (#2638)
  • [New] Implementation of PS KV DB for FBGEMM TBE operator (#2664, #2642)
  • [Improvement] Removal of D->H sync when calling lxu_cache_lookup (#2672)
  • [Improvement] Recording of functions in SSD TBE (#2670)
  • [Improvement] Added options, assertions and logs for training and inference SSD TBE (#2689, #2657)
  • [Improvement] SSD TBE backend fixes (#2645, #2671)

New Operator Groups

GenAI FP8 Operators

GenAI Support and Operators

  • [New] Integrated Gen AI ops into the build (#2512)
  • [New] Support for Triton-based operators (#2570, #2618)
  • [New] Support for CUTLASS-based operators (#2552, #2537)
  • [New] Car and All-To-All (NCCL-based) communication ops (#2606, #2667, #2631, #2624)
  • [New] Grouped query attention ops (#2673, #2504)
  • [New] CK BF16 GEMM (#2617)
  • [New] W4A8 GEMM kernels (#2558, #2607)

Pooled Embeddings

  • [Improvement] Clean up unused pooled embedding ops (#2626)
  • [Improvement] PyTorch compatibility fixes (#2619, #2629)

Sparse Operators

  • [Improvement] Increased dynamic shared memory size to support larger bucket sizes (#2500)
  • [Improvement] UINT8 support for reorder sequence embedding operator (#2531)
  • [Improvement] Fixed CPU blocking D2H in JaggedIndexSelect2dOp backward (#2510)

Benchmarks / Tests

  • [New] Unified benchmarks and unit tests for FP8 (#2609, #2699, #2666)
  • [Improvement] SSD TBE benchmarks (#2579, #2580)
  • [Improvement] SSD TBE tests (#2665, #2647)
  • [Improvement] Fixes for TBE tests and benchmarks (#2632)
  • [Improvement] nbit_cache benchmark bandwidth calculation (#2511)

Build / CI improvements and Fixes

  • [New] Support for CUDA 12.4 (#2565)
  • [Improvement] Improved AMD support (#2541, #2679)
  • [Improvement] Strengthened artifact installation process (#2491)
  • [Improvement] Memcheck added across operators (#2576, #2574, #2572, #2612, #2594, #2589, #2578)
  • [Improvement] Refactoring of large header files (#2650)
  • [Improvement] Improved build scripts to support debug flags and custom (i.e. GenAI) variants (#2702)
Assets 2
Loading

FBGEMM_GPU v0.7.0

26 Apr 17:53
@spcyppt spcyppt
v0.7.0
8c06a63
Compare
Choose a tag to compare
FBGEMM_GPU v0.7.0

Release Note

Highlights

  • New optimizer and output type supports for Table Batched Embedding (TBE) training
  • Improvement and bug fixes for TBE variable batch size
  • Enhanced TBE pipeline prefetching for UVM caching
  • Many improvements on TBE CPU kernels
  • New and enhanced low-precision operators
  • Code refactoring and reorganization for faster builds
  • New tests and benchmarks
  • PyTorch 2 support for various operators
  • Clang compilation support

Software Requirements

FBGEMM_GPU v0.6.0 has been tested and known to work on the following setups:

  • PyTorch: v2.3
  • CUDA: v11.8, 12.1
  • Python: v3.8, 3.9, 3.10, 3.11, 3.12

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 (only CUDA 12.1 variant is available)
pip install fbgemm-gpu==0.7.0

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

Alternatively, it can be fetched from PyTorch PIP:

# FBGEMM_GPU CUDA variant
pip install fbgemm-gpu==0.7.0 --index-url https://download.pytorch.org/whl/cu118/
pip install fbgemm-gpu==0.7.0 --index-url https://download.pytorch.org/whl/cu121/

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

Changes

Table batched embedding (TBE) operators

  • [New] Added BF16 output support in TBE training (#2382)
  • [New] Added Support int8 output for sequence embeddings (#2316)
  • [New] Added an auto-vectorization implementation for CPU TBE-NBit kernel with user selection (#2182, #2299)
  • [New] Added CowClip optimizer (#2226, #2243)
  • [Improvement] Extended support and bug fixes for variable batch size TBE (#2256, #2388, #2394, #2333)
  • [Improvement] Optimized cache fetch for forward split (#2216, #2282, #2289, #2262, #2218)
  • [Improvement] Caching and cache lookup for pipeline prefetching fixes and enhancements (#2164, #2309, #2287, #2308)
  • [Improvement] Built hip rules by default (#2380)
  • [New] Added a method to TBE module to recompute buffers (#2338)
  • [New] Added meta functions for PyTorch 2 support (#2347)
  • [New] Added support for MTIA in TBE modules (#2273, #2286)
  • [Improvement] Improved TBE logging and stats report (#2379, #2378, #2377, #2386, #2337)
  • [Improvement] General fixes and enhancements (#2235, #2398, #2212, #2269, #1782, #2270, #2265, #2385, #2370, #2349, #2312, #2411, #2400)
  • [Deprecation] Optimizers deprecated (#2253, #2252)
  • [Deprecation] Removed double type support from fbgemm_cuda_utils.cuh (#2335)
  • [Deprecation] Removed INT8 weight/output support from TBE GPU training

Jagged Tensor Operators

  • [Improvement] Removed device-host synchronization from keyed jagged index select (#2315)
  • [Improvement] Fixed half->int build error (#2240)

Index Select Operators

  • [Improvement] Fixed BF16 group_index_select_2d on AMD GPU (#2321)

Low-precision operators

  • [New] CPU implementation of per-channel quantize operator (#2341)
  • [New] CPU implementation for qlinear_channelwise operator (#2343)
  • [New] Enabled CPU int8 output to dequantization to bf16 on CUDA (#2242)
  • [New] Enabled dequantization for bf16 (#2241)

Pooled Embedding

  • [Improvement] Used gpu_library_selector for permute_pooled_embedding_ops_gpu (#2340)

Misc

  • [New] Implementation of CPU version of all_to_one_device (#2251)
  • [Improvement] Performance improvement of _block_bucketize_sparse_features_cuda_kernel1 (#2331)
  • [New] Created cumem_utils_cpu and added to all_deps_cpu (#2215)
  • [New] Added float support to asynchronous_complete_cumsum_cpu (#2383)
  • [Improvement] Added early exit to sparse ops (#2277, #2276, #2213, #2259)
  • [New] STBE GPU coalescing kernel (#2275)
  • [Improvement] Removed symint from tbe_input_combine_with_length_abstract (#2336)
  • [New] GPU timing and basic reporting framework (#2314)
  • [Improvement] Fixes and FBGEMM PT2 compliance (#2223, #2224, #2225, #2231, #2327)

Benchmarks / Tests

Build / CI improvements and Fixes

Assets 2
Loading

FBGEMM_GPU v0.6.0

31 Jan 19:40
@spcyppt spcyppt
v0.6.0
e0d208e
Compare
Choose a tag to compare
FBGEMM_GPU v0.6.0

Release Note

Highlights

  • Improvement and bug fixes for TBE variable batch size
  • Many TBE extensions and benchmarks
  • Enhanced TBE pipeline prefetching for UVM caching
  • Code refactoring and reorganization for faster builds
  • Many improvements and new sparse ops added
  • Improved low precision ops
  • Support for Python 3.12
  • PyTorch 2 support for various operators

Software Requirements

FBGEMM_GPU v0.6.0 has been tested and known to work on the following setups:

  • PyTorch: v2.2
  • CUDA: v11.8, 12.1
  • Python: v3.8, 3.9, 3.10, 3.11, 3.12

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 (only CUDA 12.1 variant is available)
pip install fbgemm-gpu==0.6.0

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

Alternatively, it can be fetched from PyTorch PIP:

# FBGEMM_GPU CUDA variant
pip install fbgemm-gpu==0.6.0 --index-url https://download.pytorch.org/whl/cu118/
pip install fbgemm-gpu==0.6.0 --index-url https://download.pytorch.org/whl/cu121/

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

Changes

Table batched embedding (TBE) operators

  • [Improvement] Extended support and bug fixes for variable batch size (#2012, #2043, #2107, #2150, #2188)
  • [Improvement] caching and cache lookup for pipeline prefetching (#2147, #2154, #2151)
  • [New] Support MTIA device type in FBGEMM TBE training (#1994)
  • [New] Enable sequence TBE CPU via AVX (#2195)
  • [New] Enable subwarp only for unweighted (#2051)
  • [New] Add meta functions (#2094, #2102)
  • [New] Add reverse qparam option for MTIA (#2109)
  • [New] uvm_cache_stats for direct mapped (#1951, #1952)
  • [Improvement] use memcpy for cpu emb inplace update (#2166)
  • [Improvement] Remove indices and offsets copying from prefetch (#2186)
  • [Improvement] Improve perf for L=0 cases for TBE v2 (#2046)
  • [Improvement] General fixes and enhancements (#2030, #2009)

Jagged Tensor Operators

  • [Improvement] Fix incorrect SymInt signature on dense_to_jagged (#2039)
  • [Improvement] Fix non-contiguous tensor problem in jagged_index_select (#2060, #2061)

Index Select Operators

  • [Improvement] Get total D from CPU buffer in batch_index_select_dim0 (#2079)

Low-precision operators

  • [New] Add BF16 in padded FP8 quantize ops (#2010)
  • [Improvement] Improve quantize_comm error message (#2018)
  • [Improvement] Fix illegal memory access error and initialize empty values on fp8 quantize kernel (#2131, #2176)

Pooled Embedding

  • [New] Add permute_duplicate_pooled_embeddings op for CPU (#1939)
  • [Improvement] Use PyTorch's p2p access enable function (#2000)
  • [New] Add support for duplicate in permutations for permute_pooled_embs_split (#1940)
  • [Improvement] Improve all_to_one error message (#2019)
  • [New] Add meta function for fbgemm::merge_pooled_embeddings operator (#2069)
  • [New] Add variable batch per feature support to EBC (tw/cw only) (#1986)

Misc

Benchmarks / Tests

  • [New] Benchmark block_bucketize_sparse_features uneven sharding (#2140, #2169)
  • [New] Add unit test for unique cache lookup (#2160)
  • [New] Add autogenerated opcheck tests (#2050, #2069, #2073, #2092, #2118, #2139, #2152, #2173, #2193)
  • [New] Add test for fbgemm ops. (#2136, #2082)
  • [Improvement] Modified TBE testbench to use FBGEMM generate_rquests function to generate indices and offsets (#1882)
  • [Improvement] Remove FP64 from TBE CPU tests (#2049)
  • [Improvement] Add warmup_runs to TBE benchmarks and run at least 1 warmup iter #2163
  • [Improvement] Add --pooling in TBE nbit_cpu benchmark (#2200)
  • [Improvement] Fill embedding tables with randomized scales and bias in split-TBE benchmarks (#2031)

Build / CI improvements and Fixes

Assets 2
Loading

FBGEMM_GPU v0.5.0

05 Oct 23:52
@spcyppt spcyppt
v0.5.0
a4151dd
Compare
Choose a tag to compare
FBGEMM_GPU v0.5.0

Release Notes

Highlights

  • TBE training v2 (optimized TBE forward: up to 4x kernel performance improvement)
  • Many TBE extensions including defused TBE backward-optimizer, variable batch size support, pipeline prefetching support for UVM caching
  • Many improvements and new sparse ops added
  • ARM support
  • SM 9.0 support for CUDA 12.1 for H100 GPUs
  • PyTorch 2 support for various operators, i.e., jagged tensor, pooled embedding ops

Software Requirements

FBGEMM_GPU v0.5.0 has been tested and known to work on the following setups:

  • PyTorch: v2.1
  • CUDA: v11.8, 12.1
  • Python: v3.8, 3.9, 3.10, 3.11

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 (only CUDA 12.1 variant is available)
pip install fbgemm-gpu==0.5.0

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

Alternatively, it can be fetched from PyTorch PIP:

# FBGEMM_GPU CUDA variant
pip install fbgemm-gpu==0.5.0 --index-url https://download.pytorch.org/whl/cu118/
pip install fbgemm-gpu==0.5.0 --index-url https://download.pytorch.org/whl/cu121/

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

Changes

Table batched embedding (TBE) operators

Jagged Tensor Operators

Index Select Operators

  • [New] batch_index_select_dim0 with TBE backend (#1897)
  • [New] Variable input sizes support for group_index_select_dim0 (#1968)
  • [Improvement] Improve group_index_select(#1764, #1884)

Low-precision operators

  • [New] Meta Backend FP8RowwiseQuantizedToFloat (#1890)
  • [New] Column-wise parallel quantization/dequantization (#1743)
  • [New] BF16 Support in FP8 quantize ops (#1961)
  • [Improvement] FP8 row-wise quantization optimization/improvement (#1729, #1858, #1981, #1909)

Pooled Embedding

  • [New] reduce_to_one (#1571)
  • [New] permute_duplicate_pooled_embeddings op (#1912)
  • [New] BF16 support for permute_pooled_embeddings op 1937
  • [New] Variable size input-output support for permute_pooled_embs_kernel (#1913)
  • [New] Backends (Meta) (#1853)
  • [Improvement] multi-gpu all_to_one enhancements (#1674, #1962)

Misc

  • [New] CUB kernel for 2D asynchronous_complete_cumsum (#1707)
  • [New] Backends (Meta) (#1709, #1905, #1970, #1971)
  • [New] BF16 support in permute_indices_weights_kernel_2 (#1852)
  • [New] FP16 and BF16 support in pack_segments (#1708)
  • [New] BF16 support for HBC ops. (#1744)
  • [New] BFloat16 support (#1832, #1865)
  • [Improvement] Speedup reorder_batched_ad_indices (#1901, #1902, #1932, #1933, 1711)

Benchmarks / Tests

  • [New] CLI support to GEMMsBenchmark (#1721, #1725)
  • [New] Benchmark for variable batch on TBE (#1559)
  • [New] BF16 output test coverage (#1835, #1838)
  • [New] Benchmark for reorder_batched_ad_indices (#1895)
  • [New] CPU support (#1874, #1926)
  • [Improvement] GroupIndexSelect Benchmark with zero_grad (#1559)
  • [Improvement] Add nbit-cpu-with-spec benchmark in FBGEMM-GPU's TBE benchmark suite (#1892)

Build / CI improvements and Fixes

Assets 2
Loading

FBGEMM_GPU v0.4.1

24 Mar 23:37
@q10 q10
v0.4.1
64833b5
Compare
Choose a tag to compare
FBGEMM_GPU v0.4.1

Release Notes

Software Requirements

FBGEMM_GPU v0.4.1 has been tested and known to work on the following setups:

  • PyTorch: v2.0
  • CUDA: v11.7, 11.8
  • Python: v3.8, 3.9, 3.10, 3.11

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

Availability

FBGEMM_GPU may be fetched directly from PyPI:

# FBGEMM_GPU (CUDA variant)
pip install fbgemm-gpu==0.4.1

# FBGEMM_GPU (CPU variant)
pip install fbgemm-gpu-cpu==0.4.1

Changes

This is a minor release whose main purpose is to deliver Python 3.11 support.

  • [New] Add support for Python 3.11 (#1646)
  • [Improvement] Add support for group size > 54 in group_index_select (#1611)
  • [Improvement] Implement cache miss emulation in UVM_CACHING (#1637)))
  • [Improvement] Add TensorAccessor with memcheck (#1602)
Assets 2
Loading

FBGEMM_GPU v0.4.0

15 Mar 17:08
@q10 q10
v0.4.0
ea96ea3
Compare
Choose a tag to compare
FBGEMM_GPU v0.4.0

Release Notes

Software Requirements

FBGEMM_GPU v0.4.0 has been tested and known to work on the following setups:

  • PyTorch: v2.0
  • CUDA: v11.7, 11.8
  • Python: v3.8, 3.9, 3.10 (3.11 not supported yet)

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

Availability

FBGEMM_GPU may be fetched directly from PyPI:

# FBGEMM_GPU (CUDA variant)
pip install fbgemm-gpu==0.4.0

# FBGEMM_GPU (CPU variant)
pip install fbgemm-gpu-cpu==0.4.0

Changes

Table batched embedding (TBE) operators

UVM cache improvement

  • [New] Delta in-place update (#1436)
  • [New] UVM caching stats report (#1623, #1462, #1433, #1623, #1570)
  • [Improvement] [lfu|lru]_cache_insert_byte_kernel vectorization (#1475)

Jagged Tensor Operators

Index Select Operators

  • [New] group_index_select (#1421, #1592)
  • [New] index_select for selecting KeyJaggedTensor dim 1 (previously support only dim 0) (#1429)
  • [New] jagged_index_select for CPU (#1586)

Low-precision operators

  • [New] FP8 rowwise quantized communication (#1423)

Misc

  • Support 2D inputs for asynchronous_complete_cumsum (#1573)

Benchmarks / Tests

  • [New] nbit_device_with_spec for table batched embedding inference benchmark (#1455, #1465)
  • [New] Variable bag sizes for TBE benchmark (#1450)
  • [Improvement] Parallel bottom_unique_k_per_row for faster Zipf data generation (for FBGEMM benchmarks) (#1447)

Build / CI improvements and Fixes

Assets 2
Loading

v0.3.2

19 Jan 22:22
@mjanderson09 mjanderson09
v0.3.2
b2be702
Compare
Choose a tag to compare
v0.3.2

Minor release

Assets 2
Loading

v0.3.0

28 Oct 21:29
@mjanderson09 mjanderson09
v0.3.0
9c4fa58
Compare
Choose a tag to compare
v0.3.0

New Features

Table Batched Embedding enhancements:

AMD Support (beta) (#1102, #1193)

  • FBGEMM previously supported only NVIDIA accelerators, but FBGEMM 0.3.0 started to support AMD GPUs in collaboration with AMD. Although its support is still beta (e.g., we don't have a stable release build for AMD GPUs yet), the AMD GPU implementation covers almost all the FBGEMM operators supported by NVIDIA GPUs. AMD GPU support is tested using CI with AMD MI250 GPUs.

Quantized Communication Primitives (#1219, #1337)

Sparse kernel enhancements

  • New kernel: invert_permute (#1403)
  • New kernel: truncate_jagged_1d (#1345)
  • New kernel: jagged_index_select (#1157)
  • Jagged Tensor optimization for inference use cases (#1236)

Improved documentation for Jagged Tensors and SplitTableBatchedEmbeddingBagsCodegen

Optimized 2x2 kernel for AVX2 (#1280)

Full Changelog: https://github.com/pytorch/FBGEMM/commits/v0.3.0

Assets 2
Loading