NaplesPU或NPU技术开发文档（草稿）

NaplesPU或NPU技术开发文档（草稿）

http://www.naplespu.com/doc/index.php?title=Main_Page

http://www.naplespu.com/

https://github.com/AlessandroCilardo/NaplesPU

https://github.com/AlessandroCilardo/NaplesPU-toolchain

主页

那不勒斯处理单元，被称为NaplesPU或NPU，是一个全面的开源多核加速器，涵盖了从计算核心到片上互连、一致性存储器层次结构和编译工具链的所有架构层。NaplesPU完全用系统Verilog HDL编写，利用了现代计算架构中通常存在的三种并行形式，特别是在GPU设备等异构加速器中：向量并行、硬件多线程和多核组织。配备了一个完整的基于LLVM的编译器，针对NaplesPU矢量ISA，NPU开源项目，可体验许多核心技术的所有风格。

NPU多核架构基于片上网络（NoC）连接，可配置块参数化网格。每个图块都有一个缓存控制器和一个目录控制器，处理不同图块中不同内核之间的数据一致性。计算核基于轻量级控制单元的向量流水线，以便将大部分硬件资源用于加速数据并行内核。利用硬件多线程掩盖了内存操作和长延迟指令。每个硬件线程（大致相当于OpenCL术语中的波前或NVIDIA术语中的CUDA扭曲）都有自己的PC、寄存器文件和控制寄存器。NaplesPU系统中的线程数是用户可配置的。

 

目录

1开始

1.1所需软件

1.2建造过程

1.3模拟内核

1.3.1 test.sh脚本

1.3.2 setup_project.sh脚本

1.3.3模拟.sh脚本

2完整文档

3更多关于MediaWiki的信息

开始使用

本节将展示如何接近该项目，以模拟或实现NaplesPU架构的内核。内核是指用高级编程语言（如C/C++）编写的复杂应用程序，如矩阵乘法、矩阵转置或类似应用程序。

所需软件

任何内核的模拟或实现都依赖于以下依赖关系：

Git

Xilinx Vivado 2018.2或ModelSim（例如Questa Sim-64 vsim 10.6c_1）

NaplesPU工具链

构建过程

第一步是通过克隆[1]从官方存储库中获取NaplesPU架构的源代码

在Ubuntu Linux环境中，通过启动以下命令来完成此步骤：

$git克隆https://github.com/AlessandroCilardo/NaplesPU

在NaplesPU存储库中，工具链是存储库的git子模块，因此需要创建和更新。在Ubuntu Linux环境中，只需在存储库的根文件夹中键入以下命令：

$git子模块更新--init

然后，第三步是安装工具链。[此处]描述了这一过程。

模拟内核

以下文件夹对此特别感兴趣：

软件，存储所有内核；

工具，存储所有用于模拟的脚本。

模拟内核有三种方法：

启动test.sh脚本

如果模拟器软件是Vivado，则从存储库的根文件夹启动setup_project.sh；

如果模拟器软件是ModelSim，则从存储库的根文件夹启动simulate.sh。

首先，在shell中源代码Vivado或ModelSim。这一步对所有方式都是强制性的。在Ubuntu Linux环境中：

$source Vivado/文件夹/位置/settingXX.sh

其中XX取决于安装的Vivado版本（32或64位）。

test.sh脚本

test.sh脚本位于npu/tools文件夹中，运行其中列出的所有内核，并将npu的输出与标准x86架构产生的预期结果进行比较：

$ ./test.sh[选项]

选项包括：

-h、 --帮助显示此帮助

-t、 --tool=vsim或vivado指定要使用的工具，默认值：vsim

-cn，--core麻木=VALUE指定核心编号，默认值：1

-tn，--thread-numb=VALUE指定线程号，默认值：8

test.sh脚本会自动编译内核，并在NaplesPU和x86架构上运行它们。一旦模拟终止，对于每个内核，Python脚本会比较两次执行的结果，以验证其正确性。

在tools文件夹中，文件cosim.log存储模拟器的输出。

setup_project.sh脚本

setup_project.sh脚本可以从项目的根目录按如下方式运行：

$tools/vivado/setup_project.sh[选项]

选项包括：

-h、 --帮助显示此帮助

-k、 --kernel=kernel_NAME指定要使用的内核

-s、 --单核选择单核配置，默认情况下选择多核

-c、 --core mask=VALUE指定核心激活掩码，默认值：1

-t、 --thread mask=VALUE指定线程激活掩码，默认为FF

-m、 --mode=gui或batch指定工具模式，它可以在gui或batch模式下运行，默认值：gui

此脚本启动命令中指定的内核。内核在NaplesPU架构上运行之前应该已经编译好了：

tools/vivado/setup_project.sh-k mmsc 3-t$（（16#F））-m gui

参数-c 3传递了核心激活的一个热掩码：3是（11）2，因此图块0和1将启动它们的核心。参数-t$（（16#F））表示每个核心的活动线程掩码，它是一个单热掩码，表示每个核心中哪个线程是活动的：F是（00001111）2，因此线程0到3正在运行。参数-m gui表示模拟器执行的模式。

simulate.sh脚本

simulate.sh脚本可以从项目的根目录按如下方式运行：

$tools/modelsim/simulate.sh[选项]

选项：

-h、 --帮助显示此帮助

-k、 --kernel=kernel_NAME指定要使用的内核

-s、 --单核选择单核配置，默认情况下选择多核

-c、 --core mask=VALUE指定核心激活掩码，默认值：1

-t、 --thread mask=VALUE指定线程激活掩码，默认为FF

-m、 --mode=gui或batch指定工具模式，它可以在gui或batch模式下运行，默认值：gui

此脚本启动命令中指定的内核。内核在NaplesPU架构上运行之前应该已经编译好了：

完整文档

NaplesPU硬件架构

NaplesPU工具链

NaplesPU指令集架构

加长型NaplesPU

异质瓷砖

编程模型

更多关于MediaWiki的信息

NaplesPU项目文档将基于MediaWiki。有关使用MediaWiki的信息和指南，请参阅以下链接：

用户指南（有关使用维基软件的信息）

配置设置列表

MediaWiki常见问题

MediaWiki发布邮件列表

将MediaWiki本地化为语言

了解如何在wiki上打击垃圾邮件

NaplesPU硬件架构

（重定向自nu+硬件架构）

NaplesPU多核是一个基于可配置瓦片的规则网状片上网络的模块化和深度可定制系统，旨在从头开始成为一个可扩展和参数化的平台，适合探索先进的构建解决方案。NaplesPU是在MANGO FETHPC项目的框架内开发的，其主要目标是基于专用定制硬件实现资源高效的HPC。这导致了模块化设计和指令布局，提供了足够的自由度来扩展标准指令集和基线NPU硬件设计，下文将对此进行讨论。该项目旨在提供一个完全可定制且易于扩展的多核心系统，适用于探索未来系统的软件和硬件高级解决方案。

NaplesPU的主要目标是基于专用定制硬件实现资源高效的HPC。目标是构建一个应用程序驱动的架构，为任何数据并行内核实现最佳的硬件/软件配置。众所周知，对于具有大量常规数据级并行性（DLP）的代码，专用数据并行加速器比通用处理器提供更高的效率。然而，每个并行内核都有自己的理想配置。

每个NPU瓦片都有相同的基本组件，它提供了一个可配置的GPU类开源软核，旨在用作可配置的FPGA覆盖层。这个面向HPC的加速器将SIMT范式与向量处理器模型融合在一起。此外，每个图块都有一个缓存控制器和一个目录控制器，这些组件处理不同图块中不同内核之间的数据一致性。

 

目录

1瓷砖概述

2硬件部分

2.1通用组件

2.2多核心系统

2.3单核版本

2.4测试一致性子系统

3配置NaplesPU

3.1更改线程数量

瓷砖概述

上图显示了NaplesPU多核的简化概述。每个NPU瓦片都有相同的基本组件，它提供了一个可配置的GPU式加速器，用作可配置的FPGA覆盖层，一个独立于协议的可扩展相干子系统，以及一个基于网格的网络系统，该系统通过通信网络路由硬件消息。

该系统基于依赖于片上网络（NoC，在网络架构部分中描述）的异构块的2D网格。NoC路由器与网络接口模块紧密耦合，通过四个不同的虚拟信道提供基于分组的通信。使用两级先行路由器实现基于虫洞微片的通信。网络基础设施允许块内和块间通信。一个虚拟通道专用于服务消息流。特别是，多核系统支持基于硬件屏障的分布式同步机制（部分同步）。重要的是，基于简单的有效/就绪接口，可以在每个图块中集成任何处理单元。

加速器将SIMT范式与矢量处理器模型融合在一起。类似GPU的模型展示了提高资源效率的有前景的功能。事实上，它提供了与SIMD执行单元耦合执行的硬件线程，同时减少了控制开销并隐藏了可能的长延迟。该加速器有效地利用了多线程、SIMD操作和低开销控制流构造，以及一系列高级架构定制功能，以实现底层资源的高级利用。

为了确保可扩展性，NaplesPU manycore实现了稀疏目录方法和分布式L2缓存，如Coherence架构部分所述。每个图块都部署了一个一致性维护基础设施和加速器。缓存控制器处理本地处理单元的内存请求，将加载和存储未命中转换为网络上的目录请求。它还处理来自网络的响应和转发请求，根据给定的一致性协议更新块状态，而加速器完全不知道这一点。

缓存和目录控制器的架构在设计时考虑了灵活性，不受任何特定一致性协议的约束。它们配备了一个可配置的协议ROM，为一致性协议扩展提供了一种简单的方法，因为它精确地描述了基于当前块状态对每个请求采取的行动。

硬件部分

本节介绍NaplesPU项目中的主要硬件组件及其交互。

常见组件

下文将描述整个设计中使用的基本组件，以及提供的自定义类型和GPGPU内核。这些组件也是项目单核版本的一部分（稍后描述）。

包括

基本组件

NaplesPU GPGPU核心架构

多核心系统

下文将介绍许多核心高级功能，如相干子系统和同步机制。

系统

项目界面

一致性架构

同步架构

网络架构

单核版本

该项目提供了一个单核版本，具有GPGPU内核和简化的缓存系统，下文将对此进行描述。

系统

项目界面

缓存控制器

同步

记录仪

FPGA上的系统部署

测试一致性子系统

一致性子系统附带了一个专用的测试台，如下所述。

相干注入

配置NaplesPU

NaplesPU最重要的方面之一是参数化。通过更改头文件中的相应值，可以扩展许多功能，例如缓存维度。用户设计可以为每种需求设置大量参数，例如：

NoC拓扑和Tile编号相关功能（在npu_user_defines.sv头文件中）：

NoC_X_WIDTH-X维度上的图块数量，必须是2的幂

NoC_Y_WIDTH-Y维度上的图块数量，必须是2的幂

TILE_MEMORY_ID-内存控制器的图块ID，定义NoC中的位置

TILE_H2C_ID-主机接口的瓦片ID，定义NoC中的位置

TILE_NPU-系统中NPU磁贴的数量

TILE_HT-系统中异构图块的数量

核心相关参数：

THREAD_NUMB-每个核心的线程数，位于npu_user_defines.sv头文件中。

NPU_FPU-在NPU核心中分配浮点单位。

NPU_SPM-在每个NPU内核中分配一个Scratchpad内存。

HW_LANE-在npu_Defines.sv头文件中定义SIMD扩展的宽度。

REGISTER_NUMBER-npu_defines.sv头文件中标量和矢量寄存器文件中的寄存器数。注意，改变寄存器的数量会改变特殊用途寄存器（如PC或SP）上的位置，编译器必须相应地进行修改。

缓存相关参数（在npu_user_defines.sv头文件中）：

USER_ICACHE_SET-指令缓存中的集合数，必须是2的幂次方

USER_ICACHE_WAY-指令缓存中的方式数，必须是2的幂

USER_DCACHE_SET-数据缓存中的集数，必须是2的幂

USER_DCACHE_WAY-数据缓存中的路径数必须是2的幂

USER_L2CACHE_SET-L2缓存中的集合数，必须是2的幂次方

USER_L2CACHE_WAY-L2缓存中的方式数，必须是2的幂

系统相关配置（在npu_user_defines.sv头文件中）：

IO_MAP_BASE_ADDR-开始为IO操作分配内存空间（绕过一致性）。

IO_MAP_SIZE-IO操作的内存空间宽度。

DIRECTORY_BARRIER-定义后，系统支持所有图块上的分布式目录。否则，它将分配一个同步主机。

CENTRAL_SYNC_ID-单一同步主ID，仅在DIRECTORY_BARIER未定义时使用。

更改线程数

如上所述，每个核心的线程数都是相同的，可以通过更改npu_user_defines.sv头文件中的THREAD_NUMB值轻松修改此参数。每个线程共享L1数据和指令缓存。

将参数值更改为所需值就足以在硬件侧获得线程号修改。虽然编译器有关于线程数量的信息，但链接器使用这些信息来正确管理内存布局中的堆栈。工具链存储库中的misc/lnkrscrpt.ld文件包含此信息，特别是以下行必须一致更新：

threads_per_core=0x8；

该值在crt0.s文件中用于计算堆栈尺寸和位置。

工具链

NaplesPU工具链是编译NaplesPU可执行应用程序所需的工具集合。它基于LLVM项目，利用Clang前端的自定义版本和后端的从头开始实现。该工具链附带了LLD的修改版本，这是一个用于生成与NaplesPU兼容的内存映像的elf2hex工具，最后是用于调试目的的objdump工具的自定义版本。提供了libc库的自定义实现。

该工具链支持基于C和OpenCL C内核的编译。

目录

1在Ubuntu Linux环境下构建工具链

1.1所需软件

1.2构建工艺

2 NaplesPU LLVM结构

3扩展

在Ubuntu Linux环境下构建工具链

本节将展示如何在Ubuntu Linux环境中构建NaplesPU工具链。尽管存在与包管理器相关的差异，但以下步骤在任何其他基于Unix的系统中仍然有效。

所需软件

NaplesPU工具链安装依赖于以下依赖关系：

Git

GCC

CMake

Python 2.7

libxml

Zlib

野牛

弯曲

Libedit

Swig

ncurses库

以下终端命令可用于安装所有必需的软件：

$sudo apt安装libxml2开发git cmake gcc g++python bison flex zlib1g开发swig python开发libedit开发libncurses5开发ninja build

构建工艺

首先，必须通过键入以下命令从官方存储库获取源代码：

$git克隆https://github.com/AlessandroCilardo/NaplesPU-toolchain<clone_directory>

该存储库包含一个帮助脚本setup.sh，以简化安装过程。要构建NaplesPU工具链的新实现，只需键入：

$ ./setup.sh-n

此命令启动编译过程，在/usr/local/llvm-npu中以发布模式安装工具链。如果需要调试版本，请在setup.sh中添加-d标志。还可以使用-t=<number_of_threads>参数选择编译过程的线程数。

在编译过程结束时，需要将库链接到安装文件夹：

$ ./setup.sh-l

现在，可以使用工具链构建自己的应用程序。

NaplesPU LLVM结构

NaplesPU工具链依赖于LLVM项目7.0版本，提供其库的自定义实现，以便生成NaplesPU内核。

Clang是NaplesPU的编译器前端，扩展到处理自定义内部函数的令牌识别。

LLVM核心库用于实现NaplesPU的自定义后端，以管理目标设备上的代码降低。

LLD经过调整，以满足NaplesPU架构对链接的要求。

objdump用于反汇编和分析生成的代码。

elf2hex是生成内存映像所需的工具。

有关如何实施工具链的详细信息，请查看以下链接。

NaplesPU Clang文档

NaplesPU LLVM文档

NaplesPU LLD链接器文档

NaplesPU工具

NaplesPU库

扩展

扩展NaplesPU以提供64位支持

扩展NaplesPU以支持OpenCL

ISA

目录

1注册文件

1.1数据类型

2指令格式

2.1 R型说明

2.2 I型说明

2.3 MOVEI类型说明

2.4 C型说明

2.5 J型说明

2.6 M型说明

注册文件

NPU寄存器堆由标量寄存器堆和矢量寄存器堆组成；每个包含64个寄存器。

标量寄存器堆有64个寄存器。前58个是通用寄存器，其余8个是专用寄存器。每个标量寄存器最多可以存储32位数据。

矢量寄存器堆有64个通用寄存器。每个矢量寄存器最多可以存储512位数据，每个矢量可以存储16 x 32位。

最后，还有一个由几个子寄存器组成的控制寄存器。一些信息在所有线程之间共享，另一些信息是特定于线程的，标记为“线程”的寄存器每个线程都有一个单独的实例。

 

Register	Read/Write	Shared/Thread	Description	ID
TILE_ID	Read	Shared	Tile ID	0
CORE_ID	Read	Shared	Core ID	1
THREAD_ID	Read	Thread	ThreadID	2
GLOBAL_ID	Read	Thread	Global ID, previous IDs merged as follow: TILE_ID, CORE_ID, THREAD_ID	3
GCOUNTER_LOW	Read	Shared	Low part of the Global counter register which counts processor cycles since reset	4
GCOUNTER_HIGH	Read	Shared	High part of the Global counter register which counts processor cycles since reset	5
THREAD_EN	Read	Shared	Thread enabled mask, 1 bit per thread	6
MISS_DATA	Read	Shared	Count of L1 Data cache misses	7
MISS_INSTR	Read	Shared	Count of L1 Instruction cache misses	8
PC	Read	Thread	Current PC	9
TRAP_REASON	Read	Thread	Trap Cause (see below)	10
THREAD_STATUS	Read/Write	Thread	Thread Status2 (see below)	11
ARGC	Read/Write	Shared	The number of strings pointed to by argv	12
ARGV	Read/Write	Shared	The address of command line arguments passed to main()	13
THREAD_NUMB	Read	Shared	The number of total hardware threads	14
THREAD_MISS_CC	Read	Thread	The per-thread clock cycles while the thread is idle due memory operations.	15
KERNEL_WORK	Read	Thread	The per-thread kernel clock cycles.	16
CPU_CTRL_REG	Read/Write	Shared	CPU mode register. At the moment only write policy used by the cache controller is implemented. 0 for write-back, 1 for write-through	17
UNCOHERENCE_MAP	Read/Write	Shared	Address the non-coherent table in the control register. It stores information about the non-coherent memory regions. User can define non-coherent regions addressing this special purpose register.	19
DEBUG_BASE_ADDR	Read/Write	Shared	Debug registers base address. The NPU is equipped with 16 debug registers. DEBUG_BASE_ADDR fetches the value of the first debug register, DEBUG_BASE_ADDR+1 the second and so on.	20

陷阱原因：在当前状态下，只有由于未对齐的内存访问导致的陷阱才会引发：

SPM_ADDR_MISALIGN：SPM单元中的内存访问未对齐。

LDST_ADDR_MISALIGN：LDST单元中的内存访问未对齐。

线程状态：每个线程可以处于以下状态之一：

THREAD_IDLE（值=0）：重置后，每个线程都以这种状态开始。

RUNNING（值=1）：线程正在运行内核。

END_MODE（值=2）：当发布的内核完成时，线程在此模式下切换。

TRAPPED（值=3）：线程处于陷阱模式。在当前状态下，当陷阱发生时，线程会跳到无限循环中。

WAITING_BARRIER（值=4）：线程正在等待同步事件。

数据类型

下表总结了NPU核心中可能使用的数据类型。Type列有C/C++类型名称，LLVM类型列显示LLVM中使用的类型名称，Register列显示存储特定类型值的寄存器类型。

在给定寄存器文件宽度的情况下，突出显示的类型是架构本机支持的类型。其他的是通过扩展获得的，因此它们可以被视为受支持的。它们的优点在于更有效地使用系统内存。

 

Type	LLVM Type	Register	Notes
bool	i1	scalar (32 bits)	It is expanded to 32 bits
char	i8	scalar (32 bits)	It is expanded to 32 bits
short	i16	scalar (32 bits)	It is expanded to 32 bits
int	i32	scalar (32 bits)
float	f32	scalar (32 bits)
vec16i8, vec16u8	v16i8	vector (16 x 32 bits)	It is expanded to 32 bits vector
vec16i16, vec16u16	v16i16	vector (16 x 32 bits)	It is expanded to 32 bits vector
vec16i32, vec16u32	v16i32	vector (16 x 32 bits)
vec16f32	v16f32	vector (16 x 32 bits)
vec8i8, vec8u8	v8i8	vector (8 x 32 bits)	It is expanded to 32 bits vector
vec8i16, vec8u16	v8i16	vector (8 x 32 bits)	It is expanded to 32 bits vector
vec8i32, vec8u32	v8i32	vector (8 x 32 bits)	It is expanded to 32 bits vector
vec8f32	v8f32	vector (16 x 32 bits)	It is considered as a 16 elements vector

说明格式

NaplesPU指令的固定长度为32位。它们分为六类：

R类型包括逻辑和算术运算以及内存运算。

I类型包括寄存器操作数和立即操作数之间的逻辑和算术运算。

MOVEI类型包括寄存器中立即操作数的加载操作。

用于控制操作和同步指令的C类型。

J类型包括跳转指令。

M类型包括用于访问内存的指令。

R型说明

这是用机器码编码的R型指令的格式。

RR（注册到注册）有一个目标寄存器和两个源寄存器。

RI（寄存器立即）有一个目标寄存器和一个源寄存器，以及一个在指令字中编码的立即寄存器。

R型指令的字段包括：

操作码（B29-24）是“操作码”的缩写。操作码是指令的二进制编码。对于R型指令，它只有6位。

rd（B23-18）是目标寄存器

rs0（B17-12）是第一个源寄存器。

rs1（B11-6）是第二个源寄存器。

位l（B4）用于“长”操作，即需要长整数或双精度数字的操作。如果操作需要64位寄存器，则l=1，否则l=0。

位fmt（B3-1）用于指定某个操作数是标量还是向量（格式中每个寄存器对应一位）。B3表示寄存器d，B2表示寄存器rs0，B1表示寄存器rs1。例如，如果目标寄存器应包含向量，B3=1，否则B3=0。

R型说明如下：

or	1	或	Rb
and	2	和	Rd = Ra & Rb
xor	3	xor	Rd = Ra ^ Rb
add	4	附加	Rd = Ra + Rb
sub	5	减法	Rd = Ra – Rb
mullo	6	乘法低位结果	Rd = Ra * Rb
mulhi	7	乘法高位结果	Rd = Ra * Rb
mulhu	8	乘法无符号高位结果	Rd = Ra * Rb
ashr	9	算术右移	Rd = Ra '>> Rb
shr	10	右移	Rd = Ra >> Rb
shl	11	左移	Rd = Ra << Rb
clz	12	计数前导零
ctz	13	计数尾随零
shuffle	24	矢量洗牌	Rd[i] = Ra[Rb[i]]
getlane	25	从矢量中获取通道	Rd = Ra[Rb]
move	32	移动寄存器	Rd = Ra
fadd	33	浮点加法	Rd = Ra + Rb
fsub	34	浮点减法	Rd = Ra – Rb
fmul	35	浮点乘法	Rd = Ra * Rb
fdiv	36	浮点除法	Rd = Ra / Rb
sext8	43	符号扩展8位
sext16	44	符号扩展16位
sext32	45	符号扩展32位
i32tof32	48	将整数转换为浮点数
f32toi32	49	将浮点数转换为整数
cmpeq	14	比较相等	Rd = Ra == Rb
cmpne	15	比较不相等	Rd = Ra != Rb
cmpgt	16	比较大于	Rd = Ra > Rb
cmpge	17	比较大于或等于	Rd = Ra >= Rb
cmplt	18	比较小于	Rd = Ra < Rb
cmple	19	比较少于或相等	Rd = Ra <= Rb
cmpugt	20	无符号比较大于	Rd = Ra > Rb
cmpuge	21	无符号比较大于或等于	Rd = Ra >= Rb
cmpult	22	无符号比较小于	Rd = Ra < Rb
cmpule	23	无符号比较小于或等于	Rd = Ra <= Rb
cmpfeq	37	浮点比较相等	Rd = Ra == Rb
cmpfne	38	浮点比较不相等	Rd = Ra != Rb
cmpfgt	39	浮点比较大于	Rd = Ra > Rb
cmpfge	40	浮点比较大于或等于	Rd = Ra >= Rb
cmpflt	41	浮点比较小于	Rd = Ra < Rb
cmpfle	42	浮点比较小于或等于	Rd = Ra <= Rb

I键入说明

这是用机器码编码的I型指令的格式。

I型指令的字段是：opcode（B28-24）是“操作码”的缩写。操作码是指令的二进制编码。对于*I型指令，它只有5位。

rd（B23-18）是目标寄存器

rs（B17-12）是第一个源寄存器。

imm（B11-3）是9位立即值。

fmt（B2-1）位用于指定某个操作数是标量还是向量（格式中每个寄存器对应一位）。B2表示寄存器d，B1表示寄存器rs。

I型说明如下：

助记符	操作码	语义	操作
ori	1	或	Imm
andi	2	与	Rd = Ra & Imm
xori	3	异或	Rd = Ra ^ Imm
addi	4	加法	Rd = Ra + Imm
subi	5	减法	Rd = Ra – Imm
mulli	6	乘法	Rd = Ra * Imm
mulhi	7	高倍数	Rd = Ra * Imm
mulhui	8	高倍乘无符号	Rd = Ra * Imm
ashri	9	算术右移	Rd = Ra ‘>> Imm
shri	10	右移	Rd = Ra >> Imm
shli	11	左移	Rd = Ra << Imm
getlane	25	从向量中获取通道	Rd = Ra[Imm]

MOVEI类型说明

MVI（Move Immediate）有一个目标寄存器和一个立即编码的16位指令。这是用机器码编码的MOVEI类型指令的格式。

MOVEI类型指令的字段包括：

•操作码（B26-24）是“操作码”的缩写。操作码是指令的二进制编码。对于MOVEI类型的指令，它只有3位。

•rd（B23-18）是目标寄存器

•imm（B17-2）是16位立即值。

•fmt（B1）用于指定目标寄存器是包含向量还是标量。

MOVEI类型说明如下：

助记符	操作码	语义	操作
moveil	0	移动16个较低有效位	Rd = Ra & 0xFFFF
moveih	1	移动16个最高有效位	Rd = (Ra >> 16) & 0xFFFF
movei	2	移动16个扩展名为零的低有效位	Rd = (Rd ^ Rd) & (Ra & 0xFFFF)

C类说明

这是用机器码编码的C型指令的格式。

C型指令的字段包括：

•操作码（B26-24）是“操作码”的缩写。操作码是指令的二进制编码。对于C型指令，它只有3位。

•rs0（B23-18）是第一个源寄存器。

•rs1（B17-12）是第二个源寄存器。

C型说明如下：

助记符	操作码	语义
barrier_core	0	内存屏障-确保屏障之前的所有显式数据内存传输，都在屏障之后开始的任何后续显式数据存储事务之前完成。寄存器rs0包含屏障标识号（BID）。BID可以是大于0的任意数字，即BID>0。不同的内存屏障需要不同的BID。rs1包含应同步的线程数。
flush	2	将缓存行刷新到主内存。
read_cr	3	读取控制寄存器的子寄存器。
write_cr	4	写入控制寄存器的子寄存器
dcache_inv	5	使L1缓存中的输入地址行无效。

J型说明书

这是用机器码编码的J型指令的格式。

J型指令的字段包括：

•操作码（B26-24）是“操作码”的缩写。操作码是指令的二进制编码。对于J型指令，它只有3位。

•rcond/rd（B23-18）是条件/目标寄存器。

•offset（B17-0）是偏移地址。

J型说明如下：

Mnemonic	Opcode	Meaning	Operation
jmp	0	jump - unconditionally jump to a specified location.	PC=rd or PC=PC+offset
jmpsr	1	jump to subroutine - unconditionally jump to a specified location and store the return address in the RA register.	RA=PC+4 PC=rd or RA=PC+4 PC=PC+addr
jret	3	Return from Subroutine - unconditionally return from a subroutine loading the return address from the RA register.	PC=RA
beqz	5	Conditional Branch. Branch if Equal to Zero - branche to PC+offset if the contents of the condition register is equal to zero.	if(rcond==0) PC=PC+offset else PC=PC+4
bnez	6	Conditional Branch, Branch if Not Equal to Zero - branches to PC+offset if the contents of the condition register is not equal to zero.	if(rcond!=0) PC=PC+offset else PC=PC+4

M型说明

这是用机器码编码的M型指令的格式。

M型指令的字段包括：

•操作码（B29-24）是“操作码”的缩写。操作码是指令的二进制编码。对于M型指令，它只有6位。

•rd/rs（B23-18）是目标寄存器或源寄存器

•rbase（B17-12）是基址寄存器。

•offset（B11-3）是偏移地址。

•位l（B2）未使用。保留64位扩展。

•位s（B1）用于指定某个加载/存储内存操作是否进入草稿行内存。例如，在从/向草稿行存储器加载/存储的情况下，B1=1，否则B1=0。

典型的M型指令是加载和存储指令。在这两种情况下，源/目标地址都是按基寄存器地址+立即偏移量计算的，即rbase+偏移量。在负载的情况下，rd=[rbase+offset]。同样，在存储的情况下，[rbase+offset]=rs。所有M类型的指令都可以用于对主存储器和草稿行存储器的存储操作。使用scratchpad内存操作的指令具有_scratchpad后缀。例如，load32_s8针对主存储器，而load32_s8_scratchpad指的是片上草稿行的加载操作。

M型指令可分为标量指令和向量指令。标量M型指令是：

Mnemonic	Opcode	Meaning	Operation
load32_s8	0	load memory byte [7:0] with sign extension into a 32 bit register	Rd = [Rbase + Offset]
load32_s16	1	oad memory half word [15:0] with sign extension into a 32 bit register	Rd = [Rbase + Offset]
load32	2	load memory word into a 32 bit register	Rd = [Rbase + Offset]
load32_u8	4	load memory byte [7:0] with zero extension into a 32 bit register	Rd = [Rbase + Offset]
load32_u16	5	load memory half word [15:0] with zero extension into a 32 bit register	Rd = [Rbase + Offset]
load_v16i8	7	load 16 byte [127:0] with sign extension into a 512 bit register	Rd = [Rbase + Offset]
load_v16i16	8	load 16 half word [255:0] with sign extension	Rd = [Rbase + Offset]
load_v16i32	9	load 16 words	Rd = [Rbase + Offset]
load_v16u8	11	load 16 byte [127:0] with no sign extension	Rd = [Rbase + Offset]
load_v16u16	12	load 16 half word [255:0] with no sign extension	Rd = [Rbase + Offset]
load_v8u32	13	load 8 word [255:0] with no sign extension	Rd = [Rbase + Offset]
loadg32	16	load 16 words from different memory addresses (only for scratchpad)	Rd[i] = [Rbase[i]]
store32_8	32	store 1 byte into the effective address	[Rbase + Offset] = Rs
store32_16	33	store 2 bytes into the effective address	[Rbase + Offset] = Rs
store32	34	store 1 word into the effective address	[Rbase + Offset] = Rs
store_v16i8	36	store 16 bytes from a vectorial register (data fecthing from register schema [487:480,...,39:32,7:0]) into effective address location	[Rbase + Offset] = Rs
store_v16i16	37	store 16 half words (data fetching from register schema [495:480,...,47:32,15:0]) into effective address location	[Rbase + Offset] = Rs
store_v16i32	38	store 16 words from a vectorial register into effective address location	[Rbase + Offset] = Rs
stores32	42	scatter store - store 16 words into 16 different addresses (only for scratchpad)	[Rbase[i]] = Rs[i]

Navigation menu

• Log in

• Page
• Discussion

• Read
• View source
• View history

Search

窗体顶端

窗体底端

• Main page
• Recent changes
• Random page
• Help

Tools

• What links here
• Related changes
• Special pages
• Printable version
• Permanent link
• Page information

加长型NaplesPU

（重定向自扩展nu+）

目录

1 SystemVerilog编码NaplesPU指南

2在NaplesPU核心中添加自定义指令

2.1定义新指令

2.2扩展编译器支持

2.2.1添加新的内部

2.2.2添加新指令

2.3延长NPU核心管线

2.3.1自定义单元接口

2.3.2扩展解码阶段

2.3.3扩展回写阶段

2.3.4在NPU管道中添加模块

SystemVerilog编码NaplesPU指南

这是扩展NaplesPU架构的简单指南。

1.模块的输出信号名称始终以助记符模块名称开头（例如回写的信号->wb_xxx）。

2.测试台文件名以tb_开头。

3.为每个不同的独立模块添加一个文件夹，并在主文件夹中插入分布在整个项目中的“通用”模块。

4.在新定义中使用括号进行算术运算。

5.使用结构体或typedefs来定义何时经常访问信号的子部分。

6.使用divide et impera哲学来提高可重用性和可理解性。

7.使用现有的信号类型；如果引入新的结构和typedef，请在include文件夹中为该组件的特定头文件中分配它们（例如writeback unit->writeback_defines.sv）。

在NaplesPU核心中添加自定义指令

本节介绍如何添加新的函数操作、扩展指令集以及将自定义组件添加到NaplesPU管道中。

定义新指令

第一步是在NaplesPU ISA中添加一条新指令，从指令格式开始。例如，新的算术运算应该是R型指令的一部分，而新的内存访问指令必须是M型指令的组成部分。在下面的例子中，引入了新的算术运算，称为crp。

扩展编译器支持

扩展编译器对自定义指令的支持涉及两个主要步骤：

增加一个新的内在；

添加新指令。

添加新的内部

添加新的内部函数涉及编译器后端和前端的三个不同文件。

在前端方面，Clang必须认识到这一新的内在因素。这是通过在工具链仓库的“compiler/tools/clang/include/clang/Basic/Buildings NaplesPU.def”文件中添加以下行来实现的：

//------ 交叉产品 ----------//

BUILTIN ( __builtin_npu_crossprodv16i32 , " V16iV16iV16i ", "n")

这样的宏定义了内在的签名：

• __builtin_npu_crossprodv16i32 - name
• V16iV16iV16i - input and output types
• n - optional attributes

有关更多信息，请参阅工具链仓库中的文件“compiler/tools/clang/include/clang/Bial/Builtins.def”。

然后，在“compiler/tools/clang/lib/CodeGen/CGBuiltin.cpp”中，扩展了EmitNPUBuiltinExpr方法，在switch构造中添加了一个新的case，如下所示：

// Cross Product

case NaplesPU :: BI__builtin_npu_crossprodv16i32 :

   F = CGM . getIntrinsic ( Intrinsic :: npu_crossprodv16i32 );

   break ;

密钥工作BI__builtin_npu_crossprodv16i32必须与BuiltinsNaplesPU.def文件中添加的内置一致，签名名称必须以BI开头。

最后，在后端扩展编译器/include/llvm/IR/IntegrissNaplesPU.td文件，如下所示：

// Cross Product Intrinsic

def int_npu_crossprodv16i32 : Intrinsic <[ llvm_v16i32_ty ], [ llvm_v16i32_ty , llvm_v16i32_ty ], [ IntrNoMem ], " llvm.npu.__builtin_npu_crossprodv16i32 ">;

此Table Gen代码在Clang中添加了新的内部函数，并生成了相应的AST节点。

在问题中，定义一个类（int_npu_crossprodv16i32）为TableGen固有类。本文首先介绍了构建物和入口的具体风险（llvm_v16i32_ty）、最终属性（IntrNoMem）和IR（“llvm.npu.__builtin_npu_crossprodv16i32”）的定义，并采用了NaplesPU内置定义的名称。

添加新指令

在编译器后端的ISA中添加一条新指令，扩展了工具链仓库中的编译器/lib/Target/NaplesPU/NaplesPUInstrInfo.td文件。这样的扩展名要求在编译器/lib/Target/NaplesPU/NaplesPUInstrFormats.td文件中定义的Table Gen类。特别是，用于crp指令的类是FR_TwoOp_Unmasked-32，这将指令定义为具有两个输入操作数（FR_TwoOp）的R类型，这两个操作数都是向量，没有掩码（Unmasked_32）：

// Cross Product Instruction

def CROSSPROD_32 : FR_TwoOp_Unmasked_32 <

( outs VR512W : $dst ), // output

( ins VR512W :$src0 , VR512W : $src1 ), // input

" crp $dst , $src0 , $src1 ", // corresponding assembly code

[( set v16i32 :$dst , ( int_npu_crossprodv16i32 v16i32 :$src0 , v16i32 : $src1 ))], // matching pattern

63, // ISA opcode (unique for instruction)

Fmt_V , // destination register format

Fmt_V , // src0 register format

Fmt_V >; // src1 register format

属性VR512W将操作目标寄存器定义为具有16个32位元素的向量。双重地，属性Fmt_V在指令字节码的Fmt字段中相应地设置。对于自定义模块，选择63作为操作码，在文本中进一步称为MY_opcode。

扩建NPU核心管线

在执行阶段使用新操作员扩展NPU核心管道涉及以下步骤：

定义新模块及其接口；

扩展NPU解码阶段；

扩展NPU回写阶段；

将模块添加到NPU管道中。

未完全流水线的模块也必须扩展Instruction_Buffer模块。

自定义单元接口

下面是一个示例界面：

`include " npu_user_defines.sv"

`include " npu_defines.sv"

module my_pipe (

   input clk ,

   input reset ,

   // To Instruction buffer

   output thread_mask_t my_stop ;

   // From Operand Fetch

   input opf_valid ,

   input instruction_decoded_t opf_inst_scheduled ,

   input vec_reg_size_t opf_fecthed_op0 ,

   input vec_reg_size_t opf_fecthed_op1 ,

   input hw_lane_mask_t opf_hw_lane_mask ,

   // To Writeback

   output logic my_valid ,

   output instruction_decoded_t my_inst_scheduled ,

   output vec_reg_size_t my_result ,

   output hw_lane_mask_t my_hw_lane_mask

);

如果新模块不能在每个时钟周期接受请求，则必须提供停止条件并将其转发给Instruction_Buffer模块，以防止自定义模块出现进一步问题。在上面的例子中，这是通过my_stop信号完成的，当模块无法处理进一步的请求时，它必须为高。然后将my_stop添加到Instruction_Buffer模块中，如下所示：

assign ib_instructions_valid[thread_id] = ~fifo_empty & ~( l1d_full[thread_id] & ib_instructions[thread_id].pipe_sel == PIPE_MEM ) & enable & ~(my_stop & ib_instructions[thread_id].op_code == MY_OPCODE);

如果自定义模块可以在每个时钟周期处理一个请求，则不需要最后一步。

输入信号由Operand_Fetch模块生成：

opf_valid，传入请求有效。

instruction_decoded_t opf_inst_scheduled，当前指令已解码。模块必须检查op_code，如果它等于新的操作码（MY_opcode），则必须详细说明发出的操作。指令可以是标量指令，也可以是矢量指令。这些信息存储在instruction_decoded_t字段中，每个寄存器都有一个专用位，即is_source0_vectorial、is_source1_vectorial和is_destination_vectorial位。

vec_reg_size_topf_fecthed_op0，输入寄存器的向量。

vec_reg_size_topf_fecthed_op1，输入寄存器的向量。

hw_lane_mask_topf_hw_lane-mask，硬件通道位掩码，第i位表示向量中的第i个元素必须详细说明。

输出信号被转发到回写模块：

my_valid，输出结果有效。

instruction_decoded_t my_inst_scheduled，模块必须将发出的指令与结果一起转发。

vec_reg_size_t my_result，按向量通道组织的输出结果。

hwlane_ask_t myhwlane_mask，模块必须转发与结果一起使用的硬件位掩码。

扩展解码阶段

首先，在include/npu_defines.sv中扩展了pipeline_deisp_t类型，并添加了一个新值来标识新模块（例如PIPE_new）。然后，在同一个文件中，向正确的指令类型添加了一条新指令，添加了一个新的R类型指令，因此在alu_op_t类型中添加了一种新的唯一操作码（注意，要是唯一的）。对于一个新的M类型，应该扩展memory_op_t，等等（它们都在同一个文件中）。

现在，在解码阶段，根据新的指令类型在开关构造中选择正确的情况，在这个例子中，再次引用了一个新的R类型指令，因此代码将被放置在以下情况中：

casez ( if_inst_scheduled.opcode )

  // RR

  8'b00_?????? : begin

  ...

确保在opcode=MY_opcode的情况下，解码阶段通过将pipe_sel值设置为pipe_new来为自定义模块发出新的请求，如下所示：

if ( if_inst_scheduled.opcode.alu_opcode <= MOVE || ( if_inst_scheduled.opcode.alu_opcode >= SEXT8 & if_inst_scheduled.opcode.alu_opcode <= SEXT32 )

   || if_inst_scheduled.opcode.alu_opcode == MY_OPCODE ) begin

  if (if_inst_scheduled.opcode.alu_opcode == MY_OPCODE) begin

     instruction_decoded_next.pipe_sel = PIPE_NEW ;

     instruction_decoded_next.is_int = 1'b0;

     instruction_decoded_next.is_fp = 1'b0;

  end

扩展回写阶段

在Writeback阶段，首先为自定义单元添加一个新的专用接口，如下所示：

// From MY costum module

input                                      my_valid,

input  instruction_decoded_t                 my_inst_scheduled,

input  hw_lane_t                           my_result,

input  hw_lane_mask_t                      my_mask_reg,

然后，添加一个新的写回请求FIFO，专门用于从自定义模块中获取传入结果。在这种情况下，更新include/npu_defines.sv标头中的`NUM_EX_PIPE参数（通过在之前的值中添加一个），并为自定义操作添加一个具有新ID的本地参数：

localparam  PIPE_FP_ID   = 0; // FP pipe FIFO index

localparam  PIPE_INT_ID  = 1; // INT pipe FIFO index

localparam  PIPE_SPM_ID  = 2; // SPM memory FIFO index

localparam  PIPE_MEM_ID  = 3; // LDST unit FIFO index

localparam  PIPE_NEW_ID  = 4; // NEW op FIFO index

接下来，将专用FIFO连接到模块的接口输入：

assign input_wb_request[PIPE_NEW_ID].pc                        = my_inst_scheduled.pc;

assign input_wb_request[PIPE_NEW_ID].writeback_valid           = my_valid;

assign input_wb_request[PIPE_NEW_ID].thread_id                 = my_inst_scheduled.thread_id;

assign input_wb_request[PIPE_NEW_ID].writeback_result          = my_result;

assign input_wb_request[PIPE_NEW_ID].writeback_hw_lane_mask    = my_mask_reg;

assign input_wb_request[PIPE_NEW_ID].destination               = my_inst_scheduled.destination;

assign input_wb_request[PIPE_NEW_ID].is_destination_vectorial  = my_inst_scheduled.is_destination_vectorial;

assign input_wb_request[PIPE_NEW_ID].op_code                   = my_inst_scheduled.op_code;

assign input_wb_request[PIPE_NEW_ID].pipe_sel                  = my_inst_scheduled.pipe_sel;

assign input_wb_request[PIPE_NEW_ID].is_memory_access          = my_inst_scheduled.is_memory_access;

assign input_wb_request[PIPE_NEW_ID].has_destination           = my_inst_scheduled.has_destination;

assign input_wb_request[PIPE_NEW_ID].is_branch                 = my_inst_scheduled.is_branch;

assign input_wb_request[PIPE_NEW_ID].is_control                = my_inst_scheduled.is_control;

assign input_wb_request[PIPE_NEW_ID].is_movei                  = my_inst_scheduled.is_movei;

assign input_wb_request[PIPE_NEW_ID].result_address            = 0;

最后，扩展了生成转发到寄存器文件的回写结果的代码，在构建output_wb_request.writeback信号的过程中添加了一个新案例：

//输出数据编辑器。wb_result_data直接转发到寄存器文件

always_comb begin : WB_OUTPUT_DATA_SELECTION

 case ( output_wb_request[selected_pipe].pipe_sel )

  PIPE_MEM : wb_next.wb_result_data = result_data_mem;

  PIPE_SPM : wb_next.wb_result_data = result_data_spm;

  PIPE_INT,

  PIPE_NEW,

  PIPE_FP  : wb_next.wb_result_data = output_wb_request[selected_pipe].writeback_result;

  default : wb_next.wb_result_data  = 0;

  endcase

end

在NPU管道中添加模块

首先，声明模块所需的信号：

//新管道平台-信号

logic                                        my_valid;

instruction_decoded_t                        my_inst_scheduled;

hw_lane_t                                    my_result;

hw_lane_mask_t                               my_hw_lane_mask;

然后，将模块实例化放入位于core/NPU_core.sv文件中的NPU管道中，如下所示：

my_pipe u_my_pipe (

 .clk                ( clk                ),

 .reset              ( reset              ),

 .enable             ( nfreeze            ),

 //From Operand Fetch

 .opf_valid          ( opf_valid          ),

 .opf_inst_scheduled ( opf_inst_scheduled ),

 .opf_fetched_op0    ( opf_fetched_op0    ),

 .opf_fetched_op1    ( opf_fetched_op1    ),

 .opf_hw_lane_mask   ( opf_hw_lane_mask   ),

 //To Writeback

 .my_valid           ( my_valid           ),

 .my_inst_scheduled  ( my_inst_scheduled  ),

 .my_result          ( my_result          ),

 .my_hw_lane_mask    ( my_hw_lane_mask    )

);

最后，将模块连接到写回阶段：

writeback #(

 .TILE_ID( TILE_ID )

)

u_writeback (

  .clk                 ( clk                ),

  .reset               ( reset              ),

  .enable              ( 1'b1               ),    

 ...

 //From NEW Pipe

 .my_valid           ( my_valid           ),

 .my_inst_scheduled  ( my_inst_scheduled  ),

 .my_result          ( my_result          ),

 .my_hw_lane_mask    ( my_hw_lane_mask    )

异质瓷砖

NaplesPU项目提供了一个集成到NoC中的异构图块，旨在由用户扩展。这种拼贴提供了如何在片上网络内集成定制模块的第一个示例。该项目附带了src/mc/tile/tile_ht.sv中的专用原型，旨在扩展自定义逻辑。

提供的tile_ht模块实例化了NPU图块的所有典型模块，GPGPU内核除外，以及封装在简化界面中的加载/存储单元的修改版本，该简化界面旨在便于自定义组件访问。

HT图块为用户提供了两个与系统交互的主要界面：内存界面和同步界面，下文将进一步解释。这种图块为用户特定的配置提供了以下参数：

`include "npu_user_defines.sv"

`include "npu_defines.sv"

`include "npu_coherence_defines.sv"

module tile_ht # (

   parameter TILE_ID        = 0,                 // Current tile ID

   parameter CORE_ID        = 0,                 // Current core ID, not used in this type of tile

   parameter TILE_MEMORY_ID = 9,                 // ID of the memory controller tile

   parameter THREAD_NUMB    = 8,                 // Supported thread number, each thread has a separate FIFO in the LSU and requests from different threads are elaborated concurrently - Must be a power of two

   parameter ADDRESS_WIDTH  = `ADDRESS_SIZE,     // Memory address width - has to be congruent with the system address width

   parameter DATA_WIDTH     = `CACHE_LINE_WIDTH, // Data bus width - has to be congruent with the system

   parameter L1_WAY_NUMB    = 4,                 // Number of way in the L1 data cache

   parameter L1_SET_NUMB    = 32,                // Number of L1 data sets

   parameter SYNCH_SUPP     = 1                  // Allocates barrier_core modules for synchronization support

)

 

目录

[隐藏]

•1内存接口

•2同步接口

•3服务消息接口

•提供4个永恒的假人

•5添加自定义逻辑

存储器接口

存储器接口提供了一种与相干系统交互的透明方式。内存接口为每个线程实现了一个简单的有效/可用握手，不同的线程可能会发出不同的内存事务，这些事务由一致性系统同时处理。

当一个线程有一个内存请求时，它首先检查与其ID相关的可用位，如果该可用位很高，则线程会发出一个内存事务，设置有效位并在内存接口上加载所有需要的信息。

支持的内存操作及其操作码如下：

LOAD_8      = 'h0  - 'b000000   

LOAD_16     = 'h1  - 'b000001

LOAD_32     = 'h2  - 'b000010

LOAD_V_8    = 'h7  - 'b000111

LOAD_V_16   = 'h8  - 'b001000

LOAD_V_32   = 'h9  - 'b001001

STORE_8     = 'h20 - 'b100000

STORE_16    = 'h21 - 'b100001

STORE_32    = 'h22 - 'b100010

STORE_V_8   = 'h24 - 'b100100

STORE_V_16  = 'h25 - 'b100101

STORE_V_32  = 'h26 - 'b100110

要集成到NaplesPU系统中的自定义内核应实现以下接口，以便与存储系统通信：

/*存储器接口*/

//到异构LSU

output logic                                 req_out_valid,     // Valid signal for issued memory requests

output logic [31 : 0]                        req_out_id,        // ID of the issued request, mainly used for debugging

output logic [THREAD_IDX_W - 1 : 0]          req_out_thread_id, // Thread ID of issued request. Requests running on different threads are dispatched to the CC conccurrently

output logic [7 : 0]                         req_out_op,        // Operation performed

output logic [ADDRESS_WIDTH - 1 : 0]         req_out_address,   // Issued request address

output logic [DATA_WIDTH - 1    : 0]         req_out_data,      // Data output

// From Heterogeneous LSU

input  logic                                 resp_in_valid,      // Valid signal for the incoming responses

input  logic [31 : 0]                        resp_in_id,         // ID of the incoming response, mainly used for debugging

input  logic [THREAD_IDX_W - 1 : 0]          resp_in_thread_id,  // Thread ID of the incoming response

input  logic [7 : 0]                         resp_in_op,         // Operation code

input  logic [DATA_WIDTH - 1 : 0]            resp_in_cache_line, // Incoming data

input  logic [BYTES_PERLINE - 1 : 0]         resp_in_store_mask, // Bitmask of the position of the requesting bytes in the incoming data bus

input  logic [ADDRESS_WIDTH - 1 : 0]         resp_in_address,    // Incoming response address

异构瓦片共享NPU瓦片的相同LSU和CC，因此LSU在内存接口上转发其背压信号，如下所示：

//来自异构加速器-背压信号

input  logic [THREAD_NUMB - 1 : 0]           lsu_het_almost_full,           // Thread bitmask, if i-th bit is high, i-th thread cannot issue requests.

input  logic [THREAD_NUMB - 1 : 0]           lsu_het_no_load_store_pending, // Thread bitmask, if i-th bit is low, i-th thread has no pending operations.

特别是，在为第i个线程发出内存请求之前，lsu_het_almost_full第i位必须为低。

内存接口提供性能计数器作为其接口的一部分：

// From Heterogeneous LSU - Performance counters

input  logic                                 resp_in_miss,  // LSU miss on resp_in_address

input  logic                                 resp_in_evict, // LSU eviction (replacement) on resp_in_address

input  logic                                 resp_in_flush, // LSU flush on resp_in_address

input  logic                                 resp_in_dinv,  // LSU data cache invalidatio on resp_in_address

这些信号表示L1数据缓存何时发生丢失、驱逐（或替换）、刷新和数据缓存无效。

异构图块中的LSU可以以两种不同的方式配置，即直写和回写：

output logic                                 lsu_het_ctrl_cache_wt,   // Enable Write-Through cache configuration.

当lsu_het_ctrl_cache_wt为高时，lsu充当直写缓存，当其为低时，lsu实现回写机制。

最后，如果发出的请求中的地址未对齐，内存接口会提供错误信号：

//异构加速器-刷新和错误信号

input  logic                                 lsu_het_error_valid,     // Error coming from LSU

input  register_t                            lsu_het_error_id,        // Error ID - Misaligned = 380

input  logic [THREAD_IDX_W - 1 : 0]          lsu_het_error_thread_id, // Thread involved in the Error

同步接口

同步接口将用户逻辑与拼贴块内分配的同步模块核心侧（即barrier_core单元）连接起来。这样的接口允许用户逻辑在线程粒度上同步，当参数SYNCH_SUPP为高时，瓦片实现同步支持，分配一个barrier_core模块来处理同步事件核心侧：

generate

 if ( SYNCH_SUPP == 1) begin

   barrier_core # (

        .TILE_ID        ( TILE_ID     ),

        .THREAD_NUMB    ( THREAD_NUMB ),

        .MANYCORE       ( 1           ),

        .DIS_SYNCMASTER ( DISTR_SYNC  )

   )

   u_barrier_core (

       .clk                            ( clk),

       .reset                          ( reset),         

       // Synch Request - Core Interface

       .opf_valid                      ( breq_valid),

       .opf_inst_scheduled             ( bc_inst_scheduled),

       .opf_fetched_op0                ( breq_barrier_id),

       .opf_fetched_op1                ( breq_thread_numb),

       .bc_release_val                 ( bc_release_val),

        ...

   );

 end else begin

  assign bc_release_val                = {THREAD_NUMB{1'b1}};

  assign c2n_account_valid             = 1'b0;

  assign c2n_account_message           = sync_account_message_t'(0);

  assign c2n_account_destination_valid = tile_mask_t'(0);

  assign n2c_mes_service_consumed      = 1'b0;

 end

endgenerate

同步机制支持块间和块内屏障同步。当线程到达同步点时，它会通过同步接口向分布式同步主机发出请求。然后，线程被暂停（直到用户逻辑），直到其释放信号再次为高。

如果需要同步，自定义核心可以实现以下接口：

/* 同步接口 */

//到屏障核心

//输出逻辑信号，发送同步请求                               

breq_valid,       // 冲击屏障信号，发送同步请求

output logic [31 : 0]                        breq_op_id,       // 同步 operation ID, mainly used for debugging

output logic [THREAD_NUMB - 1 : 0]           breq_thread_id,  

// 执行同步操作的线程ID

output logic [31 : 0]                        breq_barrier_id, 

// 屏障ID，在并发屏障的情况下必须是唯一的

output logic [31 : 0]                        breq_thread_numb,

// 总数-当前屏障ID上的1个同步线程

// 来自屏障核心

input  logic [THREAD_NUMB - 1 : 0]           bc_release_val

// 暂停线程位掩码等待释放（第i位低位暂停第i个线程）

服务消息接口

服务消息接口将用户逻辑与服务网络连接起来。异构瓦片可以通过此接口向其他节点发送消息，这通常用于主机瓦片通信。如果需要通过消息进行通信，自定义核心可以实现以下接口：

/* Service Message Interface */

// From Service Network

input  logic                                 message_in_valid,         // Valid bit for incoming Service Message

input  service_message_t                     message_in,               // Incoming message from Service Network

output logic                                 n2c_mes_service_consumed, // Service Message consumed

 

// To Service Network

output logic                                 message_out_valid, // Valid bit for outcoming Service Message

output service_message_t                     message_out,       // Outcoming Service Message data

input  logic                                 network_available, // Service Network availability bit

output tile_mask_t                           destination_valid  // One-Hot destinations bitmap

服务消息接口是一个标准的有效/可用接口，当Message_in_valid被断言时，传入的消息Message_in是有效的，用户应该在一个时钟周期内断言n2c_mes_Service_consumed位，这向网络接口发出信号，表明消息已被正确接收和处理。传入消息被声明为service_message_t类型，如以下代码片段所示：

typedef struct packed {

       service_message_type_t message_type;

       service_message_data_t data;

} service_message_t;

字段数据存储接收到的信息，而message_type表示传入消息的性质。在这种情况下，对于来自主机的消息，此值可能是HOST，对于来自另一个异构图块的消息，该值可能是HT_CORE。

另一方面，每当用户逻辑有消息要通过网络发送时，它都会构建输出消息message_out，将消息体存储在数据字段中，同时将message_type字段与HT_CORE值绑定。然后，必须在destination_valid输出位掩码中声明目标图块，使用位掩码对网络中的每个图块进行解码，每个位根据相应位的位置表示不同的图块，例如位置0中的位目标图块0，以此类推。连接的网络接口将把消息转发给在这样一个掩码中声明的每一个图块。当消息和目的地都准备就绪时，用户首先检查网络的可用性，读取network_available输入位，必要时等待。然后，控制逻辑可能会断言message_out_valid信号，将消息和目的地转发到网络接口，网络接口将从这一刻起负责传递。

提供异构假人

此FSM首先与NoC中的其他ht同步。ht图块中的每个虚拟核都需要LOCAL_BARIER_NUMB线程的同步（默认值=4）。

// Issue synchronization requests

SEND_BARRIER    : begin

  breq_valid       <= 1'b1;

  breq_barrier_id  <= 42;

  barrier_served   <= 1'b1;

  if(rem_barriers == 1)

    next_state <= WAIT_SYNCH;

  else

    next_state <= IDLE;

  end

SEND_BARRIER状态通过同步接口发送具有屏障ID 42的LOCAL_BARRIER_NUMB请求。它将屏障ID 42上同步的线程总数设置为total_barrier_NUMB（=LOCAL_barrier-NUMB x'TILE_HT，系统中异构块的数量）。当发出最后一个屏障时，SEND_barrier跳到WAIT_SYNC，等待来自同步主机的ACK。

// Synchronizes all dummy cores

WAIT_SYNCH : begin

 if(&bc_release_val)

    next_state    <= IDLE;

 end

此时，每个ht图块中的所有线程都是同步的，FSM会启动所有挂起的内存事务。

START_MEM_READ_TRANS执行LOCAL_WRITE_REQS读取操作（默认值=128），每次执行LOAD_8操作（操作码=0）。在默认配置中，对连续地址的128个LOAD_8操作分布在所有线程中，并通过内存接口发送给LSU。当读取操作结束时，FSM以类似的方式开始写入操作。

// Starting multiple read operations

START_MEM_READ_TRANS    : begin

  if ( rem_reads == 1 )

    next_state <= DONE;

  else

    next_state <= IDLE;

  if(lsu_het_almost_full[thread_id_read] == 1'b0) begin

    read_served       <= 1'b1;

    req_out_valid     <= 1'b1;

    req_out_id        <= rem_reads;

    req_out_op        <= 0; // LOAD_8

    incr_address      <= 1'b1;

    req_out_thread_id <= thread_id_read;

  end

end

START_MEM_WRITE_TRANS通过内存接口对连续地址执行LOCAL_WRITE_REQS（默认值=128）写入操作。这次执行的操作是STORE_8，所有ht图块都以透明的方式在争夺所有权的相同地址上发出相同的存储操作。一致性完全由LSU和CC处理，在核心侧，LSU_het_almost_full位图声明了LSU对每个线程的可用性（写入和读取）。

// Starting multiple write operations

START_MEM_WRITE_TRANS    : begin

  if ( pending_writes )

    next_state <= IDLE;

  else

    next_state <= DONE;

  if(lsu_het_almost_full[thread_id_write] == 1'b0 ) begin

    write_served      <= 1'b1;

    req_out_valid     <= 1'b1;

    req_out_id        <= rem_writes;

    req_out_thread_id <= thread_id_write;

    req_out_op        <= 'b100000; // STORE_8

    tmp_data_out[0]   <= 8'hee;

    incr_address      <= 1'b1;

  end

end

在这两种状态下，线程首先检查存储在与其ID（lsu_het_almost_full[thread_ID]）相等的位置的可用性，然后执行内存事务。

添加自定义逻辑

提供的异构虚拟核心可以用自定义加速器或用户逻辑替换，应该将其分配到tile_ht中，替换以下行：

//  -----------------------------------------------------------------------

//  -- Tile HT - Dummy Core

//  -----------------------------------------------------------------------

het_core_example  #(

       .TILE_ID     ( TILE_ID     ),

       .THREAD_NUMB ( THREAD_NUMB )

) u_dummy_het_core (

       .clk                                                ( clk                           ),

       .reset                                              ( reset                         ),

        /* Memory Interface */

       .req_out_valid                                      ( req_in_valid                  ),

       .req_out_id                                         ( req_in_id                     ),

       .req_out_thread_id                                  ( req_in_thread_id              ),

       .req_out_op                                         ( req_in_op                     ),

       .req_out_address                                    ( req_in_address                ),

       .req_out_data                                       ( req_in_data                   ),

       ...

       /* Synchronization Interface */

       .breq_valid                                         ( breq_valid                    ),

       .breq_op_id                                         ( breq_op_id                    ),

       .breq_thread_id                                     ( breq_thread_id                ),

       .breq_barrier_id                                    ( breq_barrier_id               ),

       .breq_thread_numb                                   ( breq_thread_numb              ),

       ...

);

Programming Model

Contents

 [hide] 

• 1NaplesPU Programming Model
• 1.1SIMD Support
• 1.2Thread Parallelism
• 1.2.1Thread Synchronization
• 1.3NaplesPU Other Aspects
• 1.3.1Flush instruction
• 1.3.2Scratchpad memory
• 1.3.3Barrier Instruction
• 1.4NaplesPU Intrinsics
• 1.4.1NaplesPU Misc Intrinsics
• 1.4.2NaplesPU Vector intrinsics
• 1.4.3Standard LLVM intrinsics
• 1.5Learning by Example
• 1.5.1Matrix Multiplication Multithreaded Example
• 1.5.2Matrix Multiplication Vectorial Example
• 1.5.3How to compile a kernel
• 2OpenCL support for NaplesPU
• 2.1Execution Model Matching
• 2.2Memory Model Matching
• 2.3Programming Model Matching
• 2.4OpenCL Runtime Design
• 2.5OpenCL Example

NaplesPU编程模型

SIMD支持

算术运算符（+，-，*，/，%）、关系运算符（==，！=，<，<=，>，>=）、位运算符（&，|，^，~，<<，>>）、逻辑运算符（&&，||，！）和赋值运算符（=，+=，-=，*=，/=，%=，<<=，>=，&=，^=，|=）可以与标量和向量类型一起使用，并分别产生标量或向量符号整数结果。在某些情况下，混合标量/向量操作是可能的。在这种情况下，标量被视为一个向量，所有元素都等于标量值。

例如，要添加两个向量：

#include <stdint.h>

int main (){

  vec16i32 a;

  vec16i32 b;

  …

  vec16i32 c = a+b;

}

或具有标量的向量：

#include <stdint.h>

int main (){

 vec16i32 a;

 int b;

 …

 vec16i32 c = a+b;

}

为了访问向量元素，可以使用运算符[]。例如

#include <stdint.h>

int main (){

 vec16i32 a;

 // assign some values:

 for (int i=0; i<16; i++) a[i]=i;

 int sum = 0;

 // calculate sum

 for (int i=0; i<16; i++) sum += a[i];

}

向量可以使用花括号语法进行初始化。例如，一个常数向量：

#include <stdint.h>

int main (){

  const vec16i32 a = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 };

}

或非常数向量

#include <stdint.h>

int main (){

 int x, y, z;

 ...

 vec16i32 a = { x, y, z, x, y, z, x, y, z, x, y, z, x, y, z, x};

}

具有相同元素数量的向量之间的转换可以使用LLVM固有的__builtin_convertvector来执行。矢量类型v16i32、v16u32、v16f32可以相互转换。同样，矢量类型v8i64、v8u64、v8F64可以相互转换。例如：

#include <stdint.h>

int main (){

 vec16f32 a;

 ...

 vec16i32 d = __builtin_convertvector(a,vec16i32);

}

也可以用不同数量的元素转换浮点向量。在这种情况下，需要使用两个NPU内部函数__builtin_NPU_v8f64to16f32或__builtin-NPU_v16f32tov8f64。第一个将8个双精度FP元素转换为8个单精度FP元素，这些元素放置在v16f32向量的前8个元素中。第二个将v16f32向量的前8个单精度FP元素转换为8个双精度FP元素。例如：

#include <stdint.h>

int main (){

 vec16f32 a;

 ...

 vec8f64 b = __builtin_npu_v16f32tov8f64(a);

}

矢量比较可以通过两种不同的方式进行。可以使用传统的关系运算符，得到两个向量大小相同的另一种向量类型。例如：

#include <stdint.h>

int main (){

 vec16i32 a;

 vec16i32 b;

 …

 vec16i32 c = a < b;

}

执行上述代码后，根据比较结果，向量c元素将等于0xFFFFFFFF或0x00000000。此外，还提供向量比较内部函数，如下所示：

#include <stdint.h>

int main (){

 vec16i32 a;

 vec16i32 b;

 …

 int c = __builtin_npu_mask_cmpi32_slt (a, b)

}

执行上述代码后，整数c将包含一个位图，如果需要，可以直接用于写入掩码寄存器。在NaplesPU中，使用向量比较内部函数是执行比较的自然方式。

请注意，在NaplesPU中，所有指令都被屏蔽，并且在开始时，所有通道都被启用。如果想处理SIMD控制流，需要显式地处理掩码操作，以便它们只应用于某些元素。例如，看看上面的代码：

#include <stdint.h>

int main (){

 vec16i32 a;

 vec16i32 b;

 …

 int c = __builtin_npu_mask_cmpi32_slt (a, b) //generate mask for a<b

 int rm_old = __builtin_npu_read_mask_reg();  //save mask register

 __builtin_npu_write_mask_reg(c);             //write mask register for a<b

 do_something();

 c = c^-1;                                       //generate mask for a greater equal b

 __builtin_npu_write_mask_reg(c);             //write mask register for a greater equal b

 do_somethingelse();

 __builtin_npu_write_mask_reg(rm_old);        //restore the old mask

}

线程平行度

在NPU核心中，每个线程都运行用户提供的相同代码，通过内置功能，开发人员可以将其并行化或区分流。每个线程都有一个私有堆栈，而内存系统是分布式的，所有线程都有相同的主内存视图。用户可以根据正在运行的线程ID来区分线程流，如控制寄存器中所述，每个线程都可以通过__builtin_npu_read_control_reg访问所需的控制寄存器从控制寄存器中读取其ID（用于获取线程ID，使用2；用于核心ID，使用0或1）：

#define CORE_ID    __builtin_npu_read_control_reg(0)

#define THREAD_ID  __builtin_npu_read_control_reg(2)

在当前版本中，每个图块都配备了1个NPU核心，因此每个核心ID与其图块ID重叠。

线程同步

NaplesPU支持线程之间的屏障同步。程序员需要知道需要同步的线程数量，即NumberOfThreads。对于每次同步，都有一个唯一的屏障ID。NaplesPU固有的__builtin_npu_barrier（BarrierID，NumberOfThreads-1）负责同步。可以利用等于Bmax的最大数量的屏障，即4 x线程数。屏障ID的范围从0到Bmax-1。请注意，相同的屏障ID不能在不同的内核中使用，并且只能在内核内由相同的线程或其子集多次使用。

在下面的示例中，有四个线程和两个屏障。执行一些操作后，所有线程在屏障1上同步。然后，只有线程0和1在屏障2上同步。请记住，在主代码中，用户必须提供同步线程的总数：

#include <stdint.h>

static vec16i32 C[4];

static vec16i32 D[2];

const vec16i32 A[4]={{...}    ,{...},{...},    {...}};

const vec16i32 B[4]={{...},{...},{...},    {...}};

int main(){

 //execution thread 0,1,2,3

 int threadId = __builtin_npu_read_control_reg(2);

 C[threadId] = A[threadId] + B[threadId];

 __builtin_npu_barrier(1,3);//Synchronization Threads:0,1,2,3.

 

if(threadId<2){

//execution thread 0,1

   D[threadId]=C[threadId*2]+C[(threadId*2)+1];

   __builtin_npu_barrier(2,1);//Synchronization Threads:0,1.

 }

 if(threadId==0){

   D[threadId]=D[threadId]+D[threadId+1];

   __builtin_npu_flush((int)(&D[threadId]));

 }

 return 0;

}

NaplesPU其他方面

冲洗说明

NaplesPU ISA有一个刷新指令，需要该指令来避免数据卡在缓存中。如果主机需要输出数据，则必须使用此指令。否则，主机将从主存储器读取与缓存不一致的数据。刷新指令接收所涉及变量的地址输入，并刷新整个512位缓存行。请记住将地址转换为整数，否则将看到以下错误：“无法用类型为（YOUR VARIABLE type）*'的右值初始化类型为'int'的参数”。例如，看看上面的代码：

#include <stdint.h>

 int main (){

  vec16i32 a;

  vec16i32 b;

  …

  vec16i32 c = a + b;

  __builtin_npu_flush((int)&c);

}

flush指令适用于单个512位缓存行。因此，如果变量类型大于512位，则需要多个刷新指令。

Scratchpad存储器

除了传统的主存储器外，NPU内核还支持具有不同地址空间的暂存存储器。为了使用scratchpad内存，在声明变量时应该使用__scratchpad关键字。请注意，草稿行内存中只能放置一个全局变量。例如，看看上面的代码：

#include <stdint.h>

__scratchpad int a;

int main (){

 ...

}

编译器将识别关键字和整数，并使用适当的加载/存储指令将变量放置在草稿栏中。

屏障说明

NPU支持同一核心内的线程之间或不同图块之间的屏障同步。程序员需要知道需要同步的线程数量，即NumberOfThreads。对于每个同步，都有一个唯一的屏障ID。内部__builtin_npu_barrier（BarrierID，NumberOfThreads-1）负责同步。可以利用等于Bmax的最大数量的屏障，即4 x线程数。屏障ID的范围从0到Bmax-1。请注意，相同的屏障ID不能在不同的内核中使用，并且只能在内核内由相同的线程或其子集多次使用。

在下面的示例中，有四个线程和两个屏障。执行完一些操作后，所有线程在屏障1上同步。然后，只有线程0和1在屏障2上同步。请记住，在主机端，需要用计数=4初始化屏障0，用计数=2初始化屏障1。

#include <stdint.h>

static vec16i32 C[4];

static vec16i32 D[2];

const vec16i32 A[4]={{...}    ,{...},{...},    {...}};

const vec16i32 B[4]={{...},{...},{...},    {...}};

 

int main(){

 //执行线程 0, 1, 2, 3

  int threadId = __builtin_npu_read_control_reg(2);

  C[threadId] = A[threadId] + B[threadId];

  __builtin_npu_barrier(1,3);   //Synchronization Threads:0,1,2,3.

 

 if(threadId<2){

 //执行线程0, 1

    D[threadId]=C[threadId*2]+C[(threadId*2)+1];

    __builtin_npu_barrier(2,1);  //同步线程：0, 1。

  }

  if(threadId==0){

    D[threadId]=D[threadId]+D[threadId+1];

    __builtin_npu_flush((int)(&D[threadId]));

  }

  return 0;

}

NaplesPU内部函数

NaplesPU 其他内在函数

下表总结了NaplesPU的主要内置功能：

Intrinsic name (when using these intrinsics you should add “__builtin_npu_” as prefix )	Operation	Corresponding Instruction
void barrier (int a, int b)	Refer to the description of the barrier instruction. a contains the barrier ID, while b contains the number of threads that should synchronize-1	barrier
void flush (int a)	Flush a cache line to the main memory. a contains the memory address of the cache line. Explicit integer conversion is required.	flush
int createmaskv16i32 (v16i32 a)	Convert the vector a that is made of all elements equal to 0 or -1 into a 32-bit mask value that can be written in the mask register. It can be used to compute the bitmask when a vector comparison operation is performed using the common C/C++ relational operators.	crt_maskv16
void write_mask_reg(int a)	Write a 32-bit bitmask inside the mask register	move
void write_mask_regv16i32 (v16i32 a)	Write a 512-bit vector mask inside the mask register.	crt_maskv16 + move
int read_mask_reg ()	Read a 32-bit bitmask from the mask register.	move
void write_control_reg (int a, int b)	Write values to the mask register. The integer a contains the ID of the sub-register to access, while the integer b contains the data.	write_cr
int read_control_reg (int a)	Read values from the mask register. The integer a contains the ID of the sub-register to access.	read_cr

 

NaplesPU矢量内部函数

Intrinsic name (when using these intrinsics you should add “__builtin_npu_” as prefix )	Operation	Corresponding Instruction
vec16i32 makevectori32 (int a)	Create a vector of 16 elements whose elements are all equal to a	move_i32
vec16f32 makevectorf32 (float a)	Create a vector of 16 elements whose elements are all equal to a	move_i32
vec16i32 shufflei32 (vec16i32 a, vec16i32 b) or vec16i32 shufflef32 (vec16f32 a, vec16i32 b)	Vector Shuffle - it allows elements of vector a to be copied to the output vector in different positions. The elements in vector b specify, for each corresponding position in the destination register, the indexes of the elements in the source vector register. (see the description of the shuffle instruction)	shuffle_i32 or shuffle_f32
int mask_cmp'w'32_'xyz' (vec16'w'32 a, vec16'w'32 b)	Return an integer bitmask where the i-th bit is equal to one if (a[i] comp b[i]) is true. Otherwise, the i-th bit is equal to zero. w = type of compare: i for interger and f for float. x = s if signed or u otherwise. yz = traditional compare conditions, namely: gt, ge, lt, le	cmp'xyz'_'w'32
int mask_cmp'w'32_eq (vec16'w'32 a, vec16'w'32 b)	Return an integer bitmask where the i-th bit is equal to one if (a[i] comp b[i]) is true. Otherwise, the i-th bit is equal to zero. w = type of compare: i for interger and f for float.	cmpeq_'w'32
int mask_cmp'w'32_ne (vec16'w'32 a, vec16'w'32 b)	Return an integer bitmask where the i-th bit is equal to one if (a[i] comp b[i]) is true. Otherwise, the i-th bit is equal to zero. w = type of compare: i for interger and f for float.	cmpne_'w'32

矢量比较可以通过两种不同的方式进行。可以使用传统的关系运算符，得到两个向量大小相同的另一种向量类型。例如：

#include <stdint.h>

int main (){

  vec16i32 a;

  vec16i32 b;

  …

  vec16i32 c = a < b;

}

执行上述代码后，根据比较结果，向量c元素将等于0xFFFFFFFF或0x00000000。此外，还提供向量比较内部函数，如下所示：

#include <stdint.h>

int main (){

  vec16i32 a;

  vec16i32 b;

  …

  int c = __builtin_npu_mask_cmpi32_slt (a, b)

}

执行上述代码后，整数c将包含一个位图，如果需要，可以直接用于写入掩码寄存器。使用向量比较内部函数是在NPU中执行比较的自然方式。

请注意，在NPU中，所有矢量操作都被屏蔽，启动后所有硬件通道都被启用。如果想处理SIMD控制流，需要显式地处理掩码操作，以便它们只应用于某些元素。例如，看看上面的代码：

#include <stdint.h>

int main (){

  volatile vec16i32 a;

  volatile vec16i32 b;

  …

  int c = __builtin_npu_mask_cmpi32_slt (a, b) //generate mask for a lt b

  int rm_old = __builtin_npu_read_mask_reg();  //save mask register

  __builtin_npu_write_mask_reg(c);             //write mask register for a lt b

  do_something();

  c = c^-1;                                       //generate mask for a ge b

  __builtin_npu_write_mask_reg(c);             //write mask register for a ge b

  do_somethingelse();

  __builtin_npu_write_mask_reg(rm_old);        //restore the old mask

}

在这些情况下，重要的是要防止编译器对所涉及的向量进行优化，并将其标记为易失性变量。编译器倾向于重新排列操作，最好仔细检查编译器objdumb以检查操作的顺序。如果volatile不能阻止重新排序，可以将代码嵌入到函数中或使用-O0作为优化标志。

标准LLVM内部函数

Intrinsic name	Operation	Corresponding Instruction
vec16X __builtin_convertvector(vec16Y a, vec16X)	It is is used to express generic vector type-conversion operations. The input vector and the output vector type must have the same number of elements. X can be equal to i32/f32, while Y can be equal to i32/f32	sext or itof or ftoi

具有相同元素数量的向量之间的转换可以使用LLVM固有的__builtin_convertvector来执行。矢量类型v16i32、v16u32、v16f32可以相互转换。例如：

#include <stdint.h>

int main (){

  vec16f32 a;

  ...

  vec16i32 d = __builtin_convertvector(a,vec16i32);

}

示例学习

矩阵乘法多线程示例

以下代码显示了在NPU内核上运行的矩阵乘法内核的多线程版本。该代码将输出计算分散到所有可用线程中。由于每个输出元素的计算都是独立的，因此线程并行性是通过划分函数的外循环来实现的。输出矩阵行计算在所有核心上均匀分布，并在线程上进一步分布。对于每个线程，该函数首先计算输出矩阵中要在核心级别计算的部分：N/core_NUMB，其中N是矩阵的维度，core_NUMB是系统中NPU核心的数量。然后，每个线程开始计算start_loop=（core_id*N/core_NUMB）+thread_id的外循环，将要计算的输出矩阵部分乘以正在运行的核心id（core_id），再加上正在运行的线程id（thread_id），每次迭代增加核心thread_NUMB中的线程数。

void matrix_mult(const int a[N][N], const int b[N][N], int mult[N][N], int core_id, int thread_id) {

 int start_loop = (core_id * N / CORE_NUMB) + thread_id;

 int end_loop = N / CORE_NUMB * (core_id + 1);

 

 for (int i = start_loop; i < end_loop; i += THREAD_NUMB){

   for (int j = 0; j < N; j++)

     for (int k = 0; k < N; k++)

       mult[i][j] += a[i][k] * b[k][j];

 }

}

参数core_id和thread_id由main函数传递，并通过内置函数从NPU控制寄存器中获取：

#define CORE_ID    __builtin_npu_read_control_reg(0)

#define THREAD_ID  __builtin_npu_read_control_reg(2)

这样，每个核心中的每个线程都会与其他线程同时计算输出矩阵的一部分。

当系统上的所有线程结束分配的任务时，最终结果就准备好了，需要系统同步，在主要功能中，这是通过编程模型提供的内置屏障来实现的：

__builtin_npu_barrier(42, CORE_NUMB * THREAD_NUMB - 1);

上述内置程序同步ID 42上的CORE_NUMB*THREAD_NUMB线程数（系统中运行的线程总数）。当所有线程都遇到障碍时，输出矩阵就准备好了，尽管其中大部分可能在私有L1缓存中。下一步是将输出线刷新到主存储器中：

if (THREAD_ID == 0 && CORE_ID == 0) {

     for (int i = 0; i < N*N; i += 64 / sizeof(int)) {

       __builtin_npu_flush((int) &C[i / N][i % N]);

     }

     __builtin_npu_write_control_reg(N*N, 12); // For cosimulation purpose

   }

通常，刷新操作由线程执行，在这种情况下，第一个内核中的第一个线程调用刷新内置程序，该程序将L1缓存中的输出结果发送到主内存。

主要功能的完整代码如下：

#define CORE_ID    __builtin_npu_read_control_reg(0)

#define THREAD_ID  __builtin_npu_read_control_reg(2)

 

int main(){

   init_matrix(A);

   init_matrix(B);

 

   matrix_mult(A, B, C, CORE_ID, THREAD_ID);

   __builtin_npu_barrier(CORE_ID + 1, THREAD_NUMB - 1);

 

   if (THREAD_ID == 0 && CORE_ID == 0) {

     for (int i = 0; i < N*N; i += 64 / sizeof(int)) {

       __builtin_npu_flush((int) &C[i / N][i % N]);

     }

     __builtin_npu_write_control_reg(N*N, 12); // For cosimulation purpose

   }

 

  return (int)&C;

}

 

矩阵乘法向量示例

另一方面，矩阵乘法函数的矢量版本以SIMD多线程方式计算输出矩阵。输入和输出矩阵都以特定于目标的向量类型组织，成为向量的向量。这种组织导致N列部分结果在16个硬件通道上分布；每个线程每个周期计算N个部分结果。矩阵的大小必须是16的倍数。

void kernel_function(vec16i32 *A, vec16i32 *B, vec16i32 *C, int N) {

   uint32_t coreId =  __builtin_npu_read_control_reg(0);

   uint32_t threadId = __builtin_npu_read_control_reg(2);

   uint32_t nT = 2; // number of threads

   uint32_t nL = 16; // number of lanes

   uint32_t nC = N/nL;

   uint32_t ndivnT = N/nT;

   uint32_t tIdndivnT = threadId*ndivnT;

   uint32_t tIdndivnTnC = tIdndivnT*nC;

 

   for (uint32_t i = coreId; i < ndivnT*nC; i+=CORE_NUMB){

       uint32_t col = (tIdndivnT+i)%nC;

       C[tIdndivnTnC+i] = 0;

       for (uint32_t j = 0; j < nC; j++){

           for (uint32_t k = 0; k < nL; k++){

               C[tIdndivnTnC+i] += A[tIdndivnTnC+i-col+j][k] * B[(nC*k)+(j*N)+col];

            }

       }

   }

}

请注意，C[tIdndivnTnC+i]+=A[tIdndivnTnC+i-col+j][k]*B[（nC*k）+（j*N）+col]一次对16个不同的数据执行16个操作。代码的组织和线程并行化等效于其标量版本。

如何编译内核

目前，NaplesPU工具链发布时，一些示例内核位于npu/software/kernes文件夹中。提供makefile来为NaplesPU编译这些内核。如果想添加一个新的内核，建议复制一个内核文件夹，并用自己的源代码替换C/CPP文件。然后，记得修改makefile，用当前的主C/CPP文件名更新SRCS变量。

NaplesPU的OpenCL支持

OpenCL将平台定义为主机连接到的一组计算设备。每个设备进一步分为几个计算单元（CU），每个计算单元都定义为一组处理元素（PE）。回想一下，目标平台是围绕单个核心进行架构设计的，其结构最多为八个硬件线程。每个硬件线程相互竞争，以访问16个硬件通道，对32位宽的整数或浮点操作数执行标量和向量操作。

计算设备抽象在物理上映射到NaplesPU多核架构上。NaplesPU多芯可以根据芯数进行配置。每个NPU核心映射到OpenCL计算单元上。在内部，NPU核心由硬件线程组成，每个线程都代表OpenCL处理元素的抽象。

执行模型匹配

从执行模型的角度来看，OpenCL依赖于一个N维索引空间，其中每个点代表一个内核实例执行。由于物理内核实例的执行是由硬件线程完成的，因此OpenCL工作项被映射到NPU单个硬件线程上。因此，工作组被定义为一组硬件线程，工作组中的所有工作项都在单个计算单元上执行。

内存模型匹配

OpenCL将内存划分为四个不同的空间：

•全局和常量空间：所有工作组中的所有工作项都可以访问这些空间中的元素。

•本地空间：仅对工作组内的工作项可见。

•私人空间：仅对单个工作项可见。

目标平台提供DDR内存，即OpenCL术语中的设备内存。因此，变量在物理上映射到此内存上。编译器本身通过查看地址空间限定符来验证是否满足OpenCL约束。

每个NPU内核还配备了一个Scratchpad存储器，这是每个内核独有的片上非相干存储器部分。此内存符合OpenCL本地内存功能。

最后，NPU内核中的每个硬件线程都有一个私有堆栈。此内存部分对每个硬件线程（即OpenCL工作项）都是私有的，其他线程无法寻址。因此，每个堆栈都充当OpenCL私有内存。

编程模型匹配

OpenCL支持两种编程模型，数据并行和任务并行。数据并行模型要求OpenCL索引空间的每个点执行一个内核实例。由于每个点代表一个工作项，并且这些工作项映射到硬件线程上，因此正确地满足了数据并行要求。请注意，实现的模型是一个宽松的版本，不需要对数据进行严格的一对一映射。

任务并行编程模型要求内核实例在索引空间的任何点上独立执行。在这种情况下，每个工作项都不受限于执行其他工作项的相同内核实例。编译器前端定义了一组可用于此目的的内置程序。此外，每个NPU核心由16个硬件通道构建，有助于实现锁步执行。因此，OpenCL支持被实现为允许使用向量类型。因此，以下数据类型支持向量执行：

•charn和ucharn分别映射到vec16i8和vec16u8上，其中n=16。不支持n的其他值。

•shortn、ushortn分别映射到vec16i16和vec16i32上，其中n=16。不支持n的其他值。

•intn、uintn分别映射到vec16i32和vec16u32上，其中n=16。不支持n的其他值。

•floatn，映射在vec16f32上，其中n=16。不支持n的其他值。

OpenCL运行时设计

OpenCL API是一组用于协调和管理设备的功能，这些功能为运行应用程序和监控其执行提供支持。这些API还提供了一种检索设备相关信息的方法。

下图描述了OpenCL规范中定义的OpenCL运行时的UML类图。灰色填充框表示由于缺乏硬件支持而无法使用功能。

自定义OpenCL运行时依赖于两个主要抽象：

•低级抽象，不完全依赖于硬件，提供设备主机通信支持。

•根据OpenCL API，高级抽象管理设备上运行的内核的生命周期。

OpenCL示例

以下代码显示了在NPU设备上运行的OpenCL中的矢量矩阵乘法。

#include <opencl_stdlib.h>

#define WORK_DIM 4

 

__kernel void kernel_function(__global int16 *A, __global int16 *B, __global int16 *C, int rows, int cols)

{

   __private uint32_t threadId = get_local_id(0);

 

   uint32_t nT = WORK_DIM; // number of threads

   uint32_t nL = 16;       // number of lanes

 

   uint32_t N = rows;

   uint32_t nC = N / nL;

   uint32_t ndivnT = N / nT;

   uint32_t tIdndivnT = threadId * ndivnT;

   uint32_t tIdndivnTnC = tIdndivnT * nC;

   for (uint32_t i = 0; i < ndivnT * nC; i++)

   {

       uint32_t col = (tIdndivnT + i) % nC;

       C[tIdndivnTnC + i] = 0;

       for (uint32_t j = 0; j < nC; j++)

       {

           for (uint32_t k = 0; k < nL; k++)

           {

               C[tIdndivnTnC + i] += A[tIdndivnTnC + i - col + j][k] * B[(nC * k) + (j * N) + col];

           }

       }

   }

}

 

 

http://www.naplespu.com/doc/index.php?title=Main_Page

http://www.naplespu.com/

https://github.com/AlessandroCilardo/NaplesPU

https://github.com/AlessandroCilardo/NaplesPU-toolchain