不久之前的一篇分享里，我介绍了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

刚开始碰到的问题是这样的：在Azure上开的一台HPC（4块 V100 16G）在运行了大概七八个小时之后，nvidia的显卡会挂掉，具体的表现为nvidia-smi会卡住十几分钟，之后输出No devices were found，但是执行lspci | grep -i nvidia还是可以看到四块显卡好好的挂在上面，这种情况应该直接reboot就可以修复，但是reboot了之后同样的程序运行一段时间之后显卡还是会掉。

Read More