updated sl to be non-constexpr (#149)

## Summary
<!--- This is a required section; please describe the main purpose of
this proposed code change. --->
updated seqlen for rope to be non-constexpr

<!---
## Details
This is an optional section; is there anything specific that reviewers
should be aware of?
--->
sl from constexpr to non-constexpr

## 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: RTX 3090
- [x] run `make test` to ensure correctness
- [x] run `make checkstyle` to ensure code style
- [x] run `make test-convergence` to ensure convergence

benchmark/rope_memory/rope-full-memory-benchmark-seq-2048.csv

@@ -1,4 +1,4 @@
 total_hidden_size,Liger,Hugging Face
-512.0,13.5,22.5
-2048.0,29.6,65.6
-8192.0,94.1,238.1
+512.000000,13.266113,22.266113
+2048.000000,28.641113,64.641113
+8192.000000,90.141113,234.141113

benchmark/rope_memory/rope-full-memory-benchmark-total_dim_8192.csv

@@ -1,6 +1,6 @@
 seq_len,Liger,Hugging Face
-1024.0,53.1,125.1
-2048.0,94.1,238.1
-4096.0,176.2,464.2
-8192.0,340.2,916.2
-16384.0,668.3,1820.3
+1024.000000,49.133301,121.133301
+2048.000000,90.141113,234.141113
+4096.000000,172.156738,460.156738
+8192.000000,336.187988,912.187988
+16384.000000,664.250488,1816.250488

benchmark/rope_speed/rope-backward-speed-benchmark-seq-2048.csv

@@ -1,4 +1,4 @@
 total_hidden_size,Liger,Hugging Face
-512.0,0.1,0.2
-2048.0,0.2,0.2
-8192.0,0.2,0.6
+512.000000,0.050176,0.121824
+2048.000000,0.048128,0.225280
+8192.000000,0.101376,0.801792

benchmark/rope_speed/rope-backward-speed-benchmark-total_dim_8192.csv

@@ -1,6 +1,6 @@
 seq_len,Liger,Hugging Face
-1024.0,0.2,0.3
-2048.0,0.1,0.6
-4096.0,0.2,1.2
-8192.0,0.2,2.3
-16384.0,0.4,4.5
+1024.000000,0.052224,0.416768
+2048.000000,0.101376,0.801984
+4096.000000,0.199680,1.553408
+8192.000000,0.396288,3.057664
+16384.000000,0.789504,6.062080

benchmark/rope_speed/rope-forward-speed-benchmark-seq-2048.csv

@@ -1,4 +1,4 @@
 total_hidden_size,Liger,Hugging Face
-512.0,0.1,0.1
-2048.0,0.1,0.2
-8192.0,0.1,0.5
+512.000000,0.010240,0.059392
+2048.000000,0.028672,0.168016
+8192.000000,0.101376,0.600064

benchmark/rope_speed/rope-forward-speed-benchmark-total_dim_8192.csv

@@ -1,6 +1,6 @@
 seq_len,Liger,Hugging Face
-1024.0,0.1,0.3
-2048.0,0.1,0.5
-4096.0,0.1,1.0
-8192.0,0.2,1.9
-16384.0,0.5,3.8
+1024.000000,0.052224,0.311296
+2048.000000,0.101376,0.600064
+4096.000000,0.199680,1.244160
+8192.000000,0.396288,2.484224
+16384.000000,0.789504,4.975504

benchmark/rope_speed/rope-full-speed-benchmark-seq-2048.csv

@@ -1,4 +1,4 @@
 total_hidden_size,Liger,Hugging Face
-512.0,0.3,0.4
-2048.0,0.3,0.4
-8192.0,0.3,1.1
+512.000000,0.059328,0.784384
+2048.000000,0.058880,0.784384
+8192.000000,0.201728,1.404928

benchmark/rope_speed/rope-full-speed-benchmark-total_dim_8192.csv

@@ -1,6 +1,6 @@
 seq_len,Liger,Hugging Face
-1024.0,0.6,0.6
-2048.0,0.6,1.1
-4096.0,0.6,2.2
-8192.0,0.6,4.2
-16384.0,0.9,8.3
+1024.000000,0.103424,0.789504
+2048.000000,0.201728,1.403904
+4096.000000,0.398336,2.801520
+8192.000000,0.791552,5.547008
+16384.000000,1.577984,11.062272

src/liger_kernel/ops/rope.py

@@ -13,8 +13,8 @@ def _triton_rope(
     cos_row_stride,
     sin,
     sin_row_stride,
+    sl,
     bs: tl.constexpr,
-    sl: tl.constexpr,
     n_qh: tl.constexpr,
     n_kh: tl.constexpr,
     hd: tl.constexpr,
@@ -168,8 +168,8 @@ def forward(ctx, q, k, cos, sin, position_ids=None, unsqueeze_dim=1):
             cos.stride(-2),
             sin,
             sin.stride(-2),
-            batch_size,
             seq_len,
+            batch_size,
             n_q_head,
             n_kv_head,
             head_dim,
@@ -219,8 +219,8 @@ def backward(ctx, dq, dk):
             cos.stride(-2),
             sin,
             sin.stride(-2),
-            batch_size,
             seq_len,
+            batch_size,
             n_q_head,
             n_kv_head,
             head_dim,