fix(cute): pass mDynamicCausal to kernel and init is_split_kv in SM80/SM120 FA4 forward (#178)

- Pass mDynamicCausal in the FlashAttentionForwardSm80/Sm120 kernel
  launch and add it to the kernel signature (was referenced in the
  kernel body but never passed, causing NameError on first call).
- Initialize self.is_split_kv = False in __init__ (read in four places
  but never assigned, causing AttributeError on first call).

Fixes vllm-project/vllm#51776

Signed-off-by: SayHelloToWorld <SayHelloToWorld@users.noreply.github.com>
Co-authored-by: SayHelloToWorld <SayHelloToWorld@users.noreply.github.com>

Author: SayHelloToWorld

flash_attn/cute/flash_fwd.py

@@ -626,6 +626,7 @@ def __init__(self, *args, **kwargs):
         assert self.output_quant_key is None, (
             f"Fused quant output not implemented for {type(self).__name__}"
         )
+        self.is_split_kv = False
 
     def _get_smem_layout_atom(self):
         sQ_layout_atom = sm80_utils.get_smem_layout_atom(self.dtype, self.tile_hdim)
@@ -793,6 +794,7 @@ def __call__(
             aux_data,
             fastdiv_mods,
             output_scale,
+            mDynamicCausal,
         ).launch(
             grid=grid_dim,
             block=[self.num_threads, 1, 1],
@@ -833,6 +835,7 @@ def kernel(
         aux_data: AuxData = AuxData(),
         fastdiv_mods=None,
         output_scale: Optional[cute.Tensor] = None,
+        mDynamicCausal: Optional[cute.Tensor] = None,
     ):
         # Thread index, block index
         tidx, _, _ = cute.arch.thread_idx()