【Cutlass】tile iterator concept

最新推荐文章于 2025-05-28 22:18:33 发布

CoCo_DY

最新推荐文章于 2025-05-28 22:18:33 发布

阅读量1k

点赞数

文章标签： c++ 计算机体系架构 cuda

本文链接：https://blog.youkuaiyun.com/qq_33146555/article/details/130492104

版权

Cutlass库在不同版本间对矩阵和张量操作进行了调整，从2.x版本的TileIterator到3.0及以后的cute::Tensor。TileIteratorConcept是一个核心概念，它定义了不同类型的迭代器行为，如Element、ContiguousMemoryTileIterator、ReadableTileIteratorConcept等，支持前向、双向和随机访问。MaskedTileIteratorConcept则引入了掩码机制，用于处理边界安全问题。

摘要生成于 C知道，由 DeepSeek-R1 满血版支持，前往体验 >

简介

在Cutlass 2.x之前将matix或者tensor切分成固定尺寸的操作都是在tile iterator当中（Cutlass 3.0之后对于matrix和tensor的操作都放到了cute::Tensor里面）。tile_iterator_concept不同于c++20里面的concept，这里不使用concept作为key word，仅仅是对这些类型的需求。tile_iterator内部是几个概念的组合，tile iterator concept几乎涵盖了所有需求。

Concept定义

struct	定义
TileIteratorConcept	基类：1. Element，tile的组成部分；2. 形状类型，取决于iter的实现
ContiguousMemoryTileIterator	在连续内存上任意Tile的iterator，每个thread将Element尺寸的数据块存储在pointer_offset的位置上
ReadableTileIteratorConcept（WriteableTileIteratorConcept类似）	每个thread从内存读取一块数据放到其相应的Fragment当中
ReadableContiguousTileIteratorConcept（WriteableContiguousTileIteratorConcept类似）	从pointer_offset的位置连续读取数据
ForwardTileIteratorConcept	前向遍历tile，具体的方向和Context相关，比如Gemm就这沿着K方向遍历
BidirectionalTileIteratorConcept	两个方向遍历tile
RandomAccessTileIteratorConcept	相较TensorRef进行TensorCoord的偏移
ReadableRandomAccessTileIteratorConcept（WriteableRandomAccessTileIteratorConcept类似）	根据TensorRef和TensorCoord load数据到fragment中
ReadableRandomAccessContiguousTileIteratorConcept（WriteableRandomAccessContiguousTileIteratorConcept类似）	根据TensorRef和TensorCoord load pointer_offset处数据到fragment中
MaskedTileIteratorConcept	为了部分获取数据设置的mask，可以作为oob保护
WriteableReadableForwardContiguousTileIteratorConcept	ForwardTileIteratorConcept + ReadableContiguousTileIteratorConcept + WriteableContiguousTileIteratorConcept
WriteableReadableRandomAccessContiguousTileIteratorConcept	ReadableRandomAccessContiguousTileIteratorConcept + WriteableRandomAccessContiguousTileIteratorConcept