Xe rearchitecture

[petercad]

This PR introduces a new architecture for Xe CuTe atoms (CUTLASS-level changes to come later).

Current status:

• MMA atoms:
• 2D copy atoms:
  • loads (regular/VNNI/transpose)
  • stores
  • prefetch
• 2D copy helpers:
  • low-level (make_block_2d_copy)
  • high-level (make_block_2d_copy_{A,B,C}).
• Reorders:
  • generic
  • optimized reorder atoms for {int4, uint4, int8, uint8, e2m1, e4m3, e5m2} -> {half, bfloat16}.

Link to rendered documentation here.

Note

This branch requires a very recent IGC version — ci-comp_igc-30311 or later. This IGC has important bug fixes/improvements to inline vISA needed to properly implement the new atoms.

[petercad]

@rolandschulz I will address the merge conflicts once review is done, to avoid rebasing.

[yuankuns]

Hi @petercad , will this PR support ST_T?

[petercad]

Hi @petercad , will this PR support ST_T?

No, since block 2D store messages don't support transposition. But, since these kinds of stores are occasionally useful, it might be a good idea to introduce an emulated transpose store operation, using D32 scattered writes. If you have some specific use cases, let me know.

[yuankuns]

Hi @petercad , will this PR support ST_T?

No, since block 2D store messages don't support transposition. But, since these kinds of stores are occasionally useful, it might be a good idea to introduce an emulated transpose store operation, using D32 scattered writes. If you have some specific use cases, let me know.

Yes there is a case in sdpa backward where dV=PtdQ might be calculated as dVt=dQtP. For dVt, it requires transpose write. For now, I just go dV=Pt*dQ for simplicity/delivery.

[rolandschulz]

Do you think the commit history is useful or do we not lose anything if we squash?

[anamikac-intel]

@petercad - there are couple of error I am seeing while running xe_gemm:

1. Got error from this using SignedAccType = ensure_signed_t; --> no template named 'ensure_signed_t' defined , I checked cute/util/type_traits it has 'make_signed_t' which worked but not ensure_signed_t

2. Also in void test_case(sycl::queue &Q, int m, int n, int k) there are some undefined identifiers like make_shared_usm_tensor, type_str, free_usm_tensor which I assume is custom and I could not find it anywhere in code defined.

[petercad]

@petercad - there are couple of error I am seeing while running xe_gemm:

1. Got error from this using SignedAccType = ensure_signed_t; --> no template named 'ensure_signed_t' defined , I checked cute/util/type_traits it has 'make_signed_t' which worked but not ensure_signed_t
2. Also in void test_case(sycl::queue &Q, int m, int n, int k) there are some undefined identifiers like make_shared_usm_tensor, type_str, free_usm_tensor which I assume is custom and I could not find it anywhere in code defined.

These are all in sycl_cute_common.hpp, which is pushed now.

[petercad]

Do you think the commit history is useful or do we not lose anything if we squash?

I personally like the more specific commit messages when looking back through Git history (e.g. blaming) when there are logically independent parts.