cute是一系列C++ CUDA的模板抽象,用于定义和操作线程和数据的层次化多维layouts。cute提供 Layout
以及 Tensor
对象;其中 Tensor
将 数据类型、形状、内存空间以及数据的layout进行一个封装,用户可以对这个 Tensor
进行复杂的索引操作。这样可以使得用户可以专注于逻辑上的描述,底层对应到内存实际的索引由cute来进行resolve;可以基于layout定义出tiling
以及partition
的操作。
前置的环境配置
C++17
library的组织形式
cute 是一个header-only C++ lib;所以不需要构建其源码。
include/cute
: 顶层目录下的每个header都对应cute中的基本构件,例如Layout
和Tensor
.include/cute/container
: 实现STL-like的对象,例如tuple,array以及aligned arrayinclude/cute/numeric
: 基础的数值类型,也包括非标准的floating-point 类型include/cute/algorithm
: 一些算法的实现例如copy
,fill
include/cute/arch
: 对架构特定的 matrix-matrix multiply以及copy指令的封装include/cute/atom
:arch
下指令的元信息以及partition
和tiling
的功能
tips
如何在host或device上打印cute 对象
cute::print
函数对于几乎所有的cute类型,包括pointers, integers, strides, shapes, layouts, and tensors,都有重载版本。除此之外,这个函数可以在host或者device上工作;但需注意在device上时,print的开销很大;注意在debug之后移除对应的print函数。
你也许只想对每个thread block的第0个线程或者grid的第一个block应用print。thread0()
只对于kernel的全局0号线程返回true
if (thread0()) {print(some_cute_object);
}
一些算法可能依赖某个线程或线程块,所以你可能需要打印别的线程或线程块;cute/util/debug.hpp
包含 bool thread(int tid, int bid)
函数。
其他的打印格式
cute::print_layout
: 以表格形式展示任意rank-2 的layoutcute::print_tensor
: 展示任意 rank-1 rank-2 rank-3 rank-4的tensor,打印出tensor中的值cute::print_latex
:Layout
TiledCopy
TiledMMA