Fix illegal memory access in Triton RMSNorm and RoPE (#804)

## Summary
<!--- This is a required section; please describe the main purpose of
this proposed code change. --->
When using very large tensors (e.g. seq_len=1e6, hidden_size=4096),
Triton’s default 32-bit `tl.program_id(0)` can overflow, leading to
out-of-bounds memory accesses. This change casts the program ID to
64-bit (`tl.int64`) to ensure all pointer arithmetic stays within the
valid address range. Fix #803
<!---
## Details
This is an optional section; is there anything specific that reviewers
should be aware of?
--->

## Testing Done
<!--- This is a required section; please describe how this change was
tested. --->

<!-- 
Replace BLANK with your device type. For example, A100-80G-PCIe

Complete the following tasks before sending your PR, and replace `[ ]`
with
`[x]` to indicate you have done them. 
-->

- Hardware Type: <BLANK>
- [ ] run `make test` to ensure correctness
- [ ] run `make checkstyle` to ensure code style
- [ ] run `make test-convergence` to ensure convergence

---------

Co-authored-by: Shao Tang <[email protected]>

Author: vvvdwbvvv

src/liger_kernel/ops/rms_norm.py

@@ -63,7 +63,7 @@ def _rms_norm_forward_kernel(
     3. https://arxiv.org/pdf/1910.07467
     """
 
-    row_idx = tl.program_id(0)
+    row_idx = tl.program_id(0).to(tl.int64)
     col_offsets = tl.arange(0, BLOCK_SIZE)
     mask = col_offsets < n_cols
 
@@ -137,7 +137,7 @@ def _rms_norm_backward_kernel(
     dw = sum(dy * (x / RMS)). summation over BxT dimension
     """
 
-    row_block_id = tl.program_id(0)
+    row_block_id = tl.program_id(0).to(tl.int64)
     row_start = row_block_id * rows_per_program
     row_end = min((row_block_id + 1) * rows_per_program, n_rows)
     col_offsets = tl.arange(0, BLOCK_SIZE)

src/liger_kernel/ops/rope.py

@@ -32,7 +32,7 @@ def _triton_rope(
 
     # cos size: (1, seq_len, head_dim) or (bsz, seq_len, head_dim)
     # stride: (seq_len * head_dim, head_dim, 1)
-    pid = tl.program_id(0)
+    pid = tl.program_id(0).to(tl.int64)
 
     # locate start address
     q_ptr = q_ptr + pid * q_row_stride