CN106164862A - 用于编译器优化的存储器参考元数据 - Google Patents

用于编译器优化的存储器参考元数据 Download PDF

Info

Publication number
CN106164862A
CN106164862A CN201580016245.6A CN201580016245A CN106164862A CN 106164862 A CN106164862 A CN 106164862A CN 201580016245 A CN201580016245 A CN 201580016245A CN 106164862 A CN106164862 A CN 106164862A
Authority
CN
China
Prior art keywords
memory
kernel
compilation device
operation time
compiler
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.)
Pending
Application number
CN201580016245.6A
Other languages
English (en)
Inventor
C·C·利姆
D·S·布拉克曼
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 CN106164862A publication Critical patent/CN106164862A/zh
Pending legal-status Critical Current

Links

Classifications

    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/44—Arrangements for executing specific programs
    • G06F9/455—Emulation; Interpretation; Software simulation, e.g. virtualisation or emulation of application or operating system execution engines
    • G06F9/45504—Abstract machines for programme code execution, e.g. Java virtual machine [JVM], interpreters, emulators
    • G06F9/45516—Runtime code conversion or optimisation
    • G06F9/45525—Optimisation or modification within the same instruction set architecture, e.g. HP Dynamo
    • 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
    • G06F8/43—Checking; Contextual analysis
    • G06F8/433—Dependency analysis; Data or control flow analysis
    • G06F8/434—Pointers; Aliasing
    • 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
    • G06F8/44—Encoding
    • G06F8/443—Optimisation
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/44—Arrangements for executing specific programs
    • G06F9/455—Emulation; Interpretation; Software simulation, e.g. virtualisation or emulation of application or operating system execution engines
    • G06F9/45504—Abstract machines for programme code execution, e.g. Java virtual machine [JVM], interpreters, emulators
    • G06F9/45516—Runtime code conversion or optimisation

Landscapes

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

Abstract

本发明涉及一种设备，其包含存储器和编译处理器，编译处理器经配置以：通过由在编译处理器上执行的编译器和运行时间组成的群组中的至少一者产生用于执行经编译内核的自变量；通过由在编译处理器上执行的群组中的至少一者确定对自变量的第一存储器区的第一存储器参考和对自变量的第二存储器区的第二存储器参考是否参考同一存储器区；通过群组中的至少一者基于确定而产生与第一存储器参考和第二存储器参考相关联的元数据，其中元数据指示第一存储器区与第二存储器区之间的关系。编译器和运行时间中的至少一者可基于元数据而重新编译内核，且指示目标处理器执行经重新编译的内核。

Description

用于编译器优化的存储器参考元数据
技术领域
本发明涉及编译内核的源代码，且更具体地说涉及用于编译用于存储器存取的内核的源代码的技术。
背景技术
已存在朝向所谓的异构计算架构的趋势。在异构计算架构中，称为内核的程序可使用框架来进行编译，使得例如CPU(中央处理单元)、GPU(图形处理单元)、FPGA(现场可编程门阵列)等多种不同类型的处理器可执行所述内核。最近的支持异构计算的框架包含OpenCL框架以及DirectCompute框架。
发明内容
本发明描述用于检测内核中的存储器参考的存储器混叠和存储器重叠以便产生用于编译优化的元数据的技术。为了执行本发明的技术，例如及时编译器(JIT)等编译器将程序的源代码(也被称作“内核”)编译为二进制文件。执行编译器的编译处理器可在运行时间(当编译处理器产生执行内核所需要的自变量时)使用例如OpenCL等异构计算框架编译所述内核。在本发明中描述的技术中，驱动程序分析在缓冲器中一起传递的将传递到将执行内核的目标处理器的自变量，而不是指示目标处理器使用所产生自变量执行内核。基于所述分析，驱动程序/运行时间产生指示第一存储器参考与第二存储器参考之间的关系(例如，第一存储器参考和第二存储器参考的存储器区是否重叠、重叠达何种范围等)的元数据。
如果存储器区是不相同的，那么编译处理器可使用编译器基于元数据且使用例如循环展开等更积极的编译技术而重新编译内核。驱动程序也可以能够确定内核的存储器存取重叠达何种范围，且可基于存储器重叠量使用更积极的技术重新编译内核。以此方式，本发明的技术可增加使用及时编译器编译的内核的执行性能。
在一个实例中，本发明描述一种方法，其包括：通过由在编译处理器上执行的编译器和运行时间组成的群组中的至少一者产生用于执行经编译内核的二进制代码的自变量；通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定对所述内核自变量的第一存储器区的第一存储器参考和对所述内核自变量的第二存储器区的第二存储器参考是否参考同一存储器区；通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述确定而产生与所述第一存储器参考和所述第二存储器参考相关联的元数据。所述元数据可指示所述第一存储器区与所述第二存储器区之间的关系。所述方法进一步包含响应于通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的第一和第二存储器参考并不参考所述同一存储器区而进行以下操作：通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者致使编译器基于所述元数据而重新编译所述内核；以及通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者指示目标处理器执行所述经重新编译的内核。
在另一实例中，本发明描述一种装置，其包含存储器和编译处理器，所述编译处理器经配置以：通过由在所述编译处理器上执行的编译器和运行时间组成的群组中的至少一者产生用于执行经编译内核的二进制代码的自变量；通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定对所述内核自变量的第一存储器区的第一存储器参考和对所述内核自变量的第二存储器区的第二存储器参考是否参考同一存储器区；通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述确定而产生与所述第一存储器参考和所述第二存储器参考相关联的元数据。所述元数据可指示所述第一存储器区与所述第二存储器区之间的关系，且响应于通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的第一和第二存储器参考并不参考所述同一存储器区，所述编译处理器进一步经配置以：通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者致使编译器基于所述元数据而重新编译所述内核；以及通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者指示目标处理器执行所述经重新编译的内核。
在另一实例中，本发明描述一种存储指令的非暂时性计算机可读存储媒体，所述指令当执行时致使编译处理器进行以下操作：通过由在所述编译处理器上执行的编译器和运行时间组成的群组中的至少一者产生用于执行经编译内核的二进制代码的自变量；通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定对所述内核自变量的第一存储器区的第一存储器参考和对所述内核自变量的第二存储器区的第二存储器参考是否参考同一存储器区；通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述确定而产生与所述第一存储器参考和所述第二存储器参考相关联的元数据。所述元数据指示第一存储器区与第二存储器区之间的关系，且
响应于通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的第一和第二存储器参考并不参考所述同一存储器区，所述编译处理器可进一步经配置以执行致使所述编译处理器进行以下操作的指令：通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者致使编译器基于所述元数据而重新编译所述内核；以及通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者指示目标处理器执行所述经重新编译的内核。
在附图和下文描述中陈述本发明的一或多个实例的细节。本发明的其它特征、目标和优点将从所述描述和图式以及权利要求书而显而易见。
附图说明
图1是说明根据本发明的技术的支持混叠分析以辅助编译优化的实例计算装置的框图。
图2是说明根据本发明的技术的可执行内核的处理器的一或多个着色器核心的多个处理元件的概念图。
图3A是说明根据本发明的技术的包含当执行时可造成混叠的代码的内核代码的概念图。
图3B是说明根据本发明的技术配置的编译器可能够检测的混叠的实例的概念图。
图3C是说明根据本发明的技术配置的编译器可能够检测的不重叠存储器参考的实例的概念图。
图3D是说明根据本发明的技术配置的驱动程序/运行时间可检测的重叠存储器参考的概念图。
图4A是说明根据本发明的技术的循环展开的概念图。
图4B是说明根据本发明的技术的代码重排序的概念图。
图4C是说明根据本发明的技术的代码向量化的概念图。
图5是根据本发明的技术用于产生编译器元数据以辅助编译器优化的实例方法的流程图。
具体实施方式
如上文简要地描述，当前在开发各种异构计算框架。异构计算框架的一些实例包含当前由Khronos集团开发的OpenCLTM框架以及当前由开发的DirectCompute框架。异构计算框架允许单个程序或“内核”在多种不同处理器上执行，例如CPU(中央处理单元)、GPU(图形处理单元)、FPGA(现场可编程门阵列)、DSP(数字信号处理器)等。
为了准备内核用于执行，在本发明中称为编译处理器的处理器编译内核源代码以产生将由目标处理器执行的二进制代码。目标处理器可为同一处理器或不同于目标处理器。编译处理器使用的编译器的一个实例称为及时编译(JIT)编译器。JIT编译器在执行时间(也被称作运行时间)编译源代码，而不是在执行之前编译(有时称为“提前”编译)或根本不需要先前编译的指令(称为“解译”)。
一旦内核已经编译，编译处理器便经由驱动程序和运行时间将内核的经编译二进制代码转移到目标处理器。内核还在运行时间接受自变量集合以用于在目标处理器上执行所述内核，编译处理器也将所述自变量集合转移到目标处理器。内核自变量包括缓冲器，即为自变量分配的存储器的区域。在大多数情况下，内核包含对自变量操作(即，对其读取或从其写入)的代码区段。以此方式，自变量包括内核可对其操作的用于内核的数据集。在将内核转移到目标处理器之后，编译处理器的驱动程序/运行时间执行函数调用，其在一些实例中在运行时间将自变量提供到内核。一旦内核已接收自变量，目标处理器便可开始内核的执行。
在许多情况下，内核包含代码段，例如循环，其执行直到目标处理器确定一些布尔型条件已满足或达到某一次数的迭代。编译器可能够采用各种技术以改善执行循环代码区段的性能，例如循环展开以及其它技术，例如代码重排序，和/或可改善循环和非循环代码区段的执行的向量化。
循环展开是一种优化过程，编译器通过所述过程扩展循环的若干迭代以减少或消除控制循环的指令，例如算术操作、循环结束测试，和/或当执行循环时改善高速缓冲存储器执行性能。代码重排序是另一优化，编译器可使用其来分组一系列类似指令(例如，一起加载或存储)。代码重排序可当在某些情况下执行循环代码区段时改善高速缓冲存储器性能。举例来说，当一起聚结许多加载指令(例如，在循环主体内)可改善具有高速缓冲存储器线宽的系统(下文更详细地论述)上的性能时，代码重排序可改善性能，所述高速缓冲存储器线宽是在标量指令中使用的操作数的大小的倍数。然而，如果编译器在编译之前确定加载/存储缓冲器并不彼此混叠，那么编译器聚结负载可仅为安全的。否则，由于经重排序的加载/存储指令，可发生数据损坏。
向量化是另一优化过程，编译器通过所述过程可将包含若干标量操作的源代码转换为向量指令，所述若干标量操作中的每一者每次处理单个对的操作数，所述向量指令每次处理多对操作数上的一个操作。向量化是一种形式的并行度，其可改善相对于同一代码的标量实施方案的性能。下文更详细地描述循环展开、代码重排序以及向量化。
内核的代码区段可含有存储器参考，也被称作“指标”，其可参考自变量的存储器区域。举例来说，代码区段可包含可参考内核自变量的部分的一系列存储器参考(即，对包含在内核自变量中的缓冲器的存储器参考)。内核可从自变量缓冲器读取值，且还可将数据写入到自变量缓冲器。
在一些情况下，不同存储器参考(例如具有不同名称的指针变量)可参考存储器中的同一数据位置。其中不同符号参考参考同一存储器区的情形称为“混叠”。编译器可尝试在编译时间使用静态分析或其它技术检测混叠。然而，当循环代码区段中的存储器参考所参考的数据(例如，内核自变量)是在运行时间供应时，编译器通常不能够检测循环中的存储器参考的混叠。
当编译器不能够明确地确定存储器参考是否参考同一存储器区(即存储器参考导致混叠)时，编译器可能不能够对循环执行优化技术，例如循环展开和向量化。本发明的技术可使得JIT编译器能够确定内核循环的存储器存取是否参考同一存储器区。另外，本发明的技术使得JIT编译器能够产生关于存储器参考之间的关系的元数据，且基于所产生元数据使用例如向量化和循环展开等优化来重新编译内核。
图1是说明根据本发明的技术的支持混叠分析以辅助编译优化的实例计算装置的框图。图1包含计算装置2。计算装置2可包括个人计算机、桌上型计算机、膝上型计算机、计算机工作站、平板计算装置、视频游戏平台或控制台、无线通信装置(例如，移动电话、蜂窝式电话、卫星电话及/或移动电话手持机)、手持式装置(例如，便携式视频游戏装置或个人数字助理(PDA))、个人音乐播放器、视频播放器、显示装置、电视、电视机顶盒、服务器、中间网络装置、主机计算机，或处理及/或显示图形数据的任何其它类型的装置。
如图1的实例中所说明，计算装置2包含CPU 16、系统存储器14、图形处理单元(GPU)12、及时(JIT)编译器18以及驱动程序/运行时间19。CPU 16可执行各种类型的应用程序。应用程序的实例包含网络浏览器、电子邮件应用程序、电子数据表、视频游戏、产生用于显示的可观看对象的应用程序及类似应用程序。用于执行一或多个应用程序的指令可存储在系统存储器14内。
CPU 16还可执行JIT编译器18。因此，CPU 16可出于实例的目的称为“编译处理器”。JIT编译器18包括当由CPU 16执行时可使用如上文所描述的例如OpenCL或DirectCompute等异构计算框架编译内核的源代码的编译器。JIT编译器18将源代码编译为原生代码或中间代码(例如，字节代码)用于由目标处理器执行。JIT编译器18在“运行时间”执行编译，即在执行时而不是在执行之前执行编译。JIT编译器18可当使用OpenCL编译时使用clBuildProgram()函数执行编译。另外，JIT编译器18可经配置以分析内核20的数据存取模式以确定在目标处理器GPU 12上执行的某些纤维(即线程)的数据存取是否是独立的以及其它条件是否保持。
驱动程序/运行时间19还与JIT编译器18交互以将内核源代码翻译为二进制指令或字节代码指令。驱动程序/运行时间19可使用驱动程序来执行内核源代码指令到用于目标处理器(在此实例中为GPU 12)的原生或目标代码的架构特定编译。举例来说，驱动程序/运行时间19可知道可用于目标处理器的特定向量指令或执行资源，且可以优化目标处理器上的执行性能的方式将源代码编译为原生代码。在一些实例中，例如如果存在多个目标处理器，例如如果内核将在CPU 16和GPU 12上执行，那么可存在不同驱动程序。
内核20由目标处理器(在此实例中为GPU 12)能够执行的原生或目标代码构成，例如二进制指令。JIT编译器18还可管理GPU 12的运行时间执行。CPU 16可将内核20传输到GPU 12用于执行。CPU 16还可产生自变量26，CPU 16可将其转移到GPU 12用于进一步处理。
在分配自变量26之前，CPU 16为自变量26分配自由存储器缓冲器，其为存储器的区。一旦缓冲器已经被分配，驱动程序/运行时间19便在缓冲器中存储自变量26。自变量26可包括GPU 12能够处理的多个数据值(例如，整数、浮点值、对象、值的阵列等)。另外，在内核20的执行期间，GPU 12可将数据写入到存储自变量26的缓冲器作为输出。输出的数据可包括GPU 12可转移返回至CPU 16的输出自变量。
CPU 16转移到GPU 12的自变量可被称为“输入自变量”。在其中异构计算框架是OpenCL框架的实例中，驱动程序/运行时间19可产生自变量且在运行时间传递(使得可用)到clSetKernelArg()函数。clSetKernelArg()函数接收内核20作为自变量以及内核自变量26中的任一者，且将自变量转移到GPU 12以使得执行可开始。
作为为自变量26分配存储器的部分，驱动程序/运行时间19确定与包含在内核中的存储器参考中的一些或全部相关联的自变量26的地址和存储器区。存储器参考可为特定代码区段的存储器参考，例如包含循环的代码区段，称为“循环代码区段”。基于所确定的存储器区，驱动程序/运行时间19可能够解析(即，确定)内核20的循环代码区段或其它代码区段的存储器参考是否参考自变量26的同一存储器区。
响应于产生内核自变量26用于GPU 12执行内核20，驱动程序/运行时间19可执行内核20。更确切地说，驱动程序/运行时间19可使用clEnqueueNDRangeKernel()函数将内核20分派到目标处理器。在运行时间，驱动程序/运行时间19分析自变量26，内核20接收所述自变量。驱动程序/运行时间19还分析存储器参考(例如，指针)等以确定存储器参考是否参考为自变量26分配的存储器区的同一存储器区。驱动程序/运行时间19可以成对方式分析存储器参考和自变量缓冲器以确定存储器参考是否参考同一存储器区。
驱动程序/运行时间19进一步基于存储器参考所参考的自变量26的存储器区之间的关系产生与存储器参考相关联的元数据。所述元数据可指示存储器参考之间的关系。举例来说，作为一些非限制性实例，所述元数据可包含重叠存储器参考的列表、与存储器区相关联的存储器区是否重叠、存储器区重叠的范围，以及所述重叠包括多少字节。
驱动程序/运行时间19将所产生元数据(如果存在)提供回到JIT编译器18。响应于基于元数据确定两个存储器参考并不共享确切同一存储器区，驱动程序/运行时间19可致使JIT编译器18可使用各种优化来重新编译内核20，例如循环展开、代码重排序和/或向量化。JIT编译器18可基于所产生元数据应用循环展开代码重排序和/或向量化的这些各种优化。
根据本发明的技术，编译处理器(例如CPU 16)可经配置以使用在编译处理器上执行的由JIT编译器18和驱动程序/运行时间19组成的群组中的至少一者产生用于执行经编译内核20的代码(例如，二进制代码或目标代码)的自变量26。由JIT编译器18和驱动程序/运行时间19组成的群组中的所述至少一者可进一步经配置以确定对内核自变量的第一存储器区的第一存储器参考和对内核自变量的第二存储器区的第二存储器参考是否参考同一存储器区。响应于通过由JIT编译器18和驱动程序/运行时间19组成的群组中的所述至少一者确定所述内核的第一和第二存储器参考并不参考同一存储器区，CPU 16可进一步经配置以：借助在CPU 16上执行的由JIT编译器18和驱动程序/运行时间19组成的群组中的所述至少一者致使由JIT编译器18和驱动程序/运行时间19组成的群组中的所述至少一者基于元数据重新编译内核20，且通过在CPU 16上执行的由JIT编译器18和驱动程序/运行时间19组成的群组中的所述至少一者指示目标处理器(例如GPU 12)执行经重新编译的内核20。
GPU 12可为允许大规模并行处理的专用硬件，所GPU 12可为允许大规模并行处理的专用硬件，所述硬件对于处理图形数据非常适合。以此方式，CPU 16卸载由GPU 12更好地处置的图形处理。CPU 16可根据特定应用处理接口(API)或异构计算框架与GPU 12通信。此些API的实例包含的API以及Khronos集团的异构计算框架的实例包含微软的DirectCompute、Khronos集团的OpenCLTM。然而，本发明的各方面不限于上述API和框架，且可延伸到其它类型的API。
CPU 16及GPU 12的实例包含(但不限于)数字信号处理器(DSP)、通用微处理器、专用集成电路(ASIC)、现场可编程逻辑阵列(FPGA)或其它等效集成或离散逻辑电路。在一些实例中，GPU 12可为包含集成及/或离散逻辑电路的专用硬件，所述电路为GPU 12提供适合于图形处理的大规模并行处理能力。在一些情况下，GPU 12还可包含通用处理，且可被称作通用GPU(GPGPU)。本发明中描述的技术适用于GPU 12为GPGPU的实例。
系统存储器14可包括一或多个计算机可读存储媒体。系统存储器14的实例包含(但不限于)随机存取存储器(RAM)、只读存储器(ROM)、电可擦除可编程只读存储器(EEPROM)、快闪存储器或可用以载运或存储呈指令及/或数据结构形式的所要程序代码及可由计算机或处理器存取的任何其它媒体。
在一些方面中，系统存储器14可包含致使CPU 16及/或GPU 12执行在本发明中归于CPU 16及GPU 12的功能的指令。因此，系统存储器14可为包括使一或多个处理器(例如，CPU 16和GPU 12)执行各种功能的指令的计算机可读存储媒体。
在一些实例中，系统存储器14可被视为非暂时性储存媒体。术语“非暂时性”可指示存储媒体不是以载波或实例来体现，非暂时性存储媒体可存储可随时间改变和数据(例如，存储在RAM中)。
使用JIT编译器18和驱动程序/运行时间19，CPU 16可将源代码编译为用于GPGPU应用程序的原生代码(例如，命令和数据)或字节代码。实例GPGPU数据和命令包含用于射线跟踪应用、物理学模拟的命令和场景数据，或用于任何其它类型的GPGPU内核的数据。GPGPU应用程序(例如，内核20)还可使用图形API(例如DirectX或OpenGL)或使用更通用的计算API(例如开放计算语言(OpenCL)或OpenCompute或DirectCompute)来编译。CPU 16可将用于内核20的数据发射到命令缓冲器以供处理。在各种实例中，命令缓冲器可为系统存储器14的部分或GPU 12的部分。在一些实例中，CPU 16可经由专用总线(例如PCI-Express总线或另一通用串行或并行总线)将用于GPU 12的内核20的命令和数据发射到程序。
为执行命令缓冲器中存储的内核20的操作，GPU 12可实施处理管线。处理管线包含执行如由执行于GPU 12上的软件或固件定义的功能，及由经硬接线以执行极特定功能的固定功能单元执行功能。可有可能绕过用于内核20的执行的固定功能单元，或者内核20的执行可使用固定功能单元。
内核20可在GPU 12的一或多个处理元件(也被称作“着色器核心”或“PE”)上执行。着色器核心22为用户提供功能灵活性，因为用户可对着色器进行编程以用任何可设想方式执行所需任务，如同任何其它处理器。然而，固定功能单元经针对固定功能单元执行任务的方式而硬接线。因此，固定功能单元可不提供大量功能灵活性。本发明的技术是针对例如内核20等内核在GPU着色器核心22上的执行。
一旦CPU 16将与再现图形场景或执行内核相关联的数据和/或命令传输到命令缓冲器，GPU 12便通过GPU 12的管线开始命令的执行。GPU 12的调度器24产生线程，所述线程执行与内核相关联的工作的基本单元。调度器24将线程指派到着色器核心22的特定处理元件。
图2是说明根据本发明的技术的可执行内核的处理器的一或多个着色器核心的多个处理元件的概念图。图2说明GPU 12或CPU 16的部分。GPU 12包含多个处理元件42A-42N(PE 42)，其可执行内核(例如内核20)的部分。在一些实例中，可在PE 42上执行的内核20的部分可被称为“线程束”或“工作单元”。PE 42可为着色器核心22(图1)中的一或多者的一部分。线程束或工作单元可包括线程的群组，也被称作“纤维”，GPU调度器24可将其指派到多个处理元件(例如，PE 42)用于执行。图2的每一PE可包括单指令多数据(SIMD)单元，其能够在特定时间(例如，在同时以用于并行执行)对多个数据值执行单个指令，例如向量指令。PE42还可支持对单一数据值执行单一指令，例如对单一浮点值的单一操作。
图2还包含GPU 12的调度器指派PE 42以供执行的指令44。在一些实例中，指令44可存储于命令缓冲器中。指令44可包含每一PE经配置以经配置以执行的内核的指令集。程序计数器(PC)50指示PE 42中的一或多者将执行的当前指令。在指令结束在PE 42上执行之后，PC 50的值可递增到内核20的下一个指令的地址。图2还包含寄存器46。寄存器46A-46N(寄存器46)可为能够保持多个数据值或单个值的通用寄存器。寄存器46可被“储备”，即，可加载和存储用于特定PE的数据。作为一实例，寄存器46A可限于存储用于PE 42A的数据，且可不加载或存储用于其它PE的数据。寄存器46中的每一者可将数据供应到PE 42中的一者和/或供应来自其的数据，PE 42可随后处理所述数据。
PE 42、指令44、寄存器46、高速缓冲存储器48和PC 50可包括GPU 12的着色器核心22的核心或部分。在各种实例中，线程束40可包括着色器的部分，例如几何着色器、像素着色器和/或顶点着色器，其可为GPU 12的图形管线的部分或包括例如内核20等内核的部分。在一些实例中，GPU 12可将由线程束产生的结果馈送到管线的另一级中用于额外处理。
图2还包含高速缓冲存储器48。高速缓冲存储器48是存储频繁存取的指令和数据以用于在执行期间的快速检索和存储的小存储器。虽然说明为单个高速缓冲存储器，但高速缓冲存储器48可表示多个高速缓冲存储器层级和/或单独的高速缓冲存储器。如上文所描述，在内核20的执行期间，GPU 12检索位于由PC 50的值指示的地址的指令44中的一者。GPU 12随后致使PE 42执行存储在PC 50的所述地址处的指令，所述PC 50在一些实例中可为寄存器。
GPU 12检查高速缓冲存储器48以确定高速缓冲存储器48当前是否包含将执行的下一指令，而不是从系统存储器获取在PC 50的地址处的指令，这将是不必要地缓慢的。高速缓冲存储器48的存储指令的部分称为指令高速缓冲存储器(“I高速缓冲存储器”)。如果将执行的下一指令存储在高速缓冲存储器48中，称为“高速缓冲存储器中”，那么GPU 12加载且执行经高速缓冲存储的指令。如果将执行的下一指令不存储在高速缓冲存储器48中，称为“高速缓冲存储器未中”，那么GPU 12从某个较慢的存储器、例如从系统存储器14加载下一指令用于执行。
在需要存储在存储器地址处的数据值(例如，操作数)的指令执行(例如，相加、相乘、加载、存储等)期间，GPU 12首先确定所述操作数是否存储在寄存器(例如，寄存器46中的一者)内。如果所请求数据值不存储在寄存器46中，那么GPU 12尝试从高速缓冲存储器48的保持数据值的部分存取所述数据值，所述部分称为数据高速缓冲存储器(“d高速缓冲存储器”)。如果数据值存储在高速缓冲存储器48内，那么GPU 12从高速缓冲存储器48加载所请求数据值。否则，GPU 12必须从较慢的存储器(例如系统存储器14)加载所请求数据值。类似地，如果指令致使PE 42存储或修改回到存储器中的数据值，那么高速缓冲存储器48可存储所述值到高速缓冲存储器48以使得如果其被再次写入或读取，那么在所述数据值不存储在寄存器46中的一者的情况下从高速缓冲存储器48快速检索或对高速缓冲存储器48快速覆写所述数据值。
GPU 12以称为高速缓冲存储器“线”的固定大小块对高速缓冲存储器48和从高速缓冲存储器48转移数据。高速缓冲存储器48可具有存储数百或数千不同线的容量。每一线与特定存储器地址相关联，且可存储多个字节的数据。举例来说，作为一个实例，高速缓冲存储器48的每一线可存储64个字节的数据。存储在每一线中的字节的数目称为高速缓冲存储器“宽度”。在其中高速缓冲存储器48具有可存储64个字节的数据的线的实例中，高速缓冲存储器48的高速缓冲存储器宽度是64字节。高速缓冲存储器宽度可影响代码重排序优化技术的性能，如下文将更详细地论述。
在从高速缓冲存储器48检索数据的加载操作期间，GPU 12可将检索的高速缓冲存储器数据加载到寄存器46中的一或多者中或未描绘的其它寄存器中。在指令的执行期间，PE 42可从寄存器46读取一或多个数据值。PE 42可对所述数据值执行一或多个操作，且将新值存储回到寄存器46。PE 42可执行流控制指令，例如分支、跳转、去往等。然而因为存在单一PC 50，所以PE 42可在给定时间在一个特定处仅执行由PC 50指示的指令44中的一者。
例如GPU 12的处理器可具有广泛量的向量寄存器和向量指令。因此，可使用例如向量化等优化来编译应用程序的例如JIT编译器18等编译器可增加支持向量指令或具有SIMD架构的处理器(例如GPU 12)的处理量或执行性能。
更确切地说，GPU 12可包含类似于图2中说明的那些着色器核心的数百或数千个着色器核心。每一着色器核心可能够执行向量指令。执行具有多个操作数的向量指令可极大地改善相对于含有标量指令而不是向量指令的未经优化代码的性能。此外，执行性能增加可在具有能够执行向量指令的较大数目的SIMD核心的架构上较大，因为较多通用处理器可具有有限数目的能够执行向量指令的寄存器和/或核心。
图3A是说明根据本发明的技术的包含当执行时可造成混叠的代码的内核代码的概念图。图3A的实例包含内核代码80。内核代码80包含行82、84、86和88。
内核代码80的行82是compute_output函数。行82的compute_output函数是当内核开始执行时目标处理器(例如，GPU 12)调用的函数。其大致等效于C编程语言中的“intmain()”函数，因为compute_output函数是驱动程序/运行时间19用来开始内核20的执行的程序进入点。如果目标处理器是CPU 16，那么C运行时间库可包括驱动程序/运行时间19的运行时间组件。如果GPU 12是目标处理器，那么驱动程序/运行时间19的驱动程序组件可包括运行时间。compute_ouput函数包含四个输入自变量：(1)inputImage，(2)global_cdf，(3)outputImage，和(4)local_cdf。inputImage是到输入自变量的缓冲器的指针。outputImage是当内核结束执行时到缓冲器的指针，将包含输出自变量。自变量global_cdf和local_cdf是到值的阵列的指针。行84可表示当执行时致使GPU 12分配且初始化变量的多个语句。作为一实例，执行行84可致使PE 42初始化且加载inputImage[i]等的值。
行86是循环初始化语句。循环初始化语句指示循环重复固定数目的迭代。循环以等于变量“start_offset”的起始索引i开始迭代，且当每一迭代结束执行时将i递增一。在每一循环迭代的完成时，GPU 12检查以查看布尔型条件“i<final_offset”是否仍为真。当i的值等于或大于值“final_offset”时，GPU 12停止执行循环。
在每一循环迭代内，GPU 12将outputImage的值设定于索引i，表示为等于local_cdf[inputImage[i]]的值的outputImage[i]。Local_cdf是阵列，其在此实例中以inputImage[i]的值作索引。inputImage[i]又以GPU 12随着每一循环迭代而递增的变量i作索引。
如上文所论述，outputImage和inputImage都是存储器参考。有可能到outputImage和inputImage的指针可指代存储器中的同一区(即outputImage和inputImage混叠或部分地混叠)。也可能outputImage和inputImage可指代存储器中的不同区或重叠区(即outputImage和inputImage并不混叠)。如果JIT编译器18不能够确定inputImage和outputImage是否并不混叠(即并不参考确切同一存储器区)，那么编译器可不能够使用某些编译器优化，例如向量化、代码重排序和/或循环展开。
图3B是说明根据本发明的技术配置的驱动程序/运行时间可能够检测的混叠的实例的概念图。然而，编译器可能不能够针对混叠进行优化。图3B的实例说明GPU 12可在存储器中存储的缓冲器100。出于实例的目的，来自图3B的指针outputImage和inputImage可参考缓冲器100的部分。在图3B的实例中，缓冲器100开始于存储器地址0x800(十六进制)。
在此实例中，inputImage和outputImage都参考存储在缓冲器100内的单个条目(例如，单个对象、变量等)。即，在此实例中，inputImage和outputImage混叠到确切同一存储器区，其以交叉散列指示。驱动程序/运行时间可能够检测到inputImage和outputImage参考同一存储器区。因为inputImage和outputImage参考同一存储器区，所以JIT编译器18不能够执行优化，例如循环展开和/或向量化。
响应于如图3B中所说明检测到两个存储器参考参考同一存储器区，驱动程序/运行时间19可不产生任何元数据。另外，JIT编译器18可不重新编译内核20，因为JIT编译器可针对图3C到3D中说明的情况来进行。因此，JIT编译器18可不执行也如图3C到3D中说明的代码优化中的任一者。
图3C是说明根据本发明的技术配置的驱动程序/运行时间可能够检测的不重叠存储器参考的实例的概念图。图3C说明作为与图3B中所说明的同一缓冲器的缓冲器120。缓冲器120类似地与图3B的缓冲器100开始于同一存储器地址0x800。
在图3C中，inputImage和outputImage是参考缓冲器120的两个不同存储器区的存储器参考。inputImage参考的存储器区由水平散列指示。outputImage参考的存储器区由垂直散列指示。在内核代码80且更具体地说行86和88的执行之前，JIT编译器18不管i的值如何都确定在循环的同一迭代期间inputImage[i]和outputImage[i]将不参考同一存储器区。
在运行时间期间，驱动程序/运行时间19可能够基于inputImage[i]和outputImage[i]的初始值且基于inputImage[i]和outputImage[i]的存储器地址在迭代通过循环86的过程中并不收敛的事实而确定inputImage[i]和outputImage[i]并不参考同一存储器区。换句话说，inputImage和outputImage的所参考索引始终由GPU 12单调增加的同一索引值i所参考。
响应于确定存储器参考inputImage和outputImage并不参考同一存储器区，驱动程序可产生指示inputImage与outputImage之间的关系的元数据。作为一实例，元数据可指示与inputImage和outputImage相关联的存储器区不重叠，且通过两个条目分离。元数据还可指示与inputImage和ouptutImage相关联的区的大小，以及inputImage与outputImage之间的字节数目。在产生元数据之后，JIT编译器18可从驱动程序/运行时间19接收元数据，且通过应用各种优化基于元数据重新编译内核20，如下文更详细描述。
图3D是说明根据本发明的技术配置的驱动程序/运行时间可检测的重叠存储器参考的概念图。图3D包含缓冲器130，其可为自变量缓冲器，例如自变量26(图1)。在此实例中缓冲器130开始于地址0x800。缓冲器130包含多个数据值，其说明为缓冲器130的包含矩形内的单独矩形。
如先前实例中，inputImage和outputImage是存储器参考，其参考缓冲器130的区。在此实例中，inputImage和outputImage参考的区重叠，但不是完全重叠。仅与inputImage相关联的存储器区以水平散列矩形指示。仅与outputImage相关联的存储器区以垂直散列矩形指示。由inputImage和outputImage两者参考的重叠存储器区以交叉影线矩形指示。
在运行时间，驱动程序确定inputImage和outputImage存储器参考是否参考同一存储器区。在此实例中，inputImage和outputImage重叠，但并不参考同一存储器区。驱动程序/运行时间19检测到inputImage和outputImage重叠但不相同，且产生用于JIT编译器18的元数据。所述元数据可指示关于与inputImage和outputImage相关联的区的信息，例如每一区的起始和结束地址。所述元数据可进一步包含关于重叠区的信息，例如重叠区的大小以及重叠区的起始和/或结束地址。JIT编译器18接收由驱动程序/运行时间19产生的元数据，且可通过应用根据本发明的优化技术而重新编译内核20。
图4A是说明根据本发明的技术的循环展开的概念图。图4A包含代码区段140，其一般对应于图3A中说明的内核代码80。在图4A的实例中，驱动程序/运行时间19和/或JIT编译器18可能已确定存储器参考inputImage和outputImage并不参考同一存储器区，如图3C和3D中所说明。因为inputImage和outputImage并不参考同一存储器区，所以JIT编译器18已对内核代码80执行循环展开。行142到150说明将一个迭代展开为四个迭代的结果。
图3A的行86和88说明执行单个迭代且在每一迭代之后将变量i递增一，而行142的展开循环在每一迭代之后将i递增四。行144将local_cdf[inputImage[i]]的值指派给outputImage[i]。行146将local_cdf[inputImage[i+1]]的值指派给outputImage[i+1]。行148将local_cdf[inputImage[i+2]]的值指派给outputImage[i+2]，且行150将local_cdf[inputImage[i+3]]的值指派给outputImage[i+3]。行144到150的结果是将local_cdf[inputImage[i+x]]的输出指派给outputImage[i+x]的对应值，其中x[0…3]。因此，当执行时，行142到150中说明的展开循环代码区段具有与图3A的行86到88的四个迭代相同的效果。
代码区段140的循环展开可具有相对于图3A的循环代码区段80的若干益处。第一优点是通过一个接一个地对指派中的每一者排序，JIT编译器18和/或驱动程序/运行时间19可能够实现相对于未经排序代码区段的在目标处理器(例如GPU 12)上的更好的高速缓冲存储器性能。
举例来说，在执行行144之后，GPU 12可能已在高速缓冲存储器(例如高速缓冲存储器48)中存储与inputImage和outputImage相关联的存储器区的数据中的一些或全部。如果执行指令所需要的数据不存储在寄存器(例如寄存器46)中，那么可能需要从高速缓冲存储器(例如高速缓冲存储器48)存取数据。更确切地说，GPU 12可在高速缓冲存储器48中存储inputImage和outputImage的条目，例如inputImage[i+1]、[i+2]等，以及outputImage[i+1]、[i+2]等。如果inputImage和outputImage的条目存储在GPU 12的高速缓冲存储器中，那么GPU 12可能够从高速缓冲存储器快速存取行144到150的inputImage和outputImage的参考索引的数据，而不是从较慢的存储器存取所述参考索引。
另外，当代码区段140展开时，inputImage[i，i+1，i+2…]和outputImage[i，i+1，等]的值可存储在单个高速缓冲存储器线中。相比之下当不展开时，inputImage和outputImage[i]的值可存储在不同高速缓冲存储器线中。可由循环展开引起的单个高速缓冲存储器读取中从单个高速缓冲存储器线检索inputImage的全部值相对于当执行展开代码时可导致的执行多次高速缓冲存储器读取来说可更快。
从GPU 12的高速缓冲存储器存取数据而不是从较慢的系统存储器(例如系统存储器14)存取数据可相对于行86到88增加执行行142到150的循环的性能。在一些实例中，GPU12也可以能够例如在支持超标量执行的处理器或假定行144到150之间不存在相依性的SIMD处理器上并行地执行行144到150，其中inputImage或outputImage的值取决于内核20中先前计算的值。
除改善高速缓冲存储器性能之外，如图4A的代码区段140中所说明的循环展开也减少GPU 12评估与循环相关联的布尔型条件的次数，以及GPU 12在完成每一循环迭代之后执行的跳转的次数。与图3A的代码区段80相比，在评估行142的布尔型条件“i<final_offset”是否为真之前，行142到150的代码每迭代执行四行。代码区段80相比之下在评估行82的布尔型条件是否为真之前仅执行一行。因此，GPU 12评估行142的布尔型条件的次数相对于代码区段80减少。
在GPU 12完成行142到150的循环的迭代之后，且如果GPU 12确定布尔型条件“i<final_offset”仍为真，那么GPU 12从行150跳转回到行144。在代码区段140中，GPU 12在执行四行之后执行跳转。当执行代码区段80时，GPU 12在每一迭代之后跳转。因此，相对于代码区段80，代码区段140的展开代码减少了布尔型条件的评估以及GPU 12执行的跳转的数目，其可改善执行代码区段140的执行性能。
图4B是说明根据本发明的技术的代码重排序的概念图。图4B包含代码区段160，其进一步包含行162、164、166、168、170和172。如上文所论述，驱动程序/运行时间19和/或JIT编译器18可确定对存储器的参考是否混叠到存储器的同一区。如上文相对于图4A所论述，JIT编译器18可响应于从驱动程序/运行时间19接收到确定特定代码区段中不存在存储器混叠的元数据而执行某些优化，例如图4A中说明的循环展开。
JIT编译器18和/或驱动程序/编译器19可响应于确定特定代码区段中的存储器参考并不参考同一存储器区而执行的另一优化是图4B说明的代码重排序。代码160可一般对应于图4B的展开代码的经重排序汇编语言表示。JIT编译器18和/或驱动程序/运行时间19可将代码重排序应用于非循环以及循环代码区段。在图4B中，JIT编译器18已经重排序图4A的加载和存储以使得全部加载和存储分组在一起。
行162和164是JIT编译器18和/或驱动程序/编译器19已经分组在一起的加载指令。在图4A中，例如行144的行包含多个加载和存储指令。举例来说，为了执行行144，JIT编译器18可产生三个单独指令。第一指令可为加载指令，其将来自inputImage[i]参考的存储器位置的值加载到表示为r0的寄存器中。第二指令可为加载指令，其加载local_cdf[inputImage[i]]的值且将加载的值存储到同一寄存器r0中，进而覆写r0的先前值。包含在行144中的最终指令可为存储指令，其将来自r0的值存储到outputImage[i]参考的存储器中。
行162到172说明相对于包括行144到150的指令经重排序的加载和存储指令。在行162中，组合代码指示GPU 12将来自inputImage[i]参考的存储器区的值加载(使用加载指令“ldg”)到寄存器r0中。类似地，行164致使GPU 12将存储器参考inputImage[i+1]参考的值加载到寄存器r1中。可在行162和164之后但在行166之前发生且出于简洁起见未说明的后续指令可包含致使GPU 12将来自inputImage参考的存储器区的数据加载到寄存器中的额外加载指令。
在行166、168和为简洁起见未说明的其它行中，JIT编译器18已经将从缓冲器的加载local_cdf分组在一起。行166包含加载local_cdf[r0]的内容(即在索引r0处来自阵列local_cdf的存储器的内容)且将local_cdf[r0]的内容存储到寄存器r0中进而覆写r0的内容的加载指令。类似地，行168的指令致使GPU 12在由当前存储在寄存器r1中的值指示的索引处将存储器参考local_cdf参考的内容存储到寄存器r1中。因此，在执行时168的指令致使GPU 12覆写r1的先前值。在行168之后且在行170之前发生且出于简洁起见未说明的其它指令可类似地包含当执行时致使GPU 12从local_cdf[rx]加载数据的指令，其中x是某个整数。
作为重排序代码区段140的指令的部分，JIT编译器18还将存储指令分组在一起。作为此一实例，在重排序之后，JIT编译器18已将行170和172分组在一起行170包含在位置outputImage[i]将r0的内容存储到存储器中的存储指令。类似地，行172当执行时致使GPU12在outputImage[i+1]参考的位置将寄存器r1的值存储到存储器中。出于简洁起见未说明的其它指令当执行时可类似地致使GPU 12在位置outputImage[i+x]将寄存器(例如寄存器rx，其中x是整数)的值存储到存储器。
重排序加载和存储可相对于图3A的代码80改善执行代码160的性能。更确切地说，重排序加载和存储可在某些情况下取决于高速缓冲存储器线宽度而改善性能。举例来说，当将许多加载指令聚结在一起可改善具有作为标量指令中使用的操作数大小的倍数的高速缓冲存储器线宽的系统上的性能时，代码重排序可改善执行性能。
图4C是说明根据本发明的技术的代码向量化的概念图。图4C包含代码区段180，其进一步包含行182、184和186。如上文所论述，JIT编译器18和/或编译器/驱动程序19可响应于确定代码区段中的存储器参考并不参考同一存储器区而执行某些优化，例如图4A中说明的循环展开。JIT编译器18和/或编译器/驱动程序19经配置以基于来自驱动程序/运行时间19的包含与代码区段的存储器参考相关的信息的元数据而向量化循环代码区段。
向量化是其中编译器(例如，JIT编译器18)和/或驱动程序/运行时间19将各自具有单个操作数的多个标量指令组合为具有多个操作数的单个向量指令的过程。向量化是一种形式的并行度，其通过减少处理器需要执行以完成特定代码区段的指令数目以及通过利用固有硬件能力在系统存储器14与GPU 12之间移动数据而改善执行性能。在图4C的代码区段180的实例中，JIT编译器18可如图4B中所说明重排序加载和存储。一旦JIT编译器18已经重排序加载和存储，JIT编译器18便接着可向量化类似指令的群组，如图4C中所说明。
在行182中，JIT编译器18已经将多个加载(ldg)指令组合为单个向量化指令。当执行时，向量化指令将索引[i]到[i+3]处的inputImage加载到寄存器r0到r3。类似地，在行184中，JIT编译器18将行166、168等的多个加载指令组合为单个向量化加载指令，其将local_cdf[r0-r3]的值加载到寄存器r0到r3中。并且，在行186中，JIT编译器18已将行170到172的存储(“stg”指令)组合为单个向量化存储指令，其将寄存器r0到r3的值存储到outputImage[i]到outputImage[i+3]中。
为了如图4B和4C中所说明重排序或向量化指令，JIT编译器18和/或驱动程序/运行时间19必须注意任何相依性。相依性是产生语句或指令之间的执行次序约束的关系。作为一实例，如果S1必须在S2之前执行，那么存在语句S2对另一语句S1的相依性。为了确定相依性是否禁止向量化和/或代码重排序，JIT编译器18和/或驱动程序/运行时间19可基于从驱动程序/运行时间19获得的元数据在根据本发明的技术重排序或向量化代码之前执行相依性分析。
图5是说明根据本发明的技术用于产生编译器元数据以辅助编译器优化的实例方法的流程图。应大体上理解，图6的方法可通过由JIT编译器18和驱动程序/运行时间19执行编译处理器(例如，CPU 16)和目标处理器(例如，GPU 12)组成的群组中的至少一者执行。在一些实例中，目标处理器和编译处理器可为相同的。另外，可存在多于一个编译处理器和/或目标处理器。
在图5的方法中，编译处理器(例如CPU 16)使用驱动程序/运行时间19和/或JIT编译器18以产生自变量(例如，内核自变量26)用于执行经编译内核20的二进制代码或字节代码(200)。驱动程序/运行时间19和/或JIT编译器18进一步确定对内核自变量26的第一存储器区的第一存储器参考和对内核自变量26的第二存储器区的第二存储器参考是否参考内核自变量26的同一存储器区(202)，或如图3B、3C和3D中所说明的可能关系的其它实例。
CPU 16使用驱动程序/运行时间19和/或JIT编译器18以产生与第一存储器参考和第二存储器参考相关联的元数据(204)。所述元数据指示第一存储器区与第二存储器区之间的关系，例如第一存储器区与第二存储器区之间的重叠区。所述元数据可进一步包含第一和第二存储器区之间的重叠的字节数目。在一些实例中，所述元数据可包含存储器重叠的起始地址以及存储器重叠区的结束地址。应理解，相对于图5描述的实例仅出于实例的目的而参考单个对的存储器参考。驱动程序/运行时间19和/或JIT编译器18可针对内核自变量26的所有对的存储器参考导出元数据。
响应于使用驱动程序/运行时间19确定第一和第二存储器参考并不参考内核自变量26的同一存储器区，在CPU 16上执行的JIT编译器18可致使CPU 16基于所述元数据使用JIT编译器18重新编译内核20(206)。最后，目标处理器(例如GPU 12)可执行经重新编译的内核(210)。在一些实例中，驱动程序/运行时间19和/或JIT编译器18可基于元数据确定第一和第二存储器参考并不参考同一存储器区，且可使用此信息借助优化来重新编译内核20。
在一些额外实例中，为了确定内核20的第一存储器参考和第二存储器参考是否参考同一存储器区，CPU 16可使用驱动程序/运行时间19以确定内核20的包含第一和第二存储器参考的循环代码区段。并且，为了重新编译内核，JIT编译器18可基于由驱动程序/运行时间19和/或JIT编译器18产生的元数据而展开所述循环代码区段。为了重新编译内核，JIT编译器18还可基于所产生元数据而重排序所述循环代码区段的加载操作和存储操作和存储操作中的至少一者或将所述循环代码区段的多个标量指令向量化为至少一个向量指令。在各种实例中，JIT编译器18可使用例如微软DirectCompute和/或Khronos集团的OpenCL等异构框架来重新编译内核20。
本发明中所描述的技术可至少部分实施于硬件、软件、固件或其任何组合中。举例来说，所描述技术的各种方面可实施于一或多个处理器中，包含一或多个微处理器、数字信号处理器(DSP)、专用集成电路(ASIC)、现场可编程门阵列(FPGA)，或任何其它等效集成或离散逻辑电路，以及此等组件的任何组合。术语“处理器”或“处理电路”可一般指代前述逻辑电路中的任一者(单独或结合其它逻辑电路)或例如执行处理的离散硬件的任何其它等效电路。
此硬件、软件和固件可实施于同一装置内或单独装置内以支持本发明中所描述的各种操作和功能。另外，所描述的单元、模块或组件中的任一者可一起或单独作为离散但可互操作逻辑装置而实施。将不同特征描绘为模块或单元意图强调不同功能方面且未必暗示此类模块或单元必须由单独硬件或软件组件实现。而是，与一或多个模块或单元相关联的功能性可由单独硬件、固件及/或软件组件执行，或集成到共用或单独硬件或软件组件内。
本发明中所描述的技术也可存储、体现或编码于计算机可读媒体(例如，存储指令的计算机可读存储媒体)中。嵌入或编码于计算机可读媒体中的指令可致使一或多个处理器执行本文中所描述的技术(例如，当由一或多个处理器执行指令时)。计算机可读存储媒体可包含随机存取存储器(RAM)、只读存储器(ROM)、可编程只读存储(PROM)、可擦除可编程只读存储器(EPROM)、电可擦除可编程只读存储器(EEPROM)、快闪存储器、硬盘、CD-ROM、软盘、卡盒、磁性媒体、光学媒体或其它有形计算机可读存储媒体。
计算机可读媒体可包含计算机可读存储媒体，其对应于例如上文所列的有形存储媒体的有形存储媒体。计算机可读媒体也可包括通信媒体，其包含促进计算机程序从一个地点到另一地点的传送(例如，根据通信协议)的任何媒体。以此方式，短语“计算机可读媒体”大体上可对应于(1)非暂时性有形计算机可读存储媒体，和(2)例如暂时性信号或载波等非有形计算机可读通信媒体。
已描述各种方面和实例。然而，可在不脱离所附权利要求书的范围的情况下对本发明的结构或技术作出修改。

Claims (30)

1.一种编译用于执行的内核的方法，其包括：
通过由在编译处理器上执行的编译器和运行时间组成的群组中的至少一者产生用于执行经编译内核的二进制代码的自变量；
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定对所述内核自变量的第一存储器区的第一存储器参考和对所述内核自变量的第二存储器区的第二存储器参考是否参考同一存储器区；
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述确定而产生与所述第一存储器参考和所述第二存储器参考相关联的元数据，其中所述元数据指示所述第一存储器区与所述第二存储器区之间的关系；以及
响应于通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的第一和第二存储器参考并不参考所述同一存储器区而进行以下操作：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者致使编译器基于所述元数据而重新编译所述内核；以及
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者指示目标处理器执行所述经重新编译的内核。
2.根据权利要求1所述的方法，
其中确定所述内核的所述第一存储器参考和所述第二存储器参考是否参考所述同一存储器区进一步包括：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的包含所述第一和第二存储器参考的循环代码区段，
其中重新编译所述内核包括：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述产生的元数据而展开所述循环代码区段，且编译所述展开的循环代码区段。
3.根据权利要求1所述的方法，
其中确定所述内核的所述第一存储器参考和所述第二存储器参考是否参考所述同一存储器区进一步包括：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的包含所述第一和第二存储器参考的代码区段，
其中重新编译所述内核进一步包括：
响应于确定所述代码区段的所述第一和第二存储器参考并不参考所述同一存储器区，通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述产生的元数据而重排序所述代码区段的加载操作和存储操作中的至少一者。
4.根据权利要求1所述的方法，
其中确定所述内核的所述第一存储器参考和所述第二存储器参考是否参考所述同一存储器区进一步包括：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的包含所述第一和第二存储器参考的代码区段，
其中重新编译所述内核进一步包括：
响应于确定所述代码区段的所述第一和第二存储器参考并不参考所述同一存储器区，通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述产生的元数据而将所述代码区段的多个标量指令向量化为至少一个向量指令。
5.根据权利要求1所述的方法，其中所述元数据进一步指示所述第一存储器区与所述第二存储器区之间的重叠区。
6.根据权利要求5所述的方法，其中所述元数据包含所述第一存储器区与所述第二存储器区之间的重叠的字节数目。
7.根据权利要求5所述的方法，其中所述元数据进一步包括所述存储器重叠区的起始地址和所述存储器重叠区的结束地址中的至少一者。
8.根据权利要求1所述的方法，其中所述编译处理器包括中央处理单元CPU，且所述目标处理器包括图形处理单元GPU。
9.根据权利要求1所述的方法，其中所述编译器使用包括微软DirectCompute和OpenCL中的至少一者的异构计算框架来重新编译所述内核。
10.根据权利要求1所述的方法，其中所述内核自变量包括为所述自变量分配的存储器的缓冲器区域。
11.一种装置，其包括：
存储器；以及
编译处理器，其经配置以：
通过由在所述编译处理器上执行的编译器和运行时间组成的群组中的至少一者产生用于执行经编译内核的二进制代码的自变量；
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定对所述内核自变量的第一存储器区的第一存储器参考和对所述内核自变量的第二存储器区的第二存储器参考是否参考同一存储器区；
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述确定而产生与所述第一存储器参考和所述第二存储器参考相关联的元数据，其中所述元数据指示所述第一存储器区与所述第二存储器区之间的关系；以及
响应于通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的第一和第二存储器参考并不参考所述同一存储器区而进行以下操作：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者致使编译器基于所述元数据而重新编译所述内核；以及
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者指示目标处理器执行所述经重新编译的内核。
12.根据权利要求11所述的装置，
其中为了确定所述内核的所述第一存储器参考和所述第二存储器参考是否参考所述同一存储器区，所述编译处理器进一步经配置以：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的包含所述第一和第二存储器参考的循环代码区段，
其中为了重新编译所述内核，所述编译处理器进一步经配置以：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述产生的元数据而展开所述循环代码区段，且编译所述展开的循环代码区段。
13.根据权利要求11所述的装置，
其中为了确定所述内核的所述第一存储器参考和所述第二存储器参考是否参考所述同一存储器区，所述编译处理器进一步经配置以：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的包含所述第一和第二存储器参考的代码区段，
其中为了重新编译所述内核，所述编译处理器进一步经配置以：
响应于确定所述代码区段的所述第一和第二存储器参考并不参考所述同一存储器区，通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述产生的元数据而重排序所述代码区段的加载操作和存储操作中的至少一者。
14.根据权利要求11所述的装置，
其中为了确定所述内核的所述第一存储器参考和所述第二存储器参考是否参考所述同一存储器区，所述编译处理器进一步经配置以：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的包含所述第一和第二存储器参考的代码区段，
其中为了重新编译所述内核，所述编译处理器进一步经配置以：
响应于确定所述代码区段的所述第一和第二存储器参考并不参考所述同一存储器区，通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述产生的元数据而将所述代码区段的多个标量指令向量化为至少一个向量指令。
15.根据权利要求11所述的装置，其中所述元数据进一步指示所述第一存储器区与所述第二存储器区之间的重叠区。
16.根据权利要求15所述的装置，其中所述元数据包含所述第一存储器区与所述第二存储器区之间的重叠的字节数目。
17.根据权利要求15所述的装置，其中所述元数据进一步包括所述存储器重叠区的起始地址和所述存储器重叠区的结束地址中的至少一者。
18.根据权利要求11所述的装置，其中所述编译处理器包括中央处理单元CPU，且所述目标处理器包括图形处理单元GPU。
19.根据权利要求11所述的装置，其中所述编译器使用包括微软DirectCompute和OpenCL中的至少一者的异构计算框架来重新编译所述内核。
20.根据权利要求11所述的装置，其中所述内核自变量包括为所述自变量分配的存储器的缓冲器区域。
21.一种非暂时性计算机可读存储媒体，其上包含指令，所述指令当执行时致使编译处理器进行以下操作：
通过由在所述编译处理器上执行的编译器和运行时间组成的群组中的至少一者产生用于执行经编译内核的二进制代码的自变量；
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定对所述内核自变量的第一存储器区的第一存储器参考和对所述内核自变量的第二存储器区的第二存储器参考是否参考同一存储器区；
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述确定而产生与所述第一存储器参考和所述第二存储器参考相关联的元数据，其中所述元数据指示所述第一存储器区与所述第二存储器区之间的关系；以及
响应于通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的第一和第二存储器参考并不参考所述同一存储器区而进行以下操作：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者致使编译器基于所述元数据而重新编译所述内核；以及
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者指示目标处理器执行所述经重新编译的内核。
22.根据权利要求21所述的非暂时性计算机可读存储媒体，
其中所述致使所述编译处理器确定所述内核的所述第一存储器参考和所述第二存储器参考是否参考所述同一存储器区的指令进一步包括当执行时致使所述编译处理器进行以下操作的指令：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的包含所述第一和第二存储器参考的循环代码区段，
其中所述致使所述编译处理器重新编译所述内核的指令进一步包括当执行时致使所述编译处理器进行以下操作的指令：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述产生的元数据而展开所述循环代码区段，且编译所述展开的循环代码区段。
23.根据权利要求21所述的非暂时性计算机可读存储媒体，
其中所述致使所述编译处理器确定所述内核的所述第一存储器参考和所述第二存储器参考是否参考所述同一存储器区的指令进一步包括当执行时致使所述编译处理器进行以下操作的指令：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的包含所述第一和第二存储器参考的代码区段，
其中所述致使所述编译处理器重新编译所述内核的指令进一步包括当执行时致使所述编译处理器进行以下操作的指令：
响应于确定所述代码区段的所述第一和第二存储器参考并不参考所述同一存储器区，通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述产生的元数据而重排序所述代码区段的加载操作和存储操作中的至少一者。
24.根据权利要求21所述的非暂时性计算机可读存储媒体，
其中所述致使所述编译处理器确定所述内核的所述第一存储器参考和所述第二存储器参考是否参考所述同一存储器区的指令进一步包括当执行时致使所述编译处理器进行以下操作的指令：
通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者确定所述内核的包含所述第一和第二存储器参考的代码区段，
其中所述致使所述编译处理器重新编译所述内核的指令进一步包括当执行时致使所述编译处理器进行以下操作的指令：
响应于确定所述代码区段的所述第一和第二存储器参考并不参考所述同一存储器区，通过由在所述编译处理器上执行的所述编译器和所述运行时间组成的所述群组中的所述至少一者基于所述产生的元数据而将所述代码区段的多个标量指令向量化为至少一个向量指令。
25.根据权利要求21所述的非暂时性计算机可读存储媒体，其中所述元数据进一步指示所述第一存储器区与所述第二存储器区之间的重叠区。
26.根据权利要求25所述的非暂时性计算机可读存储媒体，其中所述元数据包含所述第一存储器区与所述第二存储器区之间的重叠的字节数目。
27.根据权利要求25所述的非暂时性计算机可读存储媒体，其中所述元数据进一步包括所述存储器重叠区的起始地址和所述存储器重叠区的结束地址中的至少一者。
28.根据权利要求21所述的非暂时性计算机可读存储媒体，其中所述编译处理器包括中央处理单元CPU，且所述目标处理器包括图形处理单元GPU。
29.根据权利要求21所述的非暂时性计算机可读存储媒体，其中所述编译器使用包括微软DirectCompute和OpenCL中的至少一者的异构计算框架来重新编译所述内核。
30.根据权利要求21所述的非暂时性计算机可读存储媒体，其中所述内核自变量包括为所述自变量分配的存储器的缓冲器区域。
CN201580016245.6A 2014-04-04 2015-03-19 用于编译器优化的存储器参考元数据 Pending CN106164862A (zh)

Applications Claiming Priority (3)

Application Number Priority Date Filing Date Title
US14/245,946 US9710245B2 (en) 2014-04-04 2014-04-04 Memory reference metadata for compiler optimization
US14/245,946 2014-04-04
PCT/US2015/021585 WO2015153143A1 (en) 2014-04-04 2015-03-19 Memory reference metadata for compiler optimization

Publications (1)

Publication Number Publication Date
CN106164862A true CN106164862A (zh) 2016-11-23

Family

ID=52829334

Family Applications (1)

Application Number Title Priority Date Filing Date
CN201580016245.6A Pending CN106164862A (zh) 2014-04-04 2015-03-19 用于编译器优化的存储器参考元数据

Country Status (6)

Country Link
US (1) US9710245B2 (zh)
EP (1) EP3132347A1 (zh)
JP (1) JP6329274B2 (zh)
KR (1) KR101832656B1 (zh)
CN (1) CN106164862A (zh)
WO (1) WO2015153143A1 (zh)

Cited By (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN108470072A (zh) * 2018-03-30 2018-08-31 迅讯科技（北京）有限公司 一种查询编译方法和装置
CN113168399A (zh) * 2019-03-14 2021-07-23 西部数据技术公司 可执行存储器单元
CN114398011A (zh) * 2022-01-17 2022-04-26 安谋科技（中国）有限公司 数据存储方法、设备和介质

Families Citing this family (21)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US20130113809A1 (en) * 2011-11-07 2013-05-09 Nvidia Corporation Technique for inter-procedural memory address space optimization in gpu computing compiler
GB2514618B (en) * 2013-05-31 2020-11-11 Advanced Risc Mach Ltd Data processing systems
US9785413B2 (en) * 2015-03-06 2017-10-10 Intel Corporation Methods and apparatus to eliminate partial-redundant vector loads
US9824419B2 (en) * 2015-11-20 2017-11-21 International Business Machines Corporation Automatically enabling a read-only cache in a language in which two arrays in two different variables may alias each other
WO2017209876A1 (en) * 2016-05-31 2017-12-07 Brocade Communications Systems, Inc. Buffer manager
US10169009B2 (en) 2016-06-01 2019-01-01 International Business Machines Corporation Processor that detects memory aliasing in hardware and assures correct operation when memory aliasing occurs
US9934009B2 (en) 2016-06-01 2018-04-03 International Business Machines Corporation Processor that includes a special store instruction used in regions of a computer program where memory aliasing may occur
US10169010B2 (en) 2016-06-01 2019-01-01 International Business Machines Corporation Performing register promotion optimizations in a computer program in regions where memory aliasing may occur and executing the computer program on processor hardware that detects memory aliasing
JP6810380B2 (ja) * 2016-10-07 2021-01-06 日本電気株式会社 ソースプログラム変換システム、ソースプログラム変換方法、及びソースプログラム変換プログラム
US10108404B2 (en) * 2016-10-24 2018-10-23 International Business Machines Corporation Compiling optimized entry points for local-use-only function pointers
CN110121703B (zh) * 2016-12-28 2023-08-01 英特尔公司 用于向量通信的系统和方法
US10547491B2 (en) * 2017-08-28 2020-01-28 Genband Us Llc Transcoding with a vector processing unit
US10540194B2 (en) * 2017-12-21 2020-01-21 International Business Machines Corporation Runtime GPU/CPU selection
US11367160B2 (en) * 2018-08-02 2022-06-21 Nvidia Corporation Simultaneous compute and graphics scheduling
US10884720B2 (en) * 2018-10-04 2021-01-05 Microsoft Technology Licensing, Llc Memory ordering annotations for binary emulation
CN111340678A (zh) * 2018-12-19 2020-06-26 华为技术有限公司 一种数据缓存系统、图形处理器及数据缓存方法
US10872057B1 (en) * 2019-05-23 2020-12-22 Xilinx, Inc. Partitioning in a compiler flow for a heterogeneous multi-core architecture
JP7460902B2 (ja) 2020-06-09 2024-04-03 富士通株式会社 コンパイラプログラム、コンパイル方法、情報処理装置
JP7164267B2 (ja) * 2020-12-07 2022-11-01 インテル・コーポレーション ヘテロジニアスコンピューティングのためのシステム、方法及び装置
EP4276602A1 (de) * 2022-05-12 2023-11-15 Siemens Aktiengesellschaft System mit quellcodeumwandler-spezifizierten speicherbereichen
US20240095024A1 (en) * 2022-06-09 2024-03-21 Nvidia Corporation Program code versions

Citations (7)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US20040205740A1 (en) * 2001-03-29 2004-10-14 Lavery Daniel M. Method for collection of memory reference information and memory disambiguation
US20060059326A1 (en) * 2002-11-21 2006-03-16 Microsoft Corporation Dynamic data structures for tracking file system free space in a flash memory device
CN102314342A (zh) * 2010-06-18 2012-01-11 微软公司 用于数据并行编程模型的编译器生成的调用存根
CN102460376A (zh) * 2009-06-26 2012-05-16 英特尔公司 无约束事务存储器（utm）系统的优化
CN103116513A (zh) * 2012-07-13 2013-05-22 北京时代民芯科技有限公司 一种异构多核处理器编译器
US20130132684A1 (en) * 2011-11-18 2013-05-23 Microsoft Organization Automatic optimization for programming of many-core architectures
US8458671B1 (en) * 2008-02-12 2013-06-04 Tilera Corporation Method and system for stack back-tracing in computer programs

Family Cites Families (15)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US6877088B2 (en) * 2001-08-08 2005-04-05 Sun Microsystems, Inc. Methods and apparatus for controlling speculative execution of instructions based on a multiaccess memory condition
CA2372034A1 (en) 2002-02-14 2003-08-14 Cloakware Corporation Foiling buffer-overflow and alien-code attacks by encoding
AU2003286131A1 (en) 2002-08-07 2004-03-19 Pact Xpp Technologies Ag Method and device for processing data
US7565631B1 (en) 2004-07-02 2009-07-21 Northwestern University Method and system for translating software binaries and assembly code onto hardware
US8886887B2 (en) * 2007-03-15 2014-11-11 International Business Machines Corporation Uniform external and internal interfaces for delinquent memory operations to facilitate cache optimization
US8413126B2 (en) * 2007-06-15 2013-04-02 Cray Inc. Scalar code reduction using shortest path routing
US20090070753A1 (en) 2007-09-07 2009-03-12 International Business Machines Corporation Increase the coverage of profiling feedback with data flow analysis
US8131979B2 (en) 2008-08-15 2012-03-06 Apple Inc. Check-hazard instructions for processing vectors
US8527737B2 (en) * 2010-06-23 2013-09-03 Apple Inc. Using addresses to detect overlapping memory regions
WO2012142186A2 (en) * 2011-04-11 2012-10-18 Child Timothy Database acceleration using gpu and multicore cpu systems and methods
US8935683B2 (en) * 2011-04-20 2015-01-13 Qualcomm Incorporated Inline function linking
US8468507B2 (en) 2011-06-10 2013-06-18 Microsoft Corporation Binding executable code at runtime
US20130141443A1 (en) * 2011-12-01 2013-06-06 Michael L. Schmit Software libraries for heterogeneous parallel processing platforms
US9256915B2 (en) * 2012-01-27 2016-02-09 Qualcomm Incorporated Graphics processing unit buffer management
US9734333B2 (en) * 2012-04-17 2017-08-15 Heat Software Usa Inc. Information security techniques including detection, interdiction and/or mitigation of memory injection attacks

Patent Citations (7)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US20040205740A1 (en) * 2001-03-29 2004-10-14 Lavery Daniel M. Method for collection of memory reference information and memory disambiguation
US20060059326A1 (en) * 2002-11-21 2006-03-16 Microsoft Corporation Dynamic data structures for tracking file system free space in a flash memory device
US8458671B1 (en) * 2008-02-12 2013-06-04 Tilera Corporation Method and system for stack back-tracing in computer programs
CN102460376A (zh) * 2009-06-26 2012-05-16 英特尔公司 无约束事务存储器（utm）系统的优化
CN102314342A (zh) * 2010-06-18 2012-01-11 微软公司 用于数据并行编程模型的编译器生成的调用存根
US20130132684A1 (en) * 2011-11-18 2013-05-23 Microsoft Organization Automatic optimization for programming of many-core architectures
CN103116513A (zh) * 2012-07-13 2013-05-22 北京时代民芯科技有限公司 一种异构多核处理器编译器

Non-Patent Citations (2)

* Cited by examiner, † Cited by third party
Title
CHRISTOS MARGIOLAS等: "HYDA:A Hybrid Dependence Analysis for the adaptive", 《ADAPT "14 PROCEEDINGS OFINTERNATIONAL WORKSHOP ON ADAPTIVE SELF-TUNING COMPUTING SYSTEMS》 *
DMITRY MIKUSHIN等: "KernelGen the design and implementation of a next generation compiler platform for accelerating numerical models on GPUs", 《USI TECHNICAL REPORT SERIES IN INFORMATICS》 *

Cited By (6)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN108470072A (zh) * 2018-03-30 2018-08-31 迅讯科技（北京）有限公司 一种查询编译方法和装置
CN108470072B (zh) * 2018-03-30 2019-07-09 迅讯科技（北京）有限公司 一种查询编译方法和装置
CN113168399A (zh) * 2019-03-14 2021-07-23 西部数据技术公司 可执行存储器单元
CN113168399B (zh) * 2019-03-14 2023-09-19 西部数据技术公司 可执行存储器单元
CN114398011A (zh) * 2022-01-17 2022-04-26 安谋科技（中国）有限公司 数据存储方法、设备和介质
CN114398011B (zh) * 2022-01-17 2023-09-22 安谋科技（中国）有限公司 数据存储方法、设备和介质

Also Published As

Publication number Publication date
KR20160141753A (ko) 2016-12-09
US9710245B2 (en) 2017-07-18
WO2015153143A1 (en) 2015-10-08
JP6329274B2 (ja) 2018-05-23
KR101832656B1 (ko) 2018-02-26
JP2017509999A (ja) 2017-04-06
EP3132347A1 (en) 2017-02-22
US20150286472A1 (en) 2015-10-08

Similar Documents

Publication Publication Date Title
CN106164862A (zh) 用于编译器优化的存储器参考元数据
US11243816B2 (en) Program execution on heterogeneous platform
Membarth et al. Generating device-specific GPU code for local operators in medical imaging
US8561045B2 (en) Constructing runtime state for inlined code
US8799871B2 (en) Computation of elementwise expression in parallel
US9015683B2 (en) Method and apparatus for transforming program code
JP5639274B2 (ja) 計算プラットフォームのヘテロジニアスプロセッサの間で共有されるバーチャルメモリにおけるバーチャル機能の共有
CN110008009B (zh) 在运行时绑定常量以提高资源利用率
US9146759B2 (en) Assumption-based compilation
US20190121625A1 (en) Dynamic compiler parallelism techniques
US20110022817A1 (en) Mapping Processing Logic Having Data-Parallel Threads Across Processors
US8893104B2 (en) Method and apparatus for register spill minimization
EP2862066A1 (en) Adaptive portable libraries
US20100011339A1 (en) Single instruction multiple data (simd) code generation for parallel loops using versioning and scheduling
US7418699B2 (en) Method and system for performing link-time code optimization without additional code analysis
Mikushin et al. KernelGen--The Design and Implementation of a Next Generation Compiler Platform for Accelerating Numerical Models on GPUs
JP2015509249A5 (zh)
US20160179490A1 (en) Compiler
Reiche et al. Auto-vectorization for image processing DSLs
US10002401B2 (en) Method and apparatus for efficient processing of graphics commands
US10552135B1 (en) Reducing a size of an application package
US10353708B2 (en) Strided loading of non-sequential memory locations by skipping memory locations between consecutive loads
Haidl et al. High-level programming for many-cores using C++ 14 and the STL
Ragan-Kelley Technical Perspective: Reconsidering the Design of User-Schedulable Languages
US11762641B2 (en) Allocating variables to computer memory

Legal Events

Date Code Title Description
C06 Publication
PB01 Publication
C10 Entry into substantive examination
SE01 Entry into force of request for substantive examination
WD01 Invention patent application deemed withdrawn after publication
WD01 Invention patent application deemed withdrawn after publication

Application publication date: 20161123