2024年8月19日 · CuTe, 全称为 "collection of C++ CUDA template abstractions for defining and operating on hierarchically multidimensional layouts of threads and data", 是一个处理嵌套layout的模板抽象的集合, 其并不支持提供现成算子支持, 而是给出耳目一新的数据结构,使得复杂的线性代数 …

2024年3月21日 · atom tile主要是指上述的op原子计算在M和N方向上各自拓展多少次，在tensor core中，每个op都是一个warp，即有32个线程；M和N方向各拓展两次，就有4个warp，128个线程，所以也叫thr_layout; value tile主要是指拓展后的atom，在M和N方向上继续重复多少次计算，因为是重复，内部是loop操作，所以不会占用更多的线程，只会扩大处理的矩阵大小。 最下面就是最小的tensor core指令—-MMA_Atom（16816），小框是元素级别的计算（这个可以先不关 …

cute Copy抽象及其相互关系. 和MMA类似，cute对数据搬运提供了对数据搬运的数据结构抽象，主要包括 CopyOperation 、Copy_Traits、Copy_Atom、TiledCopy、ThrCopy和拷贝函数 cute::copy 。这些结构和函数共同完成对GPU各个层级存储之上的数据进行搬运的抽象和实现，具体地，

在看代码之前，要简单说几个 cute 里的概念：tensor、layout、shape 和 stride。 他们之间有 2 个非常简单的指代关系： 就是说，所谓 tensor，就是一个内存指针（ptr，其实是显存，不过为求方便，后面均称内存）和这个指针对应的内存排布的描述信息（layout）；而所谓排布，就是说一个 tensor 的形状（shape），即这个 tensor 是几维的，每一维多大，以及每一维的各个元素之间间隔多大（stride）。 啥叫元素间隔多大呢？

cute提供了MMA能力来完成D = A x B + C的矩阵乘法运算，其针对指令封装，适配层，原子能力、块状MMA、线程划分和执行进行了抽象，形成了MMAOperation、MMA_Traits、MMA_Atom、TiledMMA、ThrMMA、cute::gemm数据结构和函数，我们通过这些结构能够完成逻辑块状矩阵乘 …

从gemm的角度出发看多维thread tiling的实现。使用cute::copy把smem中的数据tCsA拷贝到寄存器中tCrA后直接使用cute::gemm做多维thread tiling的gemm计算。具体thread tiling的布局通过可以通过打印mma查看。

2024年5月8日 · 从gemm的角度出发看多维thread tiling的实现。使用cute::copy把smem中的数据tCsA拷贝到寄存器中tCrA后直接使用cute::gemm做多维thread tiling的gemm计算。具体thread tiling的布局通过可以通过打印mma查看。

2024年10月29日 · CuTe是用于内存排布和张量表示的模板工具,cutlass是从多个层面逐层分解逐层解耦的高性能模板库,CuTe有利于代码编写时不用过于关心一些数据偏移的计算,而cutlass则将高性能算子实现拆解,如果要单独实现一个特定的,就得按照它的代码规范实现一个,不然只能用它已 ...

2023年8月5日 · An atom in CUTLASS and CuTe is defined as the smallest number of threads and data that must participate together to execute an architecture intrinsic copy/math operation.

本文旨在探讨如何形成一个新的 CUTE MMA 描述，和怎么控制排布，以及配套的 SmemCopyAtom。 期望可以为拿 CUTE 来适配新架构/希望更自由地组合MMA以获得更高表述上限的读者做出一点帮助。