[hipblaslt] DTL optimizations

[b-shi]

Motivation

Addresses several DTL limitations:

• Add support for non-power of 2 MT when TLU=1
  • Add support for generalized NLC=1
• Add support for m0 padding LDSTr+DTL
• Add support for permuting perpDim to help mitigate BC

Technical Details

Add support for generalized NLC=1 (GNLC)

• Previously for TLU=1 and non-power of 2 MT, numLoadsCoalesced > 1 so the columns/rows in the non-summation dimension is fetched across multiple loads. This PR generalized NLC=1 by allowing num threads coalesced to be arbitrary.
• Uses magic div algo for offset calculations since divisors are static.
• Automatically supports DTL for non-power of 2 MT, since we can use same local read offsets for regular NLC=1 case.
• Limitations: Initially only supported for DTL, same size inputs, TLDS=1, LSU=1

Add support for m0 padding when using LDSTr+DTL

• Updated LdsBlockSizePerPad calcs for LDSTr to support padding.

Add support for permuting perpDim.

• In GNLC, support was added to permute the columns/rows in the parallel dimension (ex: permute columns of A, for A col-major)
• Permutation is done by applying a stride across consecutive columns/rows. Motivation is to have consecutive columns/rows map to different buffer_loads so that m0 padding can better mitigate bank conflicts.
• Permutation is done in blocks of columns/rows instead of across all columns/rows
• Formula for permutation: given B(block size where permutes are done), S (stride), I (index)

[Forward Mapping used in global read offset calcs]
S' = B / S
I' = S * (I % S') + I / S'

[Inverse Mapping used in local read offset calcs]
S = B / S'
I = S' * (I' % S) + I' / S

Test Plan

Test Result

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[math-ci]

perfci run on commit c3eb0ff

math-ci run

[math-ci]

perfci run on commit ab9d0c48f2e00e4709a9f698d1443c5113311e73

math-ci run

[math-ci]

perfci run on commit c6b883cb34f0e4ae0bce40a54f6afd2c997bacdc

math-ci run

[bnemanich]

All 950 tests passed.

[math-ci]

perfci run on commit 7db3030

math-ci run