In hopper, what is rs and ss strategy?

[202476410arsmart]

[Robert_Crovella]

in case anyone is wondering, these are part of cutlass

[Curefab]

Register Storage vs. Shared Storage (for operand A) maybe?

[202476410arsmart]

Hey, that sounds very interesting! Regardless of cutlass details, how to save operand A in register or shared? I guess… maybe in epiloge, the tensor core output is in register, either we directly save them to global, or we save to shared to make them coalesced and then to global?

[Curefab]

Yes, you can custom-save elements with the epilogue. See

Understanding cutlass GEMM hierarchy

I think this picture is showing what cutlass is doing. [gemm-hierarchy-with-epilogue-no-labels] But I am not understanding what is happening. Or what is the shape? Here they are defining several shape, why several and how it is going to work? cutlass::gemm::GemmShape<128, 128, 64>, cutlass::gemm::GemmShape<64, 64, 64>, cutlass::gemm::GemmShape<16, 8, 16>, Or here how can I relate each line to the picuture that I mentioned? ShapeMMAThreadBlock ShapeMMAWarp ShapeMMAOp

with links to

• https://developer.nvidia.com/blog/cutlass-linear-algebra-cuda/
• https://on-demand.gputechconf.com/gtc/2018/presentation/s8854-cutlass-software-primitives-for-dense-linear-algebra-at-all-levels-and-scales-within-cuda.pdf