Vortex GPU 模拟器学习

环境配置

vortex 环境配置总体来说比较流畅，参考官方给出的教程即可。但是经过尝试，发现这样跑不通。排查之后发现有两个环境变量需要设置：

• OSVERSION：设置为 export OSVERSION="ubuntu/bionic"
• TOOLDIR：设置为 export TOOLDIR=$HOME/tools

这两个环境变量最好写入 ~/.bashrc 中。

然后出现了 libtinfo 的报错（error while loading shared libraries: libtinfo.so.5: cannot open shared object file: No such file or directory），可以通过下面的方法解决：

sudo apt install libncurses5 # 经测试对 22.04 版本有效

Assignment 6 点积加速（SimX）

这项作业介绍了基础的 GPU 指令集的扩展方法，用以加速硬件中的 kernel。你将添加一条用于计算整数点积新的 RISC-V 自定义指令集：VX_DOT8。你还将在 SimX 这一周期精确模拟器中实现此指令。

步骤 1 ISA 扩展

VX_DOT8 计算两个由 $4 \times 4$ int8 整数组成的向量的点积。

Dot Product = (A1*B1 + A2*B2 + A3*B3 + A4*B4)

指令格式如下所示：

VX_DOT8 rd, rs1, rs2

其中每个源寄存器 rs1 和 rs2 保存四个 int8 元素。

rs1 := {A1, A2, A3, A4}
rs2 := {B1, B2, B3, B4}
rd  := destination int32 result

使用R-Type RISC-V指令格式（RISC 指令集共有 R I J 三种常见的指令格式）。

| funct7  | rs2    | rs1    | funct3 | rd    | opcode |
|  7 bits | 5 bits | 5 bits | 3 bits | 5 bit | 7 bits |

其中：

opcode: opcode reserved for custom instructions.
funct3 and funct7: opcode modifiers.

使用自定义扩展 opcode=0x0B，其中 func7=1 和 func3=0；（后面译码要用）

需要修改 kernel/include/vx_intrinsics.h 以添加新的 VX_DOT8 指令。

inline int vx_dot8(int a, int b) {
    size_t ret;
    asm volatile (".insn r 0x0B, 0x0, 0x1, %0, %1, %2" : "=r"(ret) : "r"(a), "r"(b)); # ret 为返回结果
    return ret;
}

阅读以下文档以了解 insn speudo 指令格式 https://sourceware.org/binutils/docs/as/RISC_002dV_002dFormats.html

步骤 2 矩阵乘法 kernel

实现一个使用新 H/W 扩展的简单矩阵乘法 GPU 内核。

下面是使用我们的新 VX_DOT8 指令的 kernel 的基本 C++ 实现：

void MatrixMultiply(int8_t A[][N], int8_t B[][N], int32_t C[][N], int N) {
  for (int i = 0; i < N; ++i) {
    for (int j = 0; j < N; ++j) {
      C[i][j] = 0;
      for (int k = 0; k < N; k += 4) {
        // Pack 4 int8_t elements from A and B into 32-bit integers
        uint32_t packedA = *((int*)(A[i] + k));
        uint32_t packedB = (uint8_t)B[k][j]
                         | ((uint8_t)B[k+1][j] << 8)
                         | ((uint8_t)B[k+2][j] << 16)
                         | ((uint8_t)B[k+3][j] << 24);
        // Accumulate the dot product result into the C
        C[i][j] += vx_dot8(packedA, packedB);
      }
    }
  }
}

将 https://github.com/vortexgpgpu/vortex/blob/master/tests/regression/sgemmx 下的 sgemmx 测试克隆到新文件夹 tests/regressions/dot8 中（用以建立测试 kernel）。

在 tests/regressions/dot8/Makefile 中将 PROJECT 名称设置为 dot8

ROOT_DIR := $(realpath ../../..)
include $(ROOT_DIR)/config.mk

PROJECT := vxdot8 # 设置项目名称

SRC_DIR := $(VORTEX_HOME)/tests/regression/$(PROJECT)

SRCS := $(SRC_DIR)/main.cpp

VX_SRCS := $(SRC_DIR)/kernel.cpp

OPTS ?= -n128

include ../common.mk

更新 main.cpp 中的 matmul_cpu 以对 int8_t 矩阵进行操作。

static void matmul_cpu(TYPE_DST* out, const TYPE_SRC* A, const TYPE_SRC* B, uint32_t width, uint32_t height) {
  TYPE_SRC* Bcopy = new TYPE_SRC[width * height];
  for (uint32_t row = 0; row < height; ++row) {
    for (uint32_t col = 0; col < width; ++col) {
        Bcopy[col * width + row] = B[row * height + col];
    }
  }
  for (uint32_t row = 0; row < height; ++row) {
    for (uint32_t col = 0; col < width; ++col) {
      TYPE_DST sum(0);
      for (uint32_t e = 0; e < width; ++e) {
          // B is transposed
          // sum += A[row * width + e] * B[e * width + col];
          sum += A[row * width + e] * Bcopy[e * width + col];
      }
      out[row * width + col] = sum;
    }
  }
}

更新 tests/regressions/dot8/kernel.cpp 中的 kernel_body 以使用 vx_dot8

void kernel_body(kernel_arg_t* __UNIFORM__ arg) {
    auto A = reinterpret_cast<TYPE_SRC*>(arg->A_addr);
    auto B = reinterpret_cast<TYPE_SRC*>(arg->B_addr);
    auto C = reinterpret_cast<TYPE_DST*>(arg->C_addr);
    auto size = arg->size;
    
    // N * N kernels
    // nous avons envie de savoir quel kernel on est
    auto grid_x = blockIdx.x;
    auto grid_y = blockIdx.y;

    auto r = grid_x * size;
    auto c = grid_y * size;
    auto startA = A + r;
    auto startB = B + c;

    TYPE_DST sum = 0;
    
    for (int k = 0; k < size; k += 4) {
        sum += vx_dot8(*((int*)(startA + k)), *((int*)(startB + k)));
    }

    *(C + r + grid_y) = sum;
}

这里的 kernel 和 cuda 的 kernel 类似。

步骤 3 对实现进行仿真

修改循环级模拟器以实现自定义 ISA 扩展。我们建议参考 SimX 中 VX_SPLIT 和 VX_PRED 指令的解码方式。

• 更新 sim/simx/decoder.cpp 中的 op_string() 函数，打印出新的指令。

  case Opcode::EXT1:
  switch (func7) {
  ...
  case 1:
    switch (func3) {
      case 0: return "VX_DOT8";
      default: 
        std::abort();
    }
  default:
    std::abort();
  }

• 更新 sim/simx/decoder.cpp 中的 Emulator::decode() 函数，解码新的指令格式。

  switch (func7) {
  case 1:
    switch (func3) {
    case 0:  // DOT8
      instr->setDestReg(rd, RegType::Integer);
      instr->addSrcReg(rs1, RegType::Integer);
      instr->addSrcReg(rs2, RegType::Integer);
      break;

• 更新 sim/simx/types.h 中的 AluType 枚举，添加 DOT8 类型。

  enum class AluType {
    ARITH,
    BRANCH,
    SYSCALL,
    IMUL,
    IDIV,
    DOT8
  };

• 更新 sim/simx/execute.cpp 中的 Emulator::execute()，实现实际的 VX_DOT8 仿真。你将在 ALU 功能单元上执行新的指令。

  switch (func7) {
  case 1:
    switch (func3) {
    case 0: { // DOT8
        trace->fu_type = FUType::ALU;
        trace->alu_type = AluType::DOT8;
        trace->src_regs[0] = {RegType::Integer, rsrc0};
        trace->src_regs[1] = {RegType::Integer, rsrc1};
        for (uint32_t t = thread_start; t < num_threads; ++t) {
          if (!warp.tmask.test(t))
            continue;
          // TODO: 执行 int 8 计算
          int8_t a, b;
        	int32_t r = 0;
        	for (uint32_t _i = 0; _i < 4; _i++) {
            a = (rsdata[t][0].i >> (8*_i)) & 0xFF;
            b = (rsdata[t][1].i >> (8*_i)) & 0xFF;
            r += static_cast<int32_t>(a) * static_cast<int32_t>(b);
          }
          rddata[t].i = r;  
        }
        rd_write = true;
      } break;
    } break;
  }

更新 func_unit.cpp 中的 AluUnit::tick() 函数，以实现 VX_DOT8 的时序。假设点积运算的执行延迟为 2 个周期。

case AluType::DOT8:
    // TODO: 时延 + 2
    output.push(trace, 2+delay);
 	break;

步骤 4 测试

你将在回归代码库下将新的加速 dot8 程序与现有的 sgemmx 内核进行比较。你将使用 N=128、(warps=4,threads=4) 和 (Warps=16,threads=16) 分别对应 1 核和 4 核。绘制总执行周期图以观察性能提升。