Skip to content
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

Jump to bottom

[fused_rmsnorm] Avoid querying device inside forward #301

Closed
wants to merge 1 commit into from
Closed

[fused_rmsnorm] Avoid querying device inside forward #301

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.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
9 changes: 8 additions & 1 deletion torchtitan/models/norms.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -213,6 +213,13 @@ def _rms_norm_bwd_kernel_sm(
tl.store(DW + row_block_id * N + cols, dw, mask=mask)


def get_sm_count():
# TODO(whc) initializing sm_count here prevents the need to query device during tracing using meta-device.
# But it also forces us to pick a favorite device.
# possibly, we could build a dict of device_id: sm_count here, and then index it during forward instead.
return torch.cuda.get_device_properties(0).multi_processor_count


class TritonFusedRMSNorm(torch.autograd.Function):
@staticmethod
def forward(ctx, x, weight, eps):
Expand Down Expand Up @@ -268,7 +275,7 @@ def backward(ctx, dy):
dx = torch.empty_like(x)
dw = torch.empty_like(weight)

sm_count = torch.cuda.get_device_properties(x.device).multi_processor_count
sm_count = get_sm_count()
_dw = torch.empty((sm_count, N), dtype=torch.float32, device=weight.device)

max_size = 65536 // x.element_size()
Expand Down