diff --git a/docs/src/assets/TF32F32.svg b/docs/src/assets/TF32F32.svg new file mode 100644 index 00000000..4b7643ec --- /dev/null +++ b/docs/src/assets/TF32F32.svg @@ -0,0 +1,8036 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/docs/src/assets/smem_copy_A.svg b/docs/src/assets/smem_copy_A.svg index fff647c3..49cc93b2 100644 --- a/docs/src/assets/smem_copy_A.svg +++ b/docs/src/assets/smem_copy_A.svg @@ -1,4 +1,4 @@ - + @@ -48,54 +48,6 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - @@ -144,54 +96,6 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - @@ -240,54 +144,6 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - @@ -336,54 +192,6 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - @@ -432,54 +240,6 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - @@ -528,54 +288,6 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - @@ -624,54 +336,6 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - @@ -720,80 +384,32 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + - + - + - + - + - + - + - + - + @@ -846,54 +462,6 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - @@ -3662,1735 +3230,482 @@ - + - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + - - - - - - - - - - - - - - + - - - - - - - - + - - + @@ -5398,4665 +3713,90 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + + + + - - - - - - - - - - - - - - - - - + + + + + + + + + - - - - - - - - - - - - - - - - - + + + + + + + + + - - - - - - - - - - - - - - - - - + + + + + + + + + - - - - - - - - - - - - - - - - - + + + + + + + + + - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + - + - + - @@ -10065,11 +3805,10 @@ - + - - + @@ -10084,11 +3823,10 @@ - + - - + @@ -10103,11 +3841,10 @@ - + - @@ -10122,11 +3859,10 @@ - + - - + @@ -10141,11 +3877,10 @@ - + - - + @@ -10160,11 +3895,10 @@ - + - - + @@ -10179,11 +3913,10 @@ - + - @@ -10198,11 +3931,10 @@ - + - - + @@ -10217,16 +3949,15 @@ - + - + - - + @@ -10235,11 +3966,10 @@ - + - - + @@ -10254,11 +3984,10 @@ - + - - + @@ -10273,11 +4002,10 @@ - + - - + @@ -10292,11 +4020,10 @@ - + - - + @@ -10311,11 +4038,10 @@ - + - - + @@ -10330,11 +4056,10 @@ - + - - + @@ -10349,11 +4074,10 @@ - + - - + @@ -10368,11 +4092,10 @@ - + - - + @@ -10387,16 +4110,15 @@ - + - + - - + @@ -10405,11 +4127,10 @@ - + - - + @@ -10427,8 +4148,8 @@ - - + + @@ -10446,8 +4167,8 @@ - - + + @@ -10465,8 +4186,8 @@ - - + + @@ -10481,11 +4202,10 @@ - + - - - + + @@ -10503,8 +4223,8 @@ - - + + @@ -10522,8 +4242,8 @@ - - + + @@ -10541,8 +4261,8 @@ - - + + @@ -10557,16 +4277,15 @@ - + - + - - + @@ -10578,8 +4297,8 @@ - - + + @@ -10597,8 +4316,8 @@ - - + + @@ -10616,8 +4335,8 @@ - - + + @@ -10635,8 +4354,8 @@ - - + + @@ -10654,8 +4373,8 @@ - - + + @@ -10673,8 +4392,8 @@ - - + + @@ -10692,8 +4411,8 @@ - - + + @@ -10711,8 +4430,8 @@ - - + + @@ -10727,16 +4446,15 @@ - + - + - - + @@ -10748,8 +4466,8 @@ - - + + @@ -10767,8 +4485,8 @@ - - + + @@ -10786,8 +4504,8 @@ - - + + @@ -10805,8 +4523,8 @@ - - + + @@ -10824,8 +4542,8 @@ - - + + @@ -10843,8 +4561,8 @@ - - + + @@ -10862,8 +4580,8 @@ - - + + @@ -10881,8 +4599,8 @@ - - + + @@ -10897,16 +4615,15 @@ - + - + - - + @@ -10918,8 +4635,8 @@ - - + + @@ -10937,8 +4654,8 @@ - - + + @@ -10956,8 +4673,8 @@ - - + + @@ -10975,8 +4692,8 @@ - - + + @@ -10994,8 +4711,8 @@ - - + + @@ -11013,8 +4730,8 @@ - - + + @@ -11032,8 +4749,8 @@ - - + + @@ -11051,8 +4768,8 @@ - - + + @@ -11067,16 +4784,15 @@ - + - + - - + @@ -11088,8 +4804,8 @@ - - + + @@ -11107,8 +4823,8 @@ - - + + @@ -11126,8 +4842,8 @@ - - + + @@ -11145,8 +4861,8 @@ - - + + @@ -11164,8 +4880,8 @@ - - + + @@ -11183,8 +4899,8 @@ - - + + @@ -11202,8 +4918,8 @@ - - + + @@ -11221,8 +4937,8 @@ - - + + @@ -11237,16 +4953,15 @@ - + - + - - + @@ -11258,8 +4973,8 @@ - - + + @@ -11277,8 +4992,8 @@ - - + + @@ -11296,8 +5011,8 @@ - - + + @@ -11315,8 +5030,8 @@ - - + + @@ -11334,8 +5049,8 @@ - - + + @@ -11353,8 +5068,8 @@ - - + + @@ -11372,8 +5087,8 @@ - - + + @@ -11391,8 +5106,8 @@ - - + + @@ -11407,16 +5122,15 @@ - + - + - - + @@ -11425,11 +5139,10 @@ - + - - + @@ -11444,11 +5157,10 @@ - + - - + @@ -11463,11 +5175,10 @@ - + - @@ -11482,11 +5193,10 @@ - + - - + @@ -11501,11 +5211,10 @@ - + - - + @@ -11520,11 +5229,10 @@ - + - - + @@ -11539,11 +5247,10 @@ - + - @@ -11558,11 +5265,10 @@ - + - - + @@ -11577,7 +5283,7 @@ - + @@ -11585,8 +5291,8 @@ - - + + @@ -11595,11 +5301,10 @@ - + - - + @@ -11614,11 +5319,10 @@ - + - - + @@ -11633,11 +5337,10 @@ - + - - + @@ -11652,11 +5355,10 @@ - + - - + @@ -11671,11 +5373,10 @@ - + - - + @@ -11690,11 +5391,10 @@ - + - - + @@ -11709,11 +5409,10 @@ - + - - + @@ -11728,11 +5427,10 @@ - + - - + @@ -11747,7 +5445,7 @@ - + @@ -11755,8 +5453,8 @@ - - + + @@ -11765,11 +5463,10 @@ - + - - + @@ -11787,8 +5484,8 @@ - - + + @@ -11806,8 +5503,8 @@ - - + + @@ -11825,8 +5522,8 @@ - - + + @@ -11841,11 +5538,10 @@ - + - - + @@ -11863,8 +5559,8 @@ - - + + @@ -11882,8 +5578,8 @@ - - + + @@ -11901,8 +5597,8 @@ - - + + @@ -11917,7 +5613,7 @@ - + @@ -11925,8 +5621,8 @@ - - + + @@ -11938,8 +5634,8 @@ - - + + @@ -11957,8 +5653,8 @@ - - + + @@ -11976,8 +5672,8 @@ - - + + @@ -11995,8 +5691,8 @@ - - + + @@ -12014,8 +5710,8 @@ - - + + @@ -12033,8 +5729,8 @@ - - + + @@ -12052,8 +5748,8 @@ - - + + @@ -12071,8 +5767,8 @@ - - + + @@ -12087,7 +5783,7 @@ - + @@ -12095,8 +5791,8 @@ - - + + @@ -12108,8 +5804,8 @@ - - + + @@ -12127,8 +5823,8 @@ - - + + @@ -12146,8 +5842,8 @@ - - + + @@ -12165,8 +5861,8 @@ - - + + @@ -12184,8 +5880,8 @@ - - + + @@ -12203,8 +5899,8 @@ - - + + @@ -12222,8 +5918,8 @@ - - + + @@ -12241,8 +5937,8 @@ - - + + @@ -12257,7 +5953,7 @@ - + @@ -12265,8 +5961,8 @@ - - + + @@ -12278,8 +5974,8 @@ - - + + @@ -12297,8 +5993,8 @@ - - + + @@ -12316,8 +6012,8 @@ - - + + @@ -12335,8 +6031,8 @@ - - + + @@ -12354,8 +6050,8 @@ - - + + @@ -12373,8 +6069,8 @@ - - + + @@ -12392,8 +6088,8 @@ - - + + @@ -12411,8 +6107,8 @@ - - + + @@ -12427,7 +6123,7 @@ - + @@ -12435,8 +6131,8 @@ - - + + @@ -12448,8 +6144,8 @@ - - + + @@ -12467,8 +6163,8 @@ - - + + @@ -12486,8 +6182,8 @@ - - + + @@ -12505,8 +6201,8 @@ - - + + @@ -12524,8 +6220,8 @@ - - + + @@ -12543,8 +6239,8 @@ - - + + @@ -12562,8 +6258,8 @@ - - + + @@ -12581,8 +6277,8 @@ - - + + @@ -12597,7 +6293,7 @@ - + @@ -12605,8 +6301,8 @@ - - + + @@ -12618,8 +6314,8 @@ - - + + @@ -12637,8 +6333,8 @@ - - + + @@ -12656,8 +6352,8 @@ - - + + @@ -12675,8 +6371,8 @@ - - + + @@ -12694,8 +6390,8 @@ - - + + @@ -12713,8 +6409,8 @@ - - + + @@ -12732,8 +6428,8 @@ - - + + @@ -12751,8 +6447,8 @@ - - + + diff --git a/docs/src/assets/smem_copy_B.svg b/docs/src/assets/smem_copy_B.svg new file mode 100644 index 00000000..fe638145 --- /dev/null +++ b/docs/src/assets/smem_copy_B.svg @@ -0,0 +1,3376 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/docs/src/manual/async.md b/docs/src/manual/async.md index c87e121e..792fa5a8 100644 --- a/docs/src/manual/async.md +++ b/docs/src/manual/async.md @@ -58,7 +58,7 @@ function matmul_kernel(A, sA_layout, copy_A, cp_async_wait() - @gc_preserve gemm!(mma_C, tCsA, tCsB, tCrC) + @gc_preserve gemm!(mma_C, tCrC, tCsA, tCsB, tCrC) sync_threads() end diff --git a/docs/src/manual/matmul.md b/docs/src/manual/matmul.md index 861fb163..e3d1067f 100644 --- a/docs/src/manual/matmul.md +++ b/docs/src/manual/matmul.md @@ -76,7 +76,7 @@ end ``` You can also call [`gemm!`] to perform the same operation: ```julia -gemm!(tCsA, tCsB, tCrC) +gemm!(tCrC, tCsA, tCsB, tCrC) ``` The complete kernel code is as follows: @@ -121,7 +121,7 @@ function matmul_kernel(A, sA_layout, tA, cp_async_wait() sync_threads() - @gc_preserve gemm!(tCsA, tCsB, tCrC) + @gc_preserve gemm!(tCrC, tCsA, tCsB, tCrC) sync_threads() end diff --git a/docs/src/manual/pipeline.md b/docs/src/manual/pipeline.md index 67de5ac5..ab71d26f 100644 --- a/docs/src/manual/pipeline.md +++ b/docs/src/manual/pipeline.md @@ -67,7 +67,7 @@ function matmul_kernel(A, sA_layout, copy_A, copyto!(copy_B, tBsB, view(tBgB, :, :, :, k+1)) end - @gc_preserve gemm!(mma_C, tCrA, tCrB, tCrC) + @gc_preserve gemm!(mma_C, tCrC, tCrA, tCrB, tCrC) end copyto!(tCgC, tCrC) @@ -161,7 +161,7 @@ for the next tile. We prefetch the next tile from global memory to shared memory smem_read, smem_write = smem_write, smem_read end - @gc_preserve gemm!(mma_C, tCrA[:, :, k_block], tCrB[:, :, k_block], tCrC) + @gc_preserve gemm!(mma_C, tCrC, tCrA[:, :, k_block], tCrB[:, :, k_block], tCrC) end end diff --git a/docs/src/manual/tensor_core.md b/docs/src/manual/tensor_core.md index fe7d17ee..c370491a 100644 --- a/docs/src/manual/tensor_core.md +++ b/docs/src/manual/tensor_core.md @@ -5,37 +5,138 @@ Tensor cores are specialized hardware accelerators designed to optimize matrix o Enabling tensor cores can be as straightforward as modifying a single line of code in the existing `matmul_kernel` function: ```julia mma = make_tiled_mma(MMAOP_8x8x4_F32F16F16F32_NT(), - @Layout((2,4,1))) + atom_layout, + tiler) ``` !!! note The NT in MMAOP_8x8x4_F32F16F16F32_NT indicates that A is in M-major order and B is in N-major order. -Let's explore what `TiledMMA` with a tensor core operation entails. +Let's explore a minimal example ```julia -mma = make_tiled_mma(MMAOP_16x8x8_F32TF32TF32F32_TN(), - @Layout((2,4,1))) +mma = make_tiled_mma(MMAOP_16x8x8_F32TF32TF32F32_TN()) print_typst(mma) ``` -![](../assets/tensorcore.svg) +![](../assets/TF32F32.svg) At first glance, the diagram may seem complex, but the concept is straightforward: the threads collective load data from matrices `A` and `B` according to the specified layout. During the matrix multiply-accumulate (MMA) computation, data is internally shared among threads—a process that is not transparent to the user. Once the computation is complete, each thread stores the results as dictated by the layout of matrix `C` shown in the illustration. +```julia +function matmul_kernel(A, sA_layout, copy_A, + B, sB_layout, copy_B, + C, mma) + sA = MoYeSharedArray(eltype(A), sA_layout) + sB = MoYeSharedArray(eltype(B), sB_layout) + + mA = MoYeArray(A) + mB = MoYeArray(B) + mC = MoYeArray(C) + + bM = size(sA_layout, 1) + bN = size(sB_layout, 1) + bK = size(sB_layout, 2) + + gA = @tile mA (bM, bK) (blockIdx().x, :) + gB = @tile mB (bN, bK) (blockIdx().y, :) + gC = @tile mC (bM, bN) (blockIdx().x, blockIdx().y) + + # copy partition + thr_copy_a = get_slice(copy_A, threadIdx().x) + tAgA = partition_S(thr_copy_a, gA) # (CPY, CPY_M, CPY_K, k) + tAsA = partition_D(thr_copy_a, sA) # (CPY, CPY_M, CPY_K) + + thr_copy_b = get_slice(copy_B, threadIdx().x) + tBgB = partition_S(thr_copy_b, gB) # (CPY, CPY_N, CPY_K, k) + tBsB = partition_D(thr_copy_b, sB) # (CPY, CPY_N, CPY_K) + + # mma partition + thr_mma = get_slice(mma, threadIdx().x) + tCsA = partition_A(thr_mma, sA) # (MMA, MMA_M, MMA_K) + tCsB = partition_B(thr_mma, sB) # (MMA, MMA_M, MMA_K) + tCgC = partition_C(thr_mma, gC) # (MMA, MMA_M, MMA_N) + + tCrA = make_fragment_A(thr_mma, tCsA) # (MMA, MMA_M, MMA_K) + tCrB = make_fragment_B(thr_mma, tCsB) + tCrC = make_fragment_C(thr_mma, tCgC) + zeros!(tCrC) + + # copy from global to shared + copyto!(copy_A, tAsA, view(tAgA, :, :, :, _1)) + copyto!(copy_B, tBsB, view(tBgB, :, :, :, _1)) + + cp_async_wait() + + # copy from shared to registers + copyto!(tCrA, tCsA) + copyto!(tCrB, tCsB) + + @gc_preserve gemm!(mma, tCrC, tCrA, tCrB, tCrC) + + copyto!(tCgC, tCrC) + @inbounds tCrC[1] # compiler bug, have to load after copyto! + + return nothing +end + +function matmul(A, B, C) + bM = _16 + bN = _8 + bK = _8 + + sA_layout = make_layout((bM, bK), (_1, bM)) + sB_layout = make_layout((bN, bK), (bK, _1)) + + TA = eltype(A) + TB = eltype(B) + TC = eltype(C) + + copy_A = make_tiled_copy(CopyAtom{CPOP_ASYNC_CACHEALWAYS{UInt128}, TA}(), + @Layout((4, 8)), + @Layout((4, 1))) + copy_B = make_tiled_copy(CopyAtom{CPOP_ASYNC_CACHEALWAYS{UInt64}, TB}(), + @Layout((8, 4), (4, 1)), + @Layout((1, 2))) + + mma = make_tiled_mma(MMAOP_16x8x8_F32TF32TF32F32_TN()) + + threads = Int(size(mma)) + blocks = (cld(size(A, 1), bM), cld(size(B, 1), bN)) + + @cuda threads=threads blocks=blocks matmul_kernel(A, sA_layout, copy_A, + B, sB_layout, copy_B, + C, mma) +end + +function test() + A = CuArray(reshape(collect(1:16*8) .* 1f0, (16,8))) + B = CuArray(reshape(collect(1:8*8) .* 1f0, (8,8))) + C = CuArray(ones(Float32, (16,8))) + matmul(A, B', C) + CUDA.synchronize() + @test C == A * B + CUDA.unsafe_free!(A) + CUDA.unsafe_free!(B) + CUDA.unsafe_free!(C) +end +``` + + ## LDMatrix The `ldmatrix` instruction at the warp level facilitates the loading of data from shared memory into registers and suffles them to align with a tensor core MMA operation. Given a tensor core MMA operation, the shuffling can be "inverted" to obtain a `TiledCopy` count for the shuffling. ```julia +mma = make_tiled_mma(MMAOP_16x8x8_F32TF32TF32F32_TN()) smem_copy_A = make_tiled_copy_A(CopyAtom{LDSM_U32x4_N, Float32}(), mma) print_typst(smem_copy_A) ``` ![](../assets/smem_copy_A.svg) -The resulting layout matches the layout of A in the `mma`. +The resulting layout on the right hand side matches the layout of A in the `mma`. !!! note - The `TN` in `MMAOP_16x8x8_F32TF32TF32F32_TN` specifies that both A and B are in `K`-major order. - The `N` in `LDSM_U32x4_N` indicates `K`-major order. + The `TN` in `MMAOP_16x8x8_F32TF32TF32F32_TN` means that both A and B are in `K`-major order. + The `N` in `LDSM_U32x4_N` means the source array is `K`-major order. !!! note The `ldmatrix` requires four consecutive threads to load 16 consecutive bytes, demanding that the layout of `A` in shared memory meet this specification. @@ -45,27 +146,193 @@ For B: smem_copy_B = make_tiled_copy_B(CopyAtom{LDSM_U32x2_N, Float32}(), mma) print_typst(smem_copy_B) ``` -!!! Note - However, using LDSM_U32x4_N for `B` would not be compatible with its layout in mma. - A developer is resposible to select a compatible `ldmatrix` operation when possible. +![](../assets/smem_copy_B.svg) + +!!! Note + The visualization of `B` in `mma` is draw as `(K, N)` but `(N, K)` in `smem_copy_B`. We then use `smem_copy_A` and `smem_copy_B` to re-tile the shared memory and registers ```julia -smem_thr_copy_A = get_slice(smem_copy_A, threadIdx().x) -tCsA = partition_S(smem_thr_copy_A, sA) -tCrA_copy_view = retile_D(smem_thr_copy_A, tCrA) +smem_thr_copy_A = get_slice(smem_copy_A, threadIdx().x) +smem_thr_copy_B = get_slice(smem_copy_B, threadIdx().x) +tCsA_retiled = partition_S(smem_thr_copy_A, sA) +tCsB_retiled = partition_S(smem_thr_copy_B, sB) +tCrA_retiled = retile_D(smem_thr_copy_A, tCrA) +tCrB_retiled = retile_D(smem_thr_copy_B, tCrB) +``` + +Complete code: +```julia +function matmul_kernel(A, sA_layout, gmem_copy_A, smem_copy_A, + B, sB_layout, gmem_copy_B, smem_copy_B, + C, mma) + sA = MoYeSharedArray(eltype(A), sA_layout) + sB = MoYeSharedArray(eltype(B), sB_layout) + + mA = MoYeArray(A) + mB = MoYeArray(B) + mC = MoYeArray(C) -smem_thr_copy_B = get_slice(smem_copy_B, threadIdx().x) -tCsB = partition_S(smem_thr_copy_B, sB) -tCrB_copy_view = retile_D(smem_thr_copy_B, tCrB) + bM = size(sA_layout, 1) + bN = size(sB_layout, 1) + bK = size(sB_layout, 2) + + gA = @tile mA (bM, bK) (blockIdx().x, :) + gB = @tile mB (bN, bK) (blockIdx().y, :) + gC = @tile mC (bM, bN) (blockIdx().x, blockIdx().y) + + # gmem copy partition + gmem_thr_copy_a = get_slice(gmem_copy_A, threadIdx().x) + tAgA = partition_S(gmem_thr_copy_a, gA) # (CPY, CPY_M, CPY_K, k) + tAsA = partition_D(gmem_thr_copy_a, sA) # (CPY, CPY_M, CPY_K) + + gmem_thr_copy_b = get_slice(gmem_copy_B, threadIdx().x) + tBgB = partition_S(gmem_thr_copy_b, gB) # (CPY, CPY_N, CPY_K, k) + tBsB = partition_D(gmem_thr_copy_b, sB) # (CPY, CPY_N, CPY_K) + + # copy from global to shared + copyto!(gmem_copy_A, tAsA, view(tAgA, :, :, :, _1)) + copyto!(gmem_copy_B, tBsB, view(tBgB, :, :, :, _1)) + + # mma partition + thr_mma = get_slice(mma, threadIdx().x) + tCsA = partition_A(thr_mma, sA) # (MMA, MMA_M, MMA_K) + tCsB = partition_B(thr_mma, sB) # (MMA, MMA_M, MMA_K) + tCgC = partition_C(thr_mma, gC) # (MMA, MMA_M, MMA_N) + + tCrA = make_fragment_A(thr_mma, tCsA) # (MMA, MMA_M, MMA_K) + tCrB = make_fragment_B(thr_mma, tCsB) # (MMA, MMA_N, MMA_K) + tCrC = make_fragment_C(thr_mma, tCgC) # (MMA, MMA_M, MMA_N) + zeros!(tCrC) + + # retile + smem_thr_copy_A = get_slice(smem_copy_A, threadIdx().x) + smem_thr_copy_B = get_slice(smem_copy_B, threadIdx().x) + tCsA_retiled = partition_S(smem_thr_copy_A, sA) + tCsB_retiled = partition_S(smem_thr_copy_B, sB) + tCrA_retiled = retile_D(smem_thr_copy_A, tCrA) + tCrB_retiled = retile_D(smem_thr_copy_B, tCrB) + + cp_async_wait() + + # copy from shared to registers + copyto!(smem_copy_A, tCrA_retiled, tCsA_retiled) + copyto!(smem_copy_B, tCrB_retiled, tCsB_retiled) + + @gc_preserve gemm!(mma, tCrC, tCrA, tCrB, tCrC) + + copyto!(tCgC, tCrC) + @inbounds tCrC[1] # compiler bug, have to load after copyto! + + return nothing +end + + +function matmul(A, B, C) + bM = _16 + bN = _8 + bK = _8 + + sA_layout = make_layout((bM, bK), (_1, bM)) + sB_layout = make_layout((bN, bK), (bK, _1)) + + TA = eltype(A) + TB = eltype(B) + TC = eltype(C) + + gmem_copy_A = make_tiled_copy(CopyAtom{CPOP_ASYNC_CACHEALWAYS{UInt128}, TA}(), + @Layout((4, 8)), + @Layout((4, 1))) + gmem_copy_B = make_tiled_copy(CopyAtom{CPOP_ASYNC_CACHEALWAYS{UInt64}, TB}(), + @Layout((8, 4), (4, 1)), + @Layout((1, 2))) + + mma = make_tiled_mma(MMAOP_16x8x8_F32TF32TF32F32_TN()) + + # Note: A is M-major so we can only use `UniversalCopy` + smem_copy_A = make_tiled_copy_A(CopyAtom{UniversalCopy{TA}, TA}(), mma) + smem_copy_B = make_tiled_copy_B(CopyAtom{LDSM_U32x2_N, TB}(), mma) + + threads = Int(size(mma)) + blocks = (cld(size(A, 1), bM), cld(size(B, 1), bN)) + + @cuda threads=threads blocks=blocks matmul_kernel(A, sA_layout, gmem_copy_A, smem_copy_A, + B, sB_layout, gmem_copy_B, smem_copy_B, + C, mma) +end + +``` + +## Tiled LdMatrix + +So far we have only talked about how to use a single `ldmatrix` instruction. +We can use `tile_to_shape` to tile it to a larger shape. +```julia +sB_atom_layout = make_layout((_8, _8), (_8, _1)) +sB_layout = tile_to_shape(sB_atom_layout, static((24, 16))) + +print_layout(sB_atom_layout) +print_layout(sB_layout) +``` +Note how the internal layout of `sB_atom_layout` is preserved in `sB_layout`. + +Updated code: +```julia +function matmul(A, B, C) + bM = _16 + bN = _8 + bK = _16 + + sA_atom_layout = @Layout (16, 8) (1, 16) + sB_atom_layout = @Layout (8, 8) (8, 1) + + sA_layout = MoYe.tile_to_shape(sA_atom_layout, (bM, bK)) + sB_layout = MoYe.tile_to_shape(sB_atom_layout, (bN, bK)) + + TA = eltype(A) + TB = eltype(B) + TC = eltype(C) + + gmem_copy_A = make_tiled_copy(CopyAtom{CPOP_ASYNC_CACHEALWAYS{UInt128}, TA}(), + @Layout((4, 8)), + @Layout((4, 1))) + gmem_copy_B = make_tiled_copy(CopyAtom{CPOP_ASYNC_CACHEALWAYS{UInt128}, TB}(), + @Layout((8, 4), (4, 1)), + @Layout((1, 4))) + + mma = make_tiled_mma(MMAOP_16x8x8_F32TF32TF32F32_TN()) + + # Note: A is M-major so we can only use `UniversalCopy` + smem_copy_A = make_tiled_copy_A(CopyAtom{UniversalCopy{TA}, TA}(), mma) + smem_copy_B = make_tiled_copy_B(CopyAtom{LDSM_U32x2_N, TB}(), mma) + + threads = Int(size(mma)) + blocks = (cld(size(A, 1), bM), cld(size(B, 1), bN)) + + @cuda threads=threads blocks=blocks matmul_kernel(A, sA_layout, gmem_copy_A, smem_copy_A, + B, sB_layout, gmem_copy_B, smem_copy_B, + C, mma) +end + +function test() + A = CUDA.randn(Float32, 16, 16) # M-major + B = CUDA.randn(Float32, 16, 8) # K-major + C = CUDA.randn(Float32, 16, 8) + matmul(A, B', C) + CUDA.synchronize() + @test C == A * B + CUDA.unsafe_free!(A) + CUDA.unsafe_free!(B) + CUDA.unsafe_free!(C) +end + +test() ``` -Here, `retile_D` acts as a composition of `tCrA` with the partitioner `smem_thr_copy_A`. -## MatMul +## Double buffering -This example computes C = A * B, with A in M-major and B in K-major order. ```julia @views function matmul_kernel(A, sA_layout, gmem_copy_A, smem_copy_A, @@ -101,22 +368,22 @@ This example computes C = A * B, with A in M-major and B in K-major order. # mma partition thr_mma = get_slice(mma, threadIdx().x) - tCgC = partition_C(thr_mma, gC) # (MMA, MMA_M, MMA_N) + tCsA = partition_A(thr_mma, sA) # (MMA, MMA_M, MMA_K, 2) + tCsB = partition_B(thr_mma, sB) # (MMA, MMA_M, MMA_K, 2) + tCgC = partition_C(thr_mma, gC) # (MMA, MMA_M, MMA_N) - # mma registers - tCrA = partition_fragment_A(thr_mma, sA[:, :, _1]) # (MMA, MMA_M, MMA_K) - tCrB = partition_fragment_B(thr_mma, sB[:, :, _1]) # (MMA, MMA_N, MMA_K) - tCrC = make_fragment_C(thr_mma, tCgC) # (MMA, MMA_M, MMA_N) + tCrA = make_fragment_A(thr_mma, tCsA[:, :, :, _1]) # (MMA, MMA_M, MMA_K) + tCrB = make_fragment_B(thr_mma, tCsB[:, :, :, _1]) # (MMA, MMA_N, MMA_K) + tCrC = make_fragment_C(thr_mma, tCgC) # (MMA, MMA_M, MMA_N) zeros!(tCrC) # retile smem_thr_copy_A = get_slice(smem_copy_A, threadIdx().x) - tCsA = partition_S(smem_thr_copy_A, sA) # (MMA, MMA_M, MMA_K, 2) - tCrA_copy_view = retile_D(smem_thr_copy_A, tCrA) # (MMA, MMA_M, MMA_K) - smem_thr_copy_B = get_slice(smem_copy_B, threadIdx().x) - tCsB = partition_S(smem_thr_copy_B, sB) # (MMA, MMA_N, MMA_K, 2) - tCrB_copy_view = retile_D(smem_thr_copy_B, tCrB) # (MMA, MMA_N, MMA_K) + tCsA_retiled = partition_S(smem_thr_copy_A, sA) # (MMA, MMA_M, MMA_K, 2) + tCsB_retiled = partition_S(smem_thr_copy_B, sB) # (MMA, MMA_N, MMA_K, 2) + tCrA_retiled = retile_D(smem_thr_copy_A, tCrA) # (MMA, MMA_M, MMA_K) + tCrB_retiled = retile_D(smem_thr_copy_B, tCrB) # (MMA, MMA_N, MMA_K) cp_async_wait() sync_threads() @@ -124,10 +391,10 @@ This example computes C = A * B, with A in M-major and B in K-major order. # Copy smem to rmem for k_block=1 smem_read = 1 smem_write = 2 - tCsA_p = view(tCsA, :, :, :, smem_read) - tCsB_p = view(tCsB, :, :, :, smem_read) - copyto!(smem_copy_A, tCrA_copy_view[:, :, _1], tCsA_p[:, :, _1]) - copyto!(smem_copy_B, tCrB_copy_view[:, :, _1], tCsB_p[:, :, _1]) + tCsA_p = view(tCsA_retiled, :, :, :, smem_read) + tCsB_p = view(tCsB_retiled, :, :, :, smem_read) + copyto!(smem_copy_A, tCrA_retiled[:, :, _1], tCsA_p[:, :, _1]) + copyto!(smem_copy_B, tCrB_retiled[:, :, _1], tCsB_p[:, :, _1]) k_tile_max = size(tAgA, 4) k_block_max = static_size(tCrA, 3) @@ -137,13 +404,13 @@ This example computes C = A * B, with A in M-major and B in K-major order. if k_block == k_block_max cp_async_wait() sync_threads() - tCsA_p = view(tCsA, :, :, :, smem_read) - tCsB_p = view(tCsB, :, :, :, smem_read) + tCsA_p = view(tCsA_retiled, :, :, :, smem_read) + tCsB_p = view(tCsB_retiled, :, :, :, smem_read) k_block_next = 1 end - copyto!(smem_copy_A, tCrA_copy_view[:, :, k_block_next], tCsA_p[:, :, k_block_next]) - copyto!(smem_copy_B, tCrB_copy_view[:, :, k_block_next], tCsB_p[:, :, k_block_next]) + copyto!(smem_copy_A, tCrA_retiled[:, :, k_block_next], tCsA_p[:, :, k_block_next]) + copyto!(smem_copy_B, tCrB_retiled[:, :, k_block_next], tCsB_p[:, :, k_block_next]) if k_block == _1 && k_tile intrinsic) - - llvm_struct = Symbol("LLVMStruct$dest_sz") - ret_type = @eval $llvm_struct{$dest_type} - if isone(dest_sz) - @eval @inline function (::$(Symbol(ld_type)))(src_addr::$ptr_type) - return ccall($intrinsic, llvmcall, $dest_type, ($ptr_type,), src_addr) - end - - @eval function Base.copyto!(op::$(Symbol(ld_type)), dest::LocalArray{$dest_type}, src::SharedArray{$src_type}) - src_ptr = pointer(src) - val = op(recast(UInt32, src_ptr)) - return unsafe_store!(pointer(dest), val, 1) - end - else - @eval @inline function (::$(Symbol(ld_type)))(src_addr::$ptr_type) - return ccall($intrinsic, llvmcall, $ret_type, ($ptr_type,), src_addr) - end - - @eval function Base.copyto!(op::$(Symbol(ld_type)), dest::LocalArray{$dest_type}, src::SharedArray{$src_type}) - src_ptr = pointer(src) - val = op(recast(UInt32, src_ptr)) - dest_ptr = pointer(dest) - Base.Cartesian.@nexprs $dest_sz i -> unsafe_store!(dest_ptr, getfield(val, i), i) - return dest - end - end - end - return ld_ops +@inline function (::LDSM_U32x2_N)(src_addr::LLVMPtr{UInt32, AS.Shared}) + return ccall("llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.b16", llvmcall, LLVMStruct2{UInt32}, (LLVMPtr{UInt32, AS.Shared},), src_addr) +end + +@inline function (::LDSM_U32x4_N)(src_addr::LLVMPtr{UInt32, AS.Shared}) + return ccall("llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16", llvmcall, LLVMStruct4{UInt32}, (LLVMPtr{UInt32, AS.Shared},), src_addr) +end + +@inline function (::LDSM_U16x2_T)(src_addr::LLVMPtr{UInt32, AS.Shared}) + return ccall("llvm.nvvm.ldmatrix.sync.aligned.m8n8.x1.trans.b16", llvmcall, UInt32, (LLVMPtr{UInt32, AS.Shared},), src_addr) +end + +@inline function (::LDSM_U16x4_T)(src_addr::LLVMPtr{UInt32, AS.Shared}) + return ccall("llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.trans.b16", llvmcall, LLVMStruct2{UInt32}, (LLVMPtr{UInt32, AS.Shared},), src_addr) +end + +@inline function (::LDSM_U16x8_T)(src_addr::LLVMPtr{UInt32, AS.Shared}) + return ccall("llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16", llvmcall, LLVMStruct4{UInt32}, (LLVMPtr{UInt32, AS.Shared},), src_addr) +end + +function Base.copyto!(op::LDSM_U32x1_N, dest::LocalArray{UInt32}, src::SharedArray{UInt128}) + @inline + src_ptr = pointer(src) + val = op(recast(UInt32, src_ptr)) + return unsafe_store!(pointer(dest), val, 1) +end + +@inbounds function Base.copyto!(op::LDSM_U32x2_N, dest::LocalArray{UInt32}, src::SharedArray{UInt128}) + @inline + src_ptr = pointer(src) + val = op(recast(UInt32, src_ptr)) + Base.Cartesian.@nexprs 2 i -> dest[i] = getfield(val, i) + return dest end -const ldmatrix_ops_list = get_ldmatrix_ops() +@inbounds function Base.copyto!(op::LDSM_U32x4_N, dest::LocalArray{UInt32}, src::SharedArray{UInt128}) + @inline + src_ptr = pointer(src) + val = op(recast(UInt32, src_ptr)) + Base.Cartesian.@nexprs 4 i -> dest[i] = getfield(val, i) + return dest +end + +function Base.copyto!(op::LDSM_U16x2_T, dest::LocalArray{UInt32}, src::SharedArray{UInt128}) + @inline + src_ptr = pointer(src) + val = op(recast(UInt32, src_ptr)) + return unsafe_store!(pointer(dest), val, 1) +end + +@inbounds function Base.copyto!(op::LDSM_U16x4_T, dest::LocalArray{UInt32}, src::SharedArray{UInt128}) + @inline + src_ptr = pointer(src) + val = op(recast(UInt32, src_ptr)) + dest_ptr = pointer(dest) + Base.Cartesian.@nexprs 2 i -> dest[i] = getfield(val, i) + return dest +end + +@inbounds function Base.copyto!(op::LDSM_U16x8_T, dest::LocalArray{UInt32}, src::SharedArray{UInt128}) + @inline + src_ptr = pointer(src) + val = op(recast(UInt32, src_ptr)) + dest_ptr = pointer(dest) + Base.Cartesian.@nexprs 4 i -> dest[i] = getfield(val, i) + return dest +end """ copyto!(ldmatrix::AbstractLdMatrix, dest::MoYeArray{UInt32}, src::MoYeArray{UInt128}) @@ -85,7 +112,7 @@ Load data from shared memory to registers. The available `AbstractLdMatrix`s are You can inspect the number and the type of registers used per thread by ```julia julia> LDSM_U32x4_N() -LD_U32x4_N() +LDSM_U32x4_N() julia> ans.DRegisters Registers{UInt32, 4} @@ -94,3 +121,15 @@ Registers{UInt32, 4} function Base.copyto!(ldmatrix::AbstractLdMatrix, dest::MoYeArray, src::MoYeArray) throw(MethodError(copyto!, (ldmatrix, dest, src))) end + + +const ldmatrix_ops_list = [ + "LDSM_U32x1_N" => "llvm.nvvm.ldmatrix.sync.aligned.m8n8.x1.b16" +"LDSM_U32x2_N" => "llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.b16" +"LDSM_U32x4_N" => "llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16" +"LDSM_U16x2_T" => "llvm.nvvm.ldmatrix.sync.aligned.m8n8.x1.trans.b16" +"LDSM_U16x4_T" => "llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.trans.b16" +"LDSM_U16x8_T" => "llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16" +] + +export LDSM_U32x1_N, LDSM_U32x2_N, LDSM_U32x4_N, LDSM_U16x2_T, LDSM_U16x4_T, LDSM_U16x8_T \ No newline at end of file diff --git a/src/arch/mma/make_mma_ops.jl b/src/arch/mma/make_mma_ops.jl index c67b08ad..4416af6d 100644 --- a/src/arch/mma/make_mma_ops.jl +++ b/src/arch/mma/make_mma_ops.jl @@ -398,10 +398,10 @@ function make_mma_ops(geoms, types_a, types_b, types_c, types_d) return ccall($mma_intrinsic, llvmcall, $d_types, ($(a_types...), $(b_types...), $(c_types...)), $(a_vars...), $(b_vars...), $(c_vars...)) end - @eval @inline function fma!(op::$_struct_name, d::MoYeArray, a::MoYeArray, b::MoYeArray, c::MoYeArray) + @eval @inline @inbounds function fma!(op::$_struct_name, d::MoYeArray, a::MoYeArray, b::MoYeArray, c::MoYeArray) val = op(a,b,c) ptr = pointer(d) - Base.Cartesian.@nexprs $d_sz i -> unsafe_store!(ptr, getfield(val, i), i) + Base.Cartesian.@nexprs $d_sz i -> d[i] = getfield(val, i) return d end end diff --git a/src/atom/mma.jl b/src/atom/mma.jl index 5518ce77..98a94823 100644 --- a/src/atom/mma.jl +++ b/src/atom/mma.jl @@ -59,6 +59,7 @@ end function Base.show(io::IO, m::MMAAtom) println(io, "MMAAtom") + println(io, " Shape_MNK: ", shape_mnk(m)) println(io, " Thread ID: ", thr_id(m)) println(io, " Layout_A_TV: ", layout_a(m)) println(io, " Layout_B_TV: ", layout_b(m)) diff --git a/src/traits/ldmatrix.jl b/src/traits/ldmatrix.jl index a71b910b..eb6ee63a 100644 --- a/src/traits/ldmatrix.jl +++ b/src/traits/ldmatrix.jl @@ -16,7 +16,7 @@ function CopyTraits{LDSM_U32x4_N}() threadid = @Layout 32 srclayout = @Layout (32, 128) (128, 1) dstlayout = @Layout (32, (32, 4)) (32, (1, 1024)) - return CopyTraits{LDSM_U32x2_N}(threadid, srclayout, dstlayout, dstlayout) + return CopyTraits{LDSM_U32x4_N}(threadid, srclayout, dstlayout, dstlayout) end function CopyTraits{LDSM_U16x2_T}() diff --git a/test/device/matmul.jl b/test/device/matmul.jl index 969e0b00..b8bd8f82 100644 --- a/test/device/matmul.jl +++ b/test/device/matmul.jl @@ -44,7 +44,7 @@ function matmul_kernel(A, sA_layout, tA, cp_async_wait() sync_threads() - @gc_preserve gemm!(tCsA, tCsB, tCrC) + @gc_preserve gemm!(tCrC, tCsA, tCsB, tCrC) sync_threads() end diff --git a/test2.jl b/test2.jl new file mode 100644 index 00000000..107b15be --- /dev/null +++ b/test2.jl @@ -0,0 +1,146 @@ +using MoYe, CUDA, Test + +@views function matmul_kernel(A, sA_layout, gmem_copy_A, smem_copy_A, + B, sB_layout, gmem_copy_B, smem_copy_B, + C, mma_C) + sA = MoYeSharedArray(eltype(A), sA_layout) + sB = MoYeSharedArray(eltype(B), sB_layout) + + mA = MoYeArray(A) + mB = MoYeArray(B) + mC = MoYeArray(C) + + bM = size(sA_layout, 1) + bN = size(sB_layout, 1) + bK = size(sB_layout, 2) + + gA = @tile mA (bM, bK) (blockIdx().x, :) + gB = @tile mB (bN, bK) (blockIdx().y, :) + gC = @tile mC (bM, bN) (blockIdx().x, blockIdx().y) + + # gmem copy partition + gmem_thr_copy_a = get_slice(gmem_copy_A, threadIdx().x) + tAgA = partition_S(gmem_thr_copy_a, gA) # (CPY, CPY_M, CPY_K, k) + tAsA = partition_D(gmem_thr_copy_a, sA) # (CPY, CPY_M, CPY_K) + + gmem_thr_copy_b = get_slice(gmem_copy_B, threadIdx().x) + tBgB = partition_S(gmem_thr_copy_b, gB) # (CPY, CPY_N, CPY_K, k) + tBsB = partition_D(gmem_thr_copy_b, sB) # (CPY, CPY_N, CPY_K) + + # mma partition + thr_mma = get_slice(mma_C, threadIdx().x) + tCsA = partition_A(thr_mma, sA) # (MMA, MMA_M, MMA_K) + tCsB = partition_B(thr_mma, sB) # (MMA, MMA_M, MMA_K) + tCgC = partition_C(thr_mma, gC) # (MMA, MMA_M, MMA_N) + + tCrA = make_fragment_A(thr_mma, tCsA) # (MMA, MMA_M, MMA_K) + tCrB = make_fragment_B(thr_mma, tCsB) # (MMA, MMA_N, MMA_K) + tCrC = make_fragment_C(thr_mma, tCgC) # (MMA, MMA_M, MMA_N) + zeros!(tCrC) + + # retile + smem_thr_copy_A = get_slice(smem_copy_A, threadIdx().x) + smem_thr_copy_B = get_slice(smem_copy_B, threadIdx().x) + tCsA_retiled = partition_S(smem_thr_copy_A, sA) + tCsB_retiled = partition_S(smem_thr_copy_B, sB) + tCrA_retiled = retile_D(smem_thr_copy_A, tCrA) + tCrB_retiled = retile_D(smem_thr_copy_B, tCrB) + + + k_tile_max = size(tAgA, 4) + @cuprintln "k_tile_max: $k_tile_max" + k_tile = 1 + # for k_tile in 1:k_tile_max + # copy from global to shared + copyto!(gmem_copy_A, tAsA, view(tAgA, :, :, :, k_tile)) + copyto!(gmem_copy_B, tBsB, view(tBgB, :, :, :, k_tile)) + cp_async_wait() + sync_threads() + + # copy from shared to registers + copyto!(smem_copy_A, tCrA_retiled, tCsA_retiled) + copyto!(smem_copy_B, tCrB_retiled, tCsB_retiled) + # MoYe.copyto_unpack!(MoYe.CopyTraits{LDSM_U32x4_N}(), view(tCrB_retiled, (:,_1), _1, _1), view(tCsB_retiled, (:,_1), _1, _1)) + if threadIdx().x == 1 + @cuprintln "Thread 1" + + @cuprintln Int32(sB[17,1]), Int32(sB[17,2]), Int32(sB[17,3]), Int32(sB[17,4]), Int32(sB[17,5]), Int32(sB[17,6]), Int32(sB[17,7]), Int32(sB[17,8]), Int32(sB[17,9]), Int32(sB[17,10]), Int32(sB[17,11]), Int32(sB[17,12]), Int32(sB[17,13]), Int32(sB[17,14]), Int32(sB[17,15]), Int32(sB[17,16]) + @cuprintln Int32(sB[18,1]), Int32(sB[18,2]), Int32(sB[18,3]), Int32(sB[18,4]), Int32(sB[18,5]), Int32(sB[18,6]), Int32(sB[18,7]), Int32(sB[18,8]), Int32(sB[18,9]), Int32(sB[18,10]), Int32(sB[18,11]), Int32(sB[18,12]), Int32(sB[18,13]), Int32(sB[18,14]), Int32(sB[18,15]), Int32(sB[18,16]) + + @cuprintln Int32(tCrB[1,1,1]), Int32(tCrB[2,1,1]), Int32(tCrB[1,2,1]), Int32(tCrB[2,2,1]) + @cuprintln Int32(tCrB_retiled[1,1,1]), Int32(tCrB_retiled[2,1,1]), Int32(tCrB_retiled[3,1,1]), Int32(tCrB_retiled[4,1,1]) + end + if threadIdx().x == 25 + @cuprintln "Thread 25" + @cuprintln Int32(tCsB_retiled[1]), Int32(tCsB_retiled[2]), Int32(tCsB_retiled[3]), Int32(tCsB_retiled[4]) + end + @gc_preserve gemm!(mma_C, tCrC, tCrA, tCrB, tCrC) + + + @inbounds tCrC[1] # compiler bug, have to load after copyto! + + sync_threads() +# end + + copyto!(tCgC, tCrC) + @inbounds tCrC[1] # compiler bug, have to load after copyto! + + sync_threads() + return nothing +end + + +function matmul(A, B, C) + bM = _32 + bN = _32 + bK = _16 + + sA_atom_layout = @Layout (32, 8) (1, 32) + sB_atom_layout = @Layout (8, 16) (16, 1) + + sA_layout = MoYe.tile_to_shape(sA_atom_layout, (bM, bK)) + sB_layout = MoYe.tile_to_shape(sB_atom_layout, (bN, bK)) + + TA = eltype(A) + TB = eltype(B) + TC = eltype(C) + + gmem_copy_A = make_tiled_copy(CopyAtom{CPOP_ASYNC_CACHEALWAYS{UInt128}, TA}(), + @Layout((4, 8)), + @Layout((4, 1))) + gmem_copy_B = make_tiled_copy(CopyAtom{CPOP_ASYNC_CACHEALWAYS{UInt128}, TB}(), + @Layout((8, 4), (4, 1)), + @Layout((1, 4))) + + mma = make_tiled_mma(MMAOP_16x8x8_F32TF32TF32F32_TN(), + @Layout((1,2,1)), + (_32, _32, _8)) + + # Note: A is M-major so we can only use `UniversalCopy` + smem_copy_A = make_tiled_copy_A(CopyAtom{UniversalCopy{TA}, TA}(), mma) + smem_copy_B = make_tiled_copy_B(CopyAtom{LDSM_U32x4_N, TB}(), mma) + + threads = Int(size(mma)) + blocks = (cld(size(A, 1), bM), cld(size(B, 1), bN)) + + @cuda threads=threads blocks=blocks matmul_kernel(A, sA_layout, gmem_copy_A, smem_copy_A, + B, sB_layout, gmem_copy_B, smem_copy_B, + C, mma) +end + +function test() + M = 32 + K = 16 + N = 32 + A = CuArray(reshape(collect(1:M*K) .* 1f0, (M,K))) + B = CuArray(reshape(collect(1:N*K) .* 1f0, (K,N))) # K-major + C = CuArray(ones(Float32, (M,N))) + matmul(A, B', C) + CUDA.synchronize() + @test C == A * B + CUDA.unsafe_free!(A) + CUDA.unsafe_free!(B) + CUDA.unsafe_free!(C) +end + +test() \ No newline at end of file