Skip to content

Commit 4f3fcca

Browse files
stashuk-olekfacebook-github-bot
authored andcommitted
Reduce Triton TBE unweighted backward reduction width (meta-pytorch#4488)
Summary: We could drop that, small tuning on buffer_size. Reviewed By: axeisghost Differential Revision: D114403703
1 parent 04fdfd2 commit 4f3fcca

2 files changed

Lines changed: 3 additions & 3 deletions

File tree

torchrec/distributed/triton_tbe/triton_table_batched_embeddings.py

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -541,7 +541,7 @@ def triton_tbe_backward_short_run_unweighted(
541541
) -> None:
542542
"""Backward kernel for short runs only. Each program handles one short run."""
543543
col_offsets = tl.arange(0, BLOCK_SIZE)
544-
buffer_size: tl.constexpr = 16
544+
buffer_size: tl.constexpr = 8
545545
buffer_offsets = tl.arange(0, buffer_size)
546546

547547
if USE_CLC:
@@ -1037,7 +1037,7 @@ def triton_tbe_backward_long_run_grad_accum_unweighted(
10371037
and atomically adds the partial result into a temp gradient buffer.
10381038
"""
10391039
col_offsets = tl.arange(0, BLOCK_SIZE)
1040-
buffer_size: tl.constexpr = 16
1040+
buffer_size: tl.constexpr = 8
10411041
buffer_offsets = tl.arange(0, buffer_size)
10421042

10431043
pid = tl.program_id(0)

torchrec/distributed/triton_tbe/triton_tbe_backward_long_run_fused.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -62,7 +62,7 @@ def triton_tbe_backward_long_run_fused_weighted(
6262
applies the optimizer update — eliminating a separate apply kernel launch.
6363
"""
6464
col_offsets = tl.arange(0, BLOCK_SIZE)
65-
buffer_size: tl.constexpr = 16
65+
buffer_size: tl.constexpr = 8
6666
buffer_offsets = tl.arange(0, buffer_size)
6767

6868
clc_phase_producer = 1

0 commit comments

Comments
 (0)