Are swizzling functions for reducing shared memory bank conflicts same for different architectures?

[navdeepkk]

So this is my understanding, please correct me if I am wrong. Swizzling of data in shared memory depends on the PTX instruction which will finally use the data. PTX instructions may vary with different architectures. Each PTX instruction may have a different thread data mapping and to have minimum bank conflicts the swizzling function should also be different i.e., a swizzling function for one PTX instruction may not result in the same bank conflicts for a different PTX instruction. Is that correct?

Thanks!

[hwu36]

No.

Shared memory swizzling is used to prevent bank conflict. Only shared memory load/store instructions could have this issue. I don't remember that shared memory banks have ever been changed.

[navdeepkk]

Thanks for the reply. Okay. so you are saying that swizzling functions are the same for Volta, Turing, and Ampere. The things that will change are the functions to load operands from shared memory into the registers and then to store back the results from shared memory into the registers. This is because the PTX instructions or the larger mma.syncs may have different thread data mappings.

So i.e., functions to load/store from/to shared memory have a device-specific implementation in CUTLASS? Is that correct?
Thanks!

[hwu36]

So i.e., functions to load/store from/to shared memory have a device-specific implementation in CUTLASS? Is that correct?

Correct.

Volta fp16 uses one shared memory layout, turing/ampere b1/s4/s8/fp16/bf16/tf32 use one shared memory layout, ampere fp64 uses one shared memory layout.