CN107710150A - 从包含层次子例程信息的中间代码产生目标代码 - Google Patents

从包含层次子例程信息的中间代码产生目标代码 Download PDF

Info

Publication number
CN107710150A
CN107710150A CN201680031758.9A CN201680031758A CN107710150A CN 107710150 A CN107710150 A CN 107710150A CN 201680031758 A CN201680031758 A CN 201680031758A CN 107710150 A CN107710150 A CN 107710150A
Authority
CN
China
Prior art keywords
subroutine
code
level
processor
higher levels
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
Application number
CN201680031758.9A
Other languages
English (en)
Other versions
CN107710150B (zh
Inventor
李·豪斯
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Qualcomm Inc
Original Assignee
Qualcomm Inc
Priority date (The priority date 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 date listed.)
Filing date
Publication date
Application filed by Qualcomm Inc filed Critical Qualcomm Inc
Publication of CN107710150A publication Critical patent/CN107710150A/zh
Application granted granted Critical
Publication of CN107710150B publication Critical patent/CN107710150B/zh
Expired - Fee Related legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Classifications

    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00Arrangements for software engineering
    • G06F8/40Transformation of program code
    • G06F8/41Compilation
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00Arrangements for software engineering
    • G06F8/40Transformation of program code
    • G06F8/41Compilation
    • G06F8/44Encoding
    • G06F8/443Optimisation
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00Arrangements for software engineering
    • G06F8/40Transformation of program code
    • G06F8/41Compilation
    • G06F8/44Encoding
    • G06F8/447Target code generation
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00Arrangements for software engineering
    • G06F8/40Transformation of program code
    • G06F8/41Compilation
    • G06F8/45Exploiting coarse grain parallelism in compilation, i.e. parallelism between groups of instructions
    • G06F8/451Code distribution
    • G06F8/452Loops
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00Arrangements for software engineering
    • G06F8/40Transformation of program code
    • G06F8/41Compilation
    • G06F8/45Exploiting coarse grain parallelism in compilation, i.e. parallelism between groups of instructions
    • G06F8/456Parallelism detection

Landscapes

  • Engineering & Computer Science (AREA)
  • General Engineering & Computer Science (AREA)
  • Theoretical Computer Science (AREA)
  • Software Systems (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • Devices For Executing Special Programs (AREA)

Abstract

本发明描述使装置接收从编译应用程序的源代码产生的中间代码的实例。所述中间代码包含从所述编译产生的信息,所述信息识别较高层级子例程中的较低层级子例程的层次结构,且所述较低层级子例程在所述应用程序的所述源代码中定义为比识别所述较低层级子例程的所述较高层级子例程更频繁地执行。所述装置经配置以基于识别所述较高层级子例程中的所述较低层级子例程的所述信息,编译所述中间代码以产生目标代码,并且存储所述目标代码。

Description

从包含层次子例程信息的中间代码产生目标代码
本申请案主张2015年6月15日申请的标题为“中间代码中的层次子例程信息(HIERARCHICAL SUB-ROUTINE INFORMATION IN INTERMEDIATE CODE)”的美国专利申请案第62/175,646号的优先权,所述专利申请案整个内容由此以引用的方式并入。
技术领域
本发明涉及编译器,且更确切地说,涉及产生用于在处理单元上执行的便携式中间代码的编译器。
背景技术
计算装置上的处理器执行从以高级编程语言写入的编译代码(即,指令)产生的目标代码。为了产生目标代码,编译器首先将以高级编程语言写入的代码转译成中间代码(即,中间语言的代码)。同一或某一其它系统上的另一编译器(或可为同一编译器)将中间语言的代码转译成接着由所述处理器或系统中的某一其它处理器执行的机器代码。
发明内容
本发明描述将函数调用包含到识别以高级语言写入的应用程序中(例如,应用程序的源代码中)的子例程之间的层次关系的中间代码中的技术。例如,编译以高级语言写入的应用程序的编译器可将函数调用包含到中间代码中。子例程之间的层次关系指示第一子例程取决于第二子例程的完成,且第二子例程比第一子例程更频繁地执行。在此情况下,第一子例程被视为与第二子例程相比处于较高层级处。
另一编译器(或可能地同一编译器)进一步将中间代码编译成目标代码。为了将中间代码编译成目标代码,编译器(例如,此其它编译器或同一编译器)使用中间代码中的函数调用产生利用将执行目标代码的处理器的处理能力的目标代码。以此方式,中间代码对于在不同处理器上执行来说是更便携的,这是因为中间代码包含分离出较高层级子例程与较低层级子例程的函数调用,从而允许编译器确定如何布置用于执行的子例程以最佳使用将在上面执行目标代码的处理器的处理能力。
在一个实例中,本发明描述一种编译方法,所述方法包括:用处理器接收从编译应用程序的源代码产生的中间代码,其中所述中间代码包含从所述编译产生的信息,所述信息识别较高层级子例程中的较低层级子例程的层次结构,且其中所述较低层级子例程在所述应用程序的所述源代码中定义为比识别所述较低层级子例程的所述较高层级子例程更频繁地执行;用所述处理器基于识别所述较高层级子例程中的所述较低层级子例程的所述信息,编译所述中间代码以产生目标代码;以及存储所述目标代码。
在一个实例中,本发明描述一种用于编译的装置,所述装置包括:存储器单元;以及集成电路,其包括处理器,所述处理器经配置以:接收从编译应用程序的源代码产生的中间代码,其中所述中间代码包含从所述编译产生的信息,所述信息识别较高层级子例程中的较低层级子例程的层次结构,且其中所述较低层级子例程在所述应用程序的所述源代码中定义为比识别所述较低层级子例程的所述较高层级子例程更频繁地执行;基于识别所述较高层级子例程中的所述较低层级子例程的所述信息,编译所述中间代码以产生目标代码;以及将所述目标代码存储于所述存储器单元中。
在一个实例中,本发明一种用于编译的装置,所述装置包括:用于接收从编译应用程序的源代码产生的中间代码的装置,其中所述中间代码包含从所述编译产生的信息,所述信息识别较高层级子例程中的较低层级子例程的层次结构,且其中所述较低层级子例程在所述应用程序的所述源代码中定义为比识别所述较低层级子例程的所述较高层级子例程更频繁地执行;用于基于识别所述较高层级子例程中的所述较低层级子例程的所述信息,编译所述中间代码以产生目标代码的装置;以及用于存储所述目标代码的装置。
在一个实例中,本发明描述一种包括指令的非暂时性计算机可读存储媒体,所述指令在被执行时致使一或多个处理器:接收从编译应用程序的源代码产生的中间代码,其中所述中间代码包含从所述编译产生的信息,所述信息识别较高层级子例程中的较低层级子例程的层次结构,且其中所述较低层级子例程在所述应用程序的所述源代码中定义为比识别所述较低层级子例程的所述较高层级子例程更频繁地执行;基于识别所述较高层级子例程中的所述较低层级子例程的所述信息,编译所述中间代码以产生目标代码;以及存储所述目标代码。
在附图和以下描述中阐明一或多个实例的细节。其它特征、目标和优点将从所述描述和图式以及权利要求书显而易见。
附图说明
图1是说明用于实施本发明的各方面的实例系统的框图。
图2是说明可经配置以实施本发明的各方面的装置的框图。
图3是更详细地说明图2的实例装置的框图。
图4是说明用于实施本发明的各方面的实例的流程图。
图5是说明用于实施本发明的各方面的另一实例的流程图。
具体实施方式
处理器(例如,图形处理单元(GPU)或中央处理单元(CPU))执行程序的指令,所述指令致使处理器执行特定功能。应用程序开发者使用特定标准的语法开发程序。例如,应用程序开发者可使用OpenCL规范或异机种系统架构(HSA)规范(作为两个非限制性实例)的应用程序编程接口(API)定义的标准。为易于说明,本发明中描述的技术仅关于OpenCL规范进行描述以帮助理解。然而,本发明中描述的技术不限于根据任何特定标准开发的应用程序。
OpenCL规范允许应用程序开发者以较高级编程语言写入程序(例如,写入应用程序的源代码),以指示处理器执行特定功能。所述程序接着可在例如GPU、CPU、现场可编程门阵列(FPGA)、数字信号处理器(DSP)以及类似者等多种不同处理器类型上执行。CPU有时称为主机,且这些其它实例硬件单元有时更一般性称为装置。在本发明中,术语装置可用以指包含CPU、GPU、FPGA和/或DSP的单元或指CPU、GPU、FPGA和/或DSP。因此,术语“装置”不应视为限于例如GPU、FPGA和/或DSP等实例,且包含CPU以及容纳这些实例组件的单元。
术语“程序”是指以较高级编程语言写入的指示处理器执行特定功能的任何指令集。例如,术语“程序”用以包含根据OpenCL规范开发的程序,或使用任何其它规范开发的程序,以及以例如C++等任何较高级编程语言写入的程序。较高级编程语言的另一实例是SYCL,其为在OpenCL顶部上构建的高级编程模型。
对于执行程序的指令的处理器,编译器将程序的指令转译成中间语言(例如,汇编语言)的代码。中间语言的代码也称为中间代码。此中间代码对于不同处理器可为相当便携的。举例来说,当用户致使装置(例如,移动装置、手提式计算机、台式计算机等)下载用于执行的程序时,所述装置可下载中间语言的中间代码。
下载中间代码的装置上的编译器可进一步编译中间代码以产生将在处理器(例如,CPU、GPU、FPGA和/或DSP)上执行的目标代码。所述处理器接着可执行所述目标代码。
产生中间语言的编译器称为前端编译器或高层级编译器,且产生目标代码的编译器称为后端编译器或低层级编译器。前端编译器和后端编译器可在不同装置上执行。举例来说,前端编译器可在应用程序开发者的装置(或容纳中间代码的服务器)上执行以产生中间语言的中间代码,且后端编译器可在用户的装置上执行以产生目标代码。在某一实例中,后端编译器可以是在程序执行期间产生目标代码而非预存储目标代码的实时后端编译器。
应理解,前端编译器和后端编译器不需要一定在不同装置上执行。例如,在以上实例中,用户装置下载中间代码。在一些其它实例中,用户装置可下载程序的指令,而非中间代码。在此类实例中,在用户装置上执行的编译器可经由编译器的前端模块产生中间代码,并且经由编译器的后端模块产生目标代码。在本发明中,不同编译器描述为产生中间代码以及产生目标代码(例如,在不同装置上执行的前端编译器和后端编译器),但所述技术不应被视为受到限制。
如上文所描述,OpenCL定义用于开发跨越平台(例如,在不同类型的处理器上)执行的程序的结构。在一些情况下,开发者使用OpenCL产生包含并不适用于所有处理器类型的不必要操作的程序。举例来说,在单程序多数据(SPMD)并行性中,处理器执行同一程序的多个个例,其中每个程序可对不同数据进行操作。OpenCL和HSA规范用于将SPMD并行性映射到单指令多数据(SIMD)执行单元。然而,在此类映射中,可丢失关于并行结构的信息。
为了解决这个问题,OpenCL规范允许“屏障”指令。屏障指令致使程序的个例暂停执行,使得可同步程序的所有个例(即,允许程序的较慢个例赶上)。所述同步允许程序的个例与另一个例通信。
例如,对于OpenCL工作组,为了在个例之间进行通信,开发者可使用屏障。这些屏障充当独立实体之间的通信操作。然而,在一些情况下,大多数这些实体可能映射到相同线程并且通过正在处理的所述方式受到更好服务。换句话说,胜于具有不同并行个例,串行执行可能是更好的。在一些情况下,OpenCL的工具链可能无法推断所述实体是映射到不同线程还是相同线程。
为了解决这个问题,在OpenCL顶部上构建SYCL高级编程模型,并且启用结构化形式的并行性。对于SYCL,所述代码维持更低信息使得其性能是更便携的(即,可跨越多个不同类型装置使用)。SYCL达成此类结果的一种方式是通过考虑到层次结构。在层次结构中,所述程序包含多个子例程(也称为程序块或模块)。所述子例程中的一者可取决于另一子例程,且子例程的执行频率可为不同的(例如,其它子例程执行每一次,相依性子例程可执行N次)。相依性子例程可被视为处于层次结构中的较低层级,且与其相依的子例程视为层次结构中的较高层级。
此类层次结构的基础实例是嵌套式for环路。在此实例中,较高层级for回路执行每一次,较低层级for回路执行N次。
在一些较早情况中,屏障用以允许不同个例的环赶上彼此。对于SYCL,代替试图达成屏障任一侧的环路分裂,层次结构明确地分成两个环。
然而,此类层次结构化程序应在不同处理器类型上执行的方式可为不同的。举例来说,在提供大规模并行性的GPU上,使程序的多个个例与略不同的输入和较少命令并行地执行可为有利的,但对于不提供大规模并行性的CPU,使程序的较少个例与较多命令一起执行可为有利的。
可能出现的一个问题是,编译程序以产生中间语言的中间代码的编译器(例如,前端编译器)可能不知道将在上面执行所述程序的处理器类型(例如,CPU或GPU)。因此,前端编译器选择一种用以产生中间语言的中间代码的方式,但此方式可能较适合一些处理器类型且并不很适合其它处理器类型。
例如,中间语言的一个实例是SPIRV,且使用SYCL写入的程序编译成SPIRV中的代码。在一些情况下,使用层次结构开发的程序(例如,使用SYCL写入的程序)在编译成中间语言(例如,SPIRV)时丢失层次结构信息。高级编译器(例如,前端编译器)必须选择一个映射(例如,构建特定于处理器类型的中间代码),而非使所述便携式映射直达后端编译器。
允许后端编译器确定应如何进一步编译层次结构可为有利的,这是因为后端编译器可经配置有指示将在上面执行程序的处理器的类型的信息。在一些技术中,前端编译器丢失中间代码中的层次结构信息,这意味着后端编译器无法编译中间代码以针对在上面执行目标代码的处理器类型来专门化所述目标代码。
根据本发明中描述的技术,前端编译器可经配置以添加函数类型的层次并且将指令调度给中间语言(IL)(例如,作为实例,IL SPIRV)的中间代码。后端编译器可使用层次结构的此信息编译中间代码并且产生目标代码,所述目标代码利用在上面执行程序的处理器的处理能力。
举例来说,为了识别层次结构,前端编译器可包含子例程(例如,程序块)内识别较低层级相依性子例程的函数。较低层级子例程可比与其相依的子例程(例如,较高层级子例程)更频繁地执行。一般来说,前端编译器可将函数标记为表示执行图表中的一些执行单元。执行图表可为识别程序的子例程彼此相关的方式的概念图。用以识别子例程的函数的实例包含当前SPMD工作项目、整个子组、整个工作组,并且可能甚至包含全局操作的调度入口点。
在一些实例中,用以下此类标度标记入口点:内核(Kernel)、工作组、子组、工作项目(作为最高层次层级到最低层次层级的实例)。这些都可为所有SPIR-V和HSAIL内核中目前的默认项。
对于映射到SYCL中并行性的更细粒度(例如,较低层次层级)的代码块(例如,子例程),标记为可调用对象/λ函数的主体。这可为前端编译器包含较高层级子例程中用以调用较低层级子例程的函数的一种实例方式。本质上,对于标记为可调用对象/λ函数,更细粒度的并行性使得其成为可内联的分离函数。在本发明中,内联意指可在程序中添加信息所针对的函数。此所得的中间语言(IL)版本的中间代码将所有此类块概述(例如,分离出)为分离函数,使得函数边界标记粒度转变。添加如下的额外调度函数以转变为更细粒度:OpSGFunctionCall、OpWIFunctionCall等。
在一些实例中,对于前端编译器来说可有可能还添加并行环路构造以发起一组子级子组(例如,较低层级子例程)。这可为更具动态的行为,而一些架构具有相当静态的层次。在静态层次中,所述行为是从工作组粒度到子组粒度的移动进入工作组中每个子组的子组函数。类似地,子组到工作项目的过渡会调用子组中的每个工作项目的工作项目函数。此实例实际上是SIMD单元上均匀到不均匀状态转变。在指令集架构(instructionsetarchitecture,ISA)表达不同状态或不同执行单元的概念的情况下,此类结构也可嵌入于硬件ISA中。标量单元可能通过此类状态改变转变成向量单元。
例如,在其中存在多个较低层级个例(例如,较低层级子例程)的实例中,前端编译器可将较低层级子例程一起编译为SIMD指令并且一起发布为单一线程。在SIMD结构中可存在并行地执行的线程的多个个例。并发线程的数目可能并不是在所有情况下都随着前端编译器在层次上向下行进而增加,替代地,较低层次可映射到单一线程的数据并行部分。
因此,在一些实例中,服务器装置(例如,用户装置下载应用程序所用的服务器装置)包含存储器单元和处理器。处理器可经配置以执行编译器(例如,前端编译器)以编译以高级编程语言写入的程序的代码。在此实例中,编译所述代码包含添加从所述编译产生的中间语言的中间代码中的信息,所述信息识别较高层级子例程中的较低层级子例程,且较低层级子例程比识别所述较低层级子例程的较高层级子例程更频繁地执行。处理器可输出中间代码以存储在存储器单元中。所述存储器单元可在处理器内部或在处理器外部。
此外,在本发明中,装置(例如,用户装置)包含存储器单元,其经配置以存储从编译以高级编程语言写入的程序的代码产生的中间语言的中间代码。中间代码包含从所述编译产生的信息,所述信息识别较高层级子例程中的较低层级子例程,且较低层级子例程比识别所述较低层级子例程的较高层级子例程更频繁地执行。所述装置还包含处理器,其经配置以接收中间代码,基于识别较高层级子例程的较低层级子例程的信息来编译(例如,用后端编译器)中间语言的中间代码以产生目标代码,以及执行所述目标代码。
图1是说明用于实施本发明的各方面的实例系统10的框图。如所说明,系统10包含主机装置12和用户装置14。主机装置12可包括广泛范围的装置,包含应用程序托管服务器、台式计算机、移动计算装置,以及笔记本(例如,手提式计算机)计算机、平板计算机。用户装置14的实例包含台式计算机和无线通信装置,例如移动计算装置、笔记本计算机、电话手持机(所谓的“智能”手机),以及平板计算机。用户装置14的额外实例包含机顶盒、电视、相机、显示装置、数字媒体播放器、视频游戏控制台、车载计算机,或类似者。用户装置14的这些实例中的一或多者可组合到一起(例如,在无线通信装置上执行的数字媒体播放器)。
用户装置14可经由信道16从主机装置12接收中间代码22。信道16可包括能够将数据从主机装置12移动到用户装置14的一或多个媒体或装置。在一个实例中,信道16可包括启用主机装置12以实时地将数据直接传输到用户装置14的一或多个通信媒体。在此实例中,主机装置12可根据例如无线通信协议等通信标准调制经编码视频数据,并且可将经调制视频数据传输到用户装置14。一或多个通信媒体可包含无线和/或有线通信媒体,例如射频(RF)谱或一或多个物理传输线。一或多个通信媒体可形成例如局域网、广域网或全球网络(例如,因特网)等分组网络的部分。一或多个通信媒体可包含路由器、交换机、基站,或其它促进从主机装置12到用户装置14的通信的设备。
在另一实例中,信道16可包含存储由主机装置12产生的的中间代码22的存储媒体。在此实例中,用户装置14可经由磁盘存取或卡存取来存取存储媒体。存储媒体可包含多种本地存取的数据存储媒体,例如蓝光光盘、DVD、CD-ROM、快闪存储器或用于存储数据的其它合适数字存储媒体。
在另一实例中,信道16可包含文件服务器或另一存储由主机装置12产生的的中间代码22的中间存储装置。在此实例中,用户装置14可经由流式传输或下载存取存储于文件服务器或其它中间存储装置处的中间代码22。文件服务器可为能够存储数据并且将数据传输到用户装置14的服务器类型。实例文件服务器包含网络服务器(例如,用于网站)、文件传输协议(FTP)服务器、网络附接存储(NAS)装置,以及本地磁盘驱动器。
如所说明,主机装置12包含存储器单元18和处理器24。处理器24可实施为多种适合电路中的任一者,例如一或多个微处理器、数字信号处理器(DSP)、专用集成电路(ASIC)、现场可编程门阵列(FPGA)、离散逻辑、硬件或其任何组合。存储器单元18可由多种存储器装置中的任一者形成,例如动态随机存取存储器(DRAM),包含同步DRAM(SDRAM)、磁阻式RAM(MRAM)、电阻式RAM(RRAM),或其它类型的存储器装置。在各种实例中,存储器单元18可与处理器24一起在芯片上。在其中存储器单元18在处理器24的芯片外的实例中,存储器单元18的实例包含但不限于随机存取存储器(RAM)、电可擦除可编程只读存储器(EEPROM)、快闪存储器,或其它可用以携带或存储呈指令和/或数据结构形式的所要程序代码并且可被计算机或处理器存取的媒体。
开发者开发将执行的应用程序并将所述应用程序作为源代码20存储于存储器单元18上。应用程序的实例包含根据OpenCL的内核或根据OpenGL的着色器程序,这两者都在例如图形处理单元(GPU)36的GPU上执行。为易于描述,将源代码20描述为在GPU 36上执行。然而,本发明中描述的技术不限于此,且源代码20可用于在不同处理器类型上执行的应用程序。
存储器单元18存储源代码20和中间代码22。处理器24执行接收源代码20的前端编译器26并且产生中间代码22以存储于存储器单元18中。在此实例中,开发者可开发例如SYCL的高级语言的源代码20。开发者可开发源代码20以使得应用程序的多个个例可并行地执行。
一种开发源代码20的方法包含以层次方式布置于源代码20中的工作组、子组以及工作项目。举例来说,工作组包含一或多个子组和/或一或多个工作项目,且子组包含一或多个工作项目。在此实例中,工作组处于最高层级,子组处于中间层级,且工作项目处于最低层级。
为易于描述,工作组、子组和工作项目通常称为子例程。在层次结构中,较高层级子例程取决于较低层级子例程的执行的完成,且较低层级子例程在源代码20中定义为比较高层级子例程更频繁地执行。作为实例,开发者可写入其中子组定义工作项目执行N次的源代码20。如更详细地解释,在存在并行地执行工作项目的所有个例的足够执行路径的情况下,较低层级子例程不必总是必须比较高层级子例程更频繁地执行。
在本发明中描述的技术中,主机装置12的处理器24从存储器单元18取得源代码20并且执行前端编译器26以编译源代码20从而产生中间代码22,处理器24将所述中间代码存储回到存储器单元18中。前端编译器26称为“前端”编译器,这是因为前端编译器26的输出是中间代码22,而不是应执行的目标代码。中间代码22的一个实例是标准化便携式中间表示(SPIR)汇编语言的代码。
中间代码22的益处是中间代码22是便携的。举例来说,例如用户装置14的装置可取得中间代码22并且编译中间代码22以在本地产生目标代码(例如,目标代码30)。以此方式,中间代码22可通用于各种类型的处理器(因此是便携的),且基于在上面执行目标代码的处理器类型在本地从中间代码22产生目标代码。
作为一个实例,用户装置14的处理器32可致使用户装置14经由信道16取得中间代码22并且可将中间代码22存储于存储器单元28中。处理器32和存储器单元28的实例类似于上文关于处理器24和存储器单元18所描述的实例。然而,处理器24和存储器单元18与处理器32和存储器单元28相比可为更“稳健的”(例如,更快速处理器和更大存储器),这是因为处理器32和存储器单元28可在与处理器24和存储器单元18相比功率和空间是宝贵的手持型装置上,处理器24和存储器单元18可在其中功率和空间并不太宝贵的台式计算机或服务器或计算单元上。所述技术不限于处理器32和存储器单元28没有处理器24和存储器单元18稳健,且此类描述仅出于说明的目的提供。
处理器32可取得中间代码22并且执行后端编译器34以产生目标代码30,处理器32将所述目标代码存储回到存储器单元28中。在一些实例中,处理器32可在程序执行期间产生目标代码30,这调用由目标代码30表示的应用程序以执行。举例来说,处理器32可执行程序(例如,视频游戏)并且在视频游戏执行期间,处理器32可确定将执行应用程序(例如,顶点着色器)。在此实例中,中间代码22可用于顶点着色器,且处理器32可执行后端编译器34以与执行视频游戏实时地产生目标代码30。所述技术不应被视为限于实时编译,且在执行之前编译是可能的。
如图1中所说明,用户装置14包含图形处理单元(GPU)36,且GPU 36执行目标代码30。然而,如在一些实例中所描述,处理器32或某一其它处理单元可执行目标代码30。GPU36的实例包含但不限于DSP、通用微处理器、ASIC、FPGA或其它等效集成或离散逻辑电路。在一些实例中,GPU 36可为具体地针对图形处理经设计的专用硬件。举例来说,图形处理可需要快速并行处理,且GPU 36可具体地针对此类快速并行处理经设计。GPU 36可有可能执行除了图形处理之外的其它任务,例如一般性处理任务。因此,GPU 36可被视为一般性处理GPU(GPGPU)。本发明中描述的技术可适用于其中GPU 36仅执行图形相关任务的实例或其中GPU 36是GPGPU的实例。
如上文所描述,目标代码30可在GPU 36上执行,但可替代地在处理器32上执行。然而,处理器32和GPU 36的执行能力可为不同的,这意味着目标代码30可不同地在处理器32和GPU 36上执行,造成花费不同的时间量来完成执行。
如所说明,GPU 36可经配置以提供SIMD处理以允许执行同一程序的多个个例(例如,SPMD并行性)。在SPMD处理中,同一程序的多个个例并行地(例如,同时)执行,且在SIMD处理中,同一线程的多个个例并行地(例如,同时)执行。程序或指令的个例中的每一者的数据(例如,针对SPMD或SIMD)可为不同的,但指令是相同的。
GPU 36可在GPU 36提供多个SIMD分道(例如,32或64)使得同一指令(例如,线程)的多个个例可并行地执行的意义上是专用的。然而,处理器32可不经设计以提供并行处理,且可更适合于依序处理。例如,处理器32可仅能够一次执行指令的一个个例,而非并行地执行多个个例。由于处理器32和GPU 36的处理能力的此类差异,目标代码30可不同地执行。
帮助理解的基础实例,假设for环路循环32次(例如,for(int i=0;i<32;i++)),且对于每个环路,GPU 36执行特定子例程。可存在实施此for环路的各种方式。第一方式是保持for环路并且使for环路循环依序执行子例程32次。第二方式是“展平”for环路,并且具有子例程的32个并行执行。也存在其它方式(例如,在16个SIMD分道上执行子例程,具有计数为2的for环路)。在上述实例中,执行此for环路的第一方式较适合处理器32,且不太适合GPU 36。执行此for环路的第二方式较适合GPU 36,且不太适合处理器32。因此,后端编译器34产生是用于此for环路的目标代码30的部分的目标代码的方式实现在执行目标代码30处非常高效的处理器32或GPU 36。
根据本发明中描述的技术,执行前端编译器26的处理器24可包含产生中间代码22的函数调用。这些函数调用可识别源代码20的层次结构。后端编译器34接着可使用这些函数调用确定编译中间代码22以产生目标代码30的方式。
例如,与先前实例一致,包含源代码20中的for环路的指令的子例程可被视为与是for环路的部分的子例程相比处于较高层级处。这是因为较高层级子例程(例如,包含for环路的指令的子例程)在完成较低层级子例程(例如,循环的指令)的所有个例之前无法完成,且较低层级子例程在源代码20中实施为比较高层级子例程执行更多次。
上述实例基于for环路来说明源代码20的层次结构的较高层级子例程和较低层级子例程。然而,所述技术不限于环路且还延伸到其它子例程。例如,较高层级子例程可包含指示较低层级子例程应执行N次的指令,但不必将这定义为for环路。存在其它此类实例,且所述技术不应被视为限于上述实例。
此外,在源代码20中,较低层级子例程可被定义为比较高层级子例程执行更多次。然而,在执行处,这可能并不仍然适用。举例来说,假设较高层级子例程的指令指示较低层级子例程执行N次,且在GPU 36上存在N数目个SIMD分道。在此实例中,后端编译器34可复制包含较低层级子例程的代码的中间代码22,并且复制较低层级子例程代码以及较高层级子例程代码N次。在此实例中,较高层级子例程和较低层级子例程两者都执行N次。然而,在源代码20中,较低层级子例程被定义比较高层级子例程更频繁地执行。
可存在前端编译器26可在不同层次的子例程之间进行划分的各种方式。如上文所描述,根据OpenCL,开发者可开发具有工作组、子组和工作项目(例如,子例程的不同实例)的源代码20。在开发期间,工作组、子组和工作项目中的每一者可被指配相应识别值(例如,工作组1、工作组2、子组1、子组2、工作项目1、工作项目2等)。前端编译器26可使用子例程识别值确定层次边界(例如,其中较低层级子例程嵌入于较高层级子例程中)。可存在前端编译器26可确定特定子例程属于哪个层级的其它方式,且所述技术不应被视为限于上述实例。
在一些实例中,前端编译器26可分离出不同子例程,且可通过函数调用识别每一层级的子例程。例如,最高层级子例程(例如,工作组)可被称为操作入口点函数(OpEntryPointSKernel)或操作函数工作组函数(OpFunctionWG或OpWGFunctionCall)。中间层级子例程(例如,子组)可被称为操作函数子组函数(OpFunctionSG或OpSGFunctionCall)。最低层级子例程(例如,工作项目)可被称为操作函数工作项目函数(OpFunctionWI或OpWIFunctionCall)。最低层级子例程也可简称为OpFunction。
在解析源代码20时,处理器24经由执行前端编译器26可在较高层级子例程中将较低层级子例程的代码替换为对所述较低层级子例程的函数调用。作为实例,子组名称可识别为SG1,且前端编译器26可将子组SG1的代码替换为函数调用:OpSGFunctionCall(SG1)。作为另一实例,具有识别值WI1的工作项目可嵌入于工作组中。在此实例中,在解析工作组的源代码期间,前端编译器26可将工作组中的工作项目WI1的代码替换为函数调用:OpWGFunctionCall(WI1)。在上述实例中,OpSGFunctionCall(SG1)可被视为工作组子例程中需要执行SG1(例如,较低层级子例程)的函数调用。类似地,OpWGFunctionCall(WI1)可被视为工作组子例程或子组子例程中需要执行WI1(例如,较低层级子例程)的函数调用。
如上文所描述,用户装置14可经由信道16从主机装置12取得中间代码22。在一些情况下,中间代码22可包含用于分离出的子例程的中间代码,而非具有直接嵌入于中间代码22中的子例程的中间代码的连续集。举例来说,用户装置14可分别取得工作组、子组和工作项目子例程中的每一者的中间代码,作为取得中间代码22的部分,而非中间代码22的指令的一个单个连续集。虽然子例程可划分为独立单元,但较高层级子例程中的每一者可包含对一或多个较低层级子例程的一或多个函数调用。
在取得中间代码22并将其存储于存储器单元28中之后,处理器32执行后端编译器34以编译中间代码22以产生目标代码30。后端编译器34可经配置以解析中间代码22,并且在对较低层级子例程的函数调用的每个个例处确定是否展平对较低层级子例程的函数调用或主要以依序方式执行较低层级子例程。如上文所描述,展平函数调用意指并行地执行更多的函数调用。
作为实例,处理器32可确定后端编译器34产生目标代码30所针对的应用程序是在处理器32上执行还是在GPU 36上执行。另外,处理器32可预配置有指示处理器32和GPU 36的处理能力的信息(或可在运行时间期间接收此类信息)。例如,处理器32可预配置(例如,处理器32的操作系统)有指示GPU 36上的SIMD分道的数目(例如,GPU36可并行地执行多少个指令)的信息。类似地,处理器32可预配置有指示处理器32可并行地执行多少个指令的信息。
以此方式,后端编译器34可以采用在上面执行应用程序的处理单元的处理能力的方式编译中间代码22。举例来说,如果处理器32确定应用程序在处理器32上执行,那么后端编译器34可解析中间代码22,并且对于其中中间代码22指示较高层级子例程中的同一较低层级子例程的多个个例(例如,对同一较低层级子例程的多个调用)的执行的个例,后端编译器34可将这些个例布置为在目标代码30中执行的依序指令,这是因为处理器32不提供并行处理能力。如果处理器32确定应用程序在GPU 36上执行,那么后端编译器34可解析中间代码22,并且对于其中中间代码22指示较高层级子例程中同一较低层级子例程的多个个例的执行的个例,后端编译器34可将这些个例布置为在目标代码30中执行的并行指令,这是因为GPU 36提供并行处理能力。
在一些实例中,在GPU 36中可能不存在用于完全并行执行的足够SIMD分道。对于这些情况,后端编译器34可尽可能地并行化执行,但视需要使用环路。作为实例,如果32个并行指令需求处理,但仅有16个SIMD分道,那么后端编译器34可并行地执行较低层级子例程的16个例,但包含两个执行的环路,使得较低层级子例程的所有32个例执行。
如上文所描述,工作项目是最低层级子例程。在一些实例中,后端编译器34在解析对工作项目的函数调用后可即刻确定工作项目将在每线程基础上实施(例如,工作项目的单个指令的多个个例并行地执行)。因此,前端编译器26可被视为在线程层级处编译工作项目子例程。
后端编译器34在解析对子组的函数调用后可即刻确定子组在每wave基础上实施(例如,多个线程一起形成wave)。举例来说,在每wave基础上实施意指组成wave的线程的指令以及数据是相同的。在此实例中,因为较高层级子例程对于较低层级子例程的多个执行循环是恒定的,所以子组子例程可被视为子组子例程调用的工作项目子例程的wave。
在一些实例中,前端编译器26可复制较低层级子例程的多个个例(例如,基于较低层级子例程执行的次数,使得复制数等于较低层级子例程执行的次数)。前端编译器26可将较低层级子例程的此多个个例一起编译成SIMD指令,并且提供对子例程的单个函数调用,其包含较低层级子例程的多个复本。在此类实例中,如果应用程序在GPU 36上执行,那么后端编译器34可分离出用于并行执行的较低层级子例程的多个复本。
上述实例关于for环路作为实例进行描述。然而,本发明中描述的技术可适用于其它情况,例如也适用于向量运算。举例来说,N位变量可被视为多个M位浮点(例如,128位变量是4个32位浮点)。指令可定义将针对N位变量执行的过程,且此过程接着可应用于M位浮点。
举例来说,使用float4表示128位流式传输的SIMD扩展(streaming SIMDextension,SSE),以下是实例指令集:
float a=1.f;
float4 b=(float4)a;//这已跨float4执行1.f 4次
b.x=2.f;//仅更新4个值中的第一个
float c=sum(b);//添加b的每个元素。c应具有值5(2+1+1+1)
然而,上述指令集可基于可用的SIMD分道的数目进行拆分,使得指令可并行地执行。例如,在层次代码中,指令可通过在更细粒度函数中包含b来将运算应用于b。作为实例,
在上述代码中,四个浮点中的每一者在不同分道中操作。举例来说,在第一分道中(例如,get_local_id(0)==0所针对的情况),b的值设置成等于2,且针对所有其它分道,b的值等于1(例如,归因于float b=a)。接着,c等于跨越对于2+1+1+1等于5的分道的总和。
以此方式,128位SSE的第一指令集可被视为粗粒度函数,且128位SSE的第二指令集可被视为细粒度函数。在一些情况下,可有可能组合两个格式。例如,存在其中粗粒度函数调用细粒度函数的层次代码。作为实例,
float a=1.f;
float4 b=(float4)a;
b=fine-update(b);//这是中间语言(IL)的粗到细函数调用指令中的一者
float c=sum(b);
其中b是IL的细粒度函数,其可表示为:
在此实例中,存在直接映射到细粒度代码中的32位浮点的粗粒度代码中的float4变量(128位向量)。在以上实例中,十字分道函数(cross-lane function)可应用于粗粒度代码,使得相同数据可更新为向量代码(在以上实例中,一次128位4元素)或细粒度代码。
因此,在一些实例中,用户装置14可被视为用于编译的装置。在此类实例中,处理器32可经配置以接收从编译应用程序的源代码20产生的中间代码22。如所描述,中间代码22包含从前端编译器26进行的编译产生的信息,其中中间代码22识别较高层级子例程中的较低层级子例程的层次结构。较低层级子例程在应用程序的源代码20中定义为比识别较低层级子例程的较高层级子例程更频繁地执行。处理器32经由后端编译器34基于识别较高层级子例程中的较低层级子例程的信息,编译中间代码22以产生目标代码30。处理器32将目标代码30存储于存储器单元28中。
在一个实例中,处理器32可确定处理器32应执行目标代码30并且可执行目标代码30。在此实例中,为了编译,处理器32经由后端编译器34可编译中间代码22以产生致使处理器32依序执行较低层级子例程的多个个例的目标代码30。
在另一实例中,处理器32可确定GPU 36应执行目标代码30,并且可指示GPU 36执行目标代码30。在此实例中,为了编译,处理器32经由后端编译器34可编译中间代码22以产生致使GPU 36并行地执行较低层级子例程的多个个例的目标代码30。
对于编译中间代码22的后端编译器34,后端编译器34可识别较高层级子例程中的要求执行较低层级子例程的函数调用(例如,OpFunctionSG或OpSGFunctionCall、OpFunctionWI或OpWIFunctionCall基于其相应识别(例如SG1或WI1)要求执行较低层级子例程)。后端编译器34可从较高层级子例程确定执行较低层级子例程的次数(例如,在较高层级子例程中存在定义执行较低层级子例程的次数的for环路或指令的情况下)。
后端编译器34可基于所述对将执行较低层级子例程的所述次数的确定,产生目标代码。举例来说,如果在处理器32上执行,那么后端编译器34可产生致使处理器32依序执行较低层级子例程将执行较低层级子例程的次数的目标代码30。如果在GPU 36上执行,那么后端编译器34可产生致使GPU 36并行地执行由较低层级子例程应执行的次数定义的多个个例的目标代码30。
在一些实例中(但不必所有实例),后端编译器34可实时地产生目标代码30。例如,处理器32可执行使用产生目标代码30所针对的应用程序的程序(例如,视频游戏根据本发明中描述的技术使用产生目标代码所用于的顶点着色器)。在此实例中,处理器32经由后端编译器34可在程序执行期间编译中间代码22(例如,在视频游戏执行期间编译顶点着色器的目标代码)。
图2是说明可实施本发明中所描述的一或多个实例技术的装置14的实例的框图。如图2中所说明,并且类似于图1中的实例,装置14可包含处理器32(例如,中央处理单元(CPU))、GPU 36以及存储器单元28。装置14可包含除了图1中所说明的组件之外的其它组件。
此外,在一些实例中,处理器32和GPU 36可形成为容纳于单个电路封装内(例如,形成为公共处理器)的公共集成电路。然而,本发明的各方面不限于此,且处理器32和GPU36中的一或多者可为容纳于独立电路封装中的独立集成电路。一般来说,装置14可被视为包含集成电路,其包含经配置以执行本发明中所描述的实例技术中的一或多者的处理器32。
存储器单元28可为计算机可读存储媒体的实例。举例来说,除了存储目标代码30和中间代码22之外,存储器单元28还可存储致使处理器32和GPU 36执行归于本发明中的每一者的函数的指令。存储器单元28可被视为包括致使一或多个处理器(例如,处理器32或GPU 36)执行各种函数的指令的计算机可读存储媒体。
在一些实例中,存储器单元28可被视为非暂时性存储媒体。术语“非暂时性”可指示存储媒体未体现于载波或传播信号中。然而,术语“非暂时性”不应解释为意味着存储器单元28是不可移动的。作为一个实例,存储器单元28可从装置14去除,并且移动到另一装置。作为另一实例,基本上类似于存储器单元28的存储装置可插入到装置14中。在某些实例中,非暂时性存储媒体可存储可随时间改变的数据(例如,在RAM中)。
GPU 36可包含着色器处理器42和固定功能管线44。着色器处理器42(有时称为着色器核心)可为在上面执行程序的目标代码的GPU 36的核心。举例来说,目标代码30可在着色器处理器42上执行,且在一些实例中,可经由着色器处理器42的多个处理元件以并行方式(例如,同时)执行。作为实例,处理器32可使用着色器处理器42中处理元件的数目确定后端编译器34如何编译中间代码22(例如,着色器处理器42的处理元件的数目定义GPU 36的SIMD分道的数目)。
固定功能管线44可包含执行固定功能的硬件单元。固定功能管线44和着色器处理器42一起形成用于与GPU 36一起进行处理的管线。举例来说,在处理之后,着色器处理器42可产生被固定功能管线44的一或多个单元接收以用于进一步处理的数据,且每个单元输出到下一个单元,其中着色器处理器42间歇地执行一些处理(例如,经由顶点着色器、几何着色器或像素着色器),直到GPU 36将最终产生的数据输出到存储器单元28。
处理器32可执行存储于存储器单元28中的一或多个程序,例如程序40。程序40的实例包含但不限于网络浏览器、用户接口、电子邮件应用程序、电子表格应用程序、文字处理应用程序、图形制作应用程序、视频游戏或其它产生可查看对象以供显示的应用程序。例如,程序40可为在执行时输出在显示器上进行显示的图形内容的视频游戏。作为一个实例,程序40可致使GPU 36或处理器32执行目标代码30(例如,程序40是视频游戏,目标代码30用于顶点着色器,且视频游戏致使GPU 36执行顶点着色器)。在一些此类实例中,处理器32可实时地编译(例如,在程序40执行期间)中间代码22以产生程序40所使用的目标代码30。
如上文所描述,处理器32可接收中间代码22并且执行后端编译器34。经由后端编译器34,处理器32可基于识别较高层级子例程中的较低层级子例程的信息,编译中间语言(例如,SPIR)的中间代码22以产生目标代码30,并且存储目标代码30以用于最终执行。处理器32可将目标代码30存储于存储器单元28中。虽然存储器单元28说明为在处理器32和GPU36外部,但在一些实例中,存储器单元28可部分地在处理器32和GPU 36内。举例来说,处理器32和GPU 36各自包含(分别或一起)内部存储器,例如高速缓冲存储器。此内部存储器可被视为是存储器单元的一般性术语的部分(例如,存储器单元28的部分)。
在一些实例中,后端编译器34可基于在上面执行应用程序的处理器类型,编译中间代码22。举例来说,如果目标代码30在处理器32上执行,那么后端编译器34可使用中间代码22中的信息优化目标代码30以用于在处理器32上执行(例如,用于依序执行)。如果目标代码30在GPU 36上执行,那么后端编译器34可使用中间代码22中的信息优化目标代码30以用于在GPU 36上执行(例如,用于并行执行)。
如上文所描述,中间代码22可由包含执行前端编译器26的处理器24的主机装置12产生。然而,所述实例不必受此限制。在一些实例中,处理器32可执行前端编译器26和后端编译器34或执行单个编译器,其执行前端编译器和后端编译器两者的函数。在此类实例中,处理器32可产生中间代码22。
也如在图2中所说明,处理器32可经配置以执行图形驱动器38。图形驱动器38的功能性可由硬件或固件或硬件、软件和固件的组合执行。为了简单起见,图形驱动器38描述为在处理器32上执行的软件。
图形驱动器38可根据特定标准经设计。在处理器32上执行的图形驱动器38可经配置以实施应用程序编程接口(API)(例如,OpenGL或OpenCL或两者的组合)。在此类实例中,着色器程序或内核(例如,产生目标代码30所用于的应用程序)可根据同一API经配置为图形驱动器38(例如,根据OpenCL的内核和根据OpenGL的着色器程序)。虽然未说明,存储器单元28可存储图形驱动器38的代码,处理器32从存储器单元28取得所述代码以用于执行。
图形驱动器38可经配置以执行处理器32和GPU 36之间的通信。因此,在本发明中,当处理器32描述为指示GPU 36或者与GPU 36通信时,处理器32经由图形驱动器38执行此类函数。举例来说,当GPU 36执行目标代码30时,处理器32可经由图形驱动器38指示GPU 36执行目标代码30(例如,从存储器单元28取得目标代码30并且执行目标代码30的指令)。
图3是更详细地说明图2的实例装置的框图。举例来说,如上文所指示,装置14的实例包含但不限于移动无线电话、PDA、包含视频显示器的视频游戏控制台、移动视频会议单元、手提式计算机、台式计算机、电视机顶盒以及类似者。
如图3中所说明,装置14可包含显示器46、处理器32、GPU 36、包含帧缓冲器54的存储器单元28、收发器模块48、用户接口50以及显示处理器52。处理器32、GPU36和存储器单元28可与图1中所说明的那些组件大体上类似或相同。出于简洁的目的,仅详细地描述在图3中展示但在图1或2中不展示的组件。
如图3中所说明的装置14可包含出于清楚起见未在图3中示出的额外模块或单元。举例来说,装置14可包含未在图3中示出的扬声器和麦克风,以在其中装置14是移动无线电话的实例中或在其中装置14是媒体播放器的扬声器中实现电话通信。此外,装置14中展示的各种模块和单元可能在装置14的每个实例中不是必要的。举例来说,在其中装置14是台式计算机或其它经装备以与外部用户接口或显示器接口连接的装置的实例中,用户接口50和显示器46可在装置14外部。作为另一实例,在其中显示器46是移动装置的触敏或压敏显示器的实例中,用户接口50可以是显示器46的部分。
显示器46可包括液晶显示器(LCD)、阴极射线管(CRT)显示器、等离子显示器、触敏显示器、压敏显示器或另一类型的显示装置。用户接口50的实例包含但不限于轨迹球、鼠标、键盘和其它类型的输入装置。用户接口50也可为触摸屏并且可并入为显示器46的一部分。
收发器模块48可包含允许在装置14与另一装置或网络(例如,经由图1的信道16)之间无线或有线通信的电路。收发器模块48可包含一或多个调制器、解调器、放大器、天线和用于有线或无线通信其它此类电路。
显示处理器52可经配置以将图形数据输出到显示器46。举例来说,显示处理器52从帧缓冲器54接收数据,所述帧缓冲器存储GPU 36输出的图形数据。显示处理器52输出所取得的数据并且控制显示器46显示器图形数据。
图4是说明用于实施本发明的各方面的实例的流程图。关于处理器32描述图4中所说明的实例。举例来说,用于编译的装置(例如用户装置14)可包含存储器单元28和集成电路,所述集成电路包括经配置以执行图4中所说明的实例技术的处理器32。
在图4中说明的实例中,处理器32可接收从编译应用程序的源代码20产生的中间代码22(60)。如上文所描述且在图4中所指示,中间代码22包含从所述编译产生的信息,所述信息识别较高层级子例程中的较低层级子例程的层次结构,且较低层级子例程在应用程序的源代码中定义为比识别较低层级子例程的较高层级子例程更频繁地执行。
在一些实例中,处理器32可接收较高层级子例程的中间代码,所述较高层级子例程的中间代码与较低层级子例程的中间代码分离。举例来说,中间代码22可不为包含所有子例程的中间代码的依序代码。而是,中间代码22可包含呈第一函数的形式的较高层级子例程的中间代码,且包含呈第二函数的形式的较低层级子例程的中间代码。第一函数可调用第二函数用于执行,但当处理器32接收第一函数和第二函数时,处理器32可分别接收所述第一函数和所述第二函数函数。
处理器32经由后端编译器34可基于识别较高层级子例程中的较低层级子例程的信息,编译中间代码22以产生目标代码30(62)。作为一个实例,后端编译器34可识别较高层级子例程中的要求执行较低层级子例程的函数调用。后端编译器34可从较高层级子例程确定较低层级子例程应执行的次数,并且基于对较低层级子例程应执行的次数的确定,产生目标代码30。
处理器32可存储目标代码30(64)。举例来说,处理器32可将目标代码30写入到存储器单元28以用于最终执行。例如,处理器32或GPU 36中的一者可执行目标代码30(66)。
作为一个实例,处理器32可确定处理器32应执行目标代码30,且处理器32可执行目标代码30。在此实例中,为了编译,处理器32经由后端编译器34可编译中间代码22以产生致使处理器32依序执行较低层级子例程的多个个例的目标代码30。作为另一实例,处理器32可确定GPU 36应执行目标代码30,且处理器32可经由图形驱动器38指示GPU 36执行着色器处理器42的一或多个处理元件上的目标代码30。在此实例中,为了编译,处理器32经由后端编译器34可编译中间代码22以产生致使GPU 36并行地执行较低层级子例程的多个个例的目标代码30。
图5是说明用于实施本发明的各方面的另一实例的流程图。如所说明,处理器24经由前端编译器26可编译源代码20以产生包含识别较高层级子例程中的较低层级子例程的层次结构的信息的中间代码22(70)。处理器24接着可经由信道16输出中间代码22以用于由用户装置14最终接收和执行(72)。举例来说,用户装置14可执行上文关于图4所描述的实例技术,例如基于识别较低层级子例程和较高层级子例程的信息来编译中间代码以产生目标代码,存储目标代码,以及执行目标代码。
如上文所描述,在产生中间代码22时,前端编译器26可基于源代码20中的识别信息在较高层级子例程和较低层级子例程之间做出区分。当前端编译器26解析较高层级子例程中的较低层级子例程时,前端编译器26将较高层级子例程中的较低层级子例程替换为对较低层级子例程的函数调用。这些函数调用要求较低层级子例程的执行,以及后端编译器34用于产生目标代码30所依赖的内容。
在一或多个实例中,所描述的功能可用硬件、软件、固件或其任何组合来实施。如果以软件实施,那么所述功能可作为一或多个指令或代码存储在包括非暂时性计算机可读媒体的制品上。计算机可读媒体可包含计算机数据存储媒体或通信媒体,通信媒体包含促进将计算机程序从一处传递到另一处的任何媒体。数据存储装置可为可由一或多个计算机或一或多个处理器存取以取得用于实施本发明中描述的技术的指令、代码和/或数据结构的任何可用媒体。作为实例而非限制,此类计算机可读媒体可包括RAM、ROM、EEPROM、CD-ROM或其它光盘存储装置、磁盘存储装置或其它磁性存储装置、闪存存储器或可用来携带或存储呈指令或数据结构形式的所要程序代码且可由计算机存取的任何其它媒体。如本文中所使用,磁盘和光盘包含压缩光盘(CD)、激光光盘、光学光盘、数字多功能光盘(DVD)、软性磁盘和蓝光光盘,其中磁盘通常以磁性方式再现数据,而光盘利用激光以光学方式再现数据。以上各项的组合也应包含在计算机可读媒体的范围内。
代码可由一或多个处理器执行,所述一或多个处理器例如一或多个DSP、通用微处理器、ASIC、FPGA,或其它等效集成或离散逻辑电路。另外,在一些方面中,可在专用硬件和/或软件模块内提供本文中所描述的功能性。此外,所述技术可完全实施于一或多个电路或逻辑元件中。
本发明的技术可实施于包含无线手持机、集成电路(IC)或一组IC(例如,芯片组)的多种装置或设备中。本发明中描述各种组件、模块或单元是为了强调经配置以执行所揭示的技术的装置的功能方面,但未必需要由不同硬件单元实现。实际上,如上文所描述,各种单元可以结合合适的软件和/或固件组合在硬件单元中,或者通过互操作硬件单元的集合来提供,所述硬件单元包含如上文所描述的一或多个处理器。
描述了各种实例。这些和其它实例在所附权利要求书的范围内。

Claims (21)

1.一种编译方法,所述方法包括:
用处理器接收从编译应用程序的源代码产生的中间代码,其中所述中间代码包含从所述编译产生的信息,所述信息识别较高层级子例程中的较低层级子例程的层次结构,且其中所述较低层级子例程在所述应用程序的所述源代码中定义为比识别所述较低层级子例程的所述较高层级子例程更频繁地执行;
用所述处理器基于识别所述较高层级子例程中的所述较低层级子例程的所述信息,编译所述中间代码以产生目标代码;以及
存储所述目标代码。
2.根据权利要求1所述的方法,其进一步包括:
确定所述处理器应执行所述目标代码;以及
用所述处理器执行所述目标代码,
其中编译包括编译所述中间代码以产生致使所述处理器依序执行所述较低层级子例程的多个个例的所述目标代码。
3.根据权利要求1所述的方法,其进一步包括:
确定图形处理单元GPU或数字信号处理器DSP应执行所述目标代码;以及
指示所述GPU或DSP执行所述目标代码,
其中编译包括编译所述中间代码以产生致使所述GPU或DSP并行地执行所述较低层级子例程的多个个例的所述目标代码。
4.根据权利要求1所述的方法,其中编译包括:
识别所述较高层级子例程中的要求执行所述较低层级子例程的函数调用;
从所述较高层级子例程确定所述较低层级子例程应执行的次数;以及
基于对所述较低层级子例程应执行的所述次数的所述确定,产生所述目标代码。
5.根据权利要求1所述的方法,其中接收所述中间代码包括接收所述较高层级子例程的中间代码,所述较高层级子例程的中间代码与所述较低层级子例程的中间代码分离。
6.根据权利要求1所述的方法,其进一步包括:
用所述处理器执行使用所述应用程序的程序,
其中编译包括在所述程序执行期间编译所述中间代码。
7.一种用于编译的装置,所述装置包括:
存储器单元;以及
集成电路,其包括处理器,所述处理器经配置以:
接收从编译应用程序的源代码产生的中间代码,其中所述中间代码包含从所述编译产生的信息,所述信息识别较高层级子例程中的较低层级子例程的层次结构,且其中所述较低层级子例程在所述应用程序的所述源代码中定义为比识别所述较低层级子例程的所述较高层级子例程更频繁地执行;
基于识别所述较高层级子例程中的所述较低层级子例程的所述信息,编译所述中间代码以产生目标代码;以及
将所述目标代码存储于所述存储器单元中。
8.根据权利要求7所述的装置,其中所述处理器经配置以:
确定所述处理器应执行所述目标代码;以及
执行所述目标代码,
其中为了编译,所述处理器经配置以编译所述中间代码以产生致使所述处理器依序执行所述较低层级子例程的多个个例的所述目标代码。
9.根据权利要求7所述的装置,所述装置进一步包括:
图形处理单元GPU或数字信号处理器DSP中的至少一者,
其中所述处理器经配置以:
确定所述GPU或DSP应执行所述目标代码;以及
指示所述GPU或DSP执行所述目标代码,
其中为了编译,所述处理器经配置以编译所述中间代码以产生致使所述GPU或DSP并行地执行所述较低层级子例程的多个个例的所述目标代码。
10.根据权利要求7所述的装置,其中为了编译,所述处理器经配置以:
识别所述较高层级子例程中的要求执行所述较低层级子例程的函数调用;
从所述较高层级子例程确定所述较低层级子例程应执行的次数;以及
基于对所述较低层级子例程应执行的所述次数的所述确定,产生所述目标代码。
11.根据权利要求7所述的装置,其中为了接收所述中间代码,所述处理器经配置以接收所述较高层级子例程的中间代码,所述较高层级子例程的中间代码与所述较低层级子例程的中间代码分离。
12.根据权利要求7所述的装置,其中所述处理器经配置以:
执行使用所述应用程序的程序,
其中为了编译,所述处理器经配置以在所述程序执行期间编译所述中间代码。
13.根据权利要求7所述的装置,其中所述装置包括无线通信装置。
14.一种用于编译的装置,所述装置包括:
用于接收从编译应用程序的源代码产生的中间代码的装置,其中所述中间代码包含从所述编译产生的信息,所述信息识别较高层级子例程中的较低层级子例程的层次结构,且其中所述较低层级子例程在所述应用程序的所述源代码中定义为比识别所述较低层级子例程的所述较高层级子例程更频繁地执行;
用于基于识别所述较高层级子例程中的所述较低层级子例程的所述信息,编译所述中间代码以产生目标代码的装置;以及
用于存储所述目标代码的装置。
15.根据权利要求14所述的装置,进一步包括:
用于确定包含所述用于编译的装置的处理器应执行所述目标代码的装置;以及
用于执行所述目标代码的装置,
其中所述用于编译的装置包括用于编译所述中间代码以产生致使所述处理器依序执行所述较低层级子例程的多个个例的所述目标代码的装置。
16.根据权利要求14所述的装置,进一步包括:
用于确定图形处理单元GPU或数字信号处理器DSP应执行所述目标代码的装置;以及
用于指示所述GPU或DSP执行所述目标代码的装置,
其中所述用于编译的装置包括用于编译所述中间代码以产生致使所述GPU或DSP并行地执行所述较低层级子例程的多个个例的所述目标代码的装置。
17.根据权利要求14所述的装置,其中所述用于编译的装置包括:
用于识别所述较高层级子例程中的要求执行所述较低层级子例程的函数调用的装置;
用于从所述较高层级子例程确定所述较低层级子例程应执行的次数的装置;以及
用于基于对所述较低层级子例程应执行的所述次数的所述确定,产生所述目标代码的装置。
18.根据权利要求14所述的装置,其中所述用于接收所述中间代码的装置包括用于接收所述较高层级子例程的中间代码的装置,所述较高层级子例程的中间代码与所述较低层级子例程的中间代码分离。
19.根据权利要求14所述的装置,进一步包括:
用于执行使用所述应用程序的程序的装置,
其中所述用于编译的装置包括用于在所述程序执行期间编译所述中间代码的装置。
20.一种包括指令的非暂时性计算机可读存储媒体,所述指令在被执行时致使一或多个处理器:
接收从编译应用程序的源代码产生的中间代码,其中所述中间代码包含从所述编译产生的信息,所述信息识别较高层级子例程中的较低层级子例程的层次结构,且其中所述较低层级子例程在所述应用程序的所述源代码中定义为比识别所述较低层级子例程的所述较高层级子例程更频繁地执行;
基于识别所述较高层级子例程中的所述较低层级子例程的所述信息,编译所述中间代码以产生目标代码;以及
存储所述目标代码。
21.根据权利要求20所述的计算机可读存储媒体,其中致使所述一或多个处理器进行编译的所述指令包括致使所述一或多个处理器执行以下步骤的指令:
识别所述较高层级子例程中的要求执行所述较低层级子例程的函数调用;
从所述较高层级子例程确定所述较低层级子例程应执行的次数;以及
基于对所述较低层级子例程应执行的所述次数的所述确定,产生所述目标代码。
CN201680031758.9A 2015-06-15 2016-04-26 从包含层次子例程信息的中间代码产生目标代码 Expired - Fee Related CN107710150B (zh)

Applications Claiming Priority (5)

Application Number Priority Date Filing Date Title
US201562175646P 2015-06-15 2015-06-15
US62/175,646 2015-06-15
US14/958,572 2015-12-03
US14/958,572 US9830134B2 (en) 2015-06-15 2015-12-03 Generating object code from intermediate code that includes hierarchical sub-routine information
PCT/US2016/029375 WO2016204865A1 (en) 2015-06-15 2016-04-26 Generating object code from intermediate code that includes hierarchical sub-routine information

Publications (2)

Publication Number Publication Date
CN107710150A true CN107710150A (zh) 2018-02-16
CN107710150B CN107710150B (zh) 2021-07-16

Family

ID=57515929

Family Applications (1)

Application Number Title Priority Date Filing Date
CN201680031758.9A Expired - Fee Related CN107710150B (zh) 2015-06-15 2016-04-26 从包含层次子例程信息的中间代码产生目标代码

Country Status (6)

Country Link
US (1) US9830134B2 (zh)
EP (1) EP3308267A1 (zh)
JP (1) JP2018517986A (zh)
KR (1) KR20180017020A (zh)
CN (1) CN107710150B (zh)
WO (1) WO2016204865A1 (zh)

Cited By (2)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN109144722A (zh) * 2018-07-20 2019-01-04 上海研鸥信息科技有限公司 一种多应用高效共用fpga资源的管理系统及方法
WO2021174538A1 (zh) * 2020-03-06 2021-09-10 深圳市欢太科技有限公司 应用处理方法及相关装置

Families Citing this family (12)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US10241761B2 (en) * 2014-12-29 2019-03-26 Nvidia Corporation System and method for compiler support for compile time customization of code
WO2018094087A1 (en) * 2016-11-17 2018-05-24 The Mathworks, Inc. Systems and methods for generating code for parallel processing units
US10394536B2 (en) * 2017-03-02 2019-08-27 International Business Machines Corporation Compiling a parallel loop with a complex access pattern for writing an array for GPU and CPU
US10241766B2 (en) 2017-06-22 2019-03-26 Microsoft Technology Licensing, Llc Application binary interface cross compilation
US10657698B2 (en) * 2017-06-22 2020-05-19 Microsoft Technology Licensing, Llc Texture value patch used in GPU-executed program sequence cross-compilation
US10102015B1 (en) 2017-06-22 2018-10-16 Microsoft Technology Licensing, Llc Just in time GPU executed program cross compilation
US10289393B2 (en) 2017-06-22 2019-05-14 Microsoft Technology Licensing, Llc GPU-executed program sequence cross-compilation
US10445119B2 (en) * 2017-06-30 2019-10-15 Intel Corporation Software reconfigurable mobile devices and methods
US11475197B2 (en) * 2018-09-25 2022-10-18 Synopsys, Inc. Hardware simulation systems and methods for identifying state-holding loops and oscillating loops
KR102649351B1 (ko) 2018-11-06 2024-03-20 삼성전자주식회사 세분화된 상태들에 기초한 그래픽스 프로세서 및 그래픽스 처리 방법
CN111091612B (zh) * 2019-10-09 2023-06-02 武汉凌久微电子有限公司 一种抽象目标码架构的着色语言机器码生成方法及装置
US11327733B2 (en) * 2020-05-27 2022-05-10 Blaize, Inc. Method of using multidimensional blockification to optimize computer program and device thereof

Citations (10)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US5815719A (en) * 1996-05-07 1998-09-29 Sun Microsystems, Inc. Method and apparatus for easy insertion of assembler code for optimization
US7228531B1 (en) * 2003-02-03 2007-06-05 Altera Corporation Methods and apparatus for optimizing a processor core on a programmable chip
CN101078994A (zh) * 2006-05-26 2007-11-28 松下电器产业株式会社 编译器装置、编译器方法和编译器程序
US20080256521A1 (en) * 2006-01-25 2008-10-16 International Business Machines Corporation Apparatus and Method for Partitioning Programs Between a General Purpose Core and One or More Accelerators
CN101563673A (zh) * 2006-12-14 2009-10-21 富士通株式会社 编译方法以及编译器
CN103119561A (zh) * 2010-09-23 2013-05-22 苹果公司 用于对非叶代码进行基于编译器的矢量化的系统和方法
US8468510B1 (en) * 2008-01-16 2013-06-18 Xilinx, Inc. Optimization of cache architecture generated from a high-level language description
CN103282878A (zh) * 2010-12-24 2013-09-04 英特尔公司 基于循环分割或索引阵列的循环并行化
CN103838614A (zh) * 2014-02-19 2014-06-04 华为技术有限公司 一种数据处理方法及装置
US20140196016A1 (en) * 2013-01-07 2014-07-10 Advanced Micro Devices, Inc. Layered programming for heterogeneous devices

Family Cites Families (33)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US5175856A (en) * 1990-06-11 1992-12-29 Supercomputer Systems Limited Partnership Computer with integrated hierarchical representation (ihr) of program wherein ihr file is available for debugging and optimizing during target execution
JP2818016B2 (ja) * 1990-08-09 1998-10-30 株式会社日立製作所 プロセス並列実行方法および装置
US5659753A (en) * 1991-02-27 1997-08-19 Digital Equipment Corporation Interface for symbol table construction in a multilanguage optimizing compiler
IL100990A (en) * 1991-02-27 1995-10-31 Digital Equipment Corp Multilingual optimization compiler that uses Gladi in the production of a multi-pass cipher
US5450575A (en) * 1991-03-07 1995-09-12 Digital Equipment Corporation Use of stack depth to identify machine code mistakes
US5339238A (en) * 1991-03-07 1994-08-16 Benson Thomas R Register usage tracking in translating code for different machine architectures by forward and reverse tracing through the program flow graph
US5923882A (en) * 1995-08-29 1999-07-13 Silicon Graphics, Inc. Cross-module optimization for dynamically-shared programs and libraries
JP3606654B2 (ja) 1995-10-30 2005-01-05 富士通株式会社 コンパイラ装置
US5920723A (en) * 1997-02-05 1999-07-06 Hewlett-Packard Company Compiler with inter-modular procedure optimization
US6195793B1 (en) 1998-07-22 2001-02-27 International Business Machines Corporation Method and computer program product for adaptive inlining in a computer system
US6675377B1 (en) * 1999-09-13 2004-01-06 Matsushita Electric Industrial Co., Ltd. Program conversion apparatus
US7036116B2 (en) * 2001-03-23 2006-04-25 International Business Machines Corporation Percolating hot function store/restores to colder calling functions
US7278136B2 (en) 2002-07-09 2007-10-02 University Of Massachusetts Reducing processor energy consumption using compile-time information
JP3992102B2 (ja) 2003-06-04 2007-10-17 インターナショナル・ビジネス・マシーンズ・コーポレーション コンパイラ装置、コンパイル方法、コンパイラプログラム、及び記録媒体
US7958330B2 (en) * 2004-02-02 2011-06-07 Fujitsu Limited Compiler program, compiler program recording medium, compile method, and program processing system
US8607209B2 (en) 2004-02-04 2013-12-10 Bluerisc Inc. Energy-focused compiler-assisted branch prediction
JP2006338616A (ja) * 2005-06-06 2006-12-14 Matsushita Electric Ind Co Ltd コンパイラ装置
US7730464B2 (en) * 2005-09-14 2010-06-01 Microsoft Corporation Code compilation management service
JP4923240B2 (ja) * 2006-01-17 2012-04-25 国立大学法人東京工業大学 プログラム処理装置、並列処理プログラム、プログラム処理方法、並列処理コンパイラ、並列処理コンパイラを格納した記録媒体およびマルチプロセッサシステム
US7921418B2 (en) * 2006-08-15 2011-04-05 International Business Machines Corporation Compile time evaluation of library functions
WO2008120287A2 (ja) * 2007-02-27 2008-10-09 Fujitsu Limited モジュール生成装置、モジュール生成方法、モジュール生成プログラムおよび該プログラムを記録した記録媒体
US8196126B2 (en) * 2007-10-29 2012-06-05 Sap Ag Methods and systems for dynamically generating and optimizing code for business rules
US20100153934A1 (en) * 2008-12-12 2010-06-17 Peter Lachner Prefetch for systems with heterogeneous architectures
JP4806060B2 (ja) * 2009-09-15 2011-11-02 インターナショナル・ビジネス・マシーンズ・コーポレーション コンパイラ・プログラム、コンパイル方法及びコンピュータ・システム
US8959496B2 (en) * 2010-04-21 2015-02-17 Microsoft Corporation Automatic parallelization in a tracing just-in-time compiler system
JP5148674B2 (ja) * 2010-09-27 2013-02-20 株式会社東芝 プログラム並列化装置およびプログラム
US9009689B2 (en) * 2010-11-09 2015-04-14 Intel Corporation Speculative compilation to generate advice messages
EP2649523B1 (fr) * 2010-12-06 2020-06-17 Google LLC Méthode de compilation d'un code intermédiaire d'une application
CN103339604B (zh) * 2011-01-31 2016-10-26 株式会社索思未来 程序生成装置、程序生成方法、处理器装置以及多处理器系统
US9529575B2 (en) 2012-02-16 2016-12-27 Microsoft Technology Licensing, Llc Rasterization of compute shaders
US9971576B2 (en) * 2013-11-20 2018-05-15 Nvidia Corporation Software development environment and method of compiling integrated source code
US9329844B2 (en) * 2014-05-30 2016-05-03 Apple Inc. Programming system and language for application development
US10241763B2 (en) * 2014-12-10 2019-03-26 Microsoft Technology Licensing, Llc. Inter-procedural type propagation for devirtualization

Patent Citations (10)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US5815719A (en) * 1996-05-07 1998-09-29 Sun Microsystems, Inc. Method and apparatus for easy insertion of assembler code for optimization
US7228531B1 (en) * 2003-02-03 2007-06-05 Altera Corporation Methods and apparatus for optimizing a processor core on a programmable chip
US20080256521A1 (en) * 2006-01-25 2008-10-16 International Business Machines Corporation Apparatus and Method for Partitioning Programs Between a General Purpose Core and One or More Accelerators
CN101078994A (zh) * 2006-05-26 2007-11-28 松下电器产业株式会社 编译器装置、编译器方法和编译器程序
CN101563673A (zh) * 2006-12-14 2009-10-21 富士通株式会社 编译方法以及编译器
US8468510B1 (en) * 2008-01-16 2013-06-18 Xilinx, Inc. Optimization of cache architecture generated from a high-level language description
CN103119561A (zh) * 2010-09-23 2013-05-22 苹果公司 用于对非叶代码进行基于编译器的矢量化的系统和方法
CN103282878A (zh) * 2010-12-24 2013-09-04 英特尔公司 基于循环分割或索引阵列的循环并行化
US20140196016A1 (en) * 2013-01-07 2014-07-10 Advanced Micro Devices, Inc. Layered programming for heterogeneous devices
CN103838614A (zh) * 2014-02-19 2014-06-04 华为技术有限公司 一种数据处理方法及装置

Non-Patent Citations (3)

* Cited by examiner, † Cited by third party
Title
ANNA DEREZINSKA 等: "Object-Oriented Mutation Applied in Common Intermediate Language Programs Originated from C#", 《2011 IEEE FOURTH INTERNATIONAL CONFERENCE ON SOFTWARE TESTING, VERIFICATION AND VALIDATION WORKSHOPS》 *
甘玲 等: "基于分层思想的变量类型提取方法", 《重庆邮电大学学报(自然科学版)》 *
高伟: "面向SIMD的自动向量化优化技术研究", 《中国优秀硕士学位论文全文数据库 信息科技辑》 *

Cited By (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN109144722A (zh) * 2018-07-20 2019-01-04 上海研鸥信息科技有限公司 一种多应用高效共用fpga资源的管理系统及方法
CN109144722B (zh) * 2018-07-20 2020-11-24 上海研鸥信息科技有限公司 一种多应用高效共用fpga资源的管理系统及方法
WO2021174538A1 (zh) * 2020-03-06 2021-09-10 深圳市欢太科技有限公司 应用处理方法及相关装置

Also Published As

Publication number Publication date
KR20180017020A (ko) 2018-02-20
JP2018517986A (ja) 2018-07-05
US20160364216A1 (en) 2016-12-15
CN107710150B (zh) 2021-07-16
EP3308267A1 (en) 2018-04-18
US9830134B2 (en) 2017-11-28
WO2016204865A1 (en) 2016-12-22

Similar Documents

Publication Publication Date Title
CN107710150A (zh) 从包含层次子例程信息的中间代码产生目标代码
KR101590734B1 (ko) 그래픽 프로세싱 유닛에서의 메모리 공유
US11694299B2 (en) Methods and apparatus to emulate graphics processing unit instructions
CN106462375B (zh) 图形管线状态对象和模型
CN109871936A (zh) 用于处理神经网络中的卷积运算的方法和装置
EP2805232B1 (en) Predication of control flow instructions having associated texture load instructions for a graphics processing unit
CN106575228A (zh) 图形处理中的渲染目标命令重新排序
US10180825B2 (en) System and method for using ubershader variants without preprocessing macros
US20160225118A1 (en) Optimizing compilation of shaders
CN104137070A (zh) 用于异构cpu-gpu计算的执行模型
US9442706B2 (en) Combining compute tasks for a graphics processing unit
US20120151459A1 (en) Nested communication operator
Grzmil et al. Performance analysis of native and cross-platform mobile applications
EP2988268B1 (en) Rendergraph compilation and use thereof for low-latency execution
Acosta et al. Towards a Unified Heterogeneous Development Model in Android TM
CN112445855B (zh) 用于图形处理器芯片的可视化分析方法和可视化分析装置
CN107918946A (zh) 执行指令的图形处理设备和方法
US10620980B2 (en) Techniques for native runtime of hypertext markup language graphics content
Acosta et al. Performance analysis of paralldroid generated programs
US11573777B2 (en) Method and apparatus for enabling autonomous acceleration of dataflow AI applications
CN111383164A (zh) 图形处理单元、计算系统及其操作方法
Acosta et al. Towards the optimal execution of Renderscript applications in Android devices
US9569202B2 (en) Transparent type-based object augmentation by orthogonal functionality

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
GR01 Patent grant
GR01 Patent grant
CF01 Termination of patent right due to non-payment of annual fee
CF01 Termination of patent right due to non-payment of annual fee

Granted publication date: 20210716