[QST] Any rules to follow when set instruction shape for GemmTensorOp? about cutlass

[hwu36]

as to your q1. 884 is the only one available on a100 for fp64

[hwu36]

on ampere, tf32 is using 1688. fp16/bf16 is using 16816. s8/u8 is using 16832. s4/us is using 16864. b1 is using 168256.

[YJHMITWEB]

@hwu36 Thanks for answering. If I understand correctly, it means the more bits the element type has, the fewer elements the tensor core can process. To this, I have an extended question:

In example_14_ampere_tf32_tensorop_gemm, I changed the datatype from float to cutlass::half_t, and I found m16n8k8 can still work, does this mean that there are some computing power wasted in each warp MMA? As it can do m16n8k16, but I only give it m16n8k8.

[hwu36]

1688 is turing era mma, it still can work on ampere. but the bigger, the better unless your problem size is small.

[YJHMITWEB]

@hwu36 Thanks!