[GEMM codegen] Add e2e pipeline

[Xinyu302]

Do followering things:

1. gpu-to-vector and nvvm lowering(See GPUVectorToGPU.cpp and NVVMCodegen.cpp)
2. handle dynamic shared memory. By modify GPU::launch(See GPUToNVVM.cpp and LegalizeGPULaunch.cpp and FuncToByre.cpp)
3. GPU pipelining using upstream code(Not best way, we can add ldmatrix prefetch here)
4. Debug GPUPackSharedMemory. As we use predicted cp.async, before store to C, we should also wait the last several copy zero length to sharedA and sharedB. (See GPUPackSharedMemory.cpp)
5. All pass use memref.alloc, because upstream's passes don't handle memref.alloca
6. add f32 codegen support (using mma.sync.aligned.m16n8k8.row.col.f32.tf32.tf32.f32), but has limited precision.(TODO: support TF32_3x in upstream populateMmaSyncF32ToTF32Patterns)
7. rename getForallOpMappedTo2DBlock to getForallOpMappedToBlock as we have supported bmm.
8. Build the lowering pipeling in linalg-tensor-opt(annotate lowering config and tiling) and linalg-memref-opt(mainly optimize gemm in here)
9. Add barrier in promotion.(See LinalgPromotion.cpp)
10. Add outlineDotOp option in byteir.compile, to control whether to use gemm codegen. And we outline single dot op to trigger codegen.
11. fix bug in GPUBlockSwizzle.cpp and support bmm swizzle
12. Now we have a hack way to modify linalg.fill + linalg.matmul + linalg.generic(epilogue) pattern on tensor level. Support epilogue fusion.

[CLAassistant]

Thank you for your submission! We really appreciate it. Like many open source projects, we ask that you all sign our Contributor License Agreement before we can accept your contribution.
1 out of 2 committers have signed the CLA.

✅ Xinyu302
❌ yangxinyu

---

yangxinyu seems not to be a GitHub user. You need a GitHub account to be able to sign the CLA. If you have already a GitHub account, please add the email address used for this commit to your account.
_{You have signed the CLA already but the status is still pending? Let us recheck it.}