Fix ZeroDivisionError in fused linear cross entropy when weight vocab dim is 0

[nileshpatil6]

Fixes #767

Bug

fused_linear_cross_entropy_forward computes:

V = weight.shape[0]
inc_factor = triton.cdiv(V, H)
chunk_size = triton.next_power_of_2(triton.cdiv(BT, inc_factor))

If weight has a vocab dimension of 0, inc_factor becomes 0 and the next triton.cdiv call divides by zero, raising a bare ZeroDivisionError from deep inside triton with no indication of what actually went wrong.

Root cause

This happens when weight is not actually materialized yet when it reaches the kernel. In the reported case, the user was training under DeepSpeed ZeRO-3 with parameter offload, and model.lm_head.weight was an empty tensor at the point the Liger loss function accessed it directly (the raw parameter, not through the module forward that normally triggers the all-gather hook). Same class of bug would show up for any other reason weight ends up empty.

Fix

Added an assertion right after V is computed that fails with a clear message pointing at the likely cause (unmaterialized weight, e.g. ZeRO-3 parameter accessed before it's gathered), instead of letting execution fall through to a cryptic ZeroDivisionError three calls deep in triton internals.

Testing

Added test_empty_weight_raises_clear_error in test/transformers/test_fused_linear_cross_entropy.py. It calls fused_linear_cross_entropy_forward directly with a weight tensor of shape (0, H) and checks the new assertion fires with the expected message.

I confirmed this reproduces the exact bug from the issue: reverting just the source change and running the test gives:

src\liger_kernel\ops\fused_linear_cross_entropy.py:60: in fused_linear_cross_entropy_forward
    chunk_size = triton.next_power_of_2(triton.cdiv(BT, inc_factor))
...
ZeroDivisionError: integer division or modulo by zero

With the fix applied:

test/transformers/test_fused_linear_cross_entropy.py::test_empty_weight_raises_clear_error PASSED

Also ran ruff check and ruff format --check on both changed files, both clean. Full test file collection (138 tests) still succeeds with no import errors.

Note: my dev box doesn't have the CUDA toolkit/MSVC set up for compiling triton kernels, so I could not run the full GPU correctness/convergence suite. The new test and the guard itself run entirely on the host before any kernel is launched, so this was verified directly.