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
http://www.naplespu.com/doc/index.php?title=Detailed_studies
主页
那不勒斯处理单元,被称为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++)编写的复杂应用程序,如矩阵乘法、矩阵转置或类似应用程序。
所需软件
任何内核的仿真或实现都依赖于以下依赖关系:
1)Git
2)Xilinx Vivado 2018.2或ModelSim (例如,Questa Sim-64 vsim 10.6c_1)
3)NaplesPU工具链
构建过程
第一步是通过克隆,从官方存储库中获取NaplesPU架构的源代码。
在Ubuntu Linux环境中,通过启动以下命令来完成此步骤:
$ git clone https://github.com/AlessandroCilardo/NaplesPU
在NaplesPU存储库中,工具链是存储库的git子模块,因此需要创建和更新。在Ubuntu Linux环境中,只需在存储库的根文件夹中,键入以下命令:
$ git submodule update --init
然后,第三步是安装工具链。[此处]描述了这一过程。
仿真内核
以下文件夹对此特别感兴趣:
1)软件,存储所有内核;
2)工具,存储所有用于仿真的脚本。
仿真内核有三种方法:
1)启动test.sh脚本
2)如果仿真器软件是Vivado,则从存储库的根文件夹启动setup_project.sh;
3)如果仿真器软件是ModelSim,则从存储库的根文件夹启动simulate.sh。
首先,在shell中源代码Vivado或ModelSim。这一步对所有方式都是强制性的。在Ubuntu Linux环境中:
$ source Vivado/folder/location/settingXX.sh
其中XX取决于安装的Vivado版本(32或64位)。
test.sh脚本
test.sh脚本位于npu/tools文件夹中,运行其中列出的所有内核,并将npu的输出与标准x86架构产生的预期结果进行比较:
$ ./test.sh[选项]
选项包括:
--h, --help,显示此帮助
-t, --tool=vsim或vivado,指定要使用的工具,默认值:vsim
-cn,--core-numb=VALUE,指定核编号,默认值:1
-tn,--thread-numb=VALUE,指定线程号,默认值:8
test.sh脚本会自动编译内核,并在NaplesPU和x86架构上运行。一旦仿真终止,对于每个内核,Python脚本会比较两次执行的结果,以验证其正确性。
在工具文件夹中,文件cosim.log存储仿真器的输出。
setup_project.sh脚本
setup_project.sh脚本可以从项目的根目录按如下方式运行:
$tools/vivado/setup_project.sh[选项]
选项包括:
--h, --help,显示此帮助
--k, --kernel=KERNEL_NAME,指定要使用的内核
--s, --single-core,选择单核配置,默认情况下选择多核
--c, --core-mask=VALUE,指定核激活掩码,默认值:1
--t, --thread-mask=VALUE,指定线程激活掩码,默认为FF
--m, --mode=gui i或batch,指定工具模式,可以在gui或batch模式下运行,默认值:gui。
此脚本启动命令中指定的内核。内核在NaplesPU架构上运行之前,应该已经编译好了:
tools/vivado/setup_project.sh -k mmsc -c 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, --help,显示此帮助
-k, --kernel=KERNEL_NAME ,指定要使用的内核
-s, --single-core,选择单核配置,默认情况下选择多核
-c, --core-mask=VALUE,指定核激活掩码,默认值:1
-t, --thread-mask=VALUE,指定线程激活掩码,默认为FF
-m, --mode=gui或batch,指定工具模式,可以在gui或batch模式下运行,默认值:gui
此脚本启动命令中指定的内核。内核在NaplesPU架构上运行之前,应该已经编译好了:
完整文档
1)NaplesPU硬件架构
2)NaplesPU工具链
3)NaplesPU指令集架构
4)加长型NaplesPU
5)异质分块
6)编程模型
Naplespu微体系结构内多个插件融合分析
http://www.naplespu.com/doc/index.php?title=Detailed_studies
介绍
打算展示替换链是如何相对于L2缓存发生的,观察目录控制器是如何演变以管理此类情况的发生的。主要目的是测试发生L2替换的所有情况,如图所示:
目录控制器中的内存
为了更好地理解替换机制的工作原理,有必要首先评估目录控制器管理的缓存的结构。以下描述表示用于所有实验的配置。如图所示,缓存由四种方式和64组组成,每种方式包含一个标签和一个数据(64位)。
管理地址为32位,其组织方式如图所示:
然后,每个控制器目录可以处理有限数量的地址,例如,图块0可以处理从0x00000000到0x3ffffff的地址,图块1可以处理从0x4000000到0x7ffffff的地址,以此类推。实验中使用了四块分块进行处理。
更换工具链
为了进行替换,需要填写与给定集合相关的所有四种方式:在同一集合中插入新条目时,实现替换。为了简单起见,需要编写代码,以便在使用的内核中直接在Assembly中运行,从而更好地控制执行的操作和使用的内存地址。为了确定正在进行替换,有必要观察do_replacement信号。此信号仅在以下情况下为高。
do_replacement = dc2_message_valid && ((allocate_cache | | update_cache) && !deallocate_cache) && !is_replacement && !dc2_message_cache_hit && dc2_message_cache_valid;
因此,如果已经处理了替换以外的请求,并且没有导致缓存行被释放,那么如果缓存中没有命中,那么将需要执行替换。可以注意到,如果没有缓存释放,则会触发替换操作,这可能违反直觉。事实上,替换操作的管理方式是直接用要插入的新行替换条目LRU,而不会使执行替换的缓存行无效,将无效操作(针对缓存控制器)委托给稍后的阶段。在do_replacement信号断言后,通过dc3替换入队信号,触发相对队列中替换请求的入队,在该队列中还插入要执行替换的缓存行的内容:
dc3_replacement_enqueue = dc2_message_valid && do_replacement;
assign dc3_replacement_request.source = dc2_message_source,
dc3_replacement_request.memory_address.tag = dc2_message_cache_tag,
dc3_replacement_request.memory_address.index = dc2_message_address.index,
dc3_replacement_request.memory_address.offset = 0,
dc3_replacement_request.data = dc2_message_cache_data,
dc3_replacement_request.state = dc2_message_cache_state,
dc3_replacement_request.分享s_list = dc2_message_cache_分享s_list,
dc3_replacement_request.owner = dc2_message_cache_owner;
注意,dc2消息address.index包含从第2阶段获取的LRU索引。
如目录控制器文档中所述,第1阶段包含一个固定优先级的调度器,如果可以发出替换请求的条件得到验证,则该调度器涉及替换请求的最高优先级处理。除了通常的信号外,阶段1还向阶段2发送以下信号:
dc1_replacement_state <= output_replacement_state;
dc1_replacement_分享s_list <= output_replacement_分享s_list;
dc1_replacement_owner <= output_replacement_owner;
这些信号从阶段2转发到阶段3。在实际处理请求的状态3中,根据所在的状态,按照协议ROM中的定义进行替换。
更换一条线路
为了测试M状态下块的替换,并评估其文档中描述的主要信号,就目录控制器而言,构建了以下内核:
if(tile_id==0){
asm(
"moveih s20, 0x0000"
"moveil s20, 0x0000"
"store32 s21, (s20)"
"moveil s20, 0x1000"
"store32 s21, (s20)"
"moveil s20, 0x2000"
"store32 s21, (s20)"
);
}else if(tile_id==1){
asm(
"moveih s20, 0x0000"
"moveil s20, 0x4000"
"store32 s21, (s20)"
"moveil s20, 0x8000"
"store32 s21, (s20)"
);
}
在这种情况下,图块0运行三个存储,图块1运行两个存储。为了区分不同的请求,请观察图中的dc1消息有效信号,当固定优先级调度器从相应的队列中选择一个请求时,该信号被断言。
从仿真中可以看出,目录控制器按以下顺序处理消息:
1.分块1存储0x00004000
2.分块0存储在0x00000000
3.分块1存储0x00008000
4.分块0存储在0x00001000
5.分块0存储在0x00002000
这些请求是以属于同一组的方式提出的,以便触发替换。
根据请求,缓存分配如下:
为了获得上述信息,评估了dc3_update_cache_way信号,该信号允许在L2中写入条目。此信号直接取决于selected_way,如果是hit,则包含hit_index,否则包含lru_way(参阅与伪lru相关的部分)。
指定选定way = hit ? hit idx : lru_way;
或者,可以观察dc1_message_cache_state信号是如何演变的。详细看看在一个简单请求的情况下,会发生什么:
缓存是通过dc3_update_cache_state信号完成的。从下图中,可以看到如上所述,条目是如何以方式0分配的。
在第五次存储时,L2缓存已满,替换机制开始。为了理解它的功能,可以从理论上观察哪些信息应该交换,然后获得实际反馈:
从下图可以观察到,请求被执行,缓存行被插入,而不是lru-way信号指示的条目;因此,可以检测到高do_replacement信号,从而也检测到do_replacement_enqueue信号),这涉及在替换队列中插入请求。
此时,通过观察各阶段信号的递增顺序,检查管道内的更换操作是如何进行的。
阶段1:在图中,可以看到条件can_issue_replacement_request得到满足,因此请求从其队列中取出。此请求的类型为“c”,即替换。
此时发送的消息取决于协议ROM,可以从以下代码摘录中看到:
{STATE_S, REPLACEMENT, 1'b?, 1'b?} : begin // 替换
// 发送BACKINV分享
dpr_output.message_forwarded_send = 1'b1;
dpr_output.message_forwarded_type = MESSAGE_BACKINV;
dpr_output.message_forwarded_to_shares = 1'b1;
// Send WB To Memory Controller
dpr_output.message_response_send = 1'b1;
dpr_output.message_response_type = MESSAGE_WB;
dpr_output.message_response_has_data = 1'b1;
dpr_output.message_response_to_memory = 1'b1;
// Next State N
dpr_output.next_state = STATE_N;
end
首先,需要向分享发送BACK_INV消息,向内存控制器发送WB消息:
=======================
目录控制器 - [Time 12810] [TILE 0] – 信号发送
转发目的地: 0010
源: 0
地址: 00004000
申请者: DCACHE
数据包类型: BACK_INV
不连贯: 0
响应目的地: 1000
源: 0
来自DC: 1
地址: 00004000
申请者: DCACHE
数据包类型: WB
数据: xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx
不连贯: 0
分享次数: 1
=======================
缓存控制器 - [时间 12890] [TILE 1] [Core 0]
源: 0
地址: 00004000
申请者: DCACHE
数据包类型: BACK_INV
不连贯: 0
与前面的示例不同,可以从状态S直接转换到状态N。
案例2:多个共享者
为了测试用单个分享替换S状态的块,并评估文档中描述的主要信号,就目录控制器而言,构建了以下内核:
if(tile_id==0){
asm( "moveih s20, 0x0000"
"moveil s20, 0x4000"
"load32 s22, (s20)"
"moveil s20, 0x0000"
"store32 s21, (s20)"
"moveil s20, 0x1000"
"store32 s21, (s20)"
"moveil s20, 0x2000"
"store32 s21, (s20)"
);
}else if(tile_id==1){
asm( "moveih s20, 0x0000"
"moveil s20, 0x4000"
"load32 s22, (s20)"
"moveil s20, 0x8000"
"store32 s21, (s20)"
);
}
在这种情况下,图块0执行加载和三个存储,图块1执行加载和一个存储。
为了区分不同的请求,观察dc2_message_valid信号,当固定优先级调度器从相应的队列中选择一个请求时,该信号会被断言。
从模拟中可以看出,目录控制器按以下顺序处理消息:
1.将分块1加载到地址0x00004000
2.将分块0加载到地址0x00004000
3.分块1存储0x00008000
4.分块0存储在0x00000000
5.分块0存储在0x00001000
6.分块0存储在0x00002000
随后在缓存中进行分配:
这些请求是以属于同一组的方式提出的,以便触发替换。
从理论角度来看,观察针对不同共享者的BACK_INV消息很有趣:
可以通过提高细节级别来发现,参考消息日志,由文件display_coherence.txt表示:
=======================
目录控制器 - [Time 13980] [TILE 0] – 信息已发出
转发目的地: 0011
源: 0
地址: 00004000
申请者: DCACHE
数据包类型: BACK_INV
不连贯: 0
响应目的地: 1000
源: 0
来自DC: 1
地址: 00004000
申请者: DCACHE
数据包类型: WB
数据: xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx
不连贯: 0
分享次数: 2
=======================
缓存控制器 - [Time 14080] [TILE 1] [Core 0]
源: 0
地址: 00004000
申请者: DCACHE
数据包类型: BACK_INV
不连贯: 0
=======================
缓存控制器 - [Time 15860] [TILE 0] [Core 0]
源: 0
地址: 00004000
申请者: DCACHE
数据包类型: BACK_INV
不连贯: 0
与单独分享情况的主要区别在于,需要向替换缓存行中的所有分享,发送无效消息。
替换一条线路
为了测试L2缓存状态中条目的替换,必须使用更复杂的内核,因为L1缓存大小与L2缓存大小相同,因此必须使用两个不同的目录控制器。
使用这种类型的实验,而不是增加L2缓存的大小,更具形成性,因为除了测试和理解所使用的PLRU算法外,可以观察到使用两个不同的DC来处理请求。此外,如引言中所述,有必要对所有实验保持相同的配置。
下面是使用的代码:
"moveih s20, 0x0000"
"moveil s20, 0x0000"
"store32 s21, (s20)"
"moveil s20, 0x1000"
"load32 s22, (s20)"
"moveil s20, 0x2000"
"load32 s22, (s20)"
"moveil s20, 0x4000"
"load32 s22, (s20)"
"moveih s20, 0x4000"
"moveil s20, 0x0000"
"store32 s21, (s20)"
"moveih s20, 0x0000"
"moveil s20, 0x1000"
"store32 s21, (s20)"
"moveil s20, 0x2000"
"store32 s21, (s20)"
"moveil s20, 0x4000"
"store32 s21, (s20)"
"moveil s20, 0x8000"
"store32 s21, (s20)"
由于希望完全控制收到到目录控制器的请求,因此使用了一个分块。这是因为系统使用确定性路由,并意味着请求将在发送方发送时到达目的地。
从代码中可以看到,Tile 0生成了9个请求:
1.存储在0x00000000
2.加载到地址0x00001000
3.加载到地址0x00002000
4.加载到地址0x00004000
5.存储在0x40000000
6.存储在0x00001000
7.存储在0x00002000
8.存储在0x00004000
9.存储在0x00008000
如目录控制器中的内存部分所定义的,请求5是发往图块1的目录控制器的。
前4个请求需要填充L2缓存,如下所示:
为了获得进入状态I的条目,上述缓存行的所有者缓存控制器必须替换它(L1)。对DC的额外请求将导致其中一条线路在L2级别被替换,因此为了触发整个过程,Tile 0缓存控制器将向Tile 1目录控制器发出请求(请求5)。
为了验证请求5,可以通过查看消息日志来移动到更高的抽象级别:
=======================
缓存控制器 - [Time 13700] [TILE 0] [Core 0] - 消息请求已发送
目录目标: 1
源: 0
地址: 40000000
数据包类型: GETM
数据: 00000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000
=======================
目录控制器 - [Time 13780] [TILE 1] – 消息已接收
源: 0
地址: 40000000
数据包类型: GETM
数据: 00000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000
还要注意图\ref{fig:PUTM}和消息日志提取中L1替换后,Tile 0目录控制器收到的相对PUTM(类型3请求)。
=======================
缓存控制器 - [Time 14470] [TILE 0] [Core 0] – 信息请求已发送
目录目标: 0
源: 0
地址 00000000
数据包类型: PUTM
数据: 44618020820180007600001411018640606403206164000060600330616000007608006820f3d00082f40000040000c04b0c2010600002e061000000xxxxxxxx
=======================
目录控制器 - [Time 14610] [TILE 0] – 信息已接收
源: 0
地址: 00000000
数据包类型: PUTM
数据: 44618020820180007600001411018640606403206164000060600330616000007608006820f3d00082f40000040000c04b0c2010600002e061000000xxxxxxxx
=======================
目录控制器 - [Time 14630] [TILE 0] – 信息已发送
响应目的地: 0001
源: 0
来自DC: 1
地址: 00000000
申请者: DCACHE
数据包类型: PUT_ACK
数据: xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx
不连贯: 0
分享次数: 0
从下图中,可以通过信号dc1_message_cache_state(包含缓存行的状态)观察缓存的状态是如何演变的,状态3表示状态“I”:
图中表示的请求6、7、8是必要的,以便将selected_way与状态I中包含缓存行的方式相匹配,因为它不再是lru。
请求更改LRU条目的状态如下:
下图显示了一个示例,说明在请求输入状态后,状态如何从S(2)更改为M(1),从而导致lru_way信号的演变:
此时,请求9生成替换:
从理论上讲,可以观察到消息的交换(见下图),在这种情况下,这要简单得多,因为该块仅由目录控制器拥有。
在这一点上,检查更换操作是如何在管道内详细进行的,观察各阶段信号的演变情况,然后:
阶段1:在下图中,可以看到条件can_issue_replacement_request已满足,因此请求已从队列中取出。此请求的类型为“c”,即替换。
第二阶段:请求被转发到第三阶段,可以检查所选的方式在未命中的情况下,如何与lru_way完全相等,即使它不会在第三阶段使用,因为会采用请求附带的方式。
第三阶段:is_replacement信号为高,因此它实际上是一个替换。
参阅图中触发替换的请求和替换本身的详细信息:
在下图中,可以详细看到与第一阶段进入状态I相关的状态变化。
根据ROM协议的规定,它向存储器控制器生成WB:
{STATE_I, REPLACEMENT, 1'b?, 1'b?} : begin // 替换
// 向内存控制器发送WB
dpr_output.message_response_send = 1'b1;
dpr_output.message_response_type = MESSAGE_WB;
dpr_output.message_response_has_data = 1'b1;
dpr_output.message_response_to_memory = 1'b1;
// 下一状态N
dpr_output.next_state = STATE_N;
end
可以通过提高细节级别来看到,参考消息日志,由文件display_coherence.txt表示:
=======================
目录控制器 - [Time 15940] [TILE 0] – 信号已发送
响应目的地: 1000
源: 0
来自DC: 1
地址: 00000000
申请者: DCACHE
数据包类型: WB
数据: 618020820180007600001411018640606403206164000060600330616000007608006820f3d00082f40000040000c04b0c2010600002e061000000xxxxxxxx
不连贯: 0
分享次数: 0
=======================
内存控制器 - [Time 16360] [TILE 3] – 信号已接收
源: 0
来自DC: 1
地址: 00000000
申请者: DCACHE
数据包类型: WB
数据: 618020820180007600001411018640606403206164000060600330616000007608006820f3d00082f40000040000c04b0c2010600002e061000000xxxxxxxx
不连贯: 0
分享次数: 0
不稳定状态下的更换
值得注意的是,不稳定的状态是无法测试的。这是由于观察了控制器目录的结构。为了处理替换,必须通过循环缓冲区将请求带入控制器,但由于缓存中只存储了稳定状态(N除外),因此无法评估稳定状态以外的状态。
参考文献链接
http://www.naplespu.com/doc/index.php?title=Detailed_studies
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位。
最后,还有一个由几个子寄存器组成的控制寄存器。一些信息在所有线程之间共享,另一些信息是特定于线程的,标记为“线程”的寄存器每个线程都有一个单独的实例。
寄存器
|
读/写
|
分享/线程
|
描述
|
ID
|
TILE_ID
|
读
|
分享
|
Tile ID
|
0
|
CORE_ID
|
读
|
分享
|
Core ID
|
1
|
THREAD_ID
|
读
|
分享
|
ThreadID
|
2
|
GLOBAL_ID
|
读
|
分享
|
全局ID,之前的ID合并如下:TILE_ID、CORE_ID、THREAD_ID
|
3
|
GCOUNTER_LOW
|
读
|
分享
|
全局计数寄存器的低位部分,用于对重置后的处理器周期进行计数
|
4
|
GCOUNTER_HIGH
|
读
|
分享
|
全局计数器寄存器的高位部分,用于对重置后的处理器周期进行计数
|
5
|
THREAD_EN
|
读
|
分享
|
线程启用掩码,每个线程1位
|
6
|
MISS_DATA
|
读
|
分享
|
L1数据缓存未命中次数
|
7
|
MISS_INSTR
|
读
|
分享
|
L1指令缓存未命中计数
|
8
|
PC
|
读
|
分享
|
当前PC
|
9
|
TRAP_REASON
|
读
|
分享
|
陷阱原因
|
10
|
THREAD_STATUS
|
读/写
|
分享
|
线程状态2
|
11
|
ARGC
|
读/写
|
分享
|
argv指向的字符串数
|
12
|
ARGV
|
读/写
|
分享
|
传递给main()的命令行参数地址
|
13
|
THREAD_NUMB
|
读
|
分享
|
硬件线程总数
|
14
|
THREAD_MISS_CC
|
读
|
分享
|
由于内存操作,线程处于空闲状态时,每个线程的时钟周期。
|
15
|
KERNEL_WORK
|
读
|
分享
|
每个线程的内核时钟周期。
|
16
|
CPU_CTRL_REG
|
读/写
|
分享
|
CPU模式寄存器。目前,只实现了缓存控制器使用的写入策略。0表示回写,1表示直写
|
17
|
UNCOHERENCE_MAP
|
读/写
|
分享
|
处理控制寄存器中的非相干表。它存储有关非相干存储区域的信息。用户可以定义寻址此专用寄存器的非相干区域。
|
19
|
DEBUG_BASE_ADDR
|
读/写
|
分享
|
调试寄存器基址。NPU配备了16个调试寄存器。DEBUG_BASE_ADDR获取第一个调试寄存器的值,第二个调试寄存器获取DEBUG_BASE_ADDR+1,以此类推。
|
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列显示存储特定类型值的寄存器类型。
在给定寄存器文件宽度的情况下,突出显示的类型是架构本机支持的类型。其他的是通过扩展获得的,因此它们可以被视为受支持的。它们的优点在于更有效地使用系统内存。
类型
|
LLVM 类型
|
寄存器
|
注释
|
bool
|
i1
|
scalar (32 bits)
|
扩展到32位
|
char
|
i8
|
scalar (32 bits)
|
扩展到32位
|
short
|
i16
|
scalar (32 bits)
|
扩展到32位
|
int
|
i32
|
scalar (32 bits)
|
|
float
|
f32
|
scalar (32 bits)
|
|
vec16i8, vec16u8
|
v16i8
|
vector (16 x 32 bits)
|
它被扩展为32位矢量
|
vec16i16, vec16u16
|
v16i16
|
vector (16 x 32 bits)
|
它被扩展为32位矢量
|
vec16i32, vec16u32
|
v16i32
|
vector (16 x 32 bits)
|
|
vec16f32
|
v16f32
|
vector (16 x 32 bits)
|
|
vec8i8, vec8u8
|
v8i8
|
vector (8 x 32 bits)
|
它被扩展为32位矢量
|
vec8i16, vec8u16
|
v8i16
|
vector (8 x 32 bits)
|
它被扩展为32位矢量
|
vec8i32, vec8u32
|
v8i32
|
vector (8 x 32 bits)
|
它被扩展为32位矢量
|
vec8f32
|
v8f32
|
vector (16 x 32 bits)
|
它被视为一个16元素向量
|
说明格式
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型说明如下:
或
|
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型说明如下:
助记符
|
操作码
|
语义
|
操作
|
jmp
|
0
|
jump-无条件跳转到指定位置。
|
PC=rd or PC=PC+offset
|
jmpsr
|
1
|
跳转到子例程-无条件跳转到指定位置并将返回地址存储在RA寄存器中。
|
RA=PC+4 PC=rd or RA=PC+4 PC=PC+addr
|
jret
|
3
|
从子程序返回-从加载RA寄存器返回地址的子程序无条件返回。
|
PC=RA
|
beqz
|
5
|
有条件分行。如果条件寄存器的内容等于零,则分支为PC+偏移。
|
if(rcond==0) PC=PC+offset else PC=PC+4
|
bnez
|
6
|
条件分支,如果不等于零则分支-如果条件寄存器的内容不等于零,则分支到PC+偏移。
|
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)用于指定某个加载/存储内存操作是否进入scratchpad内存。例如,在从/向scratchpad存储器加载/存储的情况下,B1=1,否则B1=0。
典型的M型指令是加载和存储指令。在这两种情况下,源/目标地址都是按基寄存器地址+立即偏移量计算的,即rbase+偏移量。在负载的情况下,rd=[rbase+offset]。同样,在存储的情况下,[rbase+offset]=rs。所有M类型的指令都可以用于对主存储器和scratchpad存储器的存储操作。使用scratchpad内存操作的指令具有_scratchpad后缀。例如,load32_s8针对主存储器,而load32_s8_scratchpad指的是片上scratchpad的加载操作。
M型指令可分为标量指令和向量指令。标量M型指令是:
Mnemonic
|
Opcode
|
Meaning
|
Operation
|
load32_s8
|
0
|
将带有符号扩展的内存字节[7:0]加载到32位寄存器中
|
Rd = [Rbase + Offset]
|
load32_s16
|
1
|
将带有符号扩展的存储器半字[15:0]加载到32位寄存器中
|
Rd = [Rbase + Offset]
|
load32
|
2
|
将存储字加载到32位寄存器中
|
Rd = [Rbase + Offset]
|
load32_u8
|
4
|
将零扩展的内存字节[7:0]加载到32位寄存器中
|
Rd = [Rbase + Offset]
|
load32_u16
|
5
|
将零扩展的内存半字[15:0]加载到32位寄存器中
|
Rd = [Rbase + Offset]
|
load_v16i8
|
7
|
将带有符号扩展的16字节[127:0]加载到512位寄存器中
|
Rd = [Rbase + Offset]
|
load_v16i16
|
8
|
加载16个带符号扩展名的半字[255:0]
|
Rd = [Rbase + Offset]
|
load_v16i32
|
9
|
加载16个单词
|
Rd = [Rbase + Offset]
|
load_v16u8
|
11
|
加载16个字节[127:0],不带符号扩展名
|
Rd = [Rbase + Offset]
|
load_v16u16
|
12
|
加载16个半字[255:0],不带符号扩展
|
Rd = [Rbase + Offset]
|
load_v8u32
|
13
|
加载8个单词[255:0],不带符号扩展名
|
Rd = [Rbase + Offset]
|
loadg32
|
16
|
从不同的内存地址加载16个单词(仅适用于scratchpad)
|
Rd[i] = [Rbase[i]]
|
store32_8
|
32
|
在有效地址中存储1个字节
|
[Rbase + Offset] = Rs
|
store32_16
|
33
|
在有效地址中存储2个字节
|
[Rbase + Offset] = Rs
|
store32
|
34
|
在有效地址中存储1个单词
|
[Rbase + Offset] = Rs
|
store_v16i8
|
36
|
将向量寄存器中的16个字节(来自寄存器模式[487:480,…,39:32,7:0]的数据)存储到有效地址位置
|
[Rbase + Offset] = Rs
|
store_v16i16
|
37
|
将16个半字(从寄存器模式[495:480,…,47:32,15:0]中获取的数据)存储到有效地址位置
|
[Rbase + Offset] = Rs
|
store_v16i32
|
38
|
将向量寄存器中的16个字存储到有效地址位置
|
[Rbase + Offset] = Rs
|
stores32
|
42
|
分散存储-将16个单词存储到16个不同的地址中(仅用于scratchpad)
|
[Rbase[i]] = Rs[i]
|
Navigation menu
- 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,如下所示:
// 交叉产品
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文件,如下所示:
// 交叉产品内置函数
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):
// 交叉产品指令
def CROSSPROD_32 : FR_TwoOp_Unmasked_32 <
( outs VR512W : $dst ), // 输出
( ins VR512W :$src0 , VR512W : $src1 ), //输入
" crp $dst , $src0 , $src1 ", // corresponding assembly code
[( set v16i32 :$dst , ( int_npu_crossprodv16i32 v16i32 :$src0 , v16i32 : $src1 ))],
// 匹配模式
63, // ISA操作码(指令唯一)
Fmt_V , // 目标寄存器格式
Fmt_V , // src0寄存器格式
Fmt_V >; // src1寄存器格式
属性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 ,
// 到指令缓冲区
output thread_mask_t my_stop ;
// 从操作对象获取
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 ,
// 写回
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阶段,首先为自定义单元添加一个新的专用接口,如下所示:
// 来自套装模块
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 ),
//从操作对象获取
.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 ),
//写回
.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 ),
...
//来自新管道
.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, // 当前分块ID
parameter CORE_ID = 0, // 当前核心ID,未用于此类型的图块
parameter TILE_MEMORY_ID = 9, // 内存控制器块的ID
parameter THREAD_NUMB = 8,
// 支持的线程数,每个线程在LSU中// 都有一个单独的FIFO,来自不同线程的请求被并发处理-必须是2的幂
parameter ADDRESS_WIDTH = `ADDRESS_SIZE,
// 内存地址宽度-必须与系统地址宽度一致
parameter DATA_WIDTH = `CACHE_LINE_WIDTH,
// 数据总线宽度-必须与系统一致
parameter L1_WAY_NUMB = 4, // L1数据缓存中的路径数
parameter L1_SET_NUMB = 32, // L1数据集的数量
parameter SYNCH_SUPP = 1
// 分配barrier_core模块以支持同步
)
目录
[隐藏]
•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, // 发出的内存请求的有效信号
output logic
[31 : 0] req_out_id, // 发出请求的id,主要用于调试
output logic
[THREAD_IDX_W - 1 : 0] req_out_thread_id,
// 发出请求的线程id。在不同线程上运行的请求被并发分派到CC
output logic
[7 : 0] req_out_op, // 已执行操作
output logic
[ADDRESS_WIDTH - 1 : 0] req_out_address, // 发出的请求地址
output logic
[DATA_WIDTH - 1 : 0] req_out_data, // 数据输出
//来自异构LSU
input logic resp_in_valid, // 传入响应的有效信号
input logic [31 : 0] resp_in_id, // 传入响应的id,主要用于调试
input logic [THREAD_IDX_W - 1 : 0] resp_in_thread_id,
// 传入响应的线程id
input logic [7 : 0] resp_in_op, // 操作代码
input logic [DATA_WIDTH - 1 : 0] resp_in_cache_line, // 输入数据
input logic [BYTES_PERLINE - 1 : 0] resp_in_store_mask,
// 传入数据总线中请求字节位置的位掩码
input logic [ADDRESS_WIDTH - 1 : 0] resp_in_address,
// 传入响应地址
异构瓦片共享NPU瓦片的相同LSU和CC,因此LSU在内存接口上转发其背压信号,如下所示:
//来自异构加速器-背压信号
input logic [THREAD_NUMB - 1 : 0] lsu_het_almost_full,
// 线程位掩码,如果第i位为高,则第i个线程无法发出请求。
input logic [THREAD_NUMB - 1 : 0] lsu_het_no_load_store_pending,
// 线程位掩码,如果第i位为低,则第i个线程没有未决操作。
特别是,在为第i个线程发出内存请求之前,lsu_het_almost_full第i位必须为低。
内存接口提供性能计数器作为其接口的一部分:
// 来自异构LSU-性能计数器
input logic
resp_in_miss, // resp_in_address上的LSU缺失
input logic
resp_in_evict, // resp_in_address上的LSU驱逐(替换)
input logic
resp_in_flush, // LSU在resp_in_address上刷新
input logic
resp_in_dinv, // resp_in_address上的LSU数据缓存无效
这些信号表示L1数据缓存何时发生丢失、驱逐(或替换)、刷新和数据缓存无效。
异构图块中的LSU可以以两种不同的方式配置,即直写和回写:
output logic lsu_het_ctrl_cache_wt, // 启用直写缓存配置。
当lsu_het_ctrl_cache_wt为高时,lsu充当直写缓存,当其为低时,lsu实现回写机制。
最后,如果发出的请求中的地址未对齐,内存接口会提供错误信号:
//异构加速器-刷新和错误信号
input logic lsu_het_error_valid, // 来自lsu的错误
input register_t lsu_het_error_id, // 错误id-未对齐=380
input logic [THREAD_IDX_W - 1 : 0] lsu_het_error_thread_id,
// 错误涉及的线程
同步接口
同步接口将用户逻辑与拼贴块内分配的同步模块核侧(即barrier_core单元)连接起来。这样的接口允许用户逻辑在线程粒度上同步,当参数SYNCH_SUPP为高时,瓦片实现同步支持,分配一个barrier_core模块来处理同步事件核侧:
// 启动生成
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),
// 同步请求-核心接口
.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
// 结束生成
同步机制支持块间和块内屏障同步。当线程到达同步点时,它会通过同步接口向分布式同步主机发出请求。然后,线程被暂停(直到用户逻辑),直到其释放信号再次为高。
如果需要同步,自定义核可以实现以下接口:
/* 同步接口 */
//到屏障核
//输出逻辑信号,发送同步请求
breq_valid, // 冲击屏障信号,发送同步请求
output logic
[31 : 0] breq_op_id,
// 同步操作ID,主要用于调试
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个线程)
服务消息接口
服务消息接口将用户逻辑与服务网络连接起来。异构瓦片可以通过此接口向其他节点发送消息,这通常用于主机瓦片通信。如果需要通过消息进行通信,自定义核可以实现以下接口:
/* 服务消息接口 */
// 来自服务网络
input logic message_in_valid,
// 传入服务消息的有效位
input service_message_t message_in,
// 来自服务网络的传入消息
output logic n2c_mes_service_consumed,
// 服务消息已消耗
// 服务网络
output logic message_out_valid, // 传出服务消息的有效位
output
service_message_t message_out, //输出服务消息数据
input logic network_available, //服务网络可用位
output
tile_mask_t destination_valid // 一个热门目标位图
服务消息接口是一个标准的有效/可用接口,当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)。
// 发出同步请求
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。
// 同步所有虚拟核
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以类似的方式开始写入操作。
// 启动多个读取操作
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对每个线程的可用性(写入和读取)。
// 启动多个写入操作
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中,替换以下行:
// -----------------------------------------------------------------------
//
-- 平铺HT-模型芯篇片
//
-----------------------------------------------------------------------
het_core_example #(
.TILE_ID ( TILE_ID),
.THREAD_NUMB ( THREAD_NUMB )
) u_dummy_het_core (
.clk ( clk),
.reset ( reset),
/* 内存接口 */
.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),
...
/* 同步接口*/
.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),
...
);
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;
// 分配一些值:
for
(int i=0; i<16; i++) a[i]=i;
int
sum = 0;
// 计算和
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) //生成a<b的掩码
int
rm_old = __builtin_npu_read_mask_reg();
//保存寄存器
__builtin_npu_write_mask_reg(c); //写a<b的掩码寄存器
do_something();
c =
c^-1;
//生成a 大于或等于b的掩码寄存器
__builtin_npu_write_mask_reg(c);
//写a大于或等于b的掩码寄存器寄存器
do_somethingelse();
__builtin_npu_write_mask_reg(rm_old); //恢复以前的寄存器
}
线程平行度
在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(){
//执行线程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);//Synchronization Threads:0,1.
}
if(threadId==0){
D[threadId]=D[threadId]+D[threadId+1];
__builtin_npu_flush((int)(&D[threadId]));
}
return 0;
}
NaplesPU其他方面
冲洗说明
NaplesPU 指令集有一个刷新指令,需要该指令来避免数据卡在缓存中。如果主机需要输出数据,则必须使用此指令。否则,主机将从主存储器中,读取与缓存不一致的数据。刷新指令接收所涉及变量的地址输入,并刷新整个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关键字。scratch行内存中只能放置一个全局变量。例如,看看上面的代码:
#include <stdint.h>
__scratchpad int a;
int main (){
...
}
编译器将识别关键字和整数,并使用适当的加载/存储指令,将变量放置在scratch中。
屏障说明
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); //同步线程: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的主要内置功能:
内部函数名称(使用这些内部函数时,应该添加“__builtin-npu_”作为前缀) |
操作 |
对应指令 |
voidbarrier (int a, int b) |
屏障指令说明。a包含屏障ID,而b包含应同步的线程数-1 |
barrier |
voidflush (int a) |
将缓存行刷新到主内存。a包含缓存行的内存地址。需要进行显式整数转换。 |
flush |
intcreatemaskv16i32 (v16i32 a) |
将由等于0或-1的所有元素组成的向量a,转换为可以写入掩码寄存器的32位掩码值。当使用常见的C/C++关系运算符执行向量比较操作时,可用于计算位掩码 |
crt_maskv16 |
voidwrite_mask_reg(int a) |
在掩码寄存器中写入32位位掩码 |
move |
voidwrite_mask_regv16i32 (v16i32 a) |
在掩码寄存器内写入512位向量掩码 |
crt_maskv16+ move |
intread_mask_reg () |
从掩码寄存器中读取32位位掩码 |
move |
voidwrite_control_reg (int a, int b) |
将值写入掩码寄存器。整数a包含要访问的子寄存器的ID,而整数b包含数据 |
write_cr |
intread_control_reg (int a) |
从掩码寄存器中读取值。整数a包含要访问的子寄存器的ID |
read_cr |
NaplesPU向量内部函数
内部名称(使用这些内部函数时,应添加“__builtin_npu_”作为前缀) |
操作 |
相应指令 |
vec16i32makevectori32 (int a) |
创建一个由16个元素组成的向量,这些元素都等于a |
move_i32 |
vec16f32makevectorf32 (float a) |
创建一个由16个元素组成的向量,这些元素都等于a |
move_i32 |
vec16i32shufflei32 (vec16i32 a, vec16i32 b) or vec16i32 shufflef32 (vec16f32 a,vec16i32 b) |
VectorShuffle-允许将向量a的元素复制到输出向量的不同位置。向量b中的元素为目标寄存器中的每个对应位置,指定了源向量寄存器中元素的索引。(参阅shuffle指令的说明) |
shuffle_i32or shuffle_f32 |
intmask_cmp'w'32_'xyz' (vec16'w'32 a, vec16'w'32 b) |
如果满足以下条件,则返回一个整数位掩码,其中第i位等于1
(a[i]compb[i])是正确的。否则,第i位等于零。w=比较类型:i表示interger,f表示float。如果有符号,则x=s,否则u。yz=传统比较条件,即:gt、ge、lt、le |
cmp'xyz'_'w'32 |
intmask_cmp'w'32_eq (vec16'w'32 a, vec16'w'32 b) |
如果满足以下条件,则返回一个整数位掩码,其中第i位等于1
(a[i]compb[i])是正确的。否则,第i位等于零。w=比较类型:i表示interger,f表示float |
cmpeq_'w'32 |
intmask_cmp'w'32_ne (vec16'w'32 a, vec16'w'32 b) |
如果满足以下条件,则返回一个整数位掩码,其中第i位等于1
(a[i]compb[i])是正确的。否则,第i位等于零。w=比较类型:i表示interger,f表示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)
//生成a<b的掩码
int rm_old =
__builtin_npu_read_mask_reg(); //保存掩码寄存器
__builtin_npu_write_mask_reg(c); //写a<b的掩码寄存器
do_something();
c = c^-1; //生成a>b的掩码
__builtin_npu_write_mask_reg(c); //写a>b的掩码寄存器
do_somethingelse();
__builtin_npu_write_mask_reg(rm_old); //恢复原先的掩码
}
在这些情况下,重要的是要防止编译器对所涉及的向量进行优化,并将其标记为易失性变量。编译器倾向于重新排列操作,最好仔细检查编译器objdumb,以便掌握操作的顺序。如果volatile不能阻止重新排序,可以将代码嵌入到函数中,或使用-O0作为优化标志。
标准LLVM内部函数
固有名称 |
操作 |
相应指令 |
vec16X__builtin_convertvector(vec16Y a, vec16X) |
用于表示泛型向量类型转换操作。输入向量和输出向量类型,必须具有相同数量的元素。X可以等于i32/f32,而Y可以等于i32/f32 |
sext或
itof或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);
// 用于联合仿真 }
通常,刷新操作由线程执行,在这种情况下,第一个内核中的第一个线程,调用刷新内置程序,该程序将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);
// 用于联合模拟
}
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; // 线程数
uint32_t nL = 16; // 通道数
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将内存划分为四个不同的空间:
1)全局和常量空间:工作组中的所有工作项,都可以访问这些空间中的元素。
2)本地空间:仅对工作组内的工作项可见。
3)私人空间:仅对单个工作项可见。
目标平台提供DDR内存,即OpenCL术语中的设备内存。因此,变量在物理上映射到此内存上。编译器本身通过查看地址空间限定符来验证是否满足OpenCL约束。
每个NPU内核还配备了一个Scratchpad存储器,这是每个内核独有的,片上非相干存储器部分。此内存符合OpenCL本地内存功能。
最后,NPU内核中的每个硬件线程,都有一个私有堆栈。此内存部分对每个硬件线程(即OpenCL工作项)都是私有的,其他线程无法寻址。因此,每个堆栈都充当OpenCL私有内存。
编程模型匹配
OpenCL支持两种编程模型,数据并行和任务并行。数据并行模型要求OpenCL索引空间的每个点执行一个内核实例。由于每个点代表一个工作项,并且这些工作项映射到硬件线程上,因此正确地满足了数据并行要求。实现的模型是一个宽松的版本,不需要对数据进行严格的一对一映射。
任务并行编程模型要求内核实例在索引空间的任何点上独立执行。在这种情况下,每个工作项都不受限于执行其他工作项的相同内核实例。编译器前端定义了一组可用于此目的的内置程序。此外,每个NPU核由16个硬件通道构建,有助于实现锁步执行。因此,OpenCL支持被实现为允许使用向量类型。因此,使用以下数据类型支持向量执行:
1)charn和ucharn分别映射到vec16i8和vec16u8上,其中n=16。不支持n的其他值。
2)shortn、ushortn分别映射到vec16i16和vec16i32上,其中n=16。不支持n的其他值。
3)intn、uintn分别映射到vec16i32和vec16u32上,其中n=16。不支持n的其他值。
4)floatn,映射在vec16f32上,其中n=16。不支持n的其他值。
OpenCL运行时设计
OpenCL API是一组用于协调和管理设备的功能,这些功能为运行应用程序和监控执行提供支持。这些API还提供了一种检索设备相关信息的方法。
下图描述了OpenCL规范中定义的OpenCL运行时的UML类图。灰色填充框表示由于缺乏硬件支持而无法使用功能。
自定义OpenCL运行时依赖于两个主要抽象:
1)低级抽象,不完全依赖于硬件,提供设备主机通信支持。
2)根据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; // 线程数
uint32_t nL = 16; // 通道数
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
http://www.naplespu.com/doc/index.php?title=Detailed_studies