[decode] [kernel] for GQA mode, why each thead compute s for bdy

[hanzz2007]

state_t<vec_size> st;
float s[bdy * tile_size_per_bdx]; // thread compute s tile [bdy*tile_size_per_bdx]

bdy here means GROUP_NUM, but the cta block dim is (bdx, bdy, bdz), seems each thread only need to compute tile of s float s[tile_size_per_idx] .

[hanzz2007]

```cpp

update_local_state<vec_size, bdx, bdy * tile_size_per_bdx>(
v_smem + (stage_idx * bdz + tz) * bdy * tile_size_per_bdx * head_dim, s, stage_idx, st, tx);

```cpp
template <uint32_t vec_size, uint32_t bdx, uint32_t tile_size, typename T>
__device__ __forceinline__ void update_local_state(const T* smem, const float* s,
                                                  uint32_t compute_stage_idx,
                                                  state_t<vec_size>& st, uint32_t tx) {
#pragma unroll
 for (uint32_t j = 0; j < tile_size; ++j) {
   vec_t<float, vec_size> v_vec;
   v_vec.cast_load(smem + (j * bdx + tx) * vec_size);
#pragma unroll
   for (uint32_t i = 0; i < vec_size; ++i) {
     st.o[i] = st.o[i] + s[j] * v_vec[i];
   }
 }
}

in update_local_state , thread loop over s bdy * tile_size_per_bdx tile to compute st.o which is group wise, but the bdy means group.

I'm quite confused, more explanation will be helpful!