bugfix: Fix inline RoPE in decode kernels (#847)

This PR fixes a bug in decode kernel that fails to process the
`q_rope_offset` provided in Params. The variable naming follows the
existing example in prefill kernel:


https://github.com/flashinfer-ai/flashinfer/blob/ab6484effafd5133514890999ff23129ce3b8a9b/include/flashinfer/attention/prefill.cuh#L39

Author: MasterJH5574

include/flashinfer/attention/decode.cuh

@@ -33,6 +33,8 @@
 
 namespace flashinfer {
 
+DEFINE_HAS_MEMBER(maybe_q_rope_offset)
+
 namespace cg = cooperative_groups;
 using cp_async::PrefetchMode;
 using cp_async::SharedMemFillMode;
@@ -438,7 +440,10 @@ __global__ void BatchDecodeWithPagedKVCacheKernel(const __grid_constant__ Params
   const uint32_t q_stride_n = params.q_stride_n;
   const uint32_t q_stride_h = params.q_stride_h;
   if constexpr (POS_ENCODING_MODE == PosEncodingMode::kRoPELlama) {
-    const IdType* q_rope_offset = nullptr;  // params.q_rope_offset;
+    const IdType* q_rope_offset = nullptr;
+    if constexpr (has_maybe_q_rope_offset_v<Params>) {
+      q_rope_offset = params.maybe_q_rope_offset;
+    }
     int32_t q_rope_offset_val = q_rope_offset == nullptr ? (kv_len - 1) : q_rope_offset[batch_idx];
     const float rope_rcp_scale = params.rope_rcp_scale;
     const float rope_rcp_theta = params.rope_rcp_theta;