different results between --cudagraph and do_bench_cudagraph

[BoyuanFeng]

--cudagraph runs triton.testing.do_bench_cudagraph. So these two should give the same latency. However, this is not true.

To repro, please check out #345 which updates the input settings.

python3 run.py --op rope --cudagraph gives 0.0020 for liger and 0.0027 for inductor.

HOWEVER, python repro.py gives 0.0017 for lilger and 0.0014 for inductor. repro.py copies code from tritonbench/operators/rope/operator.py and uses triton.testing.do_bench_cudagraph for benchmark.

[BoyuanFeng]

measurement of rope is wrong.. need to use this config:

torch._dynamo.config.automatic_dynamic_shapes = False
torch._dynamo.config.recompile_limit = 1000000

[xuzhao9]

I can reproduce the results of repro.py:

$ python repro.py
`LlamaRotaryEmbedding` can now be fully parameterized by passing the model config through the `config` argument. All other arguments will be removed in v4.46
compiler_latency: 0.0016405560234759716 liger_latency: 0.0018557837686023195

For TritonBench, I could not reproduce

$ python3 run.py --op rope --num-inputs 1 --input-id 20 --cudagraph --precision bf16
  (B, S)    apply_rotary_pos_emb-latency    liger_rotary_pos_emb-latency    inductor_rotary_pos_emb_full_op-latency
--------  ------------------------------  ------------------------------  -----------------------------------------
  (2, 1)               0.023093 (±0.03%)               0.001851 (±0.11%)                          0.001620 (±0.11%)
 average

The latency is consistent between Tritonbench and repro.py. Can I have the full command line you used to run Tritonbench?

When running multiple shapes:

     (B, S)    apply_rotary_pos_emb-latency    liger_rotary_pos_emb-latency    inductor_rotary_pos_emb_full_op-latency
-----------  ------------------------------  ------------------------------  -----------------------------------------
     (1, 1)               0.022456 (±0.04%)               0.001833 (±0.05%)                          0.001542 (±0.29%)
     (1, 2)               0.022948 (±0.05%)               0.006660 (±0.02%)                          0.001685 (±0.13%)
     (1, 4)               0.024544 (±0.04%)               0.006749 (±0.02%)                          0.001736 (±0.05%)
     (1, 8)               0.024819 (±0.05%)               0.006962 (±0.05%)                          0.001797 (±0.12%)
    (1, 16)               0.025044 (±0.09%)               0.007128 (±0.06%)                          0.001986 (±0.09%)
    (1, 32)               0.026207 (±0.09%)               0.007208 (±0.06%)                          0.002326 (±0.10%)
    (1, 64)               0.027679 (±0.13%)               0.007585 (±0.12%)                          0.002951 (±0.12%)
   (1, 128)               0.031816 (±0.16%)               0.008473 (±0.03%)                          0.004182 (±0.05%)
   (1, 256)               0.042510 (±0.06%)               0.010767 (±0.06%)                          0.006655 (±0.22%)
   (1, 512)               0.063019 (±0.04%)               0.016472 (±0.08%)                          0.011560 (±0.17%)
  (1, 1024)               0.108383 (±0.18%)               0.028512 (±0.06%)                          0.021545 (±0.25%)
  (1, 2048)               0.213070 (±0.16%)               0.067010 (±0.03%)                          0.041880 (±0.22%)
  (1, 4096)               0.466960 (±0.33%)               0.157722 (±0.03%)                          0.082078 (±0.34%)
  (1, 8192)               0.918918 (±2.49%)               0.309500 (±0.02%)                         0.162720 (±16.59%)
 (1, 16384)               1.800814 (±0.29%)               0.611743 (±0.02%)                          0.333553 (±8.23%)
 (1, 32768)               3.586293 (±0.68%)               1.222533 (±0.02%)                          0.790581 (±0.28%)
 (1, 65536)               7.593480 (±0.49%)               2.440178 (±0.02%)                          1.667460 (±0.08%)
(1, 131072)              15.313014 (±0.26%)               4.880147 (±0.01%)                          3.351184 (±0.14%)
(1, 262144)              30.562261 (±0.12%)               9.781654 (±0.04%)                          6.689198 (±0.14%)
     (1, 1)               0.023295 (±0.42%)               0.001835 (±0.06%)                          0.001543 (±0.49%)
     (2, 1)               0.023779 (±0.05%)               0.001843 (±0.04%)                          0.003034 (±0.07%)
     (4, 1)               0.024181 (±0.05%)               0.001851 (±0.08%)                          0.003065 (±0.09%)
     (8, 1)               0.024874 (±0.03%)               0.001879 (±0.06%)                          0.003180 (±0.09%)
    (16, 1)               0.025397 (±0.05%)               0.001961 (±0.10%)                          0.003224 (±0.12%)
    (32, 1)               0.026416 (±0.05%)               0.002058 (±0.10%)                          0.003564 (±0.09%)
    (64, 1)               0.027950 (±0.11%)               0.002292 (±0.08%)                          0.004202 (±0.07%)
   (128, 1)               0.032210 (±0.05%)               0.002522 (±0.07%)                          0.005465 (±0.05%)
   (256, 1)               0.042407 (±0.02%)               0.003173 (±0.26%)                          0.007926 (±0.03%)
   (512, 1)               0.061731 (±0.22%)               0.005081 (±0.14%)                          0.012843 (±0.07%)
  (1024, 1)               0.107613 (±0.19%)               0.007379 (±0.14%)                          0.023083 (±0.36%)
    average

The Liger kernel perf makes sense, but inductor perf is abnormal.

[BoyuanFeng]

The dynamic shape + inductor issue only happens when we run on multiple shapes. Could you please try python3 run.py --op rope --cudagraph without specifying any shapes?