3rdparty: upgrade cutlass to 3.9 (#997)

Update cutlass to latest version (v3.9) for incoming pull requests on
blackwell support.

Hopper workloads experienced slight performance degradation, will
investigate them later.

Author: yzh119

3rdparty/cutlass

@@ -158,7 +158,7 @@ CUTE_HOST_DEVICE constexpr auto upcast(Shape const& shape, Stride const& stride)
                             [](auto const& s, auto const& d) { return upcast<N, I>(s, d); });
   } else if constexpr (is_scaled_basis<Stride>::value) {
     if constexpr (Stride::mode() == I) {
-      return make_layout(shape_div(shape, Int<N>{}), shape_div(stride, Int<N>{}));
+      return make_layout(ceil_div(shape, Int<N>{}), ceil_div(stride, Int<N>{}));
     } else {
       return make_layout(shape, stride);
     }

include/flashinfer/attention/hopper/block_sparse_gather.cuh

@@ -656,7 +656,7 @@ __global__ __launch_bounds__(KTraits::NUM_THREADS) void BatchMLAPageAttentionHop
   pipeline_params.producer_arv_count = 128;
   pipeline_params.consumer_arv_count = 128;
   MainloopPipeline pipeline_q(smem_storage.pipeline_q, pipeline_params);
-  pipeline_params.role = warp_group_idx == 0 ? MainloopPipeline::ThreadCategory::Producer
+  pipeline_params.role = warp_group_idx == 0 ? MainloopPipeline::ThreadCategory::ProducerConsumer
                                              : MainloopPipeline::ThreadCategory::Consumer;
   pipeline_params.producer_arv_count = 128;
   pipeline_params.consumer_arv_count = 256;