CN112130848A - 一种面向便笺式存储器的带宽感知循环分块优化技术 - Google Patents
一种面向便笺式存储器的带宽感知循环分块优化技术 Download PDFInfo
- Publication number
- CN112130848A CN112130848A CN202011013688.9A CN202011013688A CN112130848A CN 112130848 A CN112130848 A CN 112130848A CN 202011013688 A CN202011013688 A CN 202011013688A CN 112130848 A CN112130848 A CN 112130848A
- Authority
- CN
- China
- Prior art keywords
- data
- access
- memory
- dma
- mode
- Prior art date
- Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
- Granted
Links
- 230000000903 blocking effect Effects 0.000 title claims abstract description 24
- 238000005457 optimization Methods 0.000 title abstract description 13
- 238000005516 engineering process Methods 0.000 title abstract description 10
- 238000000034 method Methods 0.000 claims abstract description 47
- 238000013075 data extraction Methods 0.000 claims abstract description 32
- 230000006870 function Effects 0.000 claims description 46
- 230000001788 irregular Effects 0.000 claims description 24
- 238000011068 loading method Methods 0.000 claims description 23
- 238000012546 transfer Methods 0.000 claims description 21
- 230000004927 fusion Effects 0.000 claims description 15
- 238000000638 solvent extraction Methods 0.000 claims description 6
- 238000004590 computer program Methods 0.000 claims description 4
- 238000003860 storage Methods 0.000 claims description 4
- 230000002776 aggregation Effects 0.000 claims description 2
- 238000004220 aggregation Methods 0.000 claims description 2
- 238000005206 flow analysis Methods 0.000 claims description 2
- 238000004458 analytical method Methods 0.000 abstract description 32
- 230000006399 behavior Effects 0.000 abstract description 13
- 125000004122 cyclic group Chemical group 0.000 abstract description 10
- 238000003066 decision tree Methods 0.000 abstract description 9
- 230000003068 static effect Effects 0.000 abstract description 7
- 238000012360 testing method Methods 0.000 abstract description 6
- 230000000694 effects Effects 0.000 abstract description 4
- 229920006235 chlorinated polyethylene elastomer Polymers 0.000 description 16
- 238000012545 processing Methods 0.000 description 7
- 230000001133 acceleration Effects 0.000 description 6
- 230000005540 biological transmission Effects 0.000 description 6
- 239000008186 active pharmaceutical agent Substances 0.000 description 5
- 238000004364 calculation method Methods 0.000 description 5
- 230000008569 process Effects 0.000 description 5
- 239000002699 waste material Substances 0.000 description 4
- 230000008901 benefit Effects 0.000 description 3
- 238000009826 distribution Methods 0.000 description 3
- 230000007246 mechanism Effects 0.000 description 3
- 240000000233 Melia azedarach Species 0.000 description 2
- 238000000136 cloud-point extraction Methods 0.000 description 2
- 239000000284 extract Substances 0.000 description 2
- 238000005192 partition Methods 0.000 description 2
- 238000007781 pre-processing Methods 0.000 description 2
- 241001522296 Erithacus rubecula Species 0.000 description 1
- 238000013459 approach Methods 0.000 description 1
- 238000003491 array Methods 0.000 description 1
- 238000013473 artificial intelligence Methods 0.000 description 1
- 238000006243 chemical reaction Methods 0.000 description 1
- 238000004891 communication Methods 0.000 description 1
- 230000007547 defect Effects 0.000 description 1
- 230000001419 dependent effect Effects 0.000 description 1
- 238000013461 design Methods 0.000 description 1
- 238000005259 measurement Methods 0.000 description 1
- 238000012986 modification Methods 0.000 description 1
- 230000004048 modification Effects 0.000 description 1
- 230000003287 optical effect Effects 0.000 description 1
- 230000008447 perception Effects 0.000 description 1
- 230000000644 propagated effect Effects 0.000 description 1
- 238000011160 research Methods 0.000 description 1
Images
Classifications
-
- G—PHYSICS
- G06—COMPUTING; CALCULATING OR COUNTING
- G06F—ELECTRIC DIGITAL DATA PROCESSING
- G06F8/00—Arrangements for software engineering
- G06F8/40—Transformation of program code
- G06F8/41—Compilation
-
- G—PHYSICS
- G06—COMPUTING; CALCULATING OR COUNTING
- G06F—ELECTRIC DIGITAL DATA PROCESSING
- G06F8/00—Arrangements for software engineering
- G06F8/30—Creation or generation of source code
- G06F8/35—Creation or generation of source code model driven
Landscapes
- Engineering & Computer Science (AREA)
- General Engineering & Computer Science (AREA)
- Theoretical Computer Science (AREA)
- Software Systems (AREA)
- Physics & Mathematics (AREA)
- General Physics & Mathematics (AREA)
- Memory System Of A Hierarchy Structure (AREA)
- Devices For Executing Special Programs (AREA)
Abstract
本发明提出了一种面向便笺式存储器的带宽感知循环分块优化技术，该技术通过协调考虑带宽利用率和片上存储器容量来增强传统的循环分块优化方法。根据针对DMA的测试分析得到的带宽行为模型，创建决策树以针对不同类型的数据访问模式，从而选择最佳的数据提取操作。利用运行时循环分块框架来确定最佳分块大小，并在运行时生成分块代码。并且，本发明还通过利用参数引导的IPA来寻找不规则访问的静态分块机会，并使用冗余计算来节省SPM容量，从而增强循环分块的效果。
Description
技术领域
本发明涉及编译优化和代码生成，特别是涉及一种面向便笺式存储器的带宽感知的循环分块优化技术。
背景技术
近年来，便笺式存储器(Scratchpad Memory，SPM)已被广泛用于诸多领域的特定架构和加速设备的片上存储器。与高速缓存(cache)不同， SPM是一种对编程人员可见的片上存储，可以由编程人员显式地来管理和使用并且映射到与主存储器不同的地址空间。基于SPM的软件显式管理特点，编程人员可以自主地控制数据在SPM的存取操作，因此使用SPM，编程人员可以更好的预测数据读写的耗时，即拥有更好的时间可预测性，而且SPM的结构更简单，更节能。这些优点让SPM成为了cache的一种替代方案。例如超级计算机神威·太湖之光(Sunway TaihuLight)，人工智能处理器寒武纪Diannao，CELL处理器都使用SPM作为片上存储器。
基于SPM的新型体系结构通常会提供直接存储器访问机制(Direct MemoryAccess，DMA)和一些可以绕过SPM的全局数据加载指令。其中DMA可以利用有限的带宽将数据提取到SPM，而可以绕过SPM的全局数据加载指令则能够直接操作寄存器中单独元素。虽然全局加载指令可以直接操作寄存器中的元素，但是由于处理器中流水线的停顿导致了读写数据时存在较长的访问等待时间，明显浪费了可以利用的数据传输能力，即浪费了可用带宽。作为代表性的基于SPM的加速设备，神威·太湖之光的核心处理器SW26010提供了DMA和全局加载指令这两种机制，它们通常表现出明显不同的带宽利用率，范围从0.3GB/s到峰值28GB/s。
循环分块技术是编译器中的重要优化技术，它是一种根据循环嵌套中局部数据的空间分布和执行时间来设置分块大小，把循环进行切分的优化技术。
在基于cache的体系结构上，通常通过计算cache容量将常用数据保持在给定的cache里，根据cache容量来选择循环分块的分块大小。目前在基于cache的体系结构上，分块大小选择已经涌现了大量的技术发明。
然而，对于基于SPM的体系结构，目前涌现的技术只是将最常用的数据对象完全存储在SPM中，它们都将精力集中在数据的分配和数据获取上。因此，如果某些经常使用的数组量大于SPM的容量大小，则无法在这些技术中对它们进行分配。其中在数据分配的技术里，代表性的工作是Lian等人提出的数据切片技术，该切片技术使用整数线性规划(Integer Linear Programming，ILP)求解器为整个程序中的所有循环嵌套进行分析并选择最佳的切片方案。而在数据获取的技术中，Bhatotia等人提出了一种有界方法，通过使用包含所有必需元素的边界框列表来使用DMA，从而获取诸如A[B[i]]这种由于B[i]取值的不确定性导致数组A存在的所有不规则访问的数据。但是，这种方法仅考虑带宽策略，并且强制使用DMA操作来进行数据提取，这样不仅会提取大量无用的数据，同时也会浪费大量的SPM空间。到目前为止，还没有技术发明可以同时考虑带宽和容量来指导循环分块。
因此，对于支持DMA的SPM体系结构，仅着重分析容量和仅分析带宽的技术发明都无法有效地使用带宽和SPM。当只分析容量进行循环分块时，例如Lian等人提出的技术，将使用DMA操作将规则的数据提取到S PM，并使用全局加载指令把不规则的数据提取到寄存器中。该方案不仅S PM容量的利用率可能出现不高的情况，而且整体带宽利用率将非常低。另一方面，当仅分析带宽来进行循环分块时，例如Bhatotia等人提出的方法，将会获取大量无用的数据，虽然实现了高带宽利用率，但是将导致S PM容量利用率严重不足。
发明内容
为解决上述现有技术中存在的问题，提供一种面向便笺式存储器的循环分块方法，包括：
步骤1，根据主机代码和内核代码中访问的内存对象，获取所述内存对象的数据访问模式，所述数据访问模式包括规则访问和非规则访问，所述规则访问包括连续内存访问和跨步内存访问；
步骤2，根据所述数据访问模式，确定主存储器数据提取操作的方式，其中所述数据提取操作的方式包括连续DMA方式、跨步DMA方式、全局加载指令方式；
步骤3，根据所述数据提取操作方式和SPM容量确定分块大小；
步骤4，根据所述数据提取操作方式和分块大小，生成循环分块代码。
优选的，所述步骤1包括：
步骤11，根据主机代码和内核代码生成融合中间表示；
步骤12，根据所述融合中间表示，统计函数间参数对象的传递调用关系，生成调用图；
步骤13，遍历所述调用图，执行前向数据流分析，获取主机程序和内核程序的数据传递调用关系以及内存对象的数据访问模式。
优选的，所述步骤2包括：针对连续的内存访问，当循环嵌套中访问的数据量大于等于第一阈值，选择连续的DMA操作，当循环嵌套中访问的数据量小于第一阈值，则聚合循环迭代/工作项/工作组的DMA操作为跨步DMA操作。
优选的，所述步骤2包括：针对跨步的内存访问，当有效数据占比大于第二阈值，选择连续DMA操作。
优选的，所述步骤2包括：针对不规则的内存访问，当访问数据的个数大于第三阈值时，选择连续DMA操作，当访问数据的个数小于等于第三阈值时，选择全局数据加载操作。
优选的，所述第一阈值为1KB。
一种基于带宽感知循环分块的编译框架,包括：访问模式分析器、数据提取决策引擎、参数化分块模块、分块大小选择器、内存访问监控器、动态分块引擎，其中，
所述访问模式分析器用于分析工作组循环的数据访问模式，所述数据访问模式包括规则访问和非规则访问；
所述数据提取决策引擎用于根据所述数据访问模式来选择数据提取操作，其中所述数据提取操作包括连续DMA方式、跨步DMA方式、全局加载指令方式；
所述分块大小选择器用于确定分块大小；
所述参数化分块模块用于根据数据提取操作和所述分块大小生成分块代码；
所述内存访问监控器用于针对所述非规则访问，在运行时检测工作组循环并收集数据访问的地址信息，并将所述地址信息通过数据提取决策引擎来确定运行时的数据提取操作，
所述动态分块引擎用于根据所述运行时的数据提取操作生成相应的 DMA操作和分块大小。
一种基于带宽感知循环的编译方法，包括：
步骤1，根据主机代码和内核代码分别生成主机中间表示和内核中间表示；
步骤2，根据主机中间表示和内核中间表示生成融合中间表示；
步骤3，基于循环分块方法优化所述融合中间表示；
步骤4，根据优化后的融合中间表示生成可执行文件。
一种计算机可读存储介质，其上存储有计算机程序，其中，该程序被处理器执行时实现上述循环分块方法的步骤。
一种计算机设备，包括存储器和处理器，在所述存储器上存储有能够在处理器上运行的计算机程序，其特征在于，所述处理器执行所述程序时实现上述循环分块方法的步骤。
本发明的优点是：通过协调考虑带宽利用率和片上存储器容量来增强传统的循环分块优化方法。根据针对DMA的测试分析得到的带宽行为模型，创建决策树以针对不同类型的数据访问模式，从而选择最佳的数据提取操作。利用运行时循环分块框架来确定最佳分块大小，并在运行时生成分块代码。本发明还通过利用参数引导的IPA来寻找不规则访问的静态分块机会，并使用冗余计算来节省SPM容量，从而增强循环分块的效果。
附图说明
图1示出了神威·太湖之光的核心处理器SW26010的架构。
图2示出了根据本发明一个实施例的一段工作项的循环代码。
图3示出了现有技术中仅考虑容量的数据加载结果。
图4示出了现有技术中仅考虑带宽的数据加载结果。
图5示出了根据本发明一个实施例的数据加载结果。
图6示出了现有技术中跨步DMA操作的过程。
图7示出了根据本发明一个实施例的决策树模型。
图8示出了根据本发明一个实施例的第二阈值曲线。
图9示出了根据本发明一个实施例的第三阈值曲线。
图10示出了根据本发明一个实施例的基于带宽感知循环分块的编译框架。
图11示出了根据本发明一个实施例的编译流程。
图12示出了根据本发明一个实施例的主机代码和和内核代码。
图13示出了根据本发明一个实施例的调用图。
图14示出了根据本发明一个实施例的循环分块后生成的代码示例。
具体实施方式
下面结合附图和具体实施例对本发明加以说明。应当理解，此处所描述的具体实施例仅用以解释本发明，并不用于限定本发明。
本发明涉及程序的编译优化技术，根据本发明的一个实施例，该技术可用于超级计算机神威·太湖之光的编译器SWCL，通过SWCL对运行于神威·太湖之光的程序进行优化。
为了易于理解如何优化运行于神威·太湖之光的程序，先介绍神威·太湖之光的核心处理器SW26010。图1示出了SW26010的主要结构， SW26010包括4个处理器核心组(CoreGroup，缩写为CG)，每个CG包括：
1个管理处理单元(Management Processing Element，缩写为MPE)；
1个计算处理单元簇(Computing Processing Element Cluster,简称为 CPE簇)，每个CPE簇包括64个CPE，可并行执行计算任务，每个CPE 包含1个64k的LDM(即SPM)；
1个DDR3内存控制器(Memory Controller，缩写为MC)，每个MC 连接主存储。
由于4个CG相同，本实施例以其中一个CG的编程为例。
由于本发明的目标要涉及到优化CG提取数据时的数据带宽，此处介绍CG数据的加载方式。CG在执行计算任务时，其主存储器的数据采用两种方式加载到CPE，一种是通过DMA(Direct Memory Access，直接内存访问)机制将主存储器中的数据传输到SPM，其带宽可达到30.9G/s， DMA是现有技术。另一种是通过全局加载指令将主存储器的数据加载到CPE的寄存器，其带宽只能达到1.6G/s。
以上介绍了本发明一个实施例涉及的硬件平台，以下介绍本发明涉及的编程相关的软件平台。
由于一个CG包含64个CPE，可并行执行计算任务，构成了一个异构系统。异构系统通常由通用处理器与一个或多个加速设备互连组成，若干不同架构的处理器或处理器核协同工作，其编程相对困难。为了解决编程困难的问题，通常采用OpenCL编程框架。神威·太湖之光的编译器 SWCL也采用了OpenCL编程框架。为了更为清晰地理解本发明，以下对OpenCL进行简要介绍。
OpenCL用于解决异构系统中不同处理器架构导致的编程困难的问题。 OpenCL编程框架包括三个模型，分别是平台模型、执行模型和内存模型。其中，
OpenCL的平台模型把异构系统分为一个主机和多个加速设备，每个加速设备分为若干个计算单元(Computing Unit，缩写为CU)，每个CU 又分为若干处理单元(ProcessingElement，缩写为PE)。在SW26010中，本发明将MPE对应于OpenCL的主机，将CPE簇对应于OpenCL的加速设备。
OpenCL的执行模型分为主机执行的主机代码和加速设备执行的内核代码。当内核代码执行时，生成一个N维的索引空间(NDRange)，NDRange 中的每个线程称为工作项(work-item)，每个工作项运行在PE上，执行相同的内核程序，但操作的数据不同。NDRange可将多个工作项划分为若干个工作组(work-group)，每个工作组运行在CU上，其空间维度与NDRange 的维度相同，每个工作组包含数目相同的工作项，在SWCL中，将CU映射为CPE。可以看出，一个工作组的每个工作项在对应的CPE上串行执行，构成一个循环。
OpenCL的内存模型包括两部分，分别是主机内存和设备内存。其中，设备内存分为全局内存(也称常量内存)、局部内存和私有内存。所有工作项都可访问全局内存，局部内存仅限一个工作组内的工作项访问，私有内存仅限一个工作项访问。在SW26010中，主机内存和全局内存都映射到主存储器，局部内存映射到SPM，私有内存映射到局部变量。
以上介绍了本发明涉及的硬件平台和软件平台。下面，以一段工作项的代码为例，介绍与本发明的优化编译代码有关的背景知识。
图2示出了一段工作项的循环代码。在该代码中，有两个数组，分别是Elm和Nbr。其中，Elm数组的元素参与乘加计算，Nbr数组的元素用于提供Elm数组元素的索引，即Nbr数组的元素指定需要访问Elm数组的哪些元素。从代码中可以看出，由于lid是递增的，每次lid的循环迭代中，用于访问Nbr元素的索引2*lid和2*lid+1指向Nbr中连续的两个元素，因此，对Nbr的访问是对内存地址的连续访问。然而，Nbr元素的值在实际场景中通常是不连续的，因此，采用Nbr元素的值作为索引访问Elm数组中的元素时，构成了对内存地址的不规则访问。上述工作项代码运行于 CPE时，要从主存储器中加载数据到CPE。
现有技术中，有一种仅考虑容量的数据加载方法，即只加载需要使用的数据，不加载任何不使用的数据，其加载方式如图3所示，由于对Nbr 的访问是连续内存访问，因此通过DMA方式直接将Nbr的数据从主存储器加载到SPM，通过多次全局加载指令ld将需要的数据分别加载到CPE 的寄存器，每个ld指令只加载一个数据。根据前文介绍的两种加载方式的访问速度，DMA是30.9G/s，ld是1.6G/是，综合起来，此种加载方式只能达到4.6G/s的带宽。此种方式虽然没有浪费SPM空间，但是带宽很低。
针对上述代码，现有技术中还有一种仅考虑带宽的加载方法，即尽量通过DMA加载数据，而不考虑加载的无用数据会占用SPM的容量，其加载方式如图4所示，此方式对Nbr的数据加载与上一种相同，也是通过 DMA连续加载，但是对Elm的加载是通过多次DMA加载完成，每次DMA 加载的元素数量固定，比如图4中为3个元素。从图4中可以看出，其通过4次DMA加载完成对Elm的加载，但是将不需要使用的元素加载到SPM 中，即图4中带点的元素。在这种方式中，由于没有采用低带宽的ld指令，综合带宽虽然提高到28G/s，但是，加载了5个4字节数据，即浪费了20个字节的空间。
本发明的方法考虑到前两种方法的缺陷，采用了带宽感知的加载方法，其效果如图5所示，对于Nbr，依然采用DMA方式，对于Elm，通过本发明的优化技术(后面将详细介绍本发明的加载方法)，采用一次DMA加一次ld将数据载入到CPE。其中，一次DMA加载了6个有效元素，2个无效元素，一次ld加载了1个有效元素。此种方式的综合带宽达到16.8G/s，并且只加载了2个4字节无用数据，即只浪费了8个字节的空间。
通过以上三种方法效果的介绍，可以看出，本发明的方法综合考虑带宽和SPM容量，在二者之间进行折衷平衡，既能减少SPM空间的浪费，又能达到较高的带宽。
从以上分析中还可以看出本发明要解决的问题所在，即怎样通过分析代码，确定数据的加载方式以及如何优化DMA加载数据的大小。
为了对本发明的方法有清晰的了解，以下还要对数据的获取方式做进一步的介绍。
前文中提到，从CPE访问主存储器有两种方式，一种是DMA方式，另一种是全局加载指令方式。其中，DMA方式又可细分为连续DMA和跨步DMA。连续DMA是指一次DMA操作从主存储器中获取地址连续的一块数据。跨步DMA是指一次DMA操作可获取不连续的几块数据，这几块数据大小固定，间隔固定，如图6所示，DMA操作从主存储(Main Mem) 中加载几块数据到CPE的本地存储LDM(即SPM)中，其Main Mem中带斜线的阴影数据是要加载的数据，大小都是bsize，相邻两块数据之间的间隔为stepsize，加载到LDM后，总的DMA传输数据的大小为DMAsize。
根据本发明的一个实施例，发明人对连续DMA方式、跨步DMA方式以及全局加载指令方式的带宽行为进行了分析，获得了带宽行为模型。对于连续DMA，当bsize大于1k，数据是256字节对齐时，可获得接近峰值的带宽，为28G/s。对于跨步DMA，其带宽与bize和stepsize相关，当 bsize很小时，例如4字节，带宽仅能达到约0.3G/s，当bsize为1024字节时，带宽能达到约20G/s。对于全局加载指令方式，带宽为1.6G/s。发明人通过试验后认为，如果将同时进行的多次DMA操作合并为一次跨步操作，可以获得接近峰值的带宽，约为28G/s。
根据本发明的一个实施例，通过以下公式计算有效带宽，
bweff＝bw×reff,
其中，
bweff为有效带宽，bw为测量带宽，reff为有效数据占比，volumeuseful为有效数据大小，volumeall为全部传输数据。
基于上述的带宽行为模型，发明人认识到，在编译代码时，通过分析工作项的代码，针对代码中对内存的不同访问模式(即连续访问模式、跨步访问模式、不规则访问模式)，采用基于决策树模型的方法确定从主存储器中提取数据的方式(即连续DMA方式、跨步DMA方式、全局加载指令方式),然后将相应数据获取方式的指令插入到循环中。决策树模型如图7所示，在对代码的数据访问模式进行分析后，当数据访问模式是连续访问模式，如果嵌套循环的最内一层的数据访问量Volumeinnermost大于等于第一阈值时，在嵌套循环的最内一层中插入连续DMA操作，如果 Volumeinnermost小于第一阈值时,根据前述的带宽行为模型，则可将相邻循环迭代/工作项/工作组的多次连续DMA操作合并为一个跨步DMA操作,从而提高有效带宽。根据本发明的一个实施例，第一阈值为1KB。当数据访问模式是跨步访问模式，通过判断有效数据占比reff决定采用何种提取数据的方式，当reff大于等于第二阈值时，可采用连续DMA方式，此时虽然加载了一部分无用数据，但是提高了有效带宽；当reff小于第二阈值时，可采用跨步DMA方式，其中，第二阈值可通过测试得到的DMA性能信息来确定，根据本发明的一个实施例，基于bsize和总的数据访问量来确定第二阈值，如图8所示，例如，当bsize为256B，总的数据访问量为 5120B时，第二阈值为58％。当数据访问模式是不规则访问时，当访问数据的个数(Nacc)大于第三阈值时，使用连续DMA操作，从而更有效地利用带宽。而当访问数据的个数小于等于第三阈值时，则使用全局数据加载指令，其中第三阈值可通过测试得到的DMA性能信息来确定，根据本发明的一个实施例，Nacc基于每个数据的大小Nbyte和总的传输数据大小DMAsize确定，如图9所示，例如，当Nbyte为4字节，DMAsize为5120时，第三阈值为70。
在决策树模型中，对于内存的连续访问模式和跨步访问模式，由于访问模式明确，其数据获取方式在程序编译时即可确定；对于不规则访问，则要插入条件判断指令，以便在程序运行时，根据实际运行数据来确定采用何种获取方式。
以上介绍了决策树模型，其基于带宽行为模型，根据对代码中数据访问模式的分析结果，确定最佳的数据获取方式。
在确定了最佳的数据获取方式后，还要根据SPM的容量，通过以下公式来确定最佳的分块大小，
opt_ts＝max{ts|ws(ts)≤C}，
其中，opt_ts是最佳分块大小，ts是分块大小，ws是循环嵌套的空间大小，C是SPM的容量。
ws是ts的函数，可通过将嵌套循环中的所有DMA传输数据的大小 (DMAsize)累加得到。在给定的ts下计算ws时，遍历每个全局内存对象(__global)的访问行为(数据访问模式),对于连续访问或跨步访问，其DMA传输数据大小可以直接累加作为分块的大小。但是，对于不规则访问，例如A[B[i]]，由于B[i]的取值不确定导致A[B[i]]的数据在内存中排布不规则，无法简单通过一条连续或跨步DMA操作提取数据。B的DMA 传输大小可视为规则访问行为，可以通过静态分析直接确定B的分块大小。 A的DMA传输大小不能通过静态分析确定，可先为其假定一个分块大小，并在运行时根据具体数据大小，通过迭代缩小分块大小的方式来确定，即先根据静态设定的分块大小计算是否满足SPM容量限制，如果不满足，则缩小分块大小，再判断是否满足SPM容量限制，如此迭代计算，直到获得一个满足SPM容量限制条件的分块大小。
以上基于SW26010介绍了如何确定数据获取方式以及如何确定最佳分块大小，当用于其他支持SPM的平台时，需要先通过测试获取平台的带宽行为模型，然后确定数据获取方式和最佳分块大小。
通过以上分析可以看出，在三种数据访问模式(连续访问、跨步访问、不规则访问)中，对于不规则访问的处理方式最为复杂。前述的处理方式，是在程序运行时确定内存的访问地址，再通过DMA操作提高带宽利用率。发明人研究认为，如果通过在编译时静态分析代码，获取不规则访问的内存地址，就可以避免在运行时计算内存地址重新分块，从而提高运行效率。为了说明本发明如何在编译时静态分析内存访问地址，首先介绍本发明一个实施例的编译流程。如图11所示，本发明基于编译器前端Clang完成主机代码(Host Code)和内核代码(Kernel Code)的语法分析和语义分析后，生成一个明确的低级的中间表示(Intermediate Representation，缩写为IR)。IR具有易于生成且能被轻松翻译为目标机器语言的性质，因此，本发明将在生成的IR上进行融合、分析以及优化。如图11所示，图中的矩形框表示处理工具或处理动作，圆角矩形表示处理后的结果。首先，通过主机编译器和内核编译器分别编译主机代码和内核代码得到主机代码的中间表示(主机IR)和内核代码的中间表示(内核IR)。根据本发明的一个实施例，主机编译器和内核编译器基于Clang实现，Clang是一个C语言、C++、Objective-C语言的轻量级开源编译器。得到主机IR和内核IR后，再通过主机-内核融合阶段(H-K融合)生成融合中间表示(融合的IR)，该H-K融合阶段由主机编译器完成。然后通过本发明的带宽感知循环优化技术(BW Tiling)得到优化后的中间表示 (Optimized IR)。之后通过中间表示转C工具(IR2C Tool)得到优化后的 C/C++以及满足平台特性的代码(例如针对SW26010处理器的Athread库，该库是神威·太湖之光计算机系统的加速线程库)。最后经过C/C++或者平台指定编译器编译链接得到可执行文件(EXE)。
神威·太湖之光的编译器SWCL也是采用上述编译流程。由于采用了通用的OpenCL框架，在应用计算程序编写中，仅是针对OpenCL的平台模型、执行模型和内存模型进行编程，在主机代码和内核代码中调用了 OpenCL的API，因此，本发明的方法也适用于其他异构平台。
在明确了编译流程的基础上，以下介绍本发明如何对主机代码和内核代码进行分析，并在获取主机代码中的函数与内核代码中的函数调用关系和参数传递关系的基础上，确定内存的访问地址，以使得在静态编译时，针对不规则内存访问确定数据获取方式和分块大小。根据本发明的一个实施例，本发明使用的内存地址分析方法是参数引导的过程间分析，包括：步骤11，生成融合IR。主机编译器和内核编译器协同工作以生成融合IR，对应于
图11中的编译流程。例如，图12展示的就是主机代码和内核代码(Host andkernel source codes)。内核编译器通过标准的编译流程，将读取的内核代码经过预处理、词法分析和语法分析，得到内核IR。针对OpenCL的执行模型，内核编译器将在此内核IR上添加一层循环迭代，即工作项循环，再在此循环的基础上添加一层循环迭代，即工作组循环。以此实现OpenCL 内核代码的执行模型。之后内核编译器会把此内核IR通过进程间通信的方法(即共享内存)发送给主机编译器，完成IR的融合。主机编译器的职责是将接收到的内核IR合并到主机IR中，并生成融合的IR，对应于图11中的H-K融合阶段。主机编译器同样通过预处理、词法分析和语法分析生成主机IR，并等待接受内核处理器发送的内核IR，接受完毕后得到同时拥有主机和内核代码信息的融合IR。本发明遍历融合后的IR，利用OpenCL的语义，把clEnqueueNDRangeKernel的函数调用点(b12，b15， b19)作为内核函数在主机代码中的调用点，创建出g->k1，h->k1和h->k2 的调用关系，最后创建得到调用图，如图13所示。具体来说，在图12 中，左栏第2行的main函数对应于图13中的main节点；main函数中在 b2标记处调用了f函数，对应于main节点的子节点f；main函数在b4标记处调用了g函数，对应于main节点的子节点g；main函数在b6处调用了h函数，对应于main节点的h节点；在图12的左栏的第16行是g函数的函数体，在g函数中的b19标记处，通过OpenCL的API
clEnqueueNDRangeKernel调用了内核函数k1，k1的函数名出现在clEnqueueNDRangeKernel的参数中，k1的函数体定义在图12右栏第12 行。由此，在图13中，g节点的子节点为k1节点；h函数的函数体在图12 右栏第1行，在其函数体的b12标记处通过OpenCL的API clEnqueueNDRangeKernel调用了内核函数k2，在b15标记处调用了内核函数k1，这两处调用对应于图13中h节点指向了k1节点和k2节点。以上说明了如何通过图12的代码得到图13的调用图。
步骤12，统计函数间参数对象的传递调用关系。OpenCL编程模型把所有主机程序和内核程序间的数据流动抽象为缓存对象控制，即cl_mem 对象。为了针对融合了主机信息和内核信息的IR并创建调用图，本发明需要遍历融合IR里的所有函数，并设计了摘要(Summary)的数据结构来收集统计函数间参数对象的传递调用关系。摘要(Summary)的数据结构包含4部分：函数名；调用该函数时的参数对象集合(workset)；函数的参数对象的数据流信息(live，包括三种信息：MOD，该对象可能在此函数中被修改；REF，该对象可能在此函数中被调用；KILL，该对象的数据确定会在此函数中被修改，也就是被结束)；函数的参数对象的定义点 (location)。以图12中的h函数为例，首先，将所有cl_mem对象(即图 12中的A，B，C)和指针(图12中的p)进行汇总，由于A的声明在图 12左栏第1行，位于main函数之外，A是全局变量，可以被所有函数访问。并且，从图12右栏第1行的h函数的定义可以看出，其参数为y，由此可以得出，h的对象集合包含A和y。然后，分析汇总的变量，并且遍历所有调用图的基本块来寻找各个内核函数的定义或调用指令(图12 中的b15,b19)。同时，如果调用了clEnqueueWriteBuffer等OpenCL API(即图12中b3,b7,b16,b8)，那么将利用其主机内核间数据传递的语义进行辅助分析，通过此分析可以发现对象A在h函数中被b15调用的内核函数 k1修改，所以A的数据流信息为{1，0，0}，即MOD为1，REF为0， KILL为0；对象y在b12调用的内核函数k2中使用，所以y的数据流信息为{0，1，0}，即MOD为0，REF为1，KILL为0。之后将执行典型的迭代数据流分析来计算，并更新所有呼叫主机的基本块，即h函数中的参数对象只有A在b15处被修改，所以h的定义点为{b15}。
步骤13，上下文有关分析：在上一步骤统计计算传递函数之后，通过从上到下遍历调用图来执行前向数据流分析(前向数据流分析是现有技术，是一个通用的编译分析方法)，把上下文信息从调用节点传播到被调用函数，同样以图12中的h函数为例，也就是分析调用点main到调用点h。根据上一步骤得知main->h之间有cl_mem对象A和B，所以此过程的参数对象集合包括A和B。，根据上一步骤可知，调用h函数的参数对象集合包括A和形式参数y，且A在b16和b3可能被定义，所以A的定义点集合为{b16，b3}。而h函数的形式参数y在此过程中的实际参数为指针p，该指针可能指向A(由b5可分析得到)，也可能指向B(由b1可分析得到)。所以y的定义点集合包含A的定义点集合{b16，b3}以及B的定义点集合 {b19}。通过OpenCL的语义(即b19：clEnqueueNDRangeKernel是OpenCL API,表示调用内核函数)打通主机程序和内核程序的数据传递关系。
以上介绍了本发明使用的内存地址分析方法，即参数引导的过程间分析。
基于以上对本发明各项技术的分析，将本发明的方法总结如下：
步骤1，基于参数引导的过程间分析，在获取主机代码中的函数与内核代码中的函数调用关系和参数传递关系的基础上，确定内存的访问地址，以使得在静态编译时，针对不规则内存访问确定数据获取方式和分块大小。
步骤2，创建数据提取的决策树。根据针对DMA的测试分析得到的带宽行为模型，创建决策树，以选择最佳数据移动操作以插入到循环嵌套中，从而最大化有效地利用带宽。针对连续的内存访问(continuous)，为最里面的循环嵌套中的数据访问生成连续的DMA操作(Continuous DMA)，也就是把连续DMA的中间表示节点插入到循环嵌套中。如果在最里面的循环嵌套中访问的数据量(Volume innermost)小于1KB，那么相邻循环迭代/工作项/工作组的DMA操作可以聚合在一起，也就是收集连续DMA的单条数据传输量和连续DMA条数信息，形成以单条数据传输量乘以DMA条数为数据传输量的跨步DMA操作(Strided DMA)。
针对跨步的内存访问(strided)，当DMA数据块大小(bsize)除以bsize 与跨步大小(stepsize)的和，若商大于指定的阈值(该阈值的大小在前文中已经进行了说明，这里不再赘述)，则可以通过获取无用数据(即冗余计算)来使用连续的DMA操作，由于这样得到连续DMA操作能接近平台的峰值带宽，虽然引入了无用数据，但是获得了高带宽，因此有效地利用带宽，提高了数据提取效率。
针对不规则的内存访问，当访问数据的个数(Nacc)大于决策阈值时，本发明使用连续DMA操作，从而更有效地利用带宽的好处将超过获取无用数据的成本，进而获得更好的效果。而当访问数据的个数小于等于决策阈值时，我们将使用全局数据加载指令(Globalload/store)。
步骤3，选择最佳分块大小。确定最佳数据获取操作后，本发明根据 SPM容量选择最佳的分块大小。本发明通过分析统计所有数据的移动操作来计算循环嵌套的空间大小。因此，在给定的分块大小下计算所需空间大小时，本发明遍历每个全局内存对象(__global)的访问行为。对于连续访问或跨步访问，其DMA传输数据大小可以直接累加作为分块的大小。但是，对于不规则访问，例如A[B[i]]，由于B[i]的取值不确定导致A[B[i]] 的数据在内存中排布不规则，无法简单通过一条连续或跨步DMA操作提取数据。本发明将B的DMA传输大小视为规则访问行为，可以通过静态分析直接确定B的分块大小。而不能静态分析确定的A的DMA传输大小，则为其假定一个分块大小，并在运行时通过运行的具体数据大小，通过迭代缩小分块大小的方式，不停重新设置A的分块大小，直到获得一个满足 SPM容量限制条件的分块大小。
步骤4，生成循环分块代码。对于常规访问，通过添加中间表示节点的方式直接生成用于计算分块大小的代码，并按照OpenCL执行模型生成参数化的循环分块代码，如图14第4-7行所示。
对于不规则访问A[B[i]]，静态地在SPM上为A和B保留了相应的空间。在使用由假定分块大小得到的DMA操作进行数据移动时，会出现有多余数据将被提取到SPM，占用额外的SPM空间，导致空间不足的情况。在这种情况下，分块太大而无法在SPM上保留一个分块。此时，可以通过递归的方法，将DMA传输大小减小一半直至满足SPM容量空间的限制来选择新的分块大小，具体来说，是先将DMA传输大小减小一半，然后判断是否满足SPM容量空间的限制，如果不满足，则再将DMA传输大小减小一半，判断是否满足SPM容量空间的限制，如此重复，直到DMA 传输大小满足了SPM容量空间的限制。此后，再动态插入相应的DMA操作和计算代码，图14就是生成的循环分块代码示例，展示的是经由本发明分块后的工作项循环(Tiled work-item loop)。第1，2行声明了放入SPM 的变量数组，其中ts是指分块大小(tile size)。第3行是work-item循环。第4，5行是计算DMA操作的起始地址，这里会用到参数引导的过程间分析，也就是结合主机代码进行分析(Obtained by host codesanalysis)。第6， 7行是添加的DMA操作语句。第8行是一次分块取数的计算循环部分。第9行是数据的索引地址，同样需要结合主机代码才能分析出来(Inserted according to hostcodes)。第10行就是计算部分的代码。
基于以上方法，本发明提供一种基于带宽感知循环分块的编译框架，如图10所示，该框架包括：访问模式分析器(Access Pattern Analyzer)、数据提取决策引擎(FetchingDecision Engine)、参数化分块模块 (Parameterized Tiling Module)、分块大小选择器(Tile Size Selector)、内存访问监控器(Memory Access Monitor)、动态分块引擎(Dynamic Tiling Engine)。其中，访问模式分析器对每个工作组循环进行分析，并把访问模式分为规则访问(Regular Accesses)和非规则访问(Irregular Accesses)。对于规则访问，本发明将通过数据提取决策引擎来选择最佳数据提取操作(Fetching OPs)，在通过分块大小选择器确定最佳分块大小，然后通过参数化分块模块生成分块代码。对于不规则访问，内存访问监控器会检测工作组循环并在运行时收集数据访问的地址信息(AccessAddress)，并将这些信息通过数据提取决策引擎来确定运行中的最佳获取操作，最后动态分块引擎会生成相应的DMA操作和更合适的分块大小。
至此，本发明优化生成的代码，在平台执行时可以平衡带宽和SPM容量，充分利用了平台的SPM和DMA带宽资源，减少了程序的运行时间。
应该注意到并理解，在不脱离后附的权利要求所要求的本发明的精神和范围的情况下，能够对上述详细描述的本发明做出各种修改和改进。因此，要求保护的技术方案的范围不受所给出的任何特定示范教导的限制。
Claims (10)
1.一种面向便笺式存储器的循环分块方法，包括：
步骤1，根据主机代码和内核代码中访问的内存对象，获取所述内存对象的数据访问模式，所述数据访问模式包括规则访问和非规则访问，所述规则访问包括连续内存访问和跨步内存访问；
步骤2，根据所述数据访问模式，确定主存储器数据提取操作的方式，其中所述数据提取操作的方式包括连续DMA方式、跨步DMA方式、全局加载指令方式；
步骤3，根据所述数据提取操作方式和SPM容量确定分块大小；
步骤4，根据所述数据提取操作方式和分块大小，生成循环分块代码。
2.根据权利要求1所述的方法，所述步骤1包括：
步骤11，根据主机代码和内核代码生成融合中间表示；
步骤12，根据所述融合中间表示，统计函数间参数对象的传递调用关系，生成调用图；
步骤13，遍历所述调用图，执行前向数据流分析，获取主机程序和内核程序的数据传递调用关系以及内存对象的数据访问模式。
3.根据权利要求1所述的方法，所述步骤2包括：
针对连续的内存访问，当循环嵌套中访问的数据量大于等于第一阈值，选择连续的DMA操作，当循环嵌套中访问的数据量小于第一阈值，则聚合循环迭代/工作项/工作组的DMA操作为跨步DMA操作。
4.根据权利要求3所述的方法，所述步骤2包括：
针对跨步的内存访问，当有效数据占比大于第二阈值，选择连续DMA操作。
5.根据权利要求3所述的方法，所述步骤2包括：
针对不规则的内存访问，当访问数据的个数大于第三阈值时，选择连续DMA操作，当访问数据的个数小于等于第三阈值时，选择全局数据加载操作。
6.根据权利要求5所述的方法，所述第一阈值为1KB。
7.一种基于带宽感知循环分块的编译框架,包括：访问模式分析器、数据提取决策引擎、参数化分块模块、分块大小选择器、内存访问监控器、动态分块引擎，其中，
所述访问模式分析器用于分析工作组循环的数据访问模式，所述数据访问模式包括规则访问和非规则访问；
所述数据提取决策引擎用于根据所述数据访问模式来选择数据提取操作，其中所述数据提取操作包括连续DMA方式、跨步DMA方式、全局加载指令方式；
所述分块大小选择器用于确定分块大小；
所述参数化分块模块用于根据数据提取操作和所述分块大小生成分块代码；
所述内存访问监控器用于针对所述非规则访问，在运行时检测工作组循环并收集数据访问的地址信息，并将所述地址信息通过数据提取决策引擎来确定运行时的数据提取操作，
所述动态分块引擎用于根据所述运行时的数据提取操作生成相应的DMA操作和分块大小。
8.一种基于带宽感知循环的编译方法，包括：
步骤1，根据主机代码和内核代码分别生成主机中间表示和内核中间表示；
步骤2，根据主机中间表示和内核中间表示生成融合中间表示；
步骤3，基于权利要求1所述方法优化所述融合中间表示；
步骤4，根据优化后的融合中间表示生成可执行文件。
9.一种计算机可读存储介质，其上存储有计算机程序，其中，该程序被处理器执行时实现根据权利要求1至6中任一项所述方法的步骤。
10.一种计算机设备，包括存储器和处理器，在所述存储器上存储有能够在处理器上运行的计算机程序，其特征在于，所述处理器执行所述程序时实现权利要求1至6中任一项所述的方法的步骤。
Priority Applications (1)
| Application Number | Priority Date | Filing Date | Title |
|---|---|---|---|
| CN202011013688.9A CN112130848B (zh) | 2020-09-24 | 2020-09-24 | 一种面向便笺式存储器的带宽感知循环分块优化方法、编译系统、设备及存储介质 |
Applications Claiming Priority (1)
| Application Number | Priority Date | Filing Date | Title |
|---|---|---|---|
| CN202011013688.9A CN112130848B (zh) | 2020-09-24 | 2020-09-24 | 一种面向便笺式存储器的带宽感知循环分块优化方法、编译系统、设备及存储介质 |
Publications (2)
| Publication Number | Publication Date |
|---|---|
| CN112130848A true CN112130848A (zh) | 2020-12-25 |
| CN112130848B CN112130848B (zh) | 2022-06-14 |
Family
ID=73839587
Family Applications (1)
| Application Number | Title | Priority Date | Filing Date |
|---|---|---|---|
| CN202011013688.9A Active CN112130848B (zh) | 2020-09-24 | 2020-09-24 | 一种面向便笺式存储器的带宽感知循环分块优化方法、编译系统、设备及存储介质 |
Country Status (1)
| Country | Link |
|---|---|
| CN (1) | CN112130848B (zh) |
Cited By (2)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| CN112965810A (zh) * | 2021-01-27 | 2021-06-15 | 合肥大多数信息科技有限公司 | 一种基于共享网络通道的多内核浏览器数据整合方法 |
| CN117312330A (zh) * | 2023-11-29 | 2023-12-29 | 中国人民解放军国防科技大学 | 基于便签式存储的向量数据聚集方法、装置及计算机设备 |
Citations (6)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| US20100262973A1 (en) * | 2009-04-09 | 2010-10-14 | Rolf Ernst | Method For Operating a Multiprocessor Computer System |
| CN101937343A (zh) * | 2010-09-17 | 2011-01-05 | 上海交通大学 | 异构多核虚拟执行环境的后端翻译框架实现的方法 |
| CN102929580A (zh) * | 2012-11-06 | 2013-02-13 | 无锡江南计算技术研究所 | 数组多引用访问的分块方法和装置 |
| CN103226487A (zh) * | 2013-04-25 | 2013-07-31 | 中国人民解放军信息工程大学 | 面向异构众核多级存储结构的数据分布与局部性优化方法 |
| CN105138335A (zh) * | 2015-08-28 | 2015-12-09 | 牟永敏 | 一种基于控制流图的函数调用路径提取方法及装置 |
| CN110187988A (zh) * | 2019-06-06 | 2019-08-30 | 中国科学技术大学 | 适用于虚函数和函数指针的静态函数调用图构建方法 |
-
2020
- 2020-09-24 CN CN202011013688.9A patent/CN112130848B/zh active Active
Patent Citations (6)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| US20100262973A1 (en) * | 2009-04-09 | 2010-10-14 | Rolf Ernst | Method For Operating a Multiprocessor Computer System |
| CN101937343A (zh) * | 2010-09-17 | 2011-01-05 | 上海交通大学 | 异构多核虚拟执行环境的后端翻译框架实现的方法 |
| CN102929580A (zh) * | 2012-11-06 | 2013-02-13 | 无锡江南计算技术研究所 | 数组多引用访问的分块方法和装置 |
| CN103226487A (zh) * | 2013-04-25 | 2013-07-31 | 中国人民解放军信息工程大学 | 面向异构众核多级存储结构的数据分布与局部性优化方法 |
| CN105138335A (zh) * | 2015-08-28 | 2015-12-09 | 牟永敏 | 一种基于控制流图的函数调用路径提取方法及装置 |
| CN110187988A (zh) * | 2019-06-06 | 2019-08-30 | 中国科学技术大学 | 适用于虚函数和函数指针的静态函数调用图构建方法 |
Cited By (4)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| CN112965810A (zh) * | 2021-01-27 | 2021-06-15 | 合肥大多数信息科技有限公司 | 一种基于共享网络通道的多内核浏览器数据整合方法 |
| CN112965810B (zh) * | 2021-01-27 | 2022-06-24 | 合肥大多数信息科技有限公司 | 一种基于共享网络通道的多内核浏览器数据整合方法 |
| CN117312330A (zh) * | 2023-11-29 | 2023-12-29 | 中国人民解放军国防科技大学 | 基于便签式存储的向量数据聚集方法、装置及计算机设备 |
| CN117312330B (zh) * | 2023-11-29 | 2024-02-09 | 中国人民解放军国防科技大学 | 基于便签式存储的向量数据聚集方法、装置及计算机设备 |
Also Published As
| Publication number | Publication date |
|---|---|
| CN112130848B (zh) | 2022-06-14 |
Similar Documents
| Publication | Publication Date | Title |
|---|---|---|
| US9798528B2 (en) | Software solution for cooperative memory-side and processor-side data prefetching | |
| Rotem et al. | Glow: Graph lowering compiler techniques for neural networks | |
| Talati et al. | Prodigy: Improving the memory latency of data-indirect irregular workloads using hardware-software co-design | |
| Wahib et al. | Scalable kernel fusion for memory-bound GPU applications | |
| US8180964B1 (en) | Optimization of cache configuration for application design | |
| KR101559090B1 (ko) | 이종 코어를 위한 자동 커널 마이그레이션 | |
| KR101573586B1 (ko) | 논-리프 코드의 컴파일러 기반 벡터화를 위한 시스템들 및 방법들 | |
| US8949532B1 (en) | Automatic generation of cache-optimized code | |
| Prasad et al. | Automatic compilation of MATLAB programs for synergistic execution on heterogeneous processors | |
| White et al. | Timing analysis for data and wrap-around fill caches | |
| Piccoli et al. | Compiler support for selective page migration in NUMA architectures | |
| Soliman et al. | WCET-driven dynamic data scratchpad management with compiler-directed prefetching | |
| CN112130848B (zh) | 一种面向便笺式存储器的带宽感知循环分块优化方法、编译系统、设备及存储介质 | |
| Chen et al. | Locality analysis through static parallel sampling | |
| Stawinoga et al. | Predictable thread coarsening | |
| Das et al. | Index array flattening through program transformation | |
| Neves et al. | Compiler-assisted data streaming for regular code structures | |
| Wu et al. | Bandwidth-aware loop tiling for dma-supported scratchpad memory | |
| Madsen et al. | Towards a streaming model for nested data parallelism | |
| Calvert | Parallelisation of java for graphics processors | |
| Chakraborty et al. | Integrating software caches with scratch pad memory | |
| Soliman | Automated compilation framework for scratchpad-based real-time systems | |
| Li et al. | FreshBreeze: A data flow approach for meeting DDDAS challenges | |
| Malhotra et al. | Library-based prefetching for pointer-intensive applications | |
| Yu et al. | Hierarchical Read/Write Analysis for Pointer-Based OpenCL Programs on RRAM |
Legal Events
| Date | Code | Title | Description |
|---|---|---|---|
| PB01 | Publication | ||
| PB01 | Publication | ||
| SE01 | Entry into force of request for substantive examination | ||
| SE01 | Entry into force of request for substantive examination | ||
| CB03 | Change of inventor or designer information |
Inventor after: Wu Mingchuan Inventor after: Liu Ying Inventor after: Cui Huimin Inventor after: Wei Qingfu Inventor after: Li Quanfeng Inventor after: Li Limin Inventor after: Lv Fang Inventor after: Feng Xiaobing Inventor before: Wu Mingchuan Inventor before: Liu Ying Inventor before: Cui Huimin Inventor before: Wei Qingfu Inventor before: Li Quanfeng Inventor before: Li Limin Inventor before: Lv Fang Inventor before: Feng Xiaobing |
|
| CB03 | Change of inventor or designer information | ||
| GR01 | Patent grant | ||
| GR01 | Patent grant | ||
| TR01 | Transfer of patent right |
Effective date of registration: 20231221 Address after: Room 1305, 13th Floor, No.1 Zhongguancun Street, Haidian District, Beijing, 100086 Patentee after: Zhongke Jiahe (Beijing) Technology Co.,Ltd. Address before: 100190 No. 6 South Road, Zhongguancun Academy of Sciences, Beijing, Haidian District Patentee before: Institute of Computing Technology, Chinese Academy of Sciences |
|
| TR01 | Transfer of patent right |