github – 测试代码 编译时，加上 -lineinfo 参数，Nsight Compute 分析时，可以看到具体的 C++/cu 代码。 1. 查看 Warp State Statistics Warp State Statistics 表征从执行/issue当前指令到执行/issue下一条指令之间的周期数。导致指令周期长的原因有多种，比如等待内存访问、指令流水线Stall等（...

1. navie tile GEMM 代码文件：navie tile GEMM 基于分块矩阵乘法的简单实现，按照 Thread Block 将矩阵划分为多个tile进行计算，在 Thread Block内，再次将 tile 划分为多个子块，由每个线程负责计算子块。 使用 Shared Memory 来缓存 tile 数据，减少全局内存访问次数。每个线程负责从全局内存中复制 ti...

GEMM Pipeline

TV-Layout 描述 CTA 中线程的 layout，以及每个线程可以访问到哪些数据。TV-Layout 的第一个 mode 定义线程在 CTA 中的分布，第二个 mode 定义每个线程处理的数据布局。见下面例子中的LayoutA_TV。 Inverse TV-Layout 描述的是数据的逻辑坐标 coord 到线程ID的映射关系。比如给定 layout 的逻辑坐标 (m, n)，经过...

给定 layout 范围内，swizzle 通过列异或操作（icol = irow ^ icol），周期性的 coord 重排，映射到新的物理地址 offset。swizzle 定义了三个参数： M：以 $2^M$ 个一维坐标连续的元素为单位，将其当做一个元素； S：控制行号、列号提取的低位偏移； B：参与 XOR 的位数，即掩码位数，用于提取一维 index 中的行号、列号...

tiled_copy.cu：官方示例 1. Cute TiledCopy 层次化的 copy 抽象，将 Copy_Atom 分为几个可组合的层次： CopyOperation：NVidia在不同的硬件架构、不同的存储层次之间数据搬运提供了不同的指令，如 ldmatrix 和 LDS 等，还有针对Ampere架构的 cp.async 等。 Copy_Traits：和 MMA_...

github – 测试代码 CuTe Tensors：官方文档 1. CuTe 中的 Tensor 划分 (Partitioning a Tensor) 在 GEMM 计算是，需要对矩阵进行划分（分块），以便线程块（Thread Block）和线程（Thread）能够并行处理数据。常用的有 Inner-Partitioning、Outer-Partitioning，以及 TV-la...

github – Layout 测试代码 1. Layout Cute(CUDA Tensor) 是 CUBLASS 扩展，用于简化张量 BLAS 操作和内存布局的管理。 最主要的概念是 Tensor 和 Layout： Layout<Shape, Stride>: 定义张量的内存布局，用于将一维内存地址映射到多维张量索引。 Shape：Lo...

测试用例：github – cuda_perf 1. 启动配置 使用Nsight Compute运行被测试CUDA程序，启动时，指定metrics为full： 2. 分析Bank Conflicts 被测试CUDA程序运行结束后，打开Nsight Compute的结果页面（Details），进入Memory Workload Analysis章节，在章节标题右侧，选择Memory...

使用到的测试代码：bank_conflict.cu 1. Bank Conflicts (Shared Memory) 1.1. Bank 划分 针对Shared Memory的访问，CUDA使用bank机制，将shared memory的访问（读/写）映射到不同的bank，以实现并行访问。bank以4字节为单位，共32个bank。这样，一个时钟周期内，可以并行访问32个不同的bank...