[FEA] Complete the cutlass::library::GemmDescription class to cover Hopper GEMM kernels

[manishucsd]

Issue

CUTLASS GemmDevice Operator contains compile-time attributes (functional and performance attribute). The GemmDevice Operator is consumed by GemmOperation[3xBase]. In the past, I have found some of the values in this data structure incorrect and sometimes just completely missing.

For e.g. given the kernel by its full procedural name = cutlass3x_sm90_tensorop_s64x128x32gemm_e4m3_e4m3_f32_bf16_bf16_128x128x128_1x2x1_0_tnn_align16_warpspecialized_cooperative_epi_tma.

This kernel has the following functional and performance attribute:

Functional Attribute

• dtypeA_dtypeB_dtypeAccumulation_dtypeC_dtypeD : e4m3_e4m3_f32_bf16_bf16
• RowMajor_ColumnMajor_ColumnMajor : tnn

Performance Attribute

• Instruction Shape : 64x128x32
• Threadblock Shape : 128x128x128
• Cluster Shape : 1x2x1
• Mainloop Kind : warpspecialized_cooperative // This is missing from GemmDescription
• Epilogue Kind : epi_tma // This is missing from GemmDescription
• AccumulationKind : "fastaccum" or default "" // This is missing from GemmDescription

I have added a test so someone at NVIDIA can start on this. Can you please uncomment the two lines, add whatever is needed to fix this?

You can follow any other enum that is lifted up to GemmDevice Operator from internal templates and used to set the data members of GemmDescription class.

We can then commit this test and add more for Hopper and make sure this class is also covered for Blackwell. The tests are CPU-only and should not take too much time in the CI, this will allow us to catch bugs like this one.

[manishucsd]

@hwu36, can we assign this to be fixed sometime in the next release. Thank you!

[hwu36]

@itramble

[manishucsd]

Here is the commit that shows what I am thinking of adding to solve this. Things to do:

• Make sure that KernelSchedule tag is available CUTLASS GemmDevice Operator
• For all the kernel compile-time info encoded in this KernelSchedule Tag for Hopper here needs to be unpacked into runtime enums like this.
• Also do the same for Blackwell kernel schedules.
• At the minimum, anything that is in kernel name should be reflected in the library enum types for us to be able to select the right kernel at the runtime. So check the kernel name for existing and new kernels in the cutlass library and check if something is missing from tools/library/include/cutlass/library/types.h

This is just the starting code, I had in my mind. If you have a better way to solve this, please choose that. I won't be working on it after posting this, but looking forward to see it solved soon :).

cc: @itramble, @hwu36

[manishucsd]

Any progress or update or ETA on this?

[manishucsd]

Reading the GEMM description object for the kernel: cutlass3x_sm90_tensorop_s64x128x16gemm_bf16_bf16_f32_bf16_bf16_128x256x64_1x1x1_0_tnn_align8_warpspecialized_cooperative_epi_tma

Functional Attributes {
  element_compute: f32
   element_scalar: f32
        element_A: bf16
         layout_A: row
        element_B: bf16
         layout_B: column
        element_C: bf16
         layout_C: column
        element_D: bf16
         layout_D: column
}

Performance Attributes: {
         instruction shape: {64, 256, 16} // incorrect (FIX asked in this **comment**)
         threadblock shape: {128, 256, 64} 
             cluster shape: {1, 1, 1}
         mainloop schedule: {warpspecialized cooperative} // this using a workaorund (FIX asked in this **bug**)
        epilogue schedule: {tma}  // this using a workaorund (FIX asked in this **bug**)
  }

Kernel name: cutlass3x_sm90_tensorop_s64x128x16gemm_bf16_bf16_f32_bf16_bf16_128x256x64_1x1x1_0_tnn_align8_warpspecialized_cooperative_epi_tma
Instruction shape doesn't match the expectation:

• Expected: s64x128x16gemm
• Observed when read from gemm_desc object: instruction shape: {64, 256, 16}

I am 99.9% certain something wrong with the code lines here. The way Operator::InstructionsShape compile-time arguments are lifted from the inner-most cute::arch templates to the outer most device-level Operator templates.

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

[github-actions]

This issue has been labeled inactive-90d due to no recent activity in the past 90 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed.