CN105740036B - 用于支持代码的编译时间定制的编译器的***和方法 - Google Patents
用于支持代码的编译时间定制的编译器的***和方法 Download PDFInfo
- Publication number
- CN105740036B CN105740036B CN201511004391.5A CN201511004391A CN105740036B CN 105740036 B CN105740036 B CN 105740036B CN 201511004391 A CN201511004391 A CN 201511004391A CN 105740036 B CN105740036 B CN 105740036B
- Authority
- CN
- China
- Prior art keywords
- code
- equipment
- lambda expression
- expression formula
- source code
- 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.)
- Active
Links
- 238000000034 method Methods 0.000 title claims abstract description 59
- 230000004044 response Effects 0.000 claims abstract description 3
- 230000004048 modification Effects 0.000 claims description 23
- 238000012986 modification Methods 0.000 claims description 23
- 230000006870 function Effects 0.000 description 143
- 230000008569 process Effects 0.000 description 24
- 238000003860 storage Methods 0.000 description 23
- HPTJABJPZMULFH-UHFFFAOYSA-N 12-[(Cyclohexylcarbamoyl)amino]dodecanoic acid Chemical compound OC(=O)CCCCCCCCCCCNC(=O)NC1CCCCC1 HPTJABJPZMULFH-UHFFFAOYSA-N 0.000 description 22
- 238000012545 processing Methods 0.000 description 14
- 238000007789 sealing Methods 0.000 description 12
- 239000011800 void material Substances 0.000 description 11
- 230000008859 change Effects 0.000 description 10
- 238000004891 communication Methods 0.000 description 8
- 101150071648 licB gene Proteins 0.000 description 6
- 238000012546 transfer Methods 0.000 description 4
- 241000208340 Araliaceae Species 0.000 description 3
- 235000005035 Panax pseudoginseng ssp. pseudoginseng Nutrition 0.000 description 3
- 235000003140 Panax quinquefolius Nutrition 0.000 description 3
- 230000005540 biological transmission Effects 0.000 description 3
- 235000008434 ginseng Nutrition 0.000 description 3
- 238000004364 calculation method Methods 0.000 description 2
- 238000013500 data storage Methods 0.000 description 2
- 238000010586 diagram Methods 0.000 description 2
- 238000009877 rendering Methods 0.000 description 2
- 238000006467 substitution reaction Methods 0.000 description 2
- 241000196324 Embryophyta Species 0.000 description 1
- 230000008901 benefit Effects 0.000 description 1
- 238000010276 construction Methods 0.000 description 1
- 230000001419 dependent effect Effects 0.000 description 1
- 238000013461 design Methods 0.000 description 1
- 238000011161 development Methods 0.000 description 1
- 238000009826 distribution Methods 0.000 description 1
- 238000005516 engineering process Methods 0.000 description 1
- 238000005304 joining Methods 0.000 description 1
- 238000004519 manufacturing process Methods 0.000 description 1
- 230000007246 mechanism Effects 0.000 description 1
- 239000013307 optical fiber Substances 0.000 description 1
- 230000002093 peripheral effect Effects 0.000 description 1
- 235000015170 shellfish Nutrition 0.000 description 1
- 238000000844 transformation Methods 0.000 description 1
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/40—Transformation of program code
- G06F8/51—Source to source
Landscapes
- Engineering & Computer Science (AREA)
- Theoretical Computer Science (AREA)
- General Engineering & Computer Science (AREA)
- Software Systems (AREA)
- Physics & Mathematics (AREA)
- General Physics & Mathematics (AREA)
- Stored Programmes (AREA)
- Devices For Executing Special Programs (AREA)
Abstract
本发明提供用于支持代码的编译时间定制的编译器的***和方法,用于处理源代码用于编译。所述方法包括访问主机源代码的一部分以及确定所述主机源代码的一部分是否包括设备lambda表达式。所述方法进一步包括响应于所述主机代码的一部分包括所述设备lambda表达式,基于所述设备lambda表达式确定唯一占位符类型实例化以及基于所述唯一占位符类型实例化修改所述设备lambda表达式以生成经修改的主机源代码。所述方法进一步包括将所述经修改的主机源代码传送至主机编译器。
Description
相关美国申请
本发明要求共同未决的、序列号为62/097,528、代理人案号为NVID-P-SC-14-0346-US0、名称为“COMPILER TRANSFORMATIONS AND MECHANISMS TO SUPPORT COMPILETIME CUSTOMIZATION OF CODE EXECUTING ON THE GPU WITH LAMBDA EXPRESSIONSDEFINED IN CODE MARKED FOR CPU EXECUTION”、提交日期为2014年12月29日的临时专利申请的权益和优先权,并且此处通过引用全部并入本申请。
本申请与序列号为13/735,981、代理人案号为NVID-P-BL-12-0178-US1、名称为“SYSTEM AND METHOD FOR COMPILER SUPPORT FOR KERNEL LAUNCHES IN DEVICE CODE”、提交日期为2013年1月7日的非临时专利申请相关,该非临时专利申请要求序列号为61/645,515、代理人案号为NVID-P-BL-12-0178-US0、名称为“SYSTEM AND METHOD FORCOMPILER SUPPORT FOR KERNEL LAUNCHES IN DEVICE CODE”、提交日期为2012年5月10日的临时专利申请的优先权,它们中的每一个均通过引用方式全部并入本申请。
技术领域
本发明的实施例总体上涉及主机源代码和设备源代码的编译,以及图形处理单元(GPU)。
背景技术
随着计算机***的发展,图形处理单元(GPU)在复杂性和计算能力两方面都日益提升。因此GPU被用于操控越来越大量和复杂的图形处理。这种处理能力增加的结果是,现在GPU能够执行图形处理和更多的通用处理任务二者。在GPU上执行通用计算任务的能力导致开发在GPU上执行通用计算任务的程序增加了,并且能够执行数目更多的复杂编程任务的相应需求增加了。
图形处理单元上的通用计算(GPGPU)程序具有在中央处理单元(CPU)上执行的主机部分和在GPU上执行的设备部分。GPGPU编译器编译GPU或设备代码并将主机代码提供给用于主机的本地编译器。因此用于主机的本地编译器可以依赖于执行GPGPU编译器所处的编程环境而改变。结果是,送至主机编译器的代码需要可由任何多种不同的主机编译器兼容。当新的特征和/或高级特征增加到编程语言中时这可能造成问题。
发明内容
因此,需要一种允许支持编程语言特征,同时允许使用多种主机编译器的任何一种来编译GPGPU或异构程序的主机代码的方案。
实施例允许主机代码中的拉姆达(lambda)表达式的编译,而不需要主机编译器的修改或定制。实施例能够识别与各自设备代码入口点相关联的lambda表达式(例如,CUDA C++中的__global__template实例化)并且用唯一占位符类型实例化代替每个lambda,从而允许在主机代码中定义的lambda表达式的主机编译程序进行编译而无需主机编译器的修改。
“设备代码入口函数”表示设备代码开始执行的代码位置。当设备代码入口函数模板采用如语言惯用语所允许的模板参数实例化时,设备代码入口函数可以由该模板生成。每个不同的模板实例化是一个不同的设备代码入口函数。
“设备lambda表达式”是在主机代码中定义的lambda表达式,其中lambda表达式包含设备代码,例如,在GPU上执行的代码。
在一个实施例中,本发明涉及一种用于处理源代码用于编译的方法。该方法包括访问主机源代码的一部分并确定主机源代码的一部分是否包括设备lambda表达式。设备lambda表达式关联于设备代码入口点(例如,CUDA C++中的__global__template实例化)。该方法进一步包括,响应于该一部分主机代码包括设备lambda表达式,基于该lambda表达式确定唯一占位符类型实例化,并且基于该唯一占位符类型实例化修改该设备lambda表达式,以产生经修改的主机源代码。在一些实施例中,唯一占位符类型实例化包括与封闭设备lambda表示式的函数相关联的类型。在一些实施例中,唯一占位符类型实例化包括封闭设备lambda表达式的函数的地址。在一些实施例中,唯一占位符类型实例化包括与设备lambda表达式相关联的唯一识别符。在一些实施例中,唯一占位符类型实例化包括一个或多个与设备lambda表达式相关联的变量。在一些实施例中,唯一占位符类型实例化通过将该一个或多个变量的值显式地传递至占位符类型实例化的构造器来捕获一个或多个与该设备lambda表达式相关联的变量。该方法进一步包括将经修改的主机源代码送至主机编译器。
该方法可以进一步包括基于关联于设备lambda表达式的该唯一占位符类型实例化确定设备代码入口函数模板特化,其中原始主机源代码包括设备代码入口函数模板和设备lambda表达式,其类型被用于该模板的实例化中。
在一个实施例中,本发明涉及用于编译代码的***。该***包括可操作为从数据存储中访问源代码的代码访问模块,其中源代码包括主机源代码和设备源代码,以及配置为在主机源代码的一部分中识别设备lambda表达式的设备lambda表达式识别模块。该***进一步包括配置为基于设备lambda表达式确定占位符类型实例化的占位符类型确定模块和配置为修改主机源代码以用占位符类型实例化代替设备lambda表达式的代码修改模块。在一些实施例中,占位符类型实例化包括与封闭设备lambda表达式的函数相关联的类型。在一些实施例中,占位符类型实例化包括封闭设备lambda表达式的函数的地址。在一些实施例中,占位符类型实例化包括与设备lambda表达式相关联的唯一识别符。在一些实施例中,占位符类型实例化包括一个或多个与设备lambda表达式相关联的变量。在一些实施例中,占位符类型实例化通过将一个或多个变量的值显式地传递至占位符类型实例化的构造器来捕获一个或多个与设备lambda表达式相关联的变量。在一些实施例中,代码修改模块进一步配置为基于由模板特化确定模块确定的设备代码入口函数模板特化修改主机代码。
该***可以进一步包括配置为基于设备lambda表达式确定设备代码入口函数模板特化的模板特化确定模块。该***可以进一步包括配置为将由代码修改模块修改的主机源代码发送至主机编译器的代码发送模块。
在另一个实施例中,本发明实现为用于修改代码用于编译的方法。该方法包括访问主机源代码的一部分并且确定该主机源代码的一部分是否包括关联于设备代码入口函数模板实例化的设备lambda表达式。该方法进一步包括基于设备lambda表达式确定唯一占位符类型实例化并且基于唯一占位符类型实例化确定设备代码入口函数模板特化。主机源代码(例如,原始的或所访问的主机源代码)包括设备代码入口函数模板和设备lambda表达式,其类型被用于该模板的实例化中。该方法进一步包括通过用唯一占位符类型实例化的实例代替设备lambda表达式来生成经修改的主机源代码,并且基于唯一占位符类型实例化(例如,其在它的模板参数中使用唯一占位符类型实例化)生成设备代码入口函数特化(例如,新的设备代码入口函数特化)。该方法进一步包括将经修改的主机源代码送至主机编译器。在一些实施例中,唯一占位符类型实例化包括与封闭设备lambda表达式的函数相关联的类型。在一些实施例中,唯一占位符实例化包括封闭设备lambda表达式的函数的地址和与设备lambda表达式相关联的识别的唯一整数。
附图说明
在随附附图的各图中,本发明的实施例以示例的方式,并且不以限定的方式进行说明,并且其中相同的引用编号涉及相同的元件。
图1示出了根据各种实施例的示范性计算机***。
图2示出了根据各种实施例的、主机代码和设备代码的示范性计算机控制器编译过程。
图3示出了根据各种实施例的、用于处理源代码文件的示范性计算机控制的过程的流程图。
图4示出了根据各种实施例的、用于用唯一占位符类型实例化确定和修改源代码文件的示范性计算机控制的过程的流程图。
图5示出了根据各种实施例的、示范性计算机***和相应的模块的框图。
具体实施方式
现在将详细参考本发明的优选实施例,其例子在随附的附图中进行说明。然而,本发明将连同优选实施例一起进行描述,其将被理解为,它们并不意图将本发明限制为这些实施例。相反地,本发明意图覆盖所有替代、修改和等价,其可以被包括在由附加的权利要求限定的本发明的精神和范围内。而且,在下面对本发明的实施例的详细的描述中,为了提供对本发明的详尽理解,详细阐述了大量具体细节。然而,本领域普通技术人员应认识到,本发明可以实践为没有这些具体细节。在其他实例中,公知的方法、程序、组件和电路由于不是使本发明的实施例不必要地不清楚的方面而没有详细描述。
符号和术语:
详细说明的某些部分,如下,以程序、步骤、逻辑块、过程和对计算机存储器中的数据位的操作的其他符号表示的形式呈现。对本领域技术人员来说,这些说明和表示是数据处理领域技术人员用于最有效地表达他们工作内容的手段。程序、计算机执行的步骤、逻辑块、过程等,在此处,并且通常被考虑为导向期望结果的步骤或指令的自洽序列。步骤是要求物理量的物理操作的那些。通常,虽然不是必要地,这些量采用能够被存储、转移、组合、比较以及不然在计算机***中能够***控的电或磁信号的形式。有时被证明是便利的,主要由于普遍使用的原因,作为位、值、元件、符号、字符、术语、数字等参考这些信号。
然而应铭记于心,所有这些以及相同术语与合适的物理量相关,并且仅仅是用于这些量的便利标记。除非特别说明,否则从下面的讨论中明显得到的,应理解为贯穿本发明,利用术语例如“处理”或“访问”或“执行”或“存储”或“渲染”等的讨论涉及集成电路(例如,图1中的计算***100)或类似电子计算设备的动作和过程,其将表示为计算机***的寄存器和存储器中的物理(电子)量的数据操控和变换成类似地表示为计算机***存储器或寄存器或其他这种信息存储、传输或显示设备中的物理量的其他数据。
计算机***环境
图1示出了根据本发明的一个实施例的计算机***100。计算机***100描述了根据本发明的实施例的基础计算机***的组件,其为某些基于硬件的和基于软件的功能提供了执行平台。总的来说,计算机***100包括至少一个CPU 101、***存储器115和至少一个图形处理单元(GPU)110。CPU 101可以通过桥组件/存储器控制器(未示出)耦连至***存储器115或通过CPU 101内部的存储器控制器(未示出)直接耦连至***存储器115。GPU 110可以耦连至显示器112。一个或多个附加的GPU可以可选择地耦连至***100,以进一步增加它的计算能力。GPU 110耦连至CPU 101和***存储器115。GPU 110可以实现为分离式组件、设计为通过连接器(例如,AGP槽、PCI-Express槽等)耦连至计算机***100的分离式显卡、分离式集成电路裸片(例如,直接安装到母板上)或实现为包括在计算机***芯片集组件(未示出)的集成电路裸片中的集成GPU。此外,可以包括用于GPU 110的本地图形存储器114,用于高带宽图形数据存储。
CPU 101和GPU 110也可以集成到单个集成电路裸片中并且CPU和GPU可以共享各种资源,例如,指令逻辑、缓冲区、功能单元等,或可以向图形和通用操作提供单独的资源。GPU可以进一步被集成到核逻辑组件中。因此,此处描述的与GPU 110相关联的任何或所有电路和/或功能也可以在合适装备的CPU 101中实现或由其实施。此外,然而此处的实施例可以参考GPU,应注意到,描述的电路和/或功能也可以在其他类的处理器(例如,通用或其他专用协处理器)或在CPU中实现。
***100可以实现为例如具有强大的通用CPU 101耦连至专用图形渲染GPU 110的台式计算机***或服务器计算机***。在这样的实施例中,可以包括增加***总线、专用音频/视频组件、I/O设备等的组件。类似地,***100可以实现为手持设备(例如,手机等)、直接广播卫星(DBS)/地面机顶盒或机顶盒视频游戏控制台设备,例如由华盛顿雷德蒙德的微软公司支持的或者由日本东京的索尼计算机娱乐公司支持的***100也可以实现为片上***,其中计算设备的电子元件(组件101、115、110和114等)整个包含在单个集成电路裸片中。例子包括具有显示器的手持设备、汽车导航***、便携式娱乐***等。
在一个示范性实施例中,GPU可操作用于在图形处理单元(GPGPU)计算上的通用计算。图形处理单元(GPGPU)程序或应用程序上的通用计算可以用统一计算机设备架构(Compute Unified Device Architecture)(CUDA)框架和开发计算语言(Open ComputingLanguage)(OpenCL)框架设计或写入。GPU 110可以执行统一计算机设备架构(CUDA)程序和开发计算语言(OpenCL)程序。应理解,GPU 110的并行体系架构相比CPU 101可以具有显著的性能优势。
在一些实施例中,存储器115包括用于编译异构代码(例如,GPGPU代码)或包括主机代码和设备代码的代码的编译器150。编译器150可以是支持修改主机代码以允许由各种主机编译器编译主机源代码的设备代码编译器。如此处描述的,编译器150可以修改包括与设备代码入口函数模板实例化相关联的设备拉姆达(lambda)表达式的主机源代码。编译器150可以将经修改的主机代码送至主机编译器(未示出)。
GPGPU编程环境和编译中的lambda表达式
GPGPU语言可以支持编程规范的子集。例如,CUDA语言可以支持C++规范的子集。C++标准引入“lambda表达式”,其允许可执行函数在应用点被定义为“内联”。lambda表达式没有被显式地命名的相关闭包类型。该闭包类型可以用于或参与模板实例化。
具有lambda表达式的示范性的部分代码:
表I–具有lambda表达式的示范性的代码部分
lambda表达式可以看起来像函数但不仅仅是函数。lambda表达式是可以在任何地方定义的代码的部分或片段。例如,lambda表达式可以在函数中定义或函数外面定义。
C++中的lambda表达式实际上是叫做“算符(functor)”的对象。算符是类对象,例如标准的C++类型,其实现了括号运算符或运算符()。运算符()是被叫做类的实例的函数。例如,如果有个类“T”以及宣告命名为“a”的类T的对象,那么因为该类具有运算符(),“a()”看起来像针对称为“a”的函数的函数调用,但实际上是调用对象“a”的括号运算符。lambda表达式是算符的实例。lambda表达式是由C++中编译器支持的具有特殊特点的一种类型的算符。
lambda表达式是算符的实例但类型的名称用户是未知的。参考表I,[=]或方括号开始lambda表达式。[=]之后直到声明末尾(例如,结尾的}以及;)的代码是lambda表达式。lambda表达式代表算符的实例,但类的名称未被定义。换句话说,lambda表达式是部分未命名的类。该类的名称仅对正在处理lambda表达式的编译器是已知的。换句话说,lambda表达式是未命名的算符。
lambda表达式允许在其要使用的代码中的点处定义功能。因而lambda表达式允许定义功能用于在使用点定制函数或类。lambda表达式可以进一步对在编译时利用用户定义的代码定制功能是有用的。例如,标准库排序函数std::sort可以在编译时用lambda表达式定制,其实现了用户定义的比较函数。作为另一个例子,快速排序功能可以被实例化并且lambda表达式作为被编译函数而通过。lambda表达式允许避免必须定义单独的类以及命名该类,其相比使用lambda表达式而言是非常繁复的。
因为lambda表达式是类的实例,因此模板可以基于lambda表达式实例化。例如,函数模板比如‘template<typename T>void foo(void)’可以用lambda的相关类型(相关类型称为“闭包类型”)实例化。编译器会知道闭包类型的名称但在代码中没有闭包类型的名称。lambda表达式也可以用于实例化类模板。
lambda表达式可以进一步捕获用在lambda表达式中的局部变量。这意味着在lambda表达式定义点处捕获变量的值。换句话说,变量的副本由lambda表达式中的变量制作。变量也可以通过lambda表达式中的符号(例如&)修改。
“设备lambda表达式”是在主机代码中定义的lambda表达式,其中lambda表达式包括设备代码,例如在GPU上执行的代码。
在一个示范性实施例中,使用CUDA程序,具有在CPU上执行的主机函数与在GPU上执行的设备函数。具有更加特殊的函数,称为全局入口函数(global entry function)(例如,__global__function),其为设备代码的入口点。在CUDA C++中,GPU入口函数可以模板化。这些模板被实例化并从主机代码启动。用户会希望采用主机代码中定义的lambda表达式定制__global__function模板实例化。然而,这不被一些GPGPU模型所允许,因为lambda表达式的类型不是显式已知的。实施例允许GPGPU(例如CUDA)实现方式消除这种限制,使得GPGPU(例如,CUDA或OpenCL)语言更加有用并更加接近C++。
GPGPU编程***可以使用本地主机编译器编译主机代码(例如,在CPU上执行的代码)。例如,CUDA编译器工具链可能不采用主机编译器运送。而是编译器前端可以运送,其将输入的混杂的CUDA源文件分离成主机(例如CPU)代码和设备(例如,GPU)代码。设备代码由CUDA后端编译器处理,生成可执行的GPU机器代码。主机代码被重新生成为高级C++(例如,包括模板)并且被传递或发送至主机编译器,呈现在平台上(例如,g++,cl.exe,clang++)。
__global__function代表设备代码入口点并且__global__function模板可以从主机代码实例化并在设备上开始。为了支持__global__function模板,给定工具链限制使用本地主机编译器以及需要生成高级主机端C++代码,CUDA编译器前端在被传递到主机编译器的代码中生成明确的__global__function模板特化(specialization)。例如,考虑下面的输入CUDA源文件:
__device__int result;
template<typename T>
__global__void foo(T in){result=in;}
int main(void)
{
foo<int><<<1,1>>(1);
foo<char><<<1,1>>>(2);
}
表II-示范性输入CUDA源文件
在发送至主机编译器的代码中,CUDA前端综合‘foo<int>’和‘foo<char>’的显式特化。特化实现了内核启动所必需的任务,例如,拷贝参数值以及调用CUDA运行时间以启动内核。
当CUDA编译器前端为__global__function模板生成显式特化时,编译器前端需要显式地命名参与实例化的类型和非类型参数(例如,表II的例子中的‘char’和‘int’)。当模板参数调用与lambda表达式相关联的闭包类型时,这出现了一个问题。闭包类型不是显式命名的。名称仅对编译器已知并且其名称对编译器实现方式来说不是标准化的。而且,__global__function模板特化在命名空间范围内定义,但如果lambda表达式在函数内定义,那么闭包类型是对封闭函数范围而言是本地的,并且结果是在函数外部不可访问。在需要类型来定义显式模板特化的情况下会需要访问函数外部的类型,因为显式特化需要在命名空间范围内定义。
用于支持代码的编译时间定制的编译器的示范性***和方法
“设备代码入口函数”表示设备代码开始执行的代码位置。当模板采用如语言惯用语所允许的模板参数实例化时,设备代码入口函数可以由设备代码入口函数模板生成。每个不同的模板实例化是一个不同的设备代码入口函数。
“设备lambda表达式”是在主机代码中定义的lambda表达式,其中lambda表达式包含设备代码,例如,在GPU上执行的代码。
实施例允许编译主机代码中的设备lambda表达式,而无需主机编译器的修改或定制。实施例可以识别与各自的设备代码入口函数模板实例化相关联的设备lambda表达式,并且用唯一占位符类型实例化的实例代替每个设备lambda表达式,由此允许由在主机代码中定义的设备lambda表达式的主机编译器进行编译。
实施例识别并且标记在主机代码中定义的每个设备lambda表达式。当生成的代码(例如C++)被发送至主机编译器时,所识别的lambda表达式中的每个用新的“占位符”命名空间范围类型模板实例化的实例的定义来代替。
实施例配置为确保唯一占位符类型实例化与所代替的lambda表达式的每一个相关联。在一些实施例中,唯一占位符类型实例化的支持与处于函数范围内的定义的候选lambda表达式相关联。例如,占位符类型的唯一性通过将两个模板参数增加到占位符类型实例化来确保:1)函数的地址,其封闭原始lambda表达式定义,并且2)唯一整数常量,其被分配以将该候选lambda表达式与在相同函数中定义的其他候选lambda表达式区分开。唯一整数常量可以通过维护以每个函数为基础的计数器来实现,每次候选lambda表达式被处理则对其进行增量。
在lambda表达式从封闭函数中捕获任何变量的情况下,采用所捕获的变量的类型来实例化占位符类型,并且占位符类型的构造器保存所捕获的变量的值。
当生成显式的__global__function模板实例化,其中该模板参数调用lambda表达式闭包类型时,编译器前端(例如,CUDA编译器前端)用占位符类型的实例化代替闭包类型(例如,lambda表达式闭包类型)。
__device__int result;
template<typename T>
__global__void foo(T in){result=in(4);}
int main(void)
{
int x=10;
double d=20;
auto lam1=[=]__device__(int in){return x+d*in;};
foo<<<1,1>>(lam1);
}
表III-具有包括__global__function模板的实例化的lambda表达式类型的示范性输入CUDA源代码
表III的源代码包括具有标准主函数的主机代码和关联于设备代码入口函数模板的设备代码,设备代码入口函数模板具有定制的__global__function模板的模板实例化。模板实例化与在主机代码中定义的设备lambda表达式相关联。lambda表达式的使用创造了这样的情形:通用模板在主机代码的函数内的定义点处定义的代码的特化中通过。应领会,失去lambda表达式支持,类需要用运算符()写到函数体的外部。例如,会创建类bar并且模板foo会由bar的实例来实例化。
表III包括‘lam1’,其定义设备lambda表达式。与lambda表达式相关联的闭包类型用于实例化__global__function模板‘foo’(例如,foo<<<1,1>>(lam1);)。换句话说,全局模板用关联于设备lambda表达式的闭包类型实例化。设备lambda表达式从封闭函数(例如,主函数)捕获变量‘d’和‘x’。变量‘d’和‘x’用在设备代码中,即使该变量是主机变量,因此编译器拷贝变量并且确保变量被拷贝到__global__function并且在设备上执行。因而,编译器通过拷贝主机变量来捕获主机变量,用于在将在设备上执行的__global__function中使用。
lambda表达式基于表III中的模板(例如,template<typename T>__global__voidfoo(T in){result=in(4);})被执行并且运算符()的实例将被赋予4的值。
应领会,表III中的代码是示范性的,并且实施例可以支持其他变化。例如,表III中相同的lambda可以用于实例化类模板,制作lambda表达式字段的一些结构,然后其被传递到foo函数模板。实施例允许使用具有设备lambda表达式的模板的任何功能,包括但不限于多级嵌套。注意到,当参考lambda表达式讨论某些方面时,实施例能够操作和/或处理参与到模板实例化或其他实例化中的任何实体,其中与实体相关联的类型的名称在命名空间范围内是不可用的和/或名称对用户代码是不可用的。
实施例支持生成并输出主机代码用于由可以以多种方式变化的本地编译器编译。GPGPU编译器可以拥有它自己的设备编译器(例如,基于低层虚拟机(LLVM))并且使用本地主机编译器用于编译主机代码。因而主机编译器可以不受生产GPGPU编译器(例如,g++,clang,cl.exe)的公司或实体的控制。因此,由实施例生成的主机代码能够用多种主机编译器中的任何一种进行编译,该多种主机编译器未被定制为采用GPGPU编译器。
对于没有lambda表达式的代码,设备代码入口函数模板的特化(例如,int,double和其他类)可以被送至主机编译器而不用修改。例如,对于用int类型实例化的模板,送至主机编译器的代码将具有int的模板特化以及相关参数。然后主机代码被编译并且在运行时间建立在运行时间送至GPU的命令,包括建立参数、启动内核以及等待内核等。
lambda表达式采用模板特化产生了一个问题,其中类型的名称是未知的并且主机编译器不可以被定制。在lambda是算符类型的实例,并且当lambda在函数中定义时,类型本身对函数而言是本地的情况下,lambda表达式进一步产生了问题。例如,对于主函数中的lambda声明,由于C++规则使得类型对主函数而言是本地的,lambda表达式(即使它是显式的)的类型将对于主函数外部的模板特化将不可用,连同本地变量。闭包的类型,例如算符,也是本地的,限制对主函数外部的闭包类型的引用。
因此,lambda表达式产生了两个问题:1)闭包类型不具有显式名称,并且2)如果闭包类型具有名称,闭包类型本身对函数范围而言是本地的。实施例配置为解决这个问题,通过识别lambda表达式、用占位符类型代替lambda表达式,其实际上处于定义lambda表达式的范围(例如,函数范围)之外。处于定义lambda表达式的范围之外的占位符类型允许占位符类型在任何地方引用并且从而用于模板特化。
参考图2-4,流程200-400阐明了本发明的多个实施例使用的示例性函数。虽然流程200-400公开了具体的功能块(“blocks”),这样的步骤是示例。即实施例可以很好地适用于执行多个其他块或流程200-400中叙述的块的变化。应领会,流程200-400中的块可以以不用于所展示的顺序执行,并且并不是流程200-400中的所有块都要执行。
参考图2,图2示出了根据各种实施例的、用于主机代码和设备代码的示范性计算机控制器编译过程200。图2描述了多个处理阶段以及数据流,在主机编译器和设备编译器之间传递信息用于GPGPU程序的编译。示范性的过程200包括两部分:1)设备代码编译和2)主机代码编译。
设备模式编译包含访问GPGPU源代码(例如,cu文件)和处理它以生成设备(例如GPU)机器代码。例如,设备机器代码可以是并行线程执行(PTX)机器代码或cubin(CUDA二进制)文件的一部分。在主机代码编译期间,所生成的设备机器代码被***到主机源代码中,然后被送至主机编译器。然后主机编译器编译主机源代码部分并输出可执行文件。
GPGPU源代码202包括主机源代码204和设备源代码206。设备源代码206由设备编译器222访问,其将设备源代码206编译成设备机器代码208。在编译设备源代码206期间,设备代码入口函数模板可以用包含关联于设备lambda表达式的闭包类型的模板参数实例化。注意到,对于用在设备代码入口函数模板的实例化中的设备lambda表达式,设备编译器注意到闭包类型并且能够生成符合C++规则的模板实例化。
主机源代码204由GPGPU编译器前端220访问,其处理主机源代码204以生成经修改的源代码210。在一些实施例中,GPGPU编译器前端220移除设备函数并且保持来自GPGPU源代码202的主机函数。设备机器代码208由GPGPU编译器前端220***或嵌入经修改的源代码210(例如,代替设备函数)的合适部分。对于在主机代码中定义的用于设备代码入口函数模板实例化(例如,在设备上执行的CUDA C++__global__template function实例化)的设备lambda表达式,lambda表达式用占位符类型实例化的实例代替。例如,设备lambda表达式用唯一占位符类型实例化的实例代替,如此处描述的。在一些实施例中,GPGPU编译器前端220确定对每个设备lambda表达式是否有变量被捕获,并用每个所捕获的变量的各自的字段作为模板参数创建占位符类型模板实例化。与设备lambda表达式相关联的占位符实例化的模板参数可以包括封闭设备lambda表达式的函数的类型、封闭函数的地址、与设备lambda表达式相关联的唯一识别符(例如,唯一分配的整数)以及被设备lambda表达式捕获的任何变量。
基于任何设备lambda表达式的占位符类型实例化替代,经修改的源代码210被送至主机编译器224或由其访问。在处理经修改的主机源代码210期间,主机编译器224不会看到设备lambda表达式而是构造器调用,以创建占位符类型的实例或与所代替的lambda表达式相关联的唯一占位符类型实例化。经修改的源代码210由主机编译器224访问并编译,以生成可执行文件230。
图3示出了根据各种实施例的、用于处理源代码文件的示范性计算机控制的过程的流程图。图3描述了修改主机代码(例如,通过GPGPU编译器前端220)的过程300,以允许编译包括在主机代码中定义的设备lambda表达式的主机代码,主机代码中定义的设备lambda表达式与各自的设备代码入口函数模板实例化相关联。
在块302,访问代码的一部分。所访问的代码的一部分可以是主机源代码(例如,主机源代码204)。
在块304,确定代码的该部分是否包括设备lambda表达式。在一些实施例中,确定主机源代码的一部分是否包括与设备代码入口函数模板实例化相关联的设备lambda表达式。如果代码的该部分包括设备lambda表达式,执行块306。如果代码的该部分不包括设备lambda表达式,执行块314。
在块306,确定唯一占位符类型实例化。如在此文中描述的,基于设备lambda表达式确定唯一占位符类型实例化。例如,基于封闭设备lambda表达式的函数的类型、封闭设备lambda表达式的函数的地址、与一个或多个设备lambda表达式的变量相关联的一个或多个类型,以及一个或多个设备lambda表达式的变量可以确定唯一占位符类型实例化。
在块308,基于唯一占位符类型实例化确定设备代码入口函数模板特化。如此文中描述的,设备代码入口函数模板特化可以包括与唯一占位符类型实例化相关联的唯一识别的信息,包括但不限于封闭设备lambda表达式的函数的类型、封闭设备lambda表达式的函数的地址、与一个或多个设备lambda表达式的变量相关联的一个或多个类型,以及一个或多个设备lambda表达式的变量。
在块310,lambda表达式用唯一占位符类型实例化的实例代替。如在此文中描述的,设备lambda表达式用唯一占位符类型实例化的实例代替,以提供要用于由主机编译器编译主机源代码的类型。
在块312,基于唯一占位符类型实例化修改设备代码入口函数模板特化。通过设备代码入口函数模板特化基于占位符类型实例化,设备代码入口函数模板特化被修改,以允许通过占位符类型实例化从主机代码调用以前的设备lambda表达式功能。
在块314,确定是否一部分代码有待处理。如果有一部分部分代码有待处理,则执行块302。如果不存在一部分代码有待处理,则执行块316。
在块316,代码被发送至主机编译器。如此文中描述的,代码已经被修改以允许编译代码,包括与设备代码入口函数模板实例化相关联的主机代码中的设备lambda表达式,而无需改变主机编译器。
图4示出了根据各种实施例的、用于确定并采用唯一占位符类型修改源代码文件的示范性计算机控制的过程的流程图。图4描述了对一个或多个设备lambda表达式的每一个确定唯一占位符类型实例化并代替设备lambda表达式的过程。应领会,虽然讨论涉及在主机函数中定义的设备lambda表达式时,但对于在函数外部的lambda表达式,实施例可以支持用唯一占位符类型实例化代替lambda表达式。
占位符类型是基于模板的,并且由于源代码文件(例如,主机源代码文件)中可能存在多个lambda表达式,每个lambda表达式都用唯一占位符类型代替。模板(例如,通用模板)基于lambda表达式用不同的参数实例化,以生成与lambda表达式相关联的唯一类型,如此文中描述的。
在编译过程期间,设备lambda表达式将从将要发送至主机编译器的代码中移除并且用设备机器代码(例如,GPU集合或GPU函数)和占位符类型实例化代替。因而主机编译器将看到对占位符类型函数的构造器调用,设备lambda表达式先前位于该占位符类型函数。
主机代码被编译并且基于所编译的主机代码指导设备执行与设备lambda表达式相关联的设备机器代码的部分。例如,当主机代码在运行时间执行时,基于占位符类型调用correct__global__特化,并且该特化转而激发GPU上的执行,最终导致GPU上设备lambda表达式的执行。
设备lambda表达式的变量被捕获并传递到设备代码入口函数(例如,利用字节拷贝)。然后设备代码入口函数执行用占位符类型的实例化代替的设备lambda表达式的功能。例如,占位符类型对象具有构造器函数,其保存在原始lambda表达式中捕获的变量的值。然后占位符类型实例被传递到设备代码入口函数模板实例化。设备代码入口函数模板实例化将在该设备(例如,GPU)上执行,并且执行与设备lambda表达式相关联的运算。
参考图4,过程400包括利用占位符类型模板为每个设备lambda表达式创建唯一占位符模板实例化。主机函数可以包括许多设备lambda表达式,每个设备lambda表达式具有不同的捕获列表(例如,参数或形参(parameter))。例如,参考表III,可能有捕获‘x’的lambda表达式,而另一个lambda表达式捕获‘x’和‘d’,以及另一个lambda表达式未捕获变量。
在块402,访问具有与设备代码入口函数模板实例化相关联的设备lambda表达式的主机代码。在一个实施例中,GPGPU(例如CUDA)编译器前端(例如,GPGPU编译器前端220)访问GPGPU源代码文件的主机源代码部分。
在块404,对于设备lambda表达式确定唯一识别符。在一个实施例中,对在每个主机函数中的每个设备lambda表达式确定唯一整数。唯一识别符可以是基于源代码中的设备lambda表达式的位置。例如,计数器可以用于随着主机源代码被处理而将唯一识别符分配给每个设备lambda表达式并在每次确定并分配的唯一识别符之后被增量。
在块406,确定封闭的函数地址。封闭设备lambda表达式的函数的地址被确定。参考表III,主函数对设备lambda表达式是封闭函数,并且确定主函数(例如,&main)的地址。
在块408,确定封闭函数的类型。封闭设备lambda表达式的函数的类型(例如,返回类型)被确定。参考表III,确定主函数具有int(*)(void)类型。封闭函数的类型以及地址的确定允许超负荷(例如,具有不同的类型和/或参数的多个函数被定义为相同的名称)的封闭函数中存在差异。
在块410,确定与设备lambda表达式相关联的任意变量。对每个设备lambda表达式,相关联的变量被确定并且将用作部分占位符。例如,参考表III,‘x’和‘d’变量是设备lambda表达式的相关联变量。
在块412,确定与设备lambda表达式相关联的任意变量的各自的类型。在一些实施例中,使用函数decltype(变量)。例如,参考表III,decltype(x)的结果会是int以及decltype(d)的结果会是double。
在块414,设备lambda表达式被代替。每个设备lambda表达式可以由占位符类型实例化代替。占位符类型实例化可以配置为调用占位符类型构造器代码。
根据过程400,来自表III、经过处理以后被传递到主机编译器的代码至此可以具有下列“main”的主体:
int main(void)
{
int x=10;
double d=20;
auto lam1=__placeholder_t<int(*)(void),(&main),1,<decl type(x)>,
<decl type(d)>(x,d);
foo<<<1,1>>(lam1);
}
表IV-传递到主机编译器的示范性代码部分
‘__placeholder_t’是占位符命名空间范围类型模板。实例化的参数是封闭函数的类型(例如,int(*)(void))、封闭函数的地址(例如,&main),跟着与lambda表达式相关联的唯一整数常量(例如,1),跟着捕获的变量的类型,以及捕获的变量(例如,x,d)。捕获的变量(例如,x,d)的值通过值被传递到占位符类型的构造器,并且该构造器将该值保存为成员对象。
在块416,确定设备代码入口函数模板特化。基于参数和用于设备lambda表达式的占位符模板的信息确定设备代码入口函数模板特化。
在块418,所确定的设备代码入口函数模板特化用于更新主机代码。例如,基于占位符类型实例化,设备代码入口函数模板实例化用设备代码入口函数模板特化代替。然后经修改的主机源代码可以发送至(本地)主机编译器。
根据过程400,来自表III、经过处理以后被传递到主机编译器的代码可以具有下列模板特化:
template<>
void__placeholder__device_stub_foo<__placeholder_t<int(*)(),(&main),1,int,double>>(__placeholder_t<int(*)(),(&main),1,int,double>&__cuda_0)
{
//code to do kernel launch
}
表V-传递到主机编译器的示范性设备代码入口函数模板特化
注意到,在表V中,代替试图引用与设备lambda表达式相关联的不可访问且未命名的闭包类型,而是模板特化参数引用占位符类型模板(例如,__placeholder_t)的实例化。
用这种方式,实施例能够支持设备代码入口函数模板实例化,其中用于实例化的模板参数引用在主机代码中定义的设备lambda表达式。未知的lambda表达式的闭包类型的名称用已命名的占位符类型代替,其中占位符类型是类型模板的实例化。然后已命名的占位符实例化的类型用在设备代码入口函数模板特化中。
实施例允许这种支持不需要对主机编译器的任何改变。设备lambda表达式的定义由命名空间范围“占位符类型”模板的实例代替。占位符类型实例化是每个lambda表达式唯一的(例如,基于lambda表达式的封闭函数地址和相关的唯一整数常量来将lambda表达式和在相同函数中定义的其他lambda表达式区分开)。捕获的变量的类型被添加到占位符类型的模板实例化参数中,从而允许支持从封闭原始lambda表达式的范围(例如,函数范围)中捕获变量。通过显式地将它们的值传递到占位符类型的模板实例化的构造器来捕获变量。
图5说明了由本发明的各种实施例使用的示范性部件。虽然在计算***环境500中公开了具体部件,应领会到,这些部件是示例。即,本发明的实施例很好地适合具有多个其他部件或计算***环境500中叙述的部件的变化。领会到,计算***环境500中的部件可以用呈现的那些以外的其他部件操作,并且不是需要***500的所有部件来完成计算***环境500的目标。
图5示出了根据本发明的一个实施例的、示范性计算机***和相应模块的框图。参考图5,用于实现实施例的示范性***模块包括通用计算***环境,例如计算***环境500。计算***环境500可以包括但不限于服务器、台式计算机、膝上型电脑、平板电脑、移动设备和智能手机。在它的最基础的配置中,计算***环境500典型地包括至少一个处理单元502和计算机可读存储介质504。依赖于精密的配置和计算***环境的类型,计算机可读存储介质504可以是易失的(比如RAM)、非易失的(比如ROM、闪存等)或若干二者的结合。计算机可读存储介质504的部分当被执行时对GPGPU程序(例如,过程200-500)执行编译代码。
此外,计算***环境500也可以具有附加的特征/功能。例如,计算***环境500也可以包括附加的存储器(可移除的和/或不可移除的),包括但不限于磁盘或光盘或磁带。这种附加的存储器在图中通过可移除存储器508和不可移除存储器510说明。计算机存储介质包括在任何方法或技术中实现的、用于存储诸如计算机可读指令、数据结构、程序模块或其他数据的信息的易失的和非易失的、可移除的和不可移除的介质。计算机可读介质504、可移除存储器508和不可移除存储器510都是计算机存储介质的示例。计算机存储介质包括但不限于RAM、ROM、EEPROM、闪存或其他存储技术、CD-ROM、数字多功能磁盘(DVD)或其他光存储器、磁带盒、磁带、磁盘存储器或其他磁存储设备,或可用于储存想要的信息的任何其他介质,并且其可由计算***环境500访问。任何这样的计算机存储介质可以是计算***环境500的一部分。
计算***环境500也可以包括通信连接512,其允许它与其他设备通信。通信连接512是通信介质的一个示例。通信介质典型地包括计算机可读指令、数据结构、程序模块或在调制数据信号诸如载波或其他传输机制中的其他数据,并且包括任何信息传送介质。本文中使用的术语计算机可读介质包括存储介质和通信介质两者。
通信连接512可以允许计算***环境500通过多种网络类型通信,其包括但不限于光纤通道、小型计算机***接口(SCSI)、蓝牙、以太网、Wi-fi、红外数据关联(IRDA)、局域网(LAN)、无线局域网(WLAN)、广域网(WAN)诸如互联网、串口和通用串行总线(USB)。领会到,通信连接512连接到的多种网络类型可以运行多种网络协议,包括但不限于传输控制协议(TCP)、互联网协议(IP)、实时传输协议(RTP)、实时传输控制协议(RTCP)、文件传输协议(FTP)以及超文本传输协议(HTTP)。
计算***环境500也可以具有输入设备514,比如键盘、鼠标、笔、语音输入设备、触摸输入设备、远程控制等。也可以包括输出设备516诸如显示器、扬声器等。所有这些设备都是本领域已知的并且不详细讨论。
在一个实施例中,计算机可读存储介质504包括异构编译器506、源代码数据存储550、主机编译器560和GPGPU程序580。源代码数据存储550(例如,文件、储存库、数据库等)可以包括源代码,包括主机源代码和设备源代码,用于编译成具有主机(例如CPU)部分582和设备(例如GPU)部分584的GPGPU程序。在一些实施例中,主机源代码可操作为被编译用于在CPU上执行,以及设备源代码可操作为被编译以在GPU上执行。主机编译器560配置为编译主机源代码,如本文中描述的。在编译来自源代码数据存储550中的主机源代码以后,由主机编译器560输出GPGPU程序580,如本文中描述的。注意到,计算机可读存储介质504的模块和部件也可以在硬件中实现。
异构编译器506包括编译器前端518和设备编译器540。异构编译器506配置为利用设备编译器540编译设备源代码,并且利用编译器前端518修改主机源代码。如本文中描述的,编译器前端518配置为修改主机源代码,以允许编译包括与设备代码入口函数模板实例化相关联的lambda表达式的主机源代码。设备编译器540(例如,设备编译器222)配置为将设备源代码编译成设备机器代码,如本文中描述的。
编译器前端518包括代码访问模块520、设备lambda表达式识别模块522、占位符类型实例化确定模块524、设备代码入口函数模板特化确定模块526、代码修改模块528和代码发送模块530。
代码访问模块520可操作为从数据存储例如源代码数据存储550中访问源代码。源代码包括主机源代码和设备源代码。设备lambda表达式识别模块522配置为识别与设备代码入口函数模板实例化相关联的主机源代码部分中的设备lambda表达式。占位符类型实例化确定模块524配置为基于设备lambda表达式确定占位符类型实例化,如本文中描述的。在一些实施例中,占位符类型实例化包括与封闭lambda表达式的函数相关联的类型。在一些实施例中,占位符类型实例化包括封闭lambda表达式的函数的地址。在一些实施例中,占位符类型实例化包括与lambda表达式相关联的唯一识别符。在一些实施例中,占位符类型实例化包括一个或多个与lambda表达式相关联的变量。在一些实施例中,占位符类型实例化通过将一个或多个变量的值显式地传递到占位符类型实例化的构造器来捕获一个或多个与lambda表达式相关联的变量。
代码修改模块528配置为修改主机源代码,以用占位符类型实例化代替设备lambda表达式。设备代码入口函数模板特化确定模块526配置为基于设备lambda表达式确定设备代码入口函数模板特化。在一些实施例中,代码修改模块528进一步配置为基于由设备代码入口函数模板特化确定模块526确定的设备代码入口函数模板特化来修改主机代码。代码发送模块530配置为将由代码修改模块修改的主机源代码发送至主机编译器。
本发明的具体实施例的前面的描述以说明和描述的意图呈现。它们不意图详尽阐述或限制本发明为公开的精确形式,并且根据上面的教导许多修改和变化是可能的。实施例被选择和描述是为了更好地解释本发明的原理和它的实际应用,从而使得本领域技术人员能够更好地利用本发明和具有多种修改的多个实施例适于特殊用途的考虑。意图由附加的权利要求和它们的等价物限定本发明的保护范围。
Claims (20)
1.一种处理源代码的方法,所述方法包括:
访问主机源代码的一部分;
确定所述主机源代码的一部分是否包括设备拉姆达表达式;
响应于所述主机代码的一部分包括所述设备拉姆达表达式,基于所述设备拉姆达表达式确定唯一占位符类型实例化;
基于所述唯一占位符类型实例化修改所述设备拉姆达表达式以生成经修改的主机源代码;以及
使得所述经修改的主机源代码对主机编译器可用。
2.根据权利要求1所述的方法,其中所述设备拉姆达表达式与设备代码入口函数模板实例化相关联。
3.根据权利要求1所述的方法,其中所述唯一占位符类型实例化包括与封闭所述设备拉姆达表达式的函数相关联的类型。
4.根据权利要求3所述的方法,其中所述唯一占位符类型实例化包括封闭所述设备拉姆达表达式的所述函数的地址。
5.根据权利要求1所述的方法,其中所述唯一占位符类型实例化包括与所述设备拉姆达表达式相关联的唯一识别符。
6.根据权利要求1所述的方法,其中所述唯一占位符类型实例化包括一个或多个与所述设备拉姆达表达式相关联的变量。
7.根据权利要求6所述的方法,其中所述唯一占位符类型实例化通过将所述一个或多个变量的值显式地传递到所述占位符类型实例化的构造器来捕获所述一个或多个与所述设备拉姆达表达式相关联的变量。
8.根据权利要求1所述的方法,进一步包括:
基于与所述设备拉姆达表达式相关联的所述唯一占位符类型实例化确定设备代码入口函数模板特化,其中所述主机源代码包括设备代码入口函数模板;以及
基于所述设备代码入口函数模板特化修改所述设备代码入口函数模板。
9.一种用于编译代码的***,所述***包括:
代码访问模块,其可操作为从数据存储访问源代码,其中所述源代码包括主机源代码和设备源代码;
设备拉姆达表达式识别模块,其配置为识别所述主机源代码的一部分中的设备拉姆达表达式;
占位符类型确定模块,其配置为基于所述设备拉姆达表达式确定占位符类型实例化;以及
代码修改模块,其配置为修改所述主机源代码以用所述占位符类型实例化代替所述设备拉姆达表达式。
10.根据权利要求9所述的***,进一步包括:
模板特化确定模块,其配置为基于所述设备拉姆达表达式确定设备代码入口函数模板特化。
11.根据权利要求10所述的***,其中所述代码修改模块进一步配置为基于由所述模板特化确定模块确定的所述设备代码入口函数模板特化修改所述主机源 代码。
12.根据权利要求9所述的***,进一步包括:
代码发送模块,其配置为将由所述代码修改模块修改的所述主机源代码发送至主机编译器。
13.根据权利要求9所述的***,其中所述占位符类型实例化包括与封闭所述设备拉姆达表达式的函数相关联的类型。
14.根据权利要求13所述的***,其中所述占位符类型实例化包括封闭所述设备拉姆达表达式的所述函数的地址。
15.根据权利要求9所述的***,其中所述占位符类型实例化包括与所述设备拉姆达表达式相关联的唯一识别符。
16.根据权利要求9所述的***,其中所述占位符类型实例化包括一个或多个与所述设备拉姆达表达式相关联的变量。
17.根据权利要求16所述的***,其中所述占位符类型实例化通过将所述一个或多个变量的值显式地传递到所述占位符类型实例化的构造器来捕获所述一个或多个与所述设备拉姆达表达式相关联的变量。
18.一种用于修改代码用于编译的方法,包括:
访问主机源代码的一部分;
确定所述主机源代码的一部分是否包括与设备代码入口函数模板实例化相关联的设备拉姆达表达式;
基于所述设备拉姆达表达式确定唯一占位符类型实例化;
基于所述唯一占位符类型实例化确定设备代码入口函数模板特化,其中所述主机源代码包括设备代码入口函数模板;以及
通过用所述唯一占位符类型实例化的实例代替所述设备拉姆达表达式生成经修改的主机源代码;
基于所述唯一占位符类型实例化生成设备代码入口函数特化;以及
使得所述经修改的主机源代码对主机编译器可用。
19.根据权利要求18所述的方法,其中所述唯一占位符类型实例化包括与封闭所述设备拉姆达表达式的函数相关联的类型。
20.如权利要求19所述的方法,其中所述唯一占位符类型实例化包括封闭所述设备拉姆达表达式的函数的地址以及与所述设备拉姆达表达式相关联的识别的唯一整数。
Applications Claiming Priority (4)
Application Number | Priority Date | Filing Date | Title |
---|---|---|---|
US201462097528P | 2014-12-29 | 2014-12-29 | |
US62/097,528 | 2014-12-29 | ||
US14/968,740 | 2015-12-14 | ||
US14/968,740 US10241761B2 (en) | 2014-12-29 | 2015-12-14 | System and method for compiler support for compile time customization of code |
Publications (2)
Publication Number | Publication Date |
---|---|
CN105740036A CN105740036A (zh) | 2016-07-06 |
CN105740036B true CN105740036B (zh) | 2019-04-23 |
Family
ID=56116836
Family Applications (1)
Application Number | Title | Priority Date | Filing Date |
---|---|---|---|
CN201511004391.5A Active CN105740036B (zh) | 2014-12-29 | 2015-12-29 | 用于支持代码的编译时间定制的编译器的***和方法 |
Country Status (3)
Country | Link |
---|---|
US (2) | US10241761B2 (zh) |
CN (1) | CN105740036B (zh) |
DE (1) | DE102015122908A1 (zh) |
Families Citing this family (5)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
US10087316B2 (en) | 2008-04-29 | 2018-10-02 | The Procter & Gamble Company | Polymeric compositions and articles comprising polylactic acid and polyolefin |
US10802855B2 (en) * | 2016-09-16 | 2020-10-13 | Oracle International Corporation | Producing an internal representation of a type based on the type's source representation |
CN107908407B (zh) * | 2017-12-11 | 2021-09-07 | 北京奇虎科技有限公司 | 编译方法、装置及终端设备 |
JP6839673B2 (ja) * | 2018-02-16 | 2021-03-10 | 日本電信電話株式会社 | アプリケーション分割装置、方法およびプログラム |
CN112232003B (zh) * | 2020-12-17 | 2021-03-30 | 芯华章科技股份有限公司 | 对设计进行仿真的方法、电子设备及存储介质 |
Citations (7)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
CN101233487A (zh) * | 2005-07-29 | 2008-07-30 | 微软公司 | 拉姆达表达式 |
CN101963918A (zh) * | 2010-10-26 | 2011-02-02 | 上海交通大学 | 实现cpu/gpu异构平台的虚拟执行环境的方法 |
CN102541618A (zh) * | 2010-12-29 | 2012-07-04 | ***通信集团公司 | 一种通用图形处理器虚拟化的实现方法、***及装置 |
CN103124956A (zh) * | 2010-09-27 | 2013-05-29 | 三星电子株式会社 | 在异类***中使用虚拟化编译和执行应用的方法和装置 |
CN103389908A (zh) * | 2012-05-09 | 2013-11-13 | 辉达公司 | 用于分开编译嵌入在主机代码中的设备代码的方法和*** |
CN103858099A (zh) * | 2011-08-02 | 2014-06-11 | 国际商业机器公司 | 用于在异构计算机上编译和运行高级程序的技术 |
CN104035781A (zh) * | 2014-06-27 | 2014-09-10 | 北京航空航天大学 | 一种快速开发异构并行程序的方法 |
Family Cites Families (39)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
US4558413A (en) * | 1983-11-21 | 1985-12-10 | Xerox Corporation | Software version management system |
US20020143929A1 (en) * | 2000-12-07 | 2002-10-03 | Maltz David A. | Method and system for collection and storage of traffic data from heterogeneous network elements in a computer network |
US7299419B2 (en) * | 2001-09-28 | 2007-11-20 | Business Objects, S.A. | Apparatus and method for combining discrete logic visual icons to form a data transformation block |
US7140007B2 (en) * | 2002-01-16 | 2006-11-21 | Xerox Corporation | Aspect-oriented programming with multiple semantic levels |
US20080313282A1 (en) * | 2002-09-10 | 2008-12-18 | Warila Bruce W | User interface, operating system and architecture |
US8296665B2 (en) * | 2004-05-11 | 2012-10-23 | Sap Ag | Developing and executing applications with configurable patterns |
US20070028222A1 (en) * | 2005-07-29 | 2007-02-01 | Microsoft Corporation | Free/outer variable capture |
US8291315B2 (en) * | 2006-02-28 | 2012-10-16 | Ricoh Co., Ltd. | Standardized network access to partial document imagery |
US7711546B2 (en) * | 2006-04-21 | 2010-05-04 | Microsoft Corporation | User interface for machine aided authoring and translation |
US8146066B2 (en) * | 2006-06-20 | 2012-03-27 | Google Inc. | Systems and methods for caching compute kernels for an application running on a parallel-processing computer system |
US8739137B2 (en) * | 2006-10-19 | 2014-05-27 | Purdue Research Foundation | Automatic derivative method for a computer programming language |
US8396827B2 (en) * | 2006-12-29 | 2013-03-12 | Sap Ag | Relation-based hierarchy evaluation of recursive nodes |
US7676461B2 (en) * | 2007-07-18 | 2010-03-09 | Microsoft Corporation | Implementation of stream algebra over class instances |
US20100235730A1 (en) * | 2009-03-13 | 2010-09-16 | Microsoft Corporation | Consume-first mode text insertion |
TWI393067B (zh) * | 2009-05-25 | 2013-04-11 | Inst Information Industry | 具有電源閘控功能之繪圖處理系統及電源閘控方法,及其電腦程式產品 |
EP2494454A4 (en) * | 2009-10-30 | 2013-05-15 | Intel Corp | TWO-WAY COMMUNICATION SUPPORT FOR HETEROGENIC PROCESSORS ENERGY COMPUTER PLATFORM |
US20110212761A1 (en) * | 2010-02-26 | 2011-09-01 | Igt | Gaming machine processor |
US8589867B2 (en) * | 2010-06-18 | 2013-11-19 | Microsoft Corporation | Compiler-generated invocation stubs for data parallel programming model |
US8756590B2 (en) * | 2010-06-22 | 2014-06-17 | Microsoft Corporation | Binding data parallel device source code |
US8488164B2 (en) * | 2010-08-13 | 2013-07-16 | Sap Ag | Mobile Printing solution |
US8380700B2 (en) * | 2011-04-13 | 2013-02-19 | Cisco Technology, Inc. | Ad hoc geospatial directory of users based on optimizing non-Turing complete executable application |
US8533698B2 (en) * | 2011-06-13 | 2013-09-10 | Microsoft Corporation | Optimizing execution of kernels |
US8612279B2 (en) * | 2011-06-27 | 2013-12-17 | Cisco Technology, Inc. | Ad hoc generation of work item entity for geospatial entity based on symbol manipulation language-based workflow item |
US8842124B2 (en) * | 2011-10-26 | 2014-09-23 | Google Inc. | Declarative interface for developing test cases for graphics programs |
US9514507B2 (en) * | 2011-11-29 | 2016-12-06 | Citrix Systems, Inc. | Methods and systems for maintaining state in a virtual machine when disconnected from graphics hardware |
US10705804B2 (en) * | 2012-02-26 | 2020-07-07 | Logistics Research Centre SIA | Strongly typed metadata access in object oriented programming languages with reflection support |
US9075913B2 (en) * | 2012-02-27 | 2015-07-07 | Qualcomm Incorporated | Validation of applications for graphics processing unit |
US10025643B2 (en) * | 2012-05-10 | 2018-07-17 | Nvidia Corporation | System and method for compiler support for kernel launches in device code |
US9086937B2 (en) * | 2012-05-16 | 2015-07-21 | Apple Inc. | Cloud-based application resource files |
US8996504B2 (en) * | 2012-05-24 | 2015-03-31 | Sybase, Inc. | Plan caching using density-based clustering |
US9557974B2 (en) * | 2012-07-10 | 2017-01-31 | Oracle International Corporation | System and method for supporting compatibility checking for lambda expression |
US20150009222A1 (en) * | 2012-11-28 | 2015-01-08 | Nvidia Corporation | Method and system for cloud based virtualized graphics processing for remote displays |
US9176716B2 (en) * | 2013-04-18 | 2015-11-03 | Siemens Aktiengesellschaft | Method and apparatus for exploiting data locality in dynamic task scheduling |
US9971576B2 (en) * | 2013-11-20 | 2018-05-15 | Nvidia Corporation | Software development environment and method of compiling integrated source code |
US10152312B2 (en) * | 2014-01-21 | 2018-12-11 | Nvidia Corporation | Dynamic compiler parallelism techniques |
US9558094B2 (en) * | 2014-05-12 | 2017-01-31 | Palo Alto Research Center Incorporated | System and method for selecting useful smart kernels for general-purpose GPU computing |
US10430169B2 (en) * | 2014-05-30 | 2019-10-01 | Apple Inc. | Language, function library, and compiler for graphical and non-graphical computation on a graphical processor unit |
US9830134B2 (en) * | 2015-06-15 | 2017-11-28 | Qualcomm Incorporated | Generating object code from intermediate code that includes hierarchical sub-routine information |
US11606198B2 (en) * | 2020-01-22 | 2023-03-14 | Valimail Inc. | Centrally managed PKI provisioning and rotation |
-
2015
- 2015-12-14 US US14/968,740 patent/US10241761B2/en not_active Expired - Fee Related
- 2015-12-29 DE DE102015122908.0A patent/DE102015122908A1/de not_active Ceased
- 2015-12-29 CN CN201511004391.5A patent/CN105740036B/zh active Active
-
2019
- 2019-02-27 US US16/287,392 patent/US20190196797A1/en active Pending
Patent Citations (7)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
CN101233487A (zh) * | 2005-07-29 | 2008-07-30 | 微软公司 | 拉姆达表达式 |
CN103124956A (zh) * | 2010-09-27 | 2013-05-29 | 三星电子株式会社 | 在异类***中使用虚拟化编译和执行应用的方法和装置 |
CN101963918A (zh) * | 2010-10-26 | 2011-02-02 | 上海交通大学 | 实现cpu/gpu异构平台的虚拟执行环境的方法 |
CN102541618A (zh) * | 2010-12-29 | 2012-07-04 | ***通信集团公司 | 一种通用图形处理器虚拟化的实现方法、***及装置 |
CN103858099A (zh) * | 2011-08-02 | 2014-06-11 | 国际商业机器公司 | 用于在异构计算机上编译和运行高级程序的技术 |
CN103389908A (zh) * | 2012-05-09 | 2013-11-13 | 辉达公司 | 用于分开编译嵌入在主机代码中的设备代码的方法和*** |
CN104035781A (zh) * | 2014-06-27 | 2014-09-10 | 北京航空航天大学 | 一种快速开发异构并行程序的方法 |
Also Published As
Publication number | Publication date |
---|---|
CN105740036A (zh) | 2016-07-06 |
US20160188352A1 (en) | 2016-06-30 |
DE102015122908A1 (de) | 2016-06-30 |
US10241761B2 (en) | 2019-03-26 |
US20190196797A1 (en) | 2019-06-27 |
Similar Documents
Publication | Publication Date | Title |
---|---|---|
CN105740036B (zh) | 用于支持代码的编译时间定制的编译器的***和方法 | |
US8756590B2 (en) | Binding data parallel device source code | |
CN106716361B (zh) | 用于运行时例程冗余跟踪的编译器高速缓存 | |
US10255045B2 (en) | Graphical representation of data in a program code editor | |
US9632761B2 (en) | Distribute workload of an application to a graphics processing unit | |
TWI556170B (zh) | 將作業系統之原始應用程式介面投射至其它程式語言(二) | |
CN102402451B (zh) | 用户定义类型的编译时边界检查 | |
Walls | Embedded software: the works | |
CN101937367B (zh) | Mpi源代码程序到基于mpi线程的程序的自动转换 | |
US8935686B2 (en) | Error-code and exception-based function dispatch tables | |
US10452369B2 (en) | Code profiling of executable library for pipeline parallelization | |
Johnson et al. | Linux application development | |
Mackey | Introducing. NET 4.0: With Visual Studio 2010 | |
Sandrieser et al. | Explicit platform descriptions for heterogeneous many-core architectures | |
US9830204B2 (en) | Facilitating communication between software components that use middleware | |
Freed et al. | Dynamic, Instance-based, object-oriented programming in Max/MSP using open sound control message delegation. | |
Deepika et al. | Automatic program generation for heterogeneous architectures | |
Tzanetakis et al. | Interoperability and the Marsyas 0.2 runtime | |
Ng et al. | Session types: towards safe and fast reconfigurable programming | |
Blesel et al. | Heimdallr: Improving compile time correctness checking for message passing with Rust | |
Fisher | ARM® Cortex® M4 Cookbook | |
Nestor Ribeiro et al. | An automated model based approach to mobile UI specification and development | |
US8135943B1 (en) | Method, apparatus, and computer-readable medium for generating a dispatching function | |
Browning et al. | C++ 20 Recipes: A Problem-solution Approach | |
US20240143296A1 (en) | METHODS AND APPARATUS FOR COMBINING CODE LARGE LANGUAGE MODELS (LLMs) WITH COMPILERS |
Legal Events
Date | Code | Title | Description |
---|---|---|---|
C06 | Publication | ||
PB01 | Publication | ||
C10 | Entry into substantive examination | ||
SE01 | Entry into force of request for substantive examination | ||
GR01 | Patent grant | ||
GR01 | Patent grant |