[QST]41_fused_multi_head_attention on sm89

What is your question?
when I use sm89

int run_attention(Options& options) {
  using Attention = AttentionKernel<
    cutlass::half_t,      // scalar_t
    cutlass::arch::Sm89,  // ArchTag
    true,                 // Memory is aligned
    kQueriesPerBlock,
    kKeysPerBlock,
    kMaxK,
    false,                // Supports dropout
    false                 // Supports bias
  >;

I get compile error

Building CUDA object examples/41_fused_multi_head_attention/CMakeFiles/41_fused_multi_head_attention_fixed_seqlen.dir/fused_multihead_attention_fixed_seqlen.cu.o
/cutlass/examples/41_fused_multi_head_attention/kernel_forward.h(409): error: incomplete type is not allowed
          kIsAligned ? DefaultConfig::kAlignmentA : GemmType::kMinimumAlignment;
                       ^
          detected during:
            instantiation of class "AttentionKernel<scalar_t_, ArchTag, isAligned_, kQueriesPerBlock_, kKeysPerBlock_, kMaxK_, kSupportsDropout_, kSupportsBias_, ToBatchHookType_>::MM0 [with scalar_t_=cutlass::half_t, ArchTag=cutlass::arch::Sm89, isAligned_=true, kQueriesPerBlock_=32, kKeysPerBlock_=128, kMaxK_=128, kSupportsDropout_=false, kSupportsBias_=false, ToBatchHookType_=DefaultToBatchHook]" at line 418
            instantiation of class "AttentionKernel<scalar_t_, ArchTag, isAligned_, kQueriesPerBlock_, kKeysPerBlock_, kMaxK_, kSupportsDropout_, kSupportsBias_, ToBatchHookType_>::MM0 [with scalar_t_=cutlass::half_t, ArchTag=cutlass::arch::Sm89, isAligned_=true, kQueriesPerBlock_=32, kKeysPerBlock_=128, kMaxK_=128, kSupportsDropout_=false, kSupportsBias_=false, ToBatchHookType_=DefaultToBatchHook]" at line 1161
            instantiation of class "AttentionKernel<scalar_t_, ArchTag, isAligned_, kQueriesPerBlock_, kKeysPerBlock_, kMaxK_, kSupportsDropout_, kSupportsBias_, ToBatchHookType_> [with scalar_t_=cutlass::half_t, ArchTag=cutlass::arch::Sm89, isAligned_=true, kQueriesPerBlock_=32, kKeysPerBlock_=128, kMaxK_=128, kSupportsDropout_=false, kSupportsBias_=false, ToBatchHookType_=DefaultToBatchHook]" at line 329 of /home/sunxin20/cutlass/examples/41_fused_multi_head_attention/fused_multihead_attention_fixed_seqlen.cu
            instantiation of class "TestbedAttention<Attention> [with Attention=AttentionKernel<cutlass::half_t, cutlass::arch::Sm89, true, 32, 128, 128, false, false, DefaultToBatchHook>]" at line 1032 of /home/sunxin20/cutlass/examples/41_fused_multi_head_attention/fused_multihead_attention_fixed_seqlen.cu
            instantiation of "int run_attention<kQueriesPerBlock,kKeysPerBlock,kMaxK>(Options &) [with kQueriesPerBlock=32, kKeysPerBlock=128, kMaxK=128]" at line 1110 of /home/sunxin20/cutlass/examples/41_fused_multi_head_attention/fused_multihead_attention_fixed_seqlen.cu

Reproduce the Sm89 instantiation in examples/41_fused_multi_head_attention/fused_multihead_attention_fixed_seqlen.cu and inspect the reported failure at examples/41_fused_multi_head_attention/kernel_forward.h:409. Trace the DefaultConfig and GemmType definitions involved in the incomplete type error; done means the example compiles for the shown configuration.