Support different cp.async and load stages for operands A and B

[jacobhinkle]

As suggested in this comment: #2105 (comment)

Async loads can only be used for vectorization 4, 8, or 16. If either operand only support vectorization size of 1 or 2 elements, then we disable async loads altogether (in #2105) and limit the double buffer load stages for both operands to 2. This means we do synchronous double buffering of both operands even if one of them is fully vectorizable.

Instead, we should enable different load stages and cp.async settings for each operand.

[kevinstephano]

Is this still relevant for Hopper?

[jacobhinkle]

Is this still relevant for Hopper?

Yes I believe so but it is very low priority as it only affects unaligned inputs.