翻译自: https://tilelang.tile-ai.cn/tutorials/debug_tools_for_tilelang.html

一个Tile Language程序（我们称为 program)到具体的硬件可执行文件的流程如下图所示，大致分为以下几步：1. 用户首先编写 Tile Language program。2. 程序会经过多个 Pass 的转换和优化处理（即 lower 阶段，相关代码位于 tilelang/engine/lower.py），最终生成中间代码，比如针对 CPU 的 LLVM 或 C 代码，或者针对 NVIDIA GPU 的 CUDA 代码等。3. 生成的中间代码会通过对应的编译器进一步编译，最终输出硬件可执行文件。

在这个过程中，用户可能会碰到大概三类问题:

1. Tile Language Program无法生成硬件可执行文件，也就是lower的过程中出现问题，我们可以归纳成生成问题。
1. 正确性问题，生成的可执行文件运行后，行为不符合预期。
1. 性能问题，执行文件的性能表现与硬件的理论值存在显著差距。

本文将重点讨论前两类问题的调试方法。至于性能问题的调优，则需要结合硬件厂商提供的性能分析工具（如 Nsight Compute、rocProf 等），通过分析具体的硬件指标进一步优化，我们将在后续文章中详细探讨。

接下来，我们以矩阵乘法（Matrix Multiplication）为例，使用 Tile Language 展示如何编写和调试相关程序。

Read More

不久之前的一篇分享里，我介绍了AMD CDNA架构(MI210, MI250, MI300)上的异步拷贝相关指令，在BitBLAS可以找到相关的实现，然而实际过程中发现AMD的异步拷贝指令的要求实际上要比那篇分享所写的更加苛刻，每个warp里的线程必须要求访问连续的数据，或者通过M0寄存器来控制每个线程的偏置。

一般来说，我们习惯这个指令就是明确的要load给定指针的一小块数据就行了，但是这个指令因为上述提到的两个限制就很难做到。经过笔者非常繁琐的Micro bencmark之后，笔者终于调教出了可以让每个线程Load给定数据块的写法，如下:

1
2
3
4
5
6
7
8
9
10
11
12

template <bool pre_nop = false>
CK_TILE_DEVICE void async_buffer_load_dword_v(void* smem, int32x4_t rsrc, index_t voffset) {
  auto const lds_ptr_sgpr = __builtin_amdgcn_readfirstlane((reinterpret_cast<uintptr_t>(smem)));
  asm volatile(
      "s_mov_b32 m0, %0; \n\t"
      "buffer_load_dword %1, %2, 0 offen lds;\n\t" ::"s"(lds_ptr_sgpr),
      "v"(voffset), "s"(rsrc)
      : "memory");
}
if constexpr(N == 4) {
  async_buffer_load_dword_v(lds_base_ptr, make_wave_buffer_resource(((int32_t *)global_base_ptr) - threadIdx.x), threadIdx.x * N /*assume 4 bytes*/);
}

在这篇文章里，笔者填一下AMD Matrix Core的坑，介绍一下过去一个月里BitBLAS针对AMD的的高性能Matrix Core支持，在这篇文章里笔者将介绍一下MFMA（AMD版的MMA）。如何进行AMD Kernel的性能分析，及Profile一个AMD Kernel，最后我们介绍若干种绞尽了笔者脑汁的优化方法，完全利用好硬件的带宽(全都是128bits的内存访问指令，并且没有Memory bank conflict)。

这篇文章涉及到的算子有矩阵乘法和Flash Attention。本篇文章的实现在BitBLAS里, Codegen以及Swizzle等Layout变换依托于TVM, TVM可以帮助我们显式地操作一个数据的Layout，相比Triton更加灵活和可观。虽然AMD提供的文档十分有限，但是在这一个月里笔者参考了很多AMD开发人员提供的实现，例如Composable Kernel和Triton for ROCm，笔者从这些项目中收获良多。

本文假设读者对Nvidia GPU的编程有一定的了解，熟悉最基本的Tile优化程序的方法，以及Tensor Core的基本概念。

Read More

最近给BitBLAS添加了AMD的后端，发现AMD的异步拷贝等和Nvidia有很大的不同(但是FA3在MI300上需要用到这一个Feature)，然而官方根本没有文档，只有Instruction Set，我在这里做一下自己的理解和解读，大部分内容是参考自这个Instruction Set。

Read More

合适地选择Block Size和Grid Size

Block Size是指一个Block中的thread个数，增大block的数量一方面有利于提高程序的并行性，但是如果同一个block的thread之间存在线程的同步，则过大的block size会带来同步的overhead，导致SM利用率降低，而Grid Size是指Block的数量，如何好的

一台Windows Server 2022，上面装了两块AMD Radeon™ VII的显卡，操作系统的和显卡的版本都很新，导致出现了一些错误，不过最后还是成功安装上了。

Read More

最近在使用NNFusion的时候发现Codegen出来的FP16的网络在V100上的性能打不过FP32（甚至要慢一倍以上），但是理论上FP16应该要比FP32有两倍的性能收益才对（V100 Cuda Core的half precision的最大吞吐量是single的两倍，在s9234的slides中看到直接使用half的情况下peak performance其实和single差不多，都是15Tops，但是Cuda Core提供了half2类型，可一次做两个half类型的运算，这是half在CUDA Core上的收益来源；V100卡上的Tensor Core只支持FP16，利用好Tensor Core可以获得非常强的加速，A100卡上的Tensor Core增加更多的精度支持）。

建议阅读的文章：

聊聊 GPU 峰值计算能力

A100 Tensor Float 32 性能实测

拿nvprof测试了一下发现主要的性能瓶颈是：half卷积算子的实现速度要比single慢一倍，而这部分运算又占了总体运行时间的绝大部分。

Read More