CN112817738A - 用于修改可执行图以实施与新任务图关联的工作负载的技术 - Google Patents

用于修改可执行图以实施与新任务图关联的工作负载的技术 Download PDF

Info

Publication number
CN112817738A
CN112817738A CN202011280064.3A CN202011280064A CN112817738A CN 112817738 A CN112817738 A CN 112817738A CN 202011280064 A CN202011280064 A CN 202011280064A CN 112817738 A CN112817738 A CN 112817738A
Authority
CN
China
Prior art keywords
graph
task
executable
topology
memory
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
CN202011280064.3A
Other languages
English (en)
Inventor
S·古芬克尔
S·琼斯
D·丰泰内
S·史蒂文森
P·库尔卡尼
J·盖斯亚
傅德禹
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.)
Nvidia Corp
Original Assignee
Nvidia Corp
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
Priority claimed from US16/732,014 external-priority patent/US20210149734A1/en
Application filed by Nvidia Corp filed Critical Nvidia Corp
Publication of CN112817738A publication Critical patent/CN112817738A/zh
Pending legal-status Critical Current

Links

Images

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/46—Multiprogramming arrangements
    • G06F9/50—Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5061—Partitioning or combining of resources
    • G06F9/5072—Grid computing
    • 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/46—Multiprogramming arrangements
    • G06F9/50—Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5005—Allocation of resources, e.g. of the central processing unit [CPU] to service a request
    • G06F9/5027—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resource being a machine, e.g. CPUs, Servers, Terminals
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/30—Creation or generation of source code
    • G06F8/35—Creation or generation of source code model driven
    • 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/45—Exploiting coarse grain parallelism in compilation, i.e. parallelism between groups of instructions
    • 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/46—Multiprogramming arrangements
    • G06F9/50—Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5061—Partitioning or combining of resources
    • G06F9/5066—Algorithms for mapping a plurality of inter-dependent sub-tasks onto a plurality of physical CPUs
    • 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/46—Multiprogramming arrangements
    • G06F9/50—Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5083—Techniques for rebalancing the load in a distributed system
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F2209/00—Indexing scheme relating to G06F9/00
    • G06F2209/50—Indexing scheme relating to G06F9/50
    • G06F2209/509—Offload

Abstract

本公开涉及用于修改可执行图以实施与新任务图关联的工作负载的技术。修改可执行图以实施不同工作负载的技术。在至少一个实施例中，通过将第二任务图的不可执行版本应用到第一任务图的可执行版本而修改第一任务图的可执行版本，使得第一任务图的可执行版本可以实施第二任务图的不可执行版本的第二工作负载。

Description

用于修改可执行图以实施与新任务图关联的工作负载的技术
相关申请的交叉引用
本申请要求2019年11月15日提交并且具有序列号201941046676的题为“用于修改可执行图以实施与新任务图关联的工作负载的技术(TECHNIQUES FOR MODIFYING ANEXECUTABLE GRAPH TO PERFORM A WORKLOAD ASSOCIATED WITH A NEW TASK GRAPH)”的印度临时专利申请的优先权利益。该相关申请的主题据此通过引用合并于此。
技术领域
描述了各种不同的实施例，其总体上涉及并行计算，更具体地，涉及用于修改可执行图以实施与新任务图关联的工作负载的技术。
背景技术
任务图是一种用于对与并行计算任务相应的工作负载建模的有用工具。在任务图中，将计算任务建模为节点，并且将计算任务之间的从属性(dependencies)建模为有向边。为了可用来使得一个或更多个计算资源实施工作负载，将任务图转换为可执行图。于是，工作负载通过以下方式实施：依照可执行图配置一个或更多个计算资源中的每一个以实施任务，在必要的情况下将数据传输至一个或更多个计算资源，以及从一个或更多个计算资源取回结果。可执行图可以多次使用以便让相同的工作负载多次实施。然而，一定的可执行图典型地可用于从其创建该可执行图的任务图的仅仅单个工作负载。
如前所述，本领域所需的是更加有效的使用可执行图以实施与新任务图关联的工作负载的技术。
附图说明
为了可以更详细地理解各种不同实施例的上述特征的方式，可以通过引用其中一些在附图中图示出的实施例而拥有上面简要总结的各种不同实施例的更具体描述。然而，应当指出的是，附图仅仅图示出发明构思的典型实施例，并且因此不应当被认为限制了其范围，因为本发明可以允许其他同样有效的实施例。
图1图示出依照至少一个实施例的示例性数据中心；
图2图示出依照至少一个实施例的处理系统；
图3图示出依照至少一个实施例的计算机系统；
图4图示出依照至少一个实施例的系统；
图5图示出依照至少一个实施例的示例性集成电路；
图6图示出依照至少一个实施例的计算系统；
图7图示出依照至少一个实施例的APU；
图8图示出依照至少一个实施例的CPU；
图9图示出依照至少一个实施例的示例性加速器集成分片(slice)；
图10A-10B图示出依照至少一个实施例的示例性图形处理器；
图11A图示出依照至少一个实施例的图形核心；
图11B图示出依照至少一个实施例的GPGPU；
图12A图示出依照至少一个实施例的并行处理器；
图12B图示出依照至少一个实施例的处理集群；
图12C图示出依照至少一个实施例的图形多处理器；
图13图示出依照至少一个实施例的图形处理器；
图14图示出依照至少一个实施例的处理器；
图15图示出依照至少一个实施例的处理器；
图16图示出依照至少一个实施例的图形处理器核心；
图17图示出依照至少一个实施例的PPU；
图18图示出依照至少一个实施例的GPC；
图19图示出依照至少一个实施例的流式多处理器；
图20图示出依照至少一个实施例的编程平台的软件堆栈；
图21图示出依照至少一个实施例的图STACK(堆栈)A的软件堆栈的CUDA实现方式；
图22图示出依照至少一个实施例的图STACK A的软件堆栈的ROCm实现方式；
图23图示出依照至少一个实施例的图STACK A的软件堆栈的OpenCL实现方式；
图24图示出依照至少一个实施例的受编程平台支持的软件；
图25图示出依照至少一个实施例编译在图20-23的编程平台上执行的代码；
图26更详细地图示出依照至少一个实施例编译在图20-23的编程平台上执行的代码；
图27图示出依照至少一个实施例在编译源代码之前转化源代码；
图28A图示出依照至少一个实施例的被配置为使用不同类型的处理单元编译和执行CUDA源代码的系统；
图28B图示出依照至少一个实施例的被配置为使用CPU和CUDA启用的GPU编译和执行图28A的CUDA源代码的系统；
图28C图示出依照至少一个实施例的被配置为使用CPU和非CUDA启用的GPU编译和执行图28A的CUDA源代码的系统；
图29图示出依照至少一个实施例的通过图28C的CUDA至HIP转化工具转化的示例性内核；
图30更详细地图示出依照至少一个实施例的图28C的非CUDA启用的GPU；
图31图示出依照至少一个实施例如何将示例性CUDA网格的线程映射为图30的不同计算单元；
图32图示出依照至少一个实施例的流程图；
图33A-C图示出依照至少一个实施例的示例性任务图；
图34图示出依照至少一个实施例的流程图；以及
图35图示出依照至少一个实施例的流程图。
具体实施方式
在以下描述中，阐述了许多特定的细节以提供对于至少一个实施例的更加彻底的理解。然而，对于本领域技术人员将显而易见的是，可以在没有这些特定的细节中的一个或更多个的情况下将这些发明构思付诸实践。
数据中心
图1图示出依照至少一个实施例的示例性数据中心100。在至少一个实施例中，数据中心100包括但不限于数据中心基础结构层110、框架层120、软件层130和应用层140。
在至少一个实施例中，如图1所示，数据中心基础结构层110可以包括资源协调器112、分组计算资源114和节点计算资源(“节点C.R.”)116(1)-116(N)，其中“N”表示任何整数、正整数。在至少一个实施例中，节点C.R.116(1)-116(N)可以包括但不限于任意数量的中央处理单元(“CPU”)或者其他处理器(包括加速器、现场可编程门阵列(“FPGA”)、图形处理器等等)、存储器设备(例如，动态只读存储器)、存储设备(例如固态或磁盘驱动器)、网络输入/输出(“NW I/O”)设备、网络交换机、虚拟机(“VM”)、电源模块以及冷却模块等等。在至少一个实施例中，来自节点C.R.116(1)-116(N)之间的一个或更多个节点C.R.可以是具有上面提到的计算资源中的一个或更多个的服务器。
在至少一个实施例中，分组计算资源114可以包括不同地理位置处的数据中心中容纳的许多机架(未示出)或者一个或更多个机架内容纳的节点C.R.的单独分组(也未示出)。分组计算资源114内的节点C.R.的单独分组可以包括分组的计算、网络、存储器或存储资源，其可以被配置或分配为支持一个或更多个工作负载。在至少一个实施例中，包括CPU或处理器的若干节点C.R.可以分组在一个或更多个机架内以提供支持一个或更多个工作负载的计算资源。在至少一个实施例中，一个或更多个机架也可以包括任意组合的任意数量的电源模块、冷却模块和网络交换机。
在至少一个实施例中，资源协调器112可以配置或者以其他方式控制一个或更多个节点C.R.116(1)-116(N)和/或分组计算资源114。在至少一个实施例中，资源协调器112可以包括用于数据中心100的软件设计基础结构(“SCT”)管理实体。在至少一个实施例中，资源协调器112可以包括硬件、软件或者其某种组合。
在至少一个实施例中，如图1所示，框架层120包括但不限于作业调度器132、配置管理器134、资源管理器136和分布式文件系统138。在至少一个实施例中，框架层120可以包括支持软件层130的软件152和/或应用层140的一个或更多个应用142的框架。在至少一个实施例中，软件152或应用142可以分别包括基于web的服务软件或应用，诸如亚马逊Web服务、谷歌云和微软Azure提供的软件或应用。在至少一个实施例中，框架层120可以是，但不限于，一种类型的自由开源软件web应用框架，诸如Apache SparkTM(此后为“Spark”)，其可以将分布式文件系统138用于大规模数据处理(例如“大数据”)。在至少一个实施例中，作业调度器132可以包括便于数据中心100的各层支持的工作负载的调度的Spark驱动程序。在至少一个实施例中，配置管理器134可能能够配置诸如软件层130和框架层120之类的不同层，包括用于支持大规模数据处理的Spark和分布式文件系统138。在至少一个实施例中，资源管理器136可能能够管理被映射为或者被分配用于支持分布式文件系统138和作业调度器132的集群或分组计算资源。在至少一个实施例中，集群或分组计算资源可以包括数据中心基础结构层110处的分组计算资源114。在至少一个实施例中，资源管理器136可以与资源协调器112配合以管理这些映射或分配的计算资源。
在至少一个实施例中，软件层130中包括的软件152可以包括由节点C.R.116(1)-116(N)中的至少部分、分组计算资源114和/或框架层120的分布式文件系统138使用的软件。一种或多种类型的软件可以包括但不限于因特网网页搜索软件、电子邮件病毒扫描软件、数据库软件以及流式视频内容软件。
在至少一个实施例中，应用层140中包括的应用142可以包括由节点C.R.116(1)-116(N)中的至少部分、分组计算资源114和/或框架层120的分布式文件系统138使用的一种或多种类型的应用。一种或多种类型的应用可以包括但不限于CUDA应用。
在至少一个实施例中，配置管理器134、资源管理器136和资源协调器112中的任何一个可以基于以任何技术上可行的方式获取的任意数量和类型的数据实现任意数量和类型的自修改动作。在至少一个实施例中，自修改动作可以使数据中心100的数据中心操作员免于做出可能不好的配置决策依据以及可能避免数据中心的未充分利用和/或性能不佳的部分。
基于计算机的系统
下面的图非限制性地阐述了可以用来实现至少一个实施例的示例性的基于计算机的系统。
图2图示出依照至少一个实施例的处理系统200。在至少一个实施例中，处理系统200包括一个或更多个处理器202以及一个或更多个图形处理器208，并且可以是单处理器桌面系统、多处理器工作站系统或者具有大量处理器202或处理器核心207的服务器系统。在至少一个实施例中，处理系统200为结合到用在移动、手持或嵌入式设备中的片上系统(“SoC”)内的处理平台。
在至少一个实施例中，处理系统200可以包括或者结合到基于服务器的游戏平台、游戏控制台、媒体控制台、移动游戏控制台、手持式游戏控制台或者在线游戏控制台内。在至少一个实施例中，处理系统200为移动电话、智能电话、平板计算设备或者移动因特网设备。在至少一个实施例中，处理系统200也可以包括可穿戴设备，与可穿戴设备耦合，或者集成在可穿戴设备内，所述可穿戴设备例如智能手表可穿戴设备、智能眼镜设备、增强现实设备或者虚拟现实设备。在至少一个实施例中，处理系统200为电视或机顶盒设备，其具有一个或更多个处理器202以及由一个或更多个图形处理器208生成的图形接口。
在至少一个实施例中，一个或更多个处理器202中的每一个包括处理指令的一个或更多个处理器核心207，这些指令在被执行时实施用于系统和用户软件的操作。在至少一个实施例中，一个或更多个处理器核心207中的每一个被配置为处理特定指令集209。在至少一个实施例中，指令集209可以促进复杂指令集计算(“CISC”)、精简指令集计算(“RISC”)或者经由甚长指令字(“VLIW”)计算。在至少一个实施例中，处理器核心207中的每一个可以处理不同的指令集209，所述指令集可以包括便于模仿其他指令集的指令。在至少一个实施例中，处理器核心207也可以包括其他处理设备，例如数字信号处理器(“DSP”)。
在至少一个实施例中，处理器202包括缓存存储器(“缓存”)204。在至少一个实施例中，处理器202可以具有单个内部缓存或者多级内部缓存。在至少一个实施例中，缓存存储器在处理器202的各种不同部件之间共享。在至少一个实施例中，处理器202也使用外部缓存(例如，3级(“L3”)缓存或者最后一级缓存(“LLC”))(未示出)，其可以使用已知的缓存一致性技术在处理器核心207之间共享。在至少一个实施例中，寄存器文件206附加地包括在处理器202中，其可以包括用于存储不同类型的数据的不同类型的寄存器(例如，整数寄存器、浮点寄存器、状态寄存器和指令指针寄存器)。在至少一个实施例中，寄存器文件206可以包括通用寄存器或其他寄存器。
在至少一个实施例中，一个或更多个处理器202与一个或更多个接口总线210耦合以便在处理系统200中的处理器202与其他部件之间传送诸如地址、数据或控制信号之类的通信信号。在至少一个实施例中，接口总线210在一个实施例中可以是处理器总线，例如一种版本的直接媒体接口(“DMI”)总线。在至少一个实施例中，接口总线210不限于DMI总线，并且可以包括一个或更多个外围部件互连总线(例如，“PCI”、PCI Express(“PCIe”)、存储器总线或者其他类型的接口总线。在至少一个实施例中，处理器202包括集成存储器控制器216和平台控制器集线器230。在至少一个实施例中，存储器控制器216促进处理系统200的存储器设备与其他部件之间的通信，而平台控制器集线器(“PCH”)230经由本地I/O总线提供到输入/输出(“I/O”)设备的连接。
在至少一个实施例中，存储器设备220可以是动态随机存取存储器(“DRAM”)设备、静态随机存储存储器(“SRAM”)设备、闪存设备、相变存储器设备或者用作处理器存储器的具有适当性能的某种其他存储器设备。在至少一个实施例中，存储器设备220可以作为用于处理系统200的系统存储器而操作来存储在一个或更多个处理器202执行应用或进程时使用的数据和指令221。在至少一个实施例中，存储器控制器216也与可选的外部图形处理器212耦合，该外部图形处理器可以与处理器202中的一个或更多个图形处理器208通信以实施图形和媒体操作。在至少一个实施例中，显示设备211可以连接到处理器202。在至少一个实施例中，显示设备211可以包括像移动电子设备或膝上型设备中那样的内部显示设备或者经由显示接口(例如显示端口等等)附接的外部显示设备中的一个或更多个。在至少一个实施例中，显示设备211可以包括头戴式显示器(“HMD”)，例如用在虚拟现实(“VR”)应用或增强现实(“AR”)应用中的立体显示设备。
在至少一个实施例中，平台控制器集线器230使得外设能够经由高速I/O总线连接到存储器设备220和处理器202。在至少一个实施例中，I/O包括但不限于音频控制器246、网络控制器234、固件接口228、无线收发器226、触摸传感器225、数据存储设备224(例如，硬盘驱动器、闪存等等)。在至少一个实施例中，数据存储设备224可以经由存储接口(例如SATA)或者经由诸如PCI或PCIe之类的外围总线连接。在至少一个实施例中，触摸传感器225可以包括触摸屏传感器、压力传感器或者指纹传感器。在至少一个实施例中，无线收发器226可以是Wi-Fi收发器、蓝牙收发器或者诸如3G、4G之类的移动网络收发器或者长期演化(“LTE”)收发器。在至少一个实施例中，固件接口228允许实现与系统固件通信，并且可以为例如统一可扩展固件接口(“UEFI”)。在至少一个实施例中，网络控制器234可以允许实现到有线网络的网络连接。在至少一个实施例中，高性能网络控制器(未示出)与接口总线210耦合。在至少一个实施例中，音频控制器246为多通道高清晰度音频控制器。在至少一个实施例中，处理系统200包括可选的旧式I/O控制器240，其用于将旧式(例如个人系统2(“PS/2”))设备耦合到处理系统200。在至少一个实施例中，平台控制器集线器230也可以连接到一个或更多个通用串行总线(“USB”)控制器242，连接输入设备，例如键盘和鼠标243组合、照相机244或者其他USB输入设备。
在至少一个实施例中，存储器控制器216和平台控制器集线器230的实例可以集成到诸如外部图形处理器212之类的离散外部图形处理器中。在至少一个实施例中，平台控制器集线器230和/或存储器控制器216可以在一个或更多个处理器202的外部。例如，在至少一个实施例中，处理系统200可以包括外部存储器控制器216和平台控制器集线器230，其可以被配置为与处理器202通信的系统芯片组内的存储器控制器集线器和外围控制器集线器。
图3图示出依照至少一个实施例的计算机系统300。在至少一个实施例中，计算机系统300可以为具有互连的设备和部件的系统、SOC或者某种组合。在至少一个实施例中，计算机系统300由可以包括执行指令的执行单元的处理器302形成。在至少一个实施例中，计算机系统300可以包括但不限于诸如处理器302之类的部件，以便采用包含逻辑的执行单元实施用于处理数据的算法。在至少一个实施例中，计算机系统300可以包括处理器，例如可从加州圣克拉拉英特尔公司获得的
Figure BDA0002780474710000085
处理器系列、XeonTM、
Figure BDA0002780474710000081
XScaleTM和/或StrongARMTM、
Figure BDA0002780474710000082
核心TM、
Figure BDA0002780474710000083
Figure BDA0002780474710000084
NervanaTM微处理器，但是也可以使用其他的系统(包括具有其他微处理器的PC、工程工作站、机顶盒等等)。在至少一个实施例中，计算机系统300可以执行可从华盛顿州雷德蒙德微软公司获得的一定版本的WINDOWS操作系统，但是也可以使用其他的操作系统(例如UNIX和Linux)、嵌入式软件和/或图形用户接口。
在至少一个实施例中，计算机系统300可以用在诸如手持式设备之类的其他设备和嵌入式应用中。手持式设备的一些示例包括蜂窝电话、互联网协议设备、数字照相机、个人数字助理(“PDA”)以及手持式PC。在至少一个实施例中，嵌入式应用可以包括微控制器、数字信号处理器(DSP)、SoC、网络计算机(“NetPC”)、机顶盒、网络集线器、广域网(“WAN”)交换机，或者可以实施一个或更多个指令的任何其他系统。
在至少一个实施例中，计算机系统300可以包括但不限于处理器302，所述处理器可以包括但不限于可以被配置为执行计算统一设备架构(“CUDA”)(
Figure BDA0002780474710000091
由加州圣克拉拉的英伟达公司开发)程序的一个或更多个执行单元308。在至少一个实施例中，CUDA程序至少为用CUDA编程语言书写的软件应用的一部分。在至少一个实施例中，计算机系统300为单处理器桌面或服务器系统。在至少一个实施例中，计算机系统300可以为多处理器系统。在至少一个实施例中，处理器302可以包括但不限于CISC微处理器、RISC微处理器、VLIW微处理器、实现指令集组合的处理器或者诸如例如数字信号处理器之类的任何其他处理器设备。在至少一个实施例中，处理器302可以耦合到处理器总线310，该总线可以在计算机系统300中的处理器302与其他部件之间传送数据信号。
在至少一个实施例中，处理器302可以包括但不限于1级(“L1”)内部缓存存储器(“缓存”)304。在至少一个实施例中，处理器302可以具有单个内部缓存或者多级内部缓存。在至少一个实施例中，缓存存储器可以驻留在处理器302的外部。在至少一个实施例中，处理器302也可以包括内部和外部缓存的组合。在至少一个实施例中，寄存器文件306可以将不同类型的数据存储在各种不同的寄存器中，包括但不限于整数寄存器、浮点寄存器、状态寄存器和指令指针寄存器。
在至少一个实施例中，包括但不限于实施整数和浮点运算的逻辑的执行单元308也驻留在处理器302中。处理器302也可以包括存储用于某些宏指令的微码的微码(“ucode”)只读存储器(“ROM”)。在至少一个实施例中，执行单元308可以包括处理打包的指令集309的逻辑。在至少一个实施例中，通过在通用处理器302的指令集中包括打包的指令集309，连同执行指令的关联电路系统一起，许多多媒体应用使用的操作可以在通用处理器302中使用打包的数据实施。在至少一个实施例中，许多多媒体应用可以通过使用全宽度的处理器数据总线对打包的数据实施操作而加速并且更加高效地执行，这可以消除跨处理器数据总线传输较小的数据单位以便一次对数据元素实施一个或更多个操作的需要。
在至少一个实施例中，执行单元308也可以用在微控制器、嵌入式处理器、图形设备、DSP和其他类型的逻辑电路中。在至少一个实施例中，计算机系统300可以包括但不限于存储器320。在至少一个实施例中，存储器320可以实现为DRAM设备、SRAM设备、闪存设备或者其他存储器设备。存储器320可以存储可以由处理器302执行的指令319和/或数据信号表示的数据321。
在至少一个实施例中，系统逻辑芯片可以耦合到处理器总线310和存储器320。在至少一个实施例中，系统逻辑芯片可以包括但不限于存储器控制器集线器(“MCH”)316，并且处理器302可以经由处理器总线310与MCH 316通信。在至少一个实施例中，MCH 316可以提供到存储器320的高带宽存储器路径318，所述存储器用于指令和数据存储并且用于存储图形命令、数据和纹理。在至少一个实施例中，MCH 316可以在计算机系统300中的处理器302、存储器320和其他部件之间引导数据信号，并且在处理器总线310、存储器320和系统I/O 322之间桥接数据信号。在至少一个实施例中，系统逻辑芯片可以提供用于耦合到图形控制器的图形端口。在至少一个实施例中，MCH 316可以通过高带宽存储器路径318耦合到存储器320，并且图形/视频卡312可以通过加速图形端口(“AGP”)互连线314耦合到MCH 316。
在至少一个实施例中，计算机系统300可以使用系统I/O 322将MCH 316耦合到I/O控制器集线器(“ICH”)330，所述系统I/O是一种专有集线器接口总线。在至少一个实施例中，ICH 330可以经由本地I/O总线提供到一些I/O设备的直接连接。在至少一个实施例中，本地I/O总线可以包括但不限于用于将外设连接到存储器320、芯片组和处理器302的高速I/O总线。示例可以包括但不限于音频控制器329、固件集线器(“快闪BIOS”)328、无线收发器326、数据存储装置324、包含用户输入接口325和键盘接口的旧式I/O控制器323、诸如USB之类的串行扩展端口327以及网络控制器334。数据存储装置324可以包括硬盘驱动器、软盘驱动器、CD-ROM设备、闪存设备或者其他大容量存储设备。
在至少一个实施例中，图3图示出一种系统，其包括互连的硬件设备或“芯片”。在至少一个实施例中，图3可以图示出一种示例性SoC。在至少一个实施例中，图3中图示出的设备可以利用专有互连线、标准化互连线(例如PCIe)或者其某种组合互连。在至少一个实施例中，系统300的一个或更多个部件使用计算快速链路(“CXL”)互连线互连。
图4图示出依照至少一个实施例的系统400。在至少一个实施例中，系统400为利用处理器410的电子设备。在至少一个实施例中，系统400可以例如且非限制性地为笔记本、塔式服务器、机架式服务器、刀片式服务器、膝上型电脑、台式电脑、平板电脑、移动设备、电话、嵌入式计算机或者任何其他适当的电子设备。
在至少一个实施例中，系统400可以包括但不限于通信耦合至任何适当数量或种类的部件、外设、模块或设备的处理器410。在至少一个实施例中，处理器410使用诸如I2C总线、系统管理总线(“SMBus”)、低引脚数(“LPC”)总线、串行外围接口(“SPI”)、高清晰度音频(“HDA”)总线、串行先进技术附件(“SATA”)总线、USB(版本1、2、3)或者通用异步接收器/发送器(“UART”)总线之类的总线或接口进行耦合。在至少一个实施例中，图4图示出一种包括互连的硬件设备或“芯片”的系统。在至少一个实施例中，图4可以图示出一种示例性SoC。在至少一个实施例中，图4中图示出的设备可以利用专有互连线、标准化互连线(例如PCIe)或者其某种组合互连。在至少一个实施例中，图4的一个或更多个部件使用CXL互连线互连。
在至少一个实施例中，图4可以包括显示器424、触摸屏425、触摸板430、近场通信单元(“NFC”)445、传感器集线器440、热传感器446、快速芯片组(“EC”)435、可信平台模块(“TPM”)438、BIOS/固件/闪存(“BIOS,FW Flash”)422，DSP 460，固态盘(“S23”)或硬盘驱动器(“HDD”)420、无线局域网单元(“WLAN”)450、蓝牙单元452、无线广域网单元(“WWAN”)456、全球定位系统(“GPS”)455、诸如USB 3.0照相机之类的照相机(“USB 3.0照相机”)454或者以例如LPDDR3标准实现的低功率双倍数据速率(“LPDDR”)存储器单元(“LPDDR3”)415。这些部件中的每一个可以以任何适当的方式实现。
在至少一个实施例中，其他部件可以通过上面讨论的部件通信耦合至处理器410。在至少一个实施例中，加速度计441、环境光传感器(“ALS”)442、罗盘443和陀螺仪444可以通信耦合至传感器集线器440。在至少一个实施例中，热传感器439、风扇437、键盘446和触摸板430可以通信耦合至EC 435。在至少一个实施例中，扬声器463、头戴式耳机464和麦克风(“mic”)465可以通信耦合至音频单元(“音频编解码器和d类放大器”)464，该音频单元反过来可以通信耦合至DSP 460。在至少一个实施例中，音频单元464可以例如且非限制性地包括音频编码器/解码器(“编解码器”)和D类放大器。在至少一个实施例中，SIM卡(“SIM”)457可以通信耦合至WWAN单元456。在至少一个实施例中，诸如WLAN单元450和蓝牙单元452以及WWAN单元456之类的部件可以以下一代形状因子(“NGFF”)实现。
图5图示出依照至少一个实施例的示例性集成电路500。在至少一个实施例中，示例性集成电路500为可以使用一个或更多个IP核心制造的SoC。在至少一个实施例中，集成电路500包括一个或更多个应用处理器505(例如CPU)、至少一个图形处理器510，并且可以附加地包括图像处理器515和/或视频处理器520，它们中的任何一个都可以是模块化IP核心。在至少一个实施例中，集成电路500包括外设或者总线逻辑，包括USB控制器525、UART控制器530、SPI/SDIO控制器535以及I2S/I2C控制器540。在至少一个实施例中，集成电路500可以包括耦合至高清晰度多媒体接口(“HDMI”)控制器550和移动行业处理器接口(“MIPI”)显示接口555中的一个或更多个的显示设备545。在至少一个实施例中，可以通过包括闪存和闪存控制器的闪存子系统560提供存储装置。在至少一个实施例中，可以经由存储器控制器565提供存储器接口，以便访问23RAM或SRAM存储器设备。在至少一个实施例中，一些集成电路附加地包括嵌入式安全引擎570。
图6图示出依照至少一个实施例的计算系统600；在至少一个实施例中，计算系统600包括处理子系统601，该处理子系统具有经由可以包括存储器集线器605的互连路径通信的一个或更多个处理器602和系统存储器604。在至少一个实施例中，存储器集线器605可以是芯片组部件内的单独部件，或者可以集成到一个或更多个处理器602中。在至少一个实施例中，存储器集线器605经由通信链路606与I/O子系统611耦合。在至少一个实施例中，I/O子系统611包括I/O集线器607，该I/O集线器可以使得计算系统600能够接收来自一个或更多个输入设备608的输入。在至少一个实施例中，I/O集线器607可以使得可能包括在一个或更多个处理器602中的显示控制器能够向一个或更多个显示设备610A提供输出。在至少一个实施例中，与I/O集线器607耦合的一个或更多个显示设备610A可以包括本地的、内部的或者嵌入式显示设备。
在至少一个实施例中，处理子系统601包括经由总线或其他通信链路613耦合至存储器集线器605的一个或更多个并行处理器612。在至少一个实施例中，通信链路613可以是任意数量的基于标准的通信链路技术或协议之一，例如但不限于PCIe，或者可以是特定于供应商的通信接口或通信构造(fabric)。在至少一个实施例中，一个或更多个并行处理器612形成聚焦于计算的并行或矢量处理系统，其可以包括大量处理核心和/或处理集群，例如许多集成核心处理器。在至少一个实施例中，一个或更多个并行处理器612形成图形处理子系统，其可以将像素输出到经由I/O集线器607耦合的一个或更多个显示设备610A之一。在至少一个实施例中，一个或更多个并行处理器612也可以包括显示控制器和显示接口(未示出)以便允许实现直接连接到一个或更多个显示设备610B。
在至少一个实施例中，系统存储单元614可以连接到I/O集线器607以便为计算系统600提供存储机制。在至少一个实施例中，I/O开关616可以用来提供允许实现I/O集线器607与其他部件之间的连接的接口机制，所述其他部件例如网络适配器618和/或可以集成到平台中的无线网络适配器619以及可以经由一个或更多个外接设备620添加的各种其他设备。在至少一个实施例中，网络适配器618可以是以太网适配器或者另一种有线网络适配器。在至少一个实施例中，无线网络适配器619可以包括Wi-Fi、蓝牙、NFC设备或者包括一个或更多个无线收音机的其他网络设备中的一个或更多个。
在至少一个实施例中，计算系统600可以包括也可以连接到I/O集线器607的、没有显式示出的其他部件，包括USB或其他端口连接、光存储驱动器、视频捕获设备等等。在至少一个实施例中，将图6中的各种不同部件互连的通信路径可以使用诸如基于PCI的协议(例如PCIe)之类的任何适当的协议或者诸如NVLink高速互连或互连协议之类的其他总线或点对点通信接口和/或协议实现。
在至少一个实施例中，一个或更多个并行处理器612结合了针对图形和视频处理优化的电路系统，包括例如视频输出电路系统，并且构成图形处理单元(“GPU”)。在至少一个实施例中，一个或更多个并行处理器612结合了针对通用处理优化的电路系统。在至少实施例中，计算系统600的部件可以在单个集成电路上与一个或更多个其他系统元件集成在一起。例如，在至少一个实施例中，一个或更多个并行处理器612、存储器集线器605、处理器602和I/O集线器607可以集成到SoC集成电路中。在至少一个实施例中，计算系统600的部件可以集成到单个封装中以形成封装系统(“SIP”)配置。在至少一个实施例中，计算系统600的部件的至少一部分可以集成到多芯片模块(“MCM”)中，该多芯片模块可以与其他多芯片模块互连成模块化计算系统。在至少一个实施例中，计算系统600中省略了I/O子系统611和显示设备610B。
处理系统
下面的图非限制性地阐述了可以用来实现至少一个实施例的示例性处理系统。
图7图示出依照至少一个实施例的加速处理单元(“APU”)700。在至少一个实施例中，APU 700由加州圣克拉拉AMD公司开发。在至少一个实施例中，APU 700可以被配置为执行诸如CUDA程序之类的应用程序。在至少一个实施例中，APU 700包括但不限于核心复合体710、图形复合体740、构造760、I/O接口770、存储器控制器780、显示控制器792和多媒体引擎794。在至少一个实施例中，APU 700可以包括但不限于任意组合的任意数量的核心复合体710、任意数量的图形复合体750、任意数量的显示控制器792以及任意数量的多媒体引擎794。出于解释的目的，在需要的情况下，相似对象的多个实例在这里用标识该对象的参考数字以及标识该实例的括号数字标记。
在至少一个实施例中，核心复合体710为CPU，图形复合体740为GPU，并且APU 700为非限制性地将710和740集成到单个芯片上的处理单元。在至少一个实施例中，可以将一些任务分配给核心复合体710，并且可以将其他任务分配给图形复合体740。在至少一个实施例中，核心复合体710被配置为执行与APU 700关联的主控软件，例如操作系统。在至少一个实施例中，核心复合体710为APU 700的主处理器，控制和协调其他处理器的操作。在至少一个实施例中，核心复合体710发出控制图形复合体740的操作的命令。在至少一个实施例中，核心复合体710可以被配置为执行从CUDA源代码导出的主机可执行代码，并且图形复合体740可以被配置为执行从CUDA源代码导出的设备可执行代码。
在至少一个实施例中，核心复合体710包括但不限于核心720(1)-720(4)和L3缓存730。在至少一个实施例中，核心复合体710可以包括但不限于任意组合的任意数量的核心720以及任意数量和类型的缓存。在至少一个实施例中，核心720被配置为执行特定指令集架构(“I20”)的指令。在至少一个实施例中，每个核心720为CPU核心。
在至少一个实施例中，每个核心720包括但不限于取得/解码单元722、整数执行引擎724、浮点执行引擎726和L2缓存728。在至少一个实施例中，取得/解码单元722取得指令，解码这样的指令，生成微操作，并且将单独的微指令分派给整数执行引擎724和浮点执行引擎726。在至少一个实施例中，取得/解码单元722可以并发地将一个微指令分派给整数执行引擎724并且将另一个微指令分派给浮点执行引擎726。在至少一个实施例中，整数执行引擎724非限制性地执行整数和存储器操作。在至少一个实施例中，浮点执行引擎726非限制性地执行浮点和矢量操作。在至少一个实施例中，取得-解码单元722将微指令分派给代替整数执行引擎724和浮点执行引擎726二者的单个执行引擎。
在至少一个实施例中，每个核心720(i)——其中i为表示核心720的特定实例的整数——可以访问核心720(i)中包括的L2缓存728(i)。在至少一个实施例中，核心复合体710(j)——其中j为表示核心复合体710的特定实例的整数——中包括的每个核心720经由核心复合体710(j)中包括的L3缓存730(j)连接到核心复合体710(j)中包括的其他核心720。在至少一个实施例中，核心复合体710(j)——其中j为表示核心复合体710的特定实例的整数——中包括的核心720可以访问核心复合体710(j)中包括的所有L3缓存730(j)。在至少一个实施例中，L3缓存730可以包括但不限于任意数量的分片。
在至少一个实施例中，图形复合体740可以被配置为以高度并行的方式实施计算操作。在至少一个实施例中，图形复合体740被配置为执行图形管线操作，例如绘制命令、像素操作、几何计算以及与向显示器再现图像关联的其他操作。在至少一个实施例中，图形复合体740被配置为执行与图形无关的操作。在至少一个实施例中，图形复合体740被配置为执行与图形有关的操作和与图形无关的操作二者。
在至少一个实施例中，图形复合体740包括但不限于任意数量的计算单元750以及L2缓存742。在至少一个实施例中，计算单元750共享L2缓存742。在至少一个实施例中，L2缓存742被分区。在至少一个实施例中，图形复合体740包括但不限于任意数量的计算单元750以及任意数量(包括零个)和类型的缓存。在至少一个实施例中，图形复合体740包括但不限于任意数量的专用图形硬件。
在至少一个实施例中，每个计算单元750包括但不限于任意数量的SIMD单元752和共享的存储器754。在至少一个实施例中，每个SIMD单元752实现一种SIMD架构并且被配置为并行地实施操作。在至少一个实施例中，每个计算单元750可以执行任意数量的线程块，但是每个线程块在单个计算单元750上执行。在至少一个实施例中，线程块包括但不限于任意数量的执行线程。在至少一个实施例中，工作组为线程块。在至少一个实施例中，每个SIMD单元752执行不同的线程束(warp)。在至少一个实施例中，线程束为一组线程(例如16个线程)，其中线程束中的每个线程属于单个线程块，并且被配置为基于单个指令集处理不同的数据集合。在至少一个实施例中，谓词(predication)可以用来禁用线程束中的一个或更多个线程。在至少一个实施例中，巷道(lane)为一种线程。在至少一个实施例中，工作项目为一种线程。在至少一个实施例中，波前为一种线程束。在至少一个实施例中，线程块中的不同波前可以一起同步，并且经由共享的存储器754通信。
在至少一个实施例中，构造760为促进跨核心复合体710、图形复合体740、I/O接口770、存储器控制器780、显示控制器792和多媒体引擎794的数据和控制传送的系统互连线。在至少一个实施例中，除了构造760之外或者代替构造760的是，APU 700可以包括但不限于促进跨任意数量和类型的、可以在APU 700内部或外部的直接或间接链接的部件的数据和控制传送的任意数量和类型的系统互连线。在至少一个实施例中，I/O接口770表示任意数量和类型的I/O接口(例如PCI、PCI扩展(“PCI-X”)、PCIe、千兆以太网(“GBE”)、USB等等)。在至少一个实施例中，各种不同类型的外围设备耦合到I/O接口770。在至少一个实施例中，耦合到I/O接口770的外围设备可以包括但不限于键盘、鼠标、打印机、扫描仪、操纵杆或者其他类型的游戏控制器、媒体记录设备、外部存储设备、网络接口卡等等。
在至少一个实施例中，显示控制器AMD92在诸如液晶显示器(“LCD”)设备之类的一个或更多个显示设备上显示图像。在至少一个实施例中，多媒体引擎240包括但不限于与多媒体有关的任意数量和类型的电路系统，例如视频解码器、视频编码器、图像信号处理器等等。在至少一个实施例中，存储器控制器780促进APU 700与统一系统存储器790之间的数据传输。在至少一个实施例中，核心复合体710和图形复合体740共享统一系统存储器790。
在至少一个实施例中，APU 700实现一种存储器子系统，其包括但不限于可以专用于一个部件或者在多个部件之间共享的任意数量和类型的存储器控制器780和存储器设备(例如共享的存储器754)。在至少一个实施例中，APU 700实现一种缓存子系统，其包括但不限于一个或更多个缓存存储器(例如L2缓存828、L3缓存730和L2缓存742)，所述缓存存储器中的每一个可以专用于任意数量的部件(例如核心720、核心复合体710、SIMD单元752、计算单元750以及图形复合体740)或者在这些部件之间共享。
图8图示出依照至少一个实施例的CPU 800。在至少一个实施例中，CPU 800由加州圣克拉拉AMD公司开发。在至少一个实施例中，CPU 800可以被配置为执行应用程序。在至少一个实施例中，CPU 800被配置为执行诸如操作系统之类的主控软件。在至少一个实施例中，CPU 800发出控制外部GPU(未示出)的操作的命令。在至少一个实施例中，CPU 800可以被配置为执行从CUDA源代码导出的主机可执行代码，并且外部GPU可以被配置为执行从这样的CUDA源代码导出的设备可执行代码。在至少一个实施例中，CPU 800包括但不限于任意数量的核心复合体810、构造860、I/O接口870以及存储器控制器880。
在至少一个实施例中，核心复合体810包括但不限于核心820(1)-820(4)以及L3缓存830。在至少一个实施例中，核心复合体810可以包括但不限于任意组合的任意数量的核心820以及任意数量和类型的缓存。在至少一个实施例中，核心820被配置为执行特定I20的指令。在至少一个实施例中，每个核心820是CPU核心。
在至少一个实施例中，每个核心820包括但不限于取得/解码单元822、整数执行引擎824、浮点执行引擎826和L2缓存828。在至少一个实施例中，取得/解码单元822取得指令，解码这样的指令，生成微操作，并且将单独的微指令分派给整数执行引擎824和浮点执行引擎826。在至少一个实施例中，取得/解码单元822可以并发地将一个微指令分派给整数执行引擎824并且将另一个微指令分派给浮点执行引擎826。在至少一个实施例中，整数执行引擎824非限制性地执行整数和存储器操作。在至少一个实施例中，浮点引擎826非限制性地执行浮点和矢量操作。在至少一个实施例中，取得-解码单元822将微指令分派给代替整数执行引擎824和浮点执行引擎826二者的单个执行引擎。
在至少一个实施例中，每个核心820(i)——其中i为表示核心820的特定实例的整数——可以访问核心820(i)中包括的L2缓存828(i)。在至少一个实施例中，核心复合体810(j)——其中j为表示核心复合体810的特定实例的整数——中包括的每个核心820经由核心复合体810(j)中包括的L3缓存830(j)连接到核心复合体810(j)中的其他核心820。在至少一个实施例中，核心复合体810(j)——其中j为表示核心复合体810的特定实例的整数——中包括的核心820可以访问核心复合体810(j)中包括的所有L3缓存830(j)。在至少一个实施例中，L3缓存830可以包括但不限于任意数量的分片。
在至少一个实施例中，构造860为促进跨核心复合体810(1)-810(N)(其中N为大于0的整数)、I/O接口870和存储器控制器880的数据和控制传送的系统互连线。在至少一个实施例中，除了构造860之外或者代替构造860的是，CPU 800可以包括但不限于促进跨任意数量和类型的、可以在CPU 800内部或外部的直接或间接链接的部件的数据和控制传送的任意数量和类型的系统互连线。在至少一个实施例中，I/O接口870表示任意数量和类型的I/O接口(例如PCI、PCI-X、PCIe、GBE、USB等等)。在至少一个实施例中，各种不同类型的外围设备耦合到I/O接口870。在至少一个实施例中，耦合到I/O接口870的外围设备可以包括但不限于显示器、键盘、鼠标、打印机、扫描仪、操纵杆或者其他类型的游戏控制器、媒体记录设备、外部存储设备、网络接口卡等等。
在至少一个实施例中，存储器控制器880促进CPU 800与系统存储器890之间的数据传输。在至少一个实施例中，核心复合体810和图形复合体840共享系统存储器890。在至少一个实施例中，CPU 800实现一种存储器子系统，其包括但不限于可以专用于一个部件或者在多个部件之间共享的任意数量和类型的存储器控制器880和存储器设备。在至少一个实施例中，CPU 800实现一种缓存子系统，其包括但不限于一个或更多个缓存存储器(例如L2缓存828和L3缓存830)，所述缓存存储器中的每一个可以专用于任意数量的部件(例如核心820和核心复合体810)或者在这些部件之间共享。
图9图示出依照至少一个实施例的示例性加速器集成分片990。当在本文中使用时，“分片”包括加速器集成电路的处理资源的指定部分。在至少一个实施例中，加速器集成电路代表图形加速模块中包括的多个图形处理引擎提供缓存管理、存储器访问、上下文管理和中断管理服务。这些图形处理引擎中的每一个可以包括单独的GPU。可替换地，这些图形处理引擎可以在GPU内包括不同类型的图形处理引擎，例如图形执行单元、媒体处理引擎(例如视频编码器/解码器)、采样器和位块传送(blit)引擎。在至少一个实施例中，图形加速模块可以为具有多个图形处理引擎的GPU。在至少一个实施例中，图形处理引擎可以是集成到公共封装、线卡或者芯片上的单独的GPU。
系统存储器914内的应用有效地址空间982存储进程元素983。在一个实施例中，响应于来自处理器907上执行的应用980的GPU调用981而存储进程元素983。进程元素983包含相应应用980的进程状态。进程元素983中包含的工作描述符(“WD”)984可以是应用请求的单个作业，或者可以包含指向作业队列的指针。在至少一个实施例中，WD 984为指向应用有效地址空间982中的作业请求队列的指针。
图形加速模块946和/或单独的图形处理引擎可以由系统中的全部进程或者进程子集共享。在至少一个实施例中，可以包括用于在虚拟化环境下设立进程状态并且将WD984发送至图形加速模块946以启动作业的基础结构。
在至少一个实施例中，专用进程编程模型是特定于实现方式的。在该模型中，单个进程拥有图形加速模块946或者单独的图形处理引擎。由于图形加速模块946为单个进程拥有，当分配了图形加速模块946时，超级监督者为拥有的分区初始化加速器集成电路，并且操作系统为拥有的进程初始化加速器集成电路。
在操作中，加速器集成分片990中的WD取得单元991取得下一个WD 984，该WD包括要由图形加速模块946的一个或更多个图形处理引擎完成的工作的指示。如图所示，来自WD984的数据可以存储在寄存器945中，并且由存储器管理单元(“MMU”)939、中断管理电路947和/或上下文管理电路948使用。例如，MMU 939的一个实施例包括用于访问OS虚拟地址空间985内的分段/页面表986的分段/页面漫游电路系统。中断管理电路947可以处理接收自图形加速模块946的中断事件(“INT”)992。当实施图形操作时，图形处理引擎生成的有效地址993被MMU 939转换成真实地址。
在一个实施例中，相同的寄存器945集合针对每个图形处理引擎和/或图形加速模块946被复制，并且可以由超级监督者或者操作系统初始化。这些复制的寄存器中的每一个可以包括在加速器集成分片990中。表1中示出了可以由超级监督者初始化的示例性寄存器。
表1超级监督者初始化的寄存器
1 分片控制寄存器
2 真实地址(RA)调度进程区域指针
3 权限掩码覆盖寄存器
4 中断矢量表条目偏移
5 中断矢量表条目限制
6 状态寄存器
7 逻辑分区ID
8 真实地址(RA)超级监督者加速器利用率记录指针
9 存储描述寄存器
表2中示出了可以由操作系统初始化的示例性寄存器。
表2操作系统初始化的寄存器
1 进程和线程标识
2 有效地址(EA)上下文保存/恢复指针
3 虚拟地址(VA)加速器利用率记录指针
4 虚拟地址(VA)存储分段表指针
5 权限掩码
6 工作描述符
在一个实施例中，每个WD 984特定于具体的图形加速模块946和/或具体的图形处理引擎。它包含图形处理引擎完成工作所需的所有信息，或者它可以是指向其中应用已经设立了要完成的工作命令队列的存储器位置的指针。
图10A-10B图示出依照至少一个实施例的示例性图形处理器。在至少一个实施例中，任何所述示例性图形处理器可以使用一个或更多个IP核心制作。除了所图示出的以外，也可以包括其他逻辑和电路。在至少一个实施例中，包括附加的图形处理器/核心、外围接口控制器或者通用处理器核心。在至少一个实施例中，所述示例性图形处理器用在SoC中。
图10A图示出依照至少一个实施例的可以使用一个或更多个IP核心制作的SoC集成电路的一个示例性图形处理器1010。图10B图示出依照至少一个实施例的可以使用一个或更多个IP核心制作的SoC集成电路的附加的示例性图形处理器1040。在至少一个实施例中，图10A的图形处理器1010为低功率图形处理器核心。在至少一个实施例中，图10B的图形处理器1040为较高性能图形处理器核心。在至少一个实施例中，图形处理器1010、1040中的每一个都可以是图5的图形处理器510的变体。
在至少一个实施例中，图形处理器1010包括顶点处理器1005和一个或更多个片段处理器1015A-1015N(例如1015A、1015B、1015C、1015D，直到1015N-1和1015N)。在至少一个实施例中，图形处理器1010可以经由单独的逻辑执行不同的着色器程序，使得顶点处理器1005被优化以执行用于顶点着色器程序的操作，而一个或更多个片段处理器1015A-1015N执行用于片段或像素着色器程序的片段(例如像素)着色操作。在至少一个实施例中，顶点处理器1005实施3D图形管线的顶点处理阶段，并且生成图元和顶点数据。在至少一个实施例中，片段处理器1015A-1015N使用顶点处理器1005生成的图元和顶点数据产生在显示设备上显示的帧缓冲区。在至少一个实施例中，片段处理器1015A-1015N被优化以执行如OpenGL API中所提供的片段着色器程序，该片段着色器程序可以用来实施与如Direct 3DAPI中所提供的像素着色器程序相似的操作。
在至少一个实施例中，图形处理器1010附加地包括一个或更多个MMU 1020A-1020B、缓存1025A-1025B以及电路互连线1030A-1030B。在至少一个实施例中，一个或更多个MMU 1020A-1020B为图形处理器1010，包括为顶点处理器1005和/或片段处理器1015A-1015N提供虚拟-物理地址的映射，其除了引用一个或更多个缓存1025A-1025B中存储的顶点或图像/纹理数据之外，也可以引用存储器中存储的顶点或图像/纹理数据。在至少一个实施例中，一个或更多个MMU 1020A-1020B可以与系统内的其他MMU同步，所述其他MMU包括与图5的一个或更多个应用处理器505、图像处理器515和/或视频处理器520关联的一个或更多个MMU，使得每个处理器505-520可以参与共享或统一的虚拟存储器系统。在至少一个实施例中，一个或更多个电路互连线1030A-1030B使得图形处理器1010能够经由SoC的内部总线或者经由直接连接与该SoC内的其他IP核心进行接口。
在至少一个实施例中，图形处理器1040包括图10A的图形处理器1010的一个或更多个MMU 1020A-1020B、缓存1025A-1025B以及电路互连线1030A-1030B。在至少一个实施例中，图形处理器1040包括一个或更多个着色器核心1055A-1055N(例如，1055A、1055B、1055C、1055D、1055E、1055F，直到1055N-1和1055N)，其提供统一着色器核心架构，在该架构中，单个核心或类型或核心可以执行所有类型的可编程着色器代码，包括实现顶点着色器、片段着色器和/或计算着色器的着色器程序代码。在至少一个实施例中，着色器核心的数量可以变化。在至少一个实施例中，图形处理器1040包括核心间任务管理器1045，其充当将执行线程分派给一个或更多个着色器核心1055A-1055N的线程分派器以及加速用于基于图块(tile)的再现的平铺操作的平铺单元1058，在所述平铺操作中，在图像空间中将用于场景的再现操作细分，例如以便利用场景中的局部空间连贯性或者优化内部缓存的使用。
图11A图示出依照至少一个实施例的图形核心1100。在至少一个实施例中，图形核心1100可以包括在图5的图形处理器510中。在至少一个实施例中，图形核心1100可以像在图10B中那样为统一着色器核心1055A-1055N。在至少一个实施例中，图形核心1100包括图形核心1100内的执行资源共用的共享指令缓存1102、纹理单元1118和缓存/共享的存储器1120。在至少一个实施例中，图形核心1100对于每个核心可以包括多个分片1101A-1101N或者分区，并且图形处理器可以包括图形核心1100的多个实例。分片1101A-1101N可以包括支持逻辑，包括本地指令缓存1104A-1104N、线程调度器1106A-1106N、线程分派器1108A-1108N以及寄存器集合1110A-1110N。在至少一个实施例中，分片1101A-1101N可以包括一组附加功能单元(“AFU”)1112A-1112N、浮点单元(“FPU”)1114A-1114N、整数算术逻辑单元(“ALU”)1116-1116N、地址计算单元(“ACU”)1113A-1113N、双精度浮点单元(“DPFPU”)1115A-1115N以及矩阵处理单元(“MPU”)1117A-1117N。
在至少一个实施例中，FPU 1114A-1114N可以实施单精度(32位)和半精度(16位)浮点运算，而DPFPU 1115A-1115N实施双精度(64位)浮点运算。在至少一个实施例中，ALU1116A-1116N可以实施8位、16位和32位精度下的可变精度整数运算，并且可以被配置用于混合精度运算。在至少一个实施例中，MPU 1117A-1117N也可以被配置用于混合精度矩阵运算，包括半精度浮点和8位整数运算。在至少一个实施例中，MPU 1117A-1117N可以实施各种各样的矩阵运算以加速CUDA程序，包括启用支持加速的一般矩阵-矩阵乘法(“GEMM”)。在至少一个实施例中，AFU 1112A-1112N可以实施不被浮点或整数单元支持的附加逻辑运算，包括三角函数运算(例如正弦、余弦等等)。
图11B图示出依照至少一个实施例的通用图形处理单元(“GPGPU”)1130。在至少一个实施例中，GPGPU 1130是高度并行的并且适合部署在多芯片模块上。在至少一个实施例中，GPGPU 1130可以被配置为使得高度并行的计算操作能够由GPU阵列实施。在至少一个实施例中，GPGPU 1130可以直接链接到GPGPU 1130的其他实例以便创建改善CUDA程序的执行时间的多GPU集群。在至少一个实施例中，GPGPU 1130包括允许实现与主机处理器的连接的主机接口1132。在至少一个实施例中，主机接口1132是一种PCIe接口。在至少一个实施例中，主机接口1132可以是一种特定于供应商的通信接口或通信构造。在至少一个实施例中，GPGPU 1130接收来自主机处理器的命令，并且使用全局调度器1134将与那些命令关联的执行线程分发给计算集群1136A-1136H集合。在至少一个实施例中，计算集群1136A-1136H共享缓存存储器1138。在至少一个实施例中，缓存存储器1138可以用作用于计算集群1136A-1136H内的缓存存储器的较高级别缓存。
在至少一个实施例中，GPGPU 1130包括经由存储器控制器1142A-1142B集合与计算集群1136A-1136H耦合的存储器1144A-1144B。在至少一个实施例中，存储器1144A-1144B可以包括各种不同类型的存储器设备，包括DRAM或者图形随机存取存储器，例如同步图形随机存取存储器(“SGRAM”)，包括图形双倍数据速率(“GDDR”)存储器。
在至少一个实施例中，计算集群1136A-1136H中的每一个包括图形核心集合，例如图11A的图形核心1100，其可以包括多种类型的整数和浮点逻辑单元，这些单元可以实施一系列精度下的计算操作，包括适合于与CUDA程序关联的计算。例如，在至少一个实施例中，至少每个计算集群1136A-1136H中的浮点单元子集可以被配置为实施16位或32位浮点运算，而不同的浮点单元子集可以被配置为实施64位浮点运算。
在至少一个实施例中，GPGPU 1130的多个实例可以被配置为作为计算集群操作。计算集群1136A-1136H可以实现任何技术上可行的用于同步和数据交换的通信技术。在至少一个实施例中，GPGPU 1130的多个实例通过主机接口1132通信。在至少一个实施例中，GPGPU 1130包括I/O集线器1139，其将GPGPU 1130与允许实现到GPGPU 1130的其他实例的直接连接的GPU链路1140耦合。在至少一个实施例中，GPU链路1140耦合到专用GPU-GPU桥接器，该桥接器允许实现GPGPU 1130的多个实例之间的通信和同步。在至少一个实施例中，GPU链路1140与高速互连线耦合以便向其他GPGPU 1130或并行处理器传送数据和接收数据。在至少一个实施例中，GPGPU 1130的多个实例位于单独的数据处理系统中并且经由网络设备通信，该网络设备可经由主机接口1132访问。在至少一个实施例中，除了主机接口1132之外或者作为主机接口1132的替代，GPU链路1140可以被配置为允许实现到主机处理器的连接。在至少一个实施例中，GPGPU 1130可以被配置为执行CUDA程序。
图12A图示出依照至少一个实施例的并行处理器1200。在至少一个实施例中，并行处理器1200的各种不同的部件可以使用诸如可编程处理器、专用集成电路(“ASIC”)或者FPGA之类的一个或更多个集成电路设备实现。
在至少一个实施例中，并行处理器1200包括并行处理单元1202。在至少一个实施例中，并行处理单元1202包括I/O单元1204，该I/O单元允许实现与包括并行处理单元1202的其他实例在内的其他设备的通信。在至少一个实施例中，I/O单元1204可以直接连接到其他设备。在至少一个实施例中，I/O单元1204经由集线器或交换机接口的使用与诸如存储器集线器605之类的其他设备连接。在至少一个实施例中，存储器集线器605与I/O单元1204之间的连接形成通信链路。在至少一个实施例中，I/O单元1204与主机接口1206和存储器交叉开关(crossbar)1216连接，其中主机接口1206接收针对实施处理操作的命令，并且存储器交叉开关1216接收针对实施存储器操作的命令。
在至少一个实施例中，当主机接口1206经由I/O单元1204接收命令缓冲区时，主机接口1206可以将实施那些命令的工作操作引导至前端1208。在至少一个实施例中，前端1208与调度器1210耦合，该调度器被配置为将命令或其他工作项目分发给处理阵列1212。在至少一个实施例中，调度器1210确保在任务被分发给处理阵列1212之前，处理阵列1212被适当配置并且处于有效状态。在至少一个实施例中，调度器1210经由在微控制器上执行的固件逻辑实现。在至少一个实施例中，微控制器实现的调度器1210可配置为在粗粒度和细粒度下实施复杂的调度和工作分发操作，允许实现在处理阵列1212上执行的线程的快速抢占和上下文切换。在至少一个实施例中，主机软件可以经由多个图形处理门铃之一证明在处理阵列1212上调度的工作负载。在至少一个实施例中，工作负载于是可以通过包括调度器1210的微控制器内的调度器1210逻辑自动地跨处理阵列1212分布。
在至少一个实施例中，处理阵列1212可以包括高达“N”个集群(例如集群1214A，集群1214B，直到集群1214N)。在至少一个实施例中，处理阵列1212的每个集群1214A-1214N可以执行大量的并发线程。在至少一个实施例中，调度器1210可以使用各种不同的调度和/或工作分发算法将工作分配给处理阵列1212的集群1214A-1214N，所述算法可以根据针对每种类型的程序或计算而出现的工作负载来变化。在至少一个实施例中，可以由调度器1210动态地处理调度，或者可以由编译器逻辑在被配置用于由处理阵列1212执行的程序逻辑的编译期间部分地帮助调度。在至少一个实施例中，处理阵列1212的不同集群1214A-1214N可以被分配用于处理不同类型的程序或者用于实施不同类型的计算。
在至少一个实施例中，处理阵列1212可以被配置为实施各种不同类型的并行处理操作。在至少一个实施例中，处理阵列1212被配置为实施通用并行计算操作。例如，在至少一个实施例中，处理阵列1212可以包括执行处理任务的逻辑，所述处理任务包括对视频和/或音频数据滤波、实施包括物理操作在内的建模操作以及实施数据变换。
在至少一个实施例中，处理阵列1212被配置为实施并行图形处理操作。在至少一个实施例中，处理阵列1212可以包括支持执行这样的图形处理操作的附加逻辑，包括但不限于实施纹理操作的纹理采样逻辑以及镶嵌逻辑和其他顶点处理逻辑。在至少一个实施例中，处理阵列1212可以被配置为执行图形处理有关的着色器程序，例如但不限于顶点着色器、镶嵌着色器、几何着色器和像素着色器。在至少一个实施例中，并行处理单元1202可以经由I/O单元1204传输来自系统存储器的数据以用于处理。在至少一个实施例中，在处理期间，传输的数据可以在处理期间存储到片上存储器(例如并行处理器存储器1222)，然后回写到系统存储器。
在至少一个实施例中，当并行处理单元1202用来实施图形处理时，调度器1210可以被配置为将处理工作负载划分为近似相等大小的任务，以便更好地允许实现将图形处理操作分发给处理阵列1212的多个集群1214A-1214N。在至少一个实施例中，处理阵列1212的部分可以被配置为实施不同类型的处理。例如，在至少一个实施例中，第一部分可以被配置为实施顶点着色和拓扑生成，第二部分可以被配置为实施镶嵌和几何着色，并且第三部分可以被配置为实施像素着色或者其他屏幕空间操作，以便产生供显示的再现图像。在至少一个实施例中，集群1214A-1214N中的一个或更多个产生的中间数据可以存储在缓冲区中以便允许中间数据在集群1214A-1214N之间传送以用于进一步处理。
在至少一个实施例中，处理阵列1212可以经由调度器1210接收要执行的处理任务，该调度器从前端1208接收限定处理任务的命令。在至少一个实施例中，处理任务可以包括要处理的数据的索引，所述数据例如表面(补丁)数据、图元数据、顶点数据和/或像素数据，以及限定将如何处理数据的状态参数和命令(例如要执行什么程序)。在至少一个实施例中，调度器1210可以被配置为取得与任务相应的索引，或者可以从前端1208接收索引。在至少一个实施例中，前端1208可以被配置为确保在传入命令缓冲区(例如批处理缓冲区、推缓冲区等等)指定的工作负载被启动之前处理阵列1212被配置成有效状态。
在至少一个实施例中，并行处理单元1202的一个或更多个实例中的每一个可以与并行处理器存储器1222耦合。在至少一个实施例中，并行处理器存储器1222可以经由存储器交叉开关1216进行访问，该存储器交叉开关可以接收来自处理阵列1212以及I/O单元1204的存储器请求。在至少一个实施例中，存储器交叉开关1216可以经由存储器接口1218访问并行处理器存储器1222。在至少一个实施例中，存储器接口1218可以包括多个分区单元(例如分区单元1220A、分区单元1220B，直到分区单元1220N)，这些分区单元中的每一个可以耦合到并行处理器存储器1222的部分(例如存储器单元)。在至少一个实施例中，分区单元1220A-1220N的数量被配置为等于存储器单元的数量，使得第一分区单元1220A具有相应的第一存储器单元1224A，第二分区单元1220B具有相应的存储器单元1224B，并且第N分区单元1220N具有相应的第N存储器单元1224N。在至少一个实施例中，分区单元1220A-1220N的数量可以不等于存储器设备的数量。
在至少一个实施例中，存储器单元1224A-1224N可以包括各种不同类型的存储器设备，包括DRAM或者图形随机存取存储器，例如SGRAM，包括GDDR存储器。在至少一个实施例中，存储器单元1224A-1224N也可以包括3D堆叠式存储器，包括但不限于高带宽存储器(“HBM”)。在至少一个实施例中，诸如帧缓冲区或者纹理贴图(map)之类的再现目标可以跨存储器单元1224A-1224N存储，允许分区单元1220A-1220N并行地写每个再现目标的部分以便高效地使用并行处理器存储器1222的可用带宽。在至少一个实施例中，可以排除并行处理器存储器1222的本地实例，以利于结合本地缓存存储器利用系统存储器的统一存储器设计。
在至少一个实施例中，处理阵列1212的集群1214A-1214N中的任何一个可以处理将写到并行处理器存储器1222内的任何存储器单元1224A-1224N的数据。在至少一个实施例中，存储器交叉开关1216可以被配置为将每个集群1214A-1214N的输出传输到任何分区单元1220A-1220N或者另一个集群1214A-1214N，后者可以对输出实施附加的处理操作。在至少一个实施例中，每个集群1214A-1214N可以通过存储器交叉开关1216与存储器接口1218通信，以便从各种不同的外部存储器设备读取或者写入到各种不同的外部存储器设备。在至少一个实施例中，存储器交叉开关1216具有到存储器接口1218的连接以便与I/O单元1204通信，以及到并行处理器存储器1222的本地实例的连接，使得不同集群1214A-1214N内的处理单元能够与系统存储器或者不在并行处理单元1202本地的其他存储器通信。在至少一个实施例中，存储器交叉开关1216可以使用虚拟通道分离集群1214A-1214N与分区单元1220A-1220N之间的业务流。
在至少一个实施例中，并行处理单元1202的多个实例可以在单个外接卡上提供，或者多个外接卡可以互连。在至少一个实施例中，并行处理单元1202的不同实例可以被配置为互操作，即使不同实例具有不同数量的处理核心、不同数量的本地并行处理器存储器和/或其他配置差异。例如，在至少一个实施例中，并行处理单元1202的一些实例可以包括相对于其他实例较高精度的浮点单元。在至少一个实施例中，结合了并行处理单元1202或者并行处理器1200的一个或更多个实例的系统可以以各种各样的配置和形状因子实现，包括但不限于台式、膝上型或者手持式个人计算机、服务器、工作站、游戏控制台和/或嵌入式系统。
图12B图示出依照至少一个实施例的处理集群1294。在至少一个实施例中，处理集群1294包括在并行处理单元内。在至少一个实施例中，处理集群1294是图12的处理集群1214A-1214N之一。在至少一个实施例中，处理集群1294可以被配置为并行地执行许多线程，其中术语“线程”指的是在特定输入数据集合上执行的特定程序的实例。在至少一个实施例中，单指令多数据(“SIMD”)指令发布技术用来在不提供多个独立指令单元的情况下支持大量线程的并行执行。在至少一个实施例中，单指令多线程(“SIMT”)技术用来使用公共指令单元支持大量通常同步的线程的并行执行，所述公共指令单元被配置为向每个处理集群1294内的处理引擎集合发出指令。
在至少一个实施例中，处理集群1294的操作可以经由将处理任务分发给SIMT并行处理器的管线管理器1232进行控制。在至少一个实施例中，管线管理器1232接收来自图12的调度器1210的指令，并且经由图形多处理器1234和/或纹理单元1236管理那些指令的执行。在至少一个实施例中，图形多处理器1234是SIMT并行处理器的示例性实例。然而，在至少一个实施例中，不同架构的各种不同类型的SIMT并行处理器可以包括在处理集群1294内。在至少一个实施例中，图形多处理器1234的一个或更多个实例可以包括在处理集群1294内。在至少一个实施例中，图形多处理器1234可以处理数据，并且数据交叉开关1240可以用来将处理的数据分发给多个可能目的地之一，包括其他着色器单元。在至少一个实施例中，管线管理器1232可以通过指定要经由数据交叉开关1240分发的处理的数据的目的地而促进处理的数据的分发。
在至少一个实施例中，处理集群1294内的每个图形多处理器1234可以包括相同的功能执行逻辑(例如算术逻辑单元、加载/存储单元(“LSU”)等等)集合。在至少一个实施例中，功能执行逻辑可以以管线方式配置，其中新的指令可以在先前的指令完成之前发布。在至少一个实施例中，功能执行逻辑支持各种各样的操作，包括整数和浮点算术、比较运算、布尔运算、移位以及各种不同的代数函数的计算。在至少一个实施例中，可以利用相同功能单元硬件实施不同的操作，并且可以存在功能单元的任意组合。
在至少一个实施例中，传送至处理集群1294的指令构成线程。在至少一个实施例中，跨并行处理引擎集合执行的线程集合为线程组。在至少一个实施例中，线程组在不同的输入数据上执行程序。在至少一个实施例中，可以将线程组内的每个线程分配给图形多处理器1234内的不同处理引擎。在至少一个实施例中，线程组可以包括比图形多处理器1234内的处理引擎数量更少的线程。在至少一个实施例中，当线程组包括比处理引擎数量更少的线程时，这些处理引擎中的一个或更多个可能在正在处理线程组的周期期间是空闲的。在至少一个实施例中，线程组也可以包括比图形多处理器1234内的处理引擎数量更多的线程。在至少一个实施例中，当线程组包括比图形多处理器1234内的处理引擎数量更多的线程时，处理可以在连续的时钟周期内实施。在至少一个实施例中，可以在图形多处理器1234上并发地执行多个线程组。
在至少一个实施例中，图形多处理器1234包括实施加载和存储操作的内部缓存存储器。在至少一个实施例中，图形多处理器1234可以放弃内部缓存，并且使用处理集群1294内的缓存存储器(例如L1缓存1248)。在至少一个实施例中，每个图形多处理器1234也有权访问在所有处理集群1294之间共享并且可以用来在线程之间传输数据的分区单元(例如图12A的分区单元1220A-1220N)内的2级(“L2”)缓存。在至少一个实施例中，图形多处理器1234也可以访问片外全局存储器，该片外全局存储器可以包括本地并行处理器存储器和/或系统存储器中的一个或更多个。在至少一个实施例中，任何在并行处理单元1202外部的存储器都可以用作全局存储器。在至少一个实施例中，处理集群1294包括可以共享可能存储在L1缓存1248中的公共指令和数据的图形多处理器1234的多个实例。
在至少一个实施例中，每个处理集群1294可以包括被配置为将虚拟地址映射为物理地址的MMU 1245。在至少一个实施例中，MMU 1245的一个或更多个实例可以驻留在图12的存储器接口1218内。在至少一个实施例中，MMU 1245包括用来将虚拟地址映射到图块的物理地址以及可选地缓存行索引的页表条目(“PTE”)集合。在至少一个实施例中，MMU 1245可以包括可能驻留在图形多处理器1234或者L1缓存1248或者处理集群1294内的地址转换后备缓冲区(“TLB”)或缓存。在至少一个实施例中，物理地址经过处理以分布表面数据访问位置以便允许分区单元之间的高效请求交织。在至少一个实施例中，缓存行索引可以用来确定对于缓存行的请求命中还是未命中。
在至少一个实施例中，处理集群1294可以被配置为使得每个图形多处理器1234耦合到纹理单元1236以便实施纹理映射操作，例如，确定纹理样本位置、读取纹理数据以及对纹理数据滤波。在至少一个实施例中，纹理数据读取自内部纹理L1缓存(未示出)或者读取自图形多处理器1234内的L1缓存，并且在需要的情况下从L2缓存、本地并行处理器存储器或者系统存储器取得。在至少一个实施例中，每个图形多处理器1234将处理的任务输出至数据交叉开关1240以便将该处理的任务提供给另一个处理集群1294以供进一步处理或者以便经由存储器交叉开关1216将该处理的任务存储在L2缓存、本地并行处理器存储器或者系统存储器中。在至少一个实施例中，预光栅操作单元(“preROP”)1242被配置为接收来自图形多处理器1234的数据，将数据引导至ROP单元，所述ROP单元可以位于如本文所描述的分区单元(例如图12的分区单元1220A-1220N)内。在至少一个实施例中，preROP 1242可以实施用于混色的优化，组织像素颜色数据，以及实施地址转换。
图12C图示出依照至少一个实施例的图形多处理器1296。在至少一个实施例中，图形多处理器1296为图12B的图形多处理器1234。在至少一个实施例中，图形多处理器1296与处理集群1294的管线管理器1232耦合。在至少一个实施例中，图形多处理器1296具有执行管线，包括但不限于指令缓存1252、指令单元1254、地址映射单元1256、寄存器文件1258、一个或更多个GPGPU核心1262以及一个或更多个LSU 1266。GPGPU核心1262和LSU 1266经由存储器和缓存互连线1268与缓存存储器1272和共享的存储器1270耦合。
在至少一个实施例中，指令缓存1252从管线管理器1232接收要执行的指令流。在至少一个实施例中，指令缓存在指令缓存1252中，并且被指令单元1254分派以用于执行。在至少一个实施例中，指令单元1254可以将指令作为线程组(例如线程束)分派，线程组的每个线程分配给GPGPU核心1262内的不同执行单元。在至少一个实施例中，指令可以通过指定统一地址空间内的地址而访问本地的、共享的或者全局的地址空间中的任何一个。在至少一个实施例中，地址映射单元1256可以用来将统一地址空间内的地址转换为可以由LSU1266访问的不同的存储器地址。
在至少一个实施例中，寄存器文件1258为图形多处理器1296的功能单元提供寄存器集合。在至少一个实施例中，寄存器文件1258为连接到图形多处理器1296的功能单元(例如GPGPU核心1262、LSU 1266)的数据路径的操作数提供临时存储。在至少一个实施例中，在每个功能单元之间划分寄存器文件1258，使得每个功能单元被分配寄存器文件1258的专用部分。在至少一个实施例中，在正由图形多处理器1296执行的不同线程组之间划分寄存器文件1258。
在至少一个实施例中，GPGPU核心1262中的每一个可以包括用来执行图形多处理器1296的指令的FPU和/或整数ALU。GPGPU核心1262可以在架构上相似或者可以在架构上不同。在至少一个实施例中，GPGPU核心1262的第一部分包括单精度FPU和整数ALU，而GPGPU核心1262的第二部分包括双精度FPU。在至少一个实施例中，FPU可以实现用于浮点算术的IEEE 754-2008标准，或者允许实现可变精度浮点算术。在至少一个实施例中，图形多处理器1296可以附加地包括一个或更多个固定功能或特殊功能单元以实施诸如拷贝矩形或者像素混合操作之类的特定功能。在至少一个实施例中，GPGPU核心1262中的一个或更多个也可以包括固定或特殊功能逻辑。
在至少一个实施例中，GPGPU核心1262包括能够在多个数据集合上实施单个指令的SIMD逻辑。在至少一个实施例中，GPGPU核心1262可以物理地执行SIMD4、SIMD8和SIMD16指令，并且在逻辑上执行SIMD1、SIMD2和SIMD32指令。在至少一个实施例中，用于GPGPU核心1262的SIMD指令可以在编译时间由着色器编译器生成，或者在执行针对单程序多数据(“SPMD”)或SIMT架构写入和编译的程序时自动地生成。在至少一个实施例中，被配置用于SIMT执行模型的程序的多个线程可以经由单个SIMD指令执行。例如，在至少一个实施例中，实施相同或相似操作的八个SIMT线程可以经由单个SIMD8逻辑单元并行地执行。
在至少一个实施例中，存储器和缓存互连线1268是将图形多处理器1296的每个功能单元连接到寄存器文件1258和共享的存储器1270的互连网络。在至少一个实施例中，存储器和缓存互连线1268为交叉开关互连线，其允许LSU 1266实现共享的存储器1270与寄存器文件1258之间的加载和存储操作。在至少一个实施例中，寄存器文件1258可以以与GPGPU核心1262相同的频率操作，因此GPGPU核心1262与寄存器文件1258之间的数据传输延迟非常低。在至少一个实施例中，共享的存储器1270可以用来允许实现在图形多处理器1296内的功能单元上执行的线程之间的通信。在至少一个实施例中，缓存存储器1272可以用作例如数据缓存，以便缓存功能单元与纹理单元1236之间传送的纹理数据。在至少一个实施例中，共享的存储器1270也可以用作程序管理的缓存。在至少一个实施例中，除了存储在缓存存储器1272内的自动缓存的数据之外，在GPGPU核心1262上执行的线程还可以以编程方式将数据存储到共享的存储器中。
在至少一个实施例中，如本文所描述的并行处理器或者GPGPU通信耦合至主机/处理器核心以便加速图形操作、机器学习操作、模式分析操作和各种不同的通用GPU(GPGPU)功能。在至少一个实施例中，GPU可以通过总线或其他互连线(例如，诸如PCIe或NVLink之类的高速互连线)通信耦合至主机处理器/核心。在至少一个实施例中，GPU可以集成到与核心相同的封装或芯片上，并且通过在封装或芯片内部的处理器总线/互连线通信耦合至核心。在至少一个实施例中，不管GPU连接的方式如何，处理器核心都可以以WD中包含的命令/指令的序列的形式将工作分配给GPU。在至少一个实施例中，GPU于是使用专用电路系统/逻辑以便高效地处理这些命令/指令。
图13图示出依照至少一个实施例的图形处理器1300。在至少一个实施例中，图形处理器1300包括环形互连线1302、管线前端1304、媒体引擎1337和图形核心1380A-1380N。在至少一个实施例中，环形互连线1302将图形处理器1300耦合至其他处理单元，包括其他图形处理器或者一个或更多个通用处理器核心。在至少一个实施例中，图形处理器1300是集成到多核心处理系统中的许多处理器之一。
在至少一个实施例中，图形处理器1300经由环形互连线1302接收批量命令。在至少一个实施例中，传入命令由管线前端1304中的命令流程序(streamer)1303解释。在至少一个实施例中，图形处理器1300包括经由图形核心1380A-1380N实现3D几何处理和媒体处理的可扩展执行逻辑。在至少一个实施例中，对于3D几何处理命令，命令流程序1303将命令提供给集合管线1336。在至少一个实施例中，对于至少一些媒体处理命令，命令流程序1303将命令提供给与媒体引擎1337耦合的视频前端1334。在至少一个实施例中，媒体引擎1337包括用于视频和图像后处理的视频质量引擎(“VQE”)1330以及提供硬件加速的媒体数据编码和解码的多格式编码/解码(“MFX”)引擎1333。在至少一个实施例中，几何管线1336和媒体引擎1337中的每一个为至少一个图形核心1380A提供的线程执行资源生成执行线程。
在至少一个实施例中，图形处理器1300包括以模块化图形核心1380A-1380N(有时称为核心分片)为特点的可扩展线程执行资源，每个模块化图形核心具有多个子核心1350A-550N、1360A-1360N(有时称为核心子分片)。在至少一个实施例中，图形处理器1300可以具有任意数量的图形核心1380A直到1380N。在至少一个实施例中，图形处理器1300包括至少具有第一子核心1350A和第二子核心1360A的图形核心1380A。在至少一个实施例中，图形处理器1300是具有单个子核心(例如子核心1350A)的低功率处理器。在至少一个实施例中，图形处理器1300包括多个图形核心1380A-1380N，每个图形核心包括第一子核心1350A-1350N集合和第二子核心1360A-1360N集合。在至少一个实施例中，第一子核心1350A-1350N中的每个子核心至少包括第一组执行单元(“EU”)1352A-1352N和媒体/纹理采样器1354A-1354N。在至少一个实施例中，第二子核心1360A-1360N中的每个子核心至少包括第二组执行单元1362A-1362N和采样器1364A-1364N。在至少一个实施例中，每个子核心1350A-1350N、1360A-1360N共享共享资源1370A-1370N集合。在至少一个实施例中，共享资源1370包括共享缓存存储器和像素操作逻辑。
图14图示出依照至少一个实施例的处理器1400。在至少一个实施例中，处理器1400可以包括但不限于实施指令的逻辑电路。在至少一个实施例中，处理器1400可以实施指令，包括x86指令、ARM指令、用于ASIC的专用指令等等。在至少一个实施例中，处理器1400可以包括存储打包的数据的寄存器，例如来自加州圣克拉拉英特尔公司的利用MMX技术启用的微处理器中的64位宽MMXTM寄存器。在至少一个实施例中，有整数和浮点两种形式的MMX寄存器可以用伴随有SIMD和流式SIMD扩展(“SSE”)指令的打包数据元素操作。在至少一个实施例中，与SSE2、SSE3、SSE4、AVX或更高版本(统称为“SSEx”)技术有关的128位宽XMM寄存器可以保持这样的打包数据操作数。在至少一个实施例中，处理器1410可以实施加速CUDA程序的指令。
在至少一个实施例中，处理器1400包括取得要执行的指令以及准备以后要在处理器管线中使用的指令的有序前端(“前端”)1401。在至少一个实施例中，前端1401可以包括若干单元。在至少一个实施例中，指令预取器1426从存储器中取得指令并且将指令馈送至指令解码器1428，该指令解码器反过来解码或解释指令。例如，在至少一个实施例中，指令解码器1428将接收的指令解码成称为“微指令”或“微操作”(也称为“micro op”或“uop”)的用于执行的一个或更多个操作。在至少一个实施例中，指令解码器1428将指令解析成可以由微架构用来实施操作的操作码以及相应的数据和控制字段。在至少一个实施例中，踪迹缓存1430可以将解码的uop汇集成用于执行的uop队列1434中的程序排序的序列或踪迹。在至少一个实施例中，当踪迹缓存1430遇到复杂指令时，微码ROM 1432提供完成操作所需的uop。
在至少一个实施例中，一些指令可以转换成单个micro-op，而其他指令需要若干micro-op以完成整个操作。在至少一个实施例中，如果需要超过四个micro-op完成一定指令，那么指令解码器1428可以访问微码ROM 1432以实施指令。在至少一个实施例中，可以将指令解码成少量的micro-op以便在指令解码器1428处处理。在至少一个实施例中，如果需要一定数量的micro-op以完成操作，则可以将指令存储在微码ROM 1432内。在至少一个实施例中，踪迹缓存1430参考入口点可编程逻辑阵列(“PLA”)以确定用于从微码ROM 1432读取完成一个或更多个指令的微码序列的正确的微指令指针。在至少一个实施例中，在微码ROM 1432完成指令的micro-op排序之后，机器的前端1401可以继续从踪迹缓存1430取得micro-op。
在至少一个实施例中，乱序执行引擎(“乱序引擎”)1403可以准备用于执行的指令。在至少一个实施例中，乱序执行逻辑具有一定数量的缓冲区以便在指令流沿着管线前行并且被安排用于执行时整理并重新排序该指令流以便优化性能。乱序执行引擎1403包括但不限于分配器/寄存器重命名器1440、存储器uop队列1442、整数/浮点uop队列1444、存储器调度器1446、快速调度器1402、慢速/通用浮点调度器(“慢速/通用FP调度器”)1404以及简单浮点调度器(“简单FP调度器”)1406。在至少一个实施例中，快速调度器1402、慢速/通用浮点调度器1404和简单浮点调度器1406在这里也统称为“uop调度器1402、1404、1406”。分配器/寄存器重命名器1440分配每个uop需要的机器缓冲区和资源以便执行。在至少一个实施例中，分配器/寄存器重命名器1440将逻辑寄存器重命名到寄存器文件中的条目上。在至少一个实施例中，分配器/寄存器重命名器1440也为两个uop队列之一中的每个uop分配条目，所述uop队列为存储器调度器1446和uop调度器1402、1404、1406之前的用于存储器操作的存储器uop队列1442以及用于非存储器操作的整数/浮点uop队列1444。在至少一个实施例中，uop调度器1402、1404、1406基于其从属输入寄存器操作数源准备就绪以及uop完成其操作所需的执行资源的可用性确定uop何时准备执行。在至少一个实施例中，至少一个实施例的快速调度器1402可以在每半个主时钟周期进行调度，而慢速/通用浮点调度器1404和简单浮点调度器1406可以每个主处理器时钟周期调度一次。在至少一个实施例中，uop调度器1402、1404、1406仲裁分派端口以便调度uop以供执行。
在至少一个实施例中，执行块b11包括但不限于整数寄存器文件/旁路网络1408、浮点寄存器文件/旁路网络(“FP寄存器文件/旁路网络”)1410、地址生成单元(“AGU”)1412和1414、快速ALU 1416和1418、慢速ALU 1420、浮点ALU(“FP”)1422以及浮点移动单元(“FP移动”)1424。在至少一个实施例中，整数寄存器文件/旁路网络1408和浮点寄存器文件/旁路网络1410在这里也称为“寄存器文件1408、1410”。在至少一个实施例中，AGUS 1412和1414、快速ALU 1416和1418、慢速ALU 1420、浮点ALU 1422以及浮点移动单元1424在这里也称为“执行单元1412、1414、1416、1418、1420、1422和1424”。在至少一个实施例中，执行块可以包括但不限于任意组合的任意数量(包括零个)和类型的寄存器文件、旁路网络、地址生成单元和执行单元。
在至少一个实施例中，寄存器文件1408、1410可以布置在uop调度器1402、1404、1406与执行单元1412、1414、1416、1418、1420、1422和1424之间。在至少一个实施例中，整数寄存器文件/旁路网络1408实施整数运算。在至少一个实施例中，浮点寄存器文件/旁路网络1410实施浮点运算。在至少一个实施例中，寄存器文件1408、1410中的每一个可以包括但不限于旁路网络，该旁路网络可以绕过尚未写入寄存器文件中的刚刚完成的结果或者将其转发至新的从属uop。在至少一个实施例中，寄存器文件1408、1410可以彼此传送数据。在至少一个实施例中，整数寄存器文件/旁路网络1408可以包括但不限于两个单独的寄存器文件，一个用于低阶32位数据的寄存器文件以及另一个用于高阶32位数据的寄存器文件。在至少一个实施例中，浮点寄存器文件/旁路网络1410可以包括但不限于128位宽条目，因为浮点指令典型地具有宽度从64位到128位的操作数。
在至少一个实施例中，执行单元1412、1414、1416、1418、1420、1422和1424可以执行指令。在至少一个实施例中，寄存器文件1408、1410存储微指令需要执行的整数和浮点数据操作数。在至少一个实施例中，处理器1400可以包括但不限于任意数量和组合的执行单元1412、1414、1416、1418、1420、1422、1424。在至少一个实施例中，浮点ALU 1422和浮点移动单元1424可以执行浮点、MMX、SIMD、AVX和SSE或者其他运算。在至少一个实施例中，浮点ALU 1422可以包括但不限于64位乘64位浮点除法器以执行除法、平方根和余数micro op。在至少一个实施例中，涉及浮点值的指令可以利用浮点硬件进行处理。在至少一个实施例中，可以将ALU操作传递至快速ALU 1416、1418。在至少一个实施例中，快速ALU 1416、1418可以执行具有半个时钟周期的有效延迟的快速操作。在至少一个实施例中，最复杂的整数运算转给慢速ALU 1420，因为慢速ALU 1420可以包括但不限于用于长延迟类型的操作的整数执行硬件，例如乘法器、移位、标志逻辑和分支处理。在至少一个实施例中，存储器加载/存储操作可以由AGU 1412、1414执行。在至少一个实施例中，快速ALU 1416、快速ALU 1418和慢速ALU 1420可以在64位数据操作数上实施整数运算。在至少一个实施例中，快速ALU1416、快速ALU 1418和慢速ALU 1420可以被实现为支持各种各样的数据位大小，包括16、32、128、256等等。在至少一个实施例中，浮点ALU 1422和浮点移动单元1424可以被实现为具有各种宽度的位的一系列操作数。在至少一个实施例中，浮点ALU 1422和浮点移动单元1424可以结合SIMD和多媒体指令在128位宽打包数据操作数上操作。
在至少一个实施例中，uop调度器1402、1404、1406在父加载完成执行之前分派从属操作。在至少一个实施例中，由于uop可以在处理器1400中推测性地调度和执行，因而处理器1400也可以包括处理存储器未命中的逻辑。在至少一个实施例中，如果数据缓存中的数据加载未命中，那么在管线中可能存在留给调度器暂时不正确的数据的进行中(inflight)的从属操作。在至少一个实施例中，重播机制跟踪并且重新执行使用不正确数据的指令。在至少一个实施例中，从属操作可能需要重播，并且独立操作可以被允许完成。在至少一个实施例中，处理器的至少一个实施例的调度器和重播机制也可以被设计为捕获用于文本字符串比较操作的指令序列。
在至少一个实施例中，术语“寄存器”可以指板载处理器存储位置，其可以用作标识操作数的指令的部分。在至少一个实施例中，寄存器可以是可以从处理器外部(从编程员的角度)使用的那些寄存器。在至少一个实施例中，寄存器可能不限于特定类型的电路。相反地，在至少一个实施例中，寄存器可以存储数据、提供数据以及实施本文描述的功能。在至少一个实施例中，本文描述的寄存器可以使用任意数量的不同技术通过处理器内的电路系统实现，例如专用物理寄存器、使用寄存器重命名的动态分配的物理寄存器、专用和动态分配的物理寄存器的组合等等。在至少一个实施例中，整数寄存器存储32位整数数据。至少一个实施例的寄存器文件也包含用于打包的数据的八个多媒体SIMD寄存器。
图15图示出依照至少一个实施例的处理器1500。在至少一个实施例中，处理器1500包括但不限于一个或更多个处理器核心(“核心”)1502A-1502N、集成存储器控制器1514以及集成图形处理器1508。在至少一个实施例中，处理器1500可以包括附加的核心，直到且包括虚线框所表示的附加处理器核心1502N。在至少一个实施例中，处理器核心1502A-1502N中的每一个包括一个或更多个内部缓存单元1504A-1504N。在至少一个实施例中，每个处理器核心也有权访问一个或更多个共享的缓存单元1506。
在至少一个实施例中，内部缓存单元1504A-1504N和共享的缓存单元1506表示处理器1500内的一种缓存存储器层次结构。在至少一个实施例中，缓存存储器单元1504A-1504N可以包括每个处理器核心内的至少一个层级的指令和数据缓存以及一个或更多个层级的共享的中级缓存，例如L2、L3、4级(“L4”)，或者其他层级的缓存，其中将外部存储器之前的最高层级的缓存分类为LLC。在至少一个实施例中，缓存一致性逻辑维持各种不同的缓存单元1506与1504A-1504N之间的一致性。
在至少一个实施例中，处理器1500也可以包括一组一个或更多个总线控制器单元1516和系统代理核心1510。在至少一个实施例中，一个或更多个总线控制器单元1516管理外围总线集合，例如一个或更多个PCI或PCI快速总线。在至少一个实施例中，系统代理核心1510提供用于各种不同处理器部件的管理功能。在至少一个实施例中，系统代理核心1510包括管理对各种不同外部存储器设备(未示出)的访问的一个或更多个集成存储器控制器1514。
在至少一个实施例中，处理器核心1502A-1502N中的一个或更多个包括对于同时多线程的支持。在至少一个实施例中，系统代理核心1510包括用于在多线程处理期间协调和操作处理器核心1502A-1502N的部件。在至少一个实施例中，系统代理核心1510可以附加地包括功率控制单元(“PCU”)，该功率控制单元包括调节处理器核心1502A-1502N和图形处理器1508的一个或更多个功率状态的逻辑和部件。
在至少一个实施例中，处理器1500附加地包括执行图形处理操作的图形处理器1508。在至少一个实施例中，图形处理器1508与共享的缓存单元1506和包括一个或更多个集成存储器控制器1514的系统代理核心1510耦合。在至少一个实施例中，系统代理核心1510也包括驱动图形处理器输出到一个或更多个耦合的显示器的显示控制器1511。在至少一个实施例中，显示控制器1511也可以是经由至少一个互连线与图形处理器1508耦合的单独的模块，或者可以集成在图形处理器1508内。
在至少一个实施例中，基于环形的互连单元1512用来耦合处理器1500的内部部件。在至少一个实施例中，可以使用可替换的互连单元，例如点对点互连线、交换互连线或者其他技术。在至少一个实施例中，图形处理器1508经由I/O链路1513与环形互连线1512耦合。
在至少一个实施例中，I/O链路1513表示多种I/O互连线中的至少一个，包括促进各种不同的处理器部件与诸如eDRAM模块之类的高性能嵌入式存储器模块1518之间的通信的封装I/O互连线。在至少一个实施例中，处理器核心1502A-1502N中的每一个和图形处理器1508使用嵌入式存储器模块1518作为共享的LLC。
在至少一个实施例中，处理器核心1502A-1502N为执行公共指令集架构的同构核心。在至少一个实施例中，处理器核心1502A-1502N就I20而言是异构的，其中处理器核心1502A-1502N中的一个或更多个执行公共指令集，而处理器核心1502A-1502N中的一个或更多个其他核心执行公共指令集或者不同指令集的子集。在至少一个实施例中，处理器核心1502A-1502N就微架构而言是异构的，其中具有相对较高功耗的一个或更多个核心与具有较低功耗的一个或更多个核心耦合。在至少一个实施例中，处理器1500可以在一个或更多个芯片上实现或者实现为SoC集成电路。
图16图示出依照所描述的至少一个实施例的图形处理器核心1600。在至少一个实施例中，图形处理器核心1600包括在图形核心阵列中。在至少一个实施例中，有时称为核心分片的图形处理器核心1600可以是模块化图形处理器内的一个或更多个图形核心。在至少一个实施例中，图形处理器核心1600是一个图形核心分片的示例，并且如这里所描述的图形处理器可以包括基于目标功率和性能包迹的多个图形核心分片。在至少一个实施例中，每个图形核心1600可以包括与也称为子分片的多个子核心1601A-1601F耦合的固定功能块1630，其包括通用和固定功能逻辑的模块化块。
在至少一个实施例中，固定功能块1630包括几何/固定功能管线1636，其例如在较低性能和/或较低功率图形处理器实现方式中可以由图形处理器1600中的所有子核心共享。在至少一个实施例中，几何/固定功能管线1636包括3D固定功能管线、视频前端单元、线程派生器(spawner)和线程分派器以及管理统一返回缓冲区的统一返回缓冲区管理器。
在至少一个实施例中，固定功能块1630也包括图形SoC接口1637、图形微控制器1638和媒体管线1639。图形SoC接口1637提供图形核心1600与SoC集成电路内的其他处理器核心之间的接口。在至少一个实施例中，图形微控制器1638为可编程子处理器，其可配置来管理图形处理器1600的各种不同的功能，包括线程分派、调度和抢占。在至少一个实施例中，媒体管线1639包括促进对包括图像和视频数据在内的多媒体数据的解码、编码、预处理和/或后处理的逻辑。在至少一个实施例中，媒体管线1639经由对子核心1601-1601F内的计算或采样逻辑的请求实现媒体操作。
在至少一个实施例中，SoC接口1637使得图形核心1600能够与通用应用处理器核心(例如CPU)和/或SoC内的其他部件通信，所述部件包括存储器层次结构元件，例如共享的LLC存储器、系统RAM和/或嵌入式片上或封装DRAM。在至少一个实施例中，SoC接口1637也可以允许实现与诸如照相机成像管线之类的SoC内的固定功能设备的通信，并且允许使用和/或实现可以在图形核心1600与SoC内的CPU之间共享的全局存储器原子。在至少一个实施例中，SoC接口1637也可以实现用于图形核心1600的电源管理控制，并且允许实现图形核心1600的时钟域与SoC内的其他时钟域之间的接口。在至少一个实施例中，SoC接口1637允许实现从被配置为向图形处理器内的一个或更多个图形核心中的每一个提供命令和指令的命令流程序和全局线程分派器接收命令缓冲区。在至少一个实施例中，命令和指令可以在要实施媒体操作时分派给媒体管线1639，或者在要实施图形处理操作时分派给几何和固定功能管线(例如几何和固定功能管线1636、几何和固定功能管线1614)。
在至少一个实施例中，图形微控制器1638可以被配置为实施各种不同的用于图形核心1600的调度和管理任务。在至少一个实施例中，图形微控制器1638可以对子核心1601A-1601F内的执行单元(EU)阵列1602A-1602F、1604A-1604F内的各种不同的图形并行引擎实施图形和/或计算工作负载调度。在至少一个实施例中，在包括图形核心1600的SoC的CPU核心上执行的主机软件可以提交工作负载多个图形处理器门铃之一，这在适当的图形引擎上调用调度操作。在至少一个实施例中，调度操作包括确定接下来运行哪个工作负载，将工作负载提交给命令流程序，抢占在引擎上运行的现有工作负载，监视工作负载的进度，以及在工作负载完成时通知主机软件。在至少一个实施例中，图形微控制器1638也可以促进图形核心1600的低功率或空闲状态，向图形核心1600提供独立于操作系统和/或系统上的图形驱动程序软件，跨低功率状态过渡节省和恢复图形核心1600内的寄存器的能力。
在至少一个实施例中，图形核心1600可以具有大于或少于所图示的子核心1601A-1601F，高达N个模块化子核心。在至少一个实施例中，对于每组N个子核心，图形核心1600也可以包括共享的功能逻辑1610、共享的和/或缓存存储器1612、几何/固定功能管线1614以及附加的固定功能逻辑1616，以加速各种不同的图形和计算处理操作。在至少一个实施例中，共享的功能逻辑1610可以包括可以由图形核心1600内的每N个子核心共享的逻辑单元(例如采样器、数学和/或线程间通信逻辑)。共享的和/或缓存存储器1612可以是用于图形核心1600内的N个子核心1601A-1601F的LLC，并且也可以用作可由多个子核心访问的共享的存储器。在至少一个实施例中，几何/固定功能管线1614可以代替固定功能块1630内的几何/固定功能管线1636被包括在内，并且可以包括相同或相似的逻辑单元。
在至少一个实施例中，图形核心1600包括附加的固定功能逻辑1616，其可以包括供图形核心1600使用的各种不同的固定功能加速逻辑。在至少一个实施例中，附加的固定功能逻辑1616包括用于仅位置着色的附加的几何管线。在仅位置着色中，存在至少两个几何管线，而在几何/固定功能管线1616、1636内的完整几何管线中，剔除(cull)管线是可以包括在附加的固定功能逻辑1616内的附加的几何管线。在至少一个实施例中，剔除管线是完整几何管线的精简版本。在至少一个实施例中，完整管线和剔除管线可以执行一定应用的不同实例，每个实例具有单独的上下文。在至少一个实施例中，仅位置着色可以隐藏丢弃的三角形的长时间剔除运行，使得着色在一些实例中能够更早完成。例如，在至少一个实施例中，附加的固定功能逻辑1616内的剔除管线逻辑可以与主应用并行地执行位置着色器，并且通常比完整管线更快地生成关键结果，因为剔除管线取得并着色顶点的位置属性，而不实施光栅化和将像素再现到帧缓冲区。在至少一个实施例中，剔除管线可以使用生成的关键结果计算所有三角形的可见性信息，而不考虑那些三角形是否被剔除。在至少一个实施例中，完整管线(其在本实例中可以称为重播管线)可以消耗可见性信息以跳过剔除的三角形以便仅仅着色最终传递至光栅化阶段的可见三角形。
在至少一个实施例中，附加的固定功能逻辑1616也可以包括诸如固定功能矩阵乘法逻辑之类的通用处理加速逻辑，以便加速CUDA程序。
在至少一个实施例中，每个图形子核心1601A-1601F包括可以用来响应于图形管线、媒体管线或者着色器程序的请求实施图形、媒体和计算操作的执行资源集合。在至少一个实施例中，图形子核心1601A-1601F包括多个EU阵列1602A-1602F、1604A-1604F，线程分派和线程间通信(“TD/IC”)逻辑1603A-1603F，3D(例如纹理)采样器1605A-1605F，媒体采样器1606A-1606F，着色器处理器1607A-1607F，以及共享的本地存储器(“SLM”)1608A-1608F。EU阵列1602A-1602F、1604A-1604F中的每一个包括多个执行单元，这些执行单元是能够为服务图形、媒体或计算操作而实施浮点和整数/定点逻辑运算的GPGPU，包括图形、媒体或计算着色器程序。在至少一个实施例中，TD/IC逻辑1603A-1603F为子核心内的执行单元实施本地线程分派和线程控制操作，并且促进在子核心的执行单元上执行的线程之间的通信。在至少一个实施例中，3D采样器1605A-1605F可以将纹理或者其他3D图形相关数据读入存储器中。在至少一个实施例中，3D采样器可以基于配置的样本状态和与给定纹理关联的纹理格式不同地读取纹理数据。在至少一个实施例中，媒体采样器1606A-1606F可以基于与媒体数据关联的类型和格式实施相似的读取操作。在至少一个实施例中，每个图形子核心1601A-1601F可以可替换地包括统一3D和媒体采样器。在至少一个实施例中，在子核心1601A-1601F中的每一个内的执行单元上执行的线程可以利用每个子核心内的共享的本地存储器1608A-1608F，以使得线程组内执行的线程能够使用片上存储器的公共池执行。
图17图示出依照至少一个实施例的并行处理单元(“PPU”)1700。在至少一个实施例中，PPU 1700被配置有机器可读代码，该代码在被PPU 1700执行的情况下使得PPU 1700执行本文描述的过程和技术中的一些或全部。在至少一个实施例中，PPU 1700是一种多线程化处理器，其在一个或更多个集成电路设备上实现并且利用多线程作为延迟隐藏技术，该技术被设计为在多个线程上并行地处理计算机可读指令(也称为机器可读指令或者简称为指令)。在至少一个实施例中，线程指的是执行线程并且是被配置为由PPU 1700执行的指令集的实例化。在至少一个实施例中，PPU 1700是一种GPU，其被配置为实现用于处理三维(“3D”)图形数据以便生成在诸如LCD设备之类的显示设备上显示的二维(“2D”)图像数据的图形再现管线。在至少一个实施例中，PPU 1700被利用来实施诸如线性代数运算和机器学习操作之类的计算。图17仅仅出于说明的目的图示出一种示例性并行处理器，并且应当被解释为可以在至少一个实施例中实现的处理器架构的一个非限制性示例。
在至少一个实施例中，一个或更多个PPU 1700被配置为加速高性能计算(“HPC”)、数据中心和机器学习应用。在至少一个实施例中，一个或更多个PPU 1700被配置为加速CUDA程序。在至少一个实施例中，PPU 1700包括但不限于I/O单元1706、前端单元1710、调度器单元1712、工作分发单元1714、集线器1716、交叉开关(“Xbar”)1720、一个或更多个通用处理集群(“GPC”)1718以及一个或更多个分区单元(“存储器分区单元”)1722。在至少一个实施例中，PPU 1700经由一个或更多个高速GPU互连线(“GPU互连线”)1708连接到主机处理器或者其他PPU 1700。在至少一个实施例中，PPU 1700经由互连线1702连接到主机处理器或者其他外围设备。在至少一个实施例中，PPU 1700连接到包括一个或更多个存储器设备(“存储器”)1704在内的本地存储器。在至少一个实施例中，存储器设备1704包括但不限于一个或更多个动态随机存取存储器(DRAM)设备。在至少一个实施例中，一个或更多个DRAM设备被配置为和/或可配置为高带宽存储器(“HBM”)子系统，具有堆叠在每个设备内的多个DRAM管芯。
在至少一个实施例中，高速GPU互连线1708可以指基于有线的多巷道通信链路，其由系统用来扩展和包括与一个或更多个CPU组合的一个或更多个PPU 1700，支持PPU 1700与CPU之间的缓存连贯性以及CPU主控。在至少一个实施例中，数据和/或命令由高速GPU互连线1708通过集线器1716向/从PPU 1700的其他单元传送，所述其他单元例如一个或更多个拷贝引擎、视频编码器、视频解码器、电源管理单元以及图17中可能未显式图示出的其他部件。
在至少一个实施例中，I/O单元1706被配置为通过系统总线1702从主机处理器(未在图17中图示出)传送和接收通信(例如命令、数据)。在至少一个实施例中，I/O单元1706经由系统总线1702直接地或者通过诸如存储器桥接器之类的一个或更多个中间设备与主机处理器通信。在至少一个实施例中，I/O单元1706可以经由系统总线1702与诸如PPU 1700中的一个或更多个之类的一个或更多个其他处理器通信。在至少一个实施例中，I/O单元1706实现用于通过PCIe总线进行通信的PCIe接口。在至少一个实施例中，I/O单元1706实现用于与外部设备通信的接口。
在至少一个实施例中，I/O单元1706解码经由系统总线1702接收的数据包。在至少一个实施例中，至少一些数据包表示被配置为使得PPU 1700实施各种不同的操作的命令。在至少一个实施例中，I/O单元1706将解码的命令传送至如命令所指定的PPU 1700的各种不同的其他单元。在至少一个实施例中，将命令传送至前端单元1710和/或传送至PPU 1700的集线器1716或其他单元，例如一个或更多个拷贝引擎、视频编码器、视频解码器、电源管理单元等等(未在图17中显式地图示出)。在至少一个实施例中，I/O单元1706被配置为路由PPU 1700的各种不同的逻辑单元之间和之中的通信。
在至少一个实施例中，主机处理器执行的程序编码缓冲区中将工作负载提供给PPU 1700以供处理的命令流。在至少一个实施例中，工作负载包括指令以及要由那些指令处理的数据。在至少一个实施例中，缓冲区是可由主机处理器和PPU 1700二者访问(例如读/写)的存储器中的区域——主机接口单元可以被配置为经由I/O单元1706通过系统总线1702传送的存储器请求访问连接到系统总线1702的系统存储器中的缓冲区。在至少一个实施例中，主机处理器将命令流写到缓冲区，并且然后将指向命令流的开始的指针传送至PPU1700，使得前端单元1710接收指向一个或更多个命令流的指针，并且管理一个或更多个命令流，从命令流中读取命令以及将命令转发至PPU 1700的各种不同的单元。
在至少一个实施例中，前端单元1710耦合到调度器单元1712，该调度器单元配置各个不同的GPC 1718以处理由一个或更多个命令流限定的任务。在至少一个实施例中，调度器单元1712被配置为跟踪与调度器单元1712管理的各种不同的任务有关的状态信息，其中状态信息可以指示任务被分配给哪个GPC 1718，任务是活动的还是不活动的，与任务关联的优先级，等等。在至少一个实施例中，调度器单元1712管理GPC 1718中的一个或更多个上的多个任务的执行。
在至少一个实施例中，调度器单元1712耦合到工作分发单元1714，该工作分发单元被配置为分派用于在GPC 1718上执行的任务。在至少一个实施例中，工作分发单元1714跟踪从调度器单元1714接收的一定数量的调度任务，并且工作分发单元1714为每一个GPC1718管理待决任务池和活动任务池。在至少一个实施例中，待决任务池包括包含被分配由特定GPC 1718处理的任务的一定数量的分槽(例如32个分槽)；活动任务池可以包括用于正由GPC 1718积极处理的任务的一定数量的分槽(例如4个分槽)，使得当GPC 1718之一完成任务的执行时，该任务从用于GPC 1718的活动任务池中逐出，并且来自待决任务池的其他任务之一被选择且调度用于在GPC 1718上执行。在至少一个实施例中，如果活动任务在GPC1718上是空闲，例如当等待解决数据从属性时，那么该活动任务从GPC 1718中逐出并且返回到待决任务池，同时待决任务池中的另一个任务被选择并且调度用于在GPC 1718上执行。
在至少一个实施例中，工作分发单元1714经由XBar 1720与一个或更多个GPC1718通信。在至少一个实施例中，XBar 1720是一种将PPU 1700的许多单元耦合到PPU 1700的其他单元并且可以被配置为将工作分发单元1714耦合到特定GPC 1718的互连网络。在至少一个实施例中，PPU 1700的一个或更多个其他单元也可以经由集线器1716连接到XBar1720。
在至少一个实施例中，任务由调度器单元1712管理并且由工作分发单元1714分派到GPC 1718之一。GPC 1718被配置为处理任务并且生成结果。在至少一个实施例中，结果可以由GPC 1718内的其他任务消费，经由XBar 1720路由至不同的GPC 1718，或者存储在存储器1704中。在至少一个实施例中，结果可以经由分区单元1722写到存储器1704，该分区单元实现一种用于向/从存储器1704读写数据的存储器接口。在至少一个实施例中，结果可以经由高速GPU互连线1708传送至另一个PPU 1704或CPU。在至少一个实施例中，PPU 1700包括但不限于一定数量U的分区单元1722，该数量等于耦合到PPU 1700的单独的且不同的存储器设备1704的数量。
在至少一个实施例中，主机处理器执行实现应用编程接口(“API”)的驱动程序内核，该应用编程接口使得主机处理器上执行的一个或更多个应用能够调度用于在PPU 1700上执行的操作。在至少一个实施例中，多个计算应用由PPU 1700同时执行，并且PPU 1700为多个计算应用提供隔离、服务质量(“QoS”)和独立地址空间。在至少一个实施例中，应用生成使得驱动程序内核生成供PPU 1700执行的一个或更多个任务的(例如API调用形式的)指令，并且驱动程序内核将任务输出到正由PPU 1700处理的一个或更多个流。在至少一个实施例中，每个任务包括可以称为线程束的一组或多组相关线程。在至少一个实施例中，线程束包括可以并行地执行的多个相关线程(例如32个线程)。在至少一个实施例中，合作线程可以指包括实施任务的指令并且通过共享的存储器交换数据的多个线程。
图18图示出依照至少一个实施例的GPC 1800。在至少一个实施例中，GPC 1800是图17的GPC 1718。在至少一个实施例中，每个GPC 1800包括但不限于用于处理任务的一定数量的硬件单元，并且每个GPC 1800包括但不限于管线管理器1802、预光栅操作单元(“PROP”)1804、光栅引擎1808、工作分发交叉开关(“WDX”)1816、MMU 1818、一个或更多个数据处理集群(“DPC”)1806以及任何适当的零件组合。
在至少一个实施例中，GPC 1800的操作由管线管理器1802控制。在至少一个实施例中，管线管理器1802管理用于处理分配给GPC 1800的任务的一个或更多个DPC 1806的配置。在至少一个实施例中，管线管理器1802将一个或更多个DPC 1806中的至少一个配置为实现图形再现管线的至少一部分。在至少一个实施例中，DPC 1806被配置为在可编程流式多处理器(“SM”)1814上执行顶点着色器程序。在至少一个实施例中，管线管理器1802被配置为将接收自工作分发单元的数据包路由至GPC 1800内的适当逻辑单元，并且在至少一个实施例中，一些数据包可以被路由至PROP 1804中的固定功能硬件单元和/或光栅引擎1808，而其他数据包可以被路由至DPC 1806以供图元引擎1812或SM 1814处理。在至少一个实施例中，管线管理器1802将DPC 1806中的至少一个配置为实现计算管线。在至少一个实施例中，管线管理器1802将DPC 1806中的至少一个配置为执行CUDA程序的至少一部分。
在至少一个实施例中，PROP单元1804被配置为将光栅引擎1808和DPC 1806生成的数据路由至诸如上面结合图17更详细地描述的存储器分区单元1722之类的分区单元中的光栅操作(“ROP”)单元。在至少一个实施例中，PROP单元1804被配置为针对颜色混合实施优化，组织像素数据，实施地址转换等等。在至少一个实施例中，光栅引擎1808包括但不限于被配置为实施各种不同的光栅操作的一定数量的固定功能硬件单元，并且在至少一个实施例中，光栅引擎1808包括但不限于设置引擎、粗略光栅引擎、剔除引擎、裁剪引擎、精细光栅引擎、图块合并引擎及其任何适当的组合。在至少一个实施例中，设置引擎接收变换的顶点，并且生成与顶点限定的几何图元关联的平面方程；平面方程被传送至粗略光栅引擎以便生成用于图元的覆盖范围信息(例如用于图块的x,y覆盖掩码)；粗略光栅引擎的输出被传送至其中未通过z测试的与图元关联的片段被剔除的剔除引擎，并且被传送至其中位于观察截椎体之外的片段被裁剪的裁剪引擎。在至少一个实施例中，裁剪和剔除后幸存的片段被传递至精细光栅引擎以便基于设置引擎生成的平面方程生成用于像素片段的属性。在至少一个实施例中，光栅引擎1808的输出包括要由任何适当的实体处理，例如由在DPC1806内实现的片段着色器处理的片段。
在至少一个实施例中，GPC 1800中包括的每个DPC 1806包括但不限于M型管道控制器(“MPC”)1810、图元引擎1812、一个或更多个SM 1814及其任何适当的组合。在至少一个实施例中，MPC 1810控制DPC 1806的操作，将接收自管线管理器1802的数据包路由至DPC1806中的适当单元。在至少一个实施例中，与顶点关联的数据包被路由至图元引擎1812，该图元引擎被配置为从存储器中取得与顶点关联的顶点属性；形成对照的是，可以将与着色器程序关联的数据包传送至SM 1814。
在至少一个实施例中，SM 1814包括但不限于被配置为处理由一定数量的线程表示的任务的可编程流式处理器。在至少一个实施例中，SM 1814是多线程的并且被配置为并发地执行来自特定线程组的多个线程(例如32个线程)，并且实现一种SIMD架构，其中线程组(例如线程束)中的每个线程被配置为基于相同的指令集处理不同的数据集合。在至少一个实施例中，线程组中的所有线程执行相同的指令。在至少一个实施例中，SM 1814实现一种SIMT架构，其中线程组中的每个线程被配置为基于相同的指令集处理不同的数据集合，但是其中线程组中的各个线程在执行期间被允许发散。在至少一个实施例中，为每个线程束维持程序计数器、调用堆栈和执行状态，在线程束内的线程发散时允许实现线程束之间的并发性以及线程束内的串行执行。在另一个实施例中，为每个个别的线程维持程序计数器、调用堆栈和执行状态，允许实现线程数之内和之间的所有线程之间的同样并发性。在至少一个实施例中，为每个个别线程维持执行状态，并且执行相同指令的线程可以收敛并且并行地执行以提高效率。SM 1814的至少一个实施例结合图19更详细地进行了描述。
在至少一个实施例中，MMU 1818提供GPC 1800与存储器分区单元(例如图17的分区单元1722)之间的接口，并且MMU 1818提供虚拟地址-物理地址的转换、存储器保护以及存储器请求的仲裁。在至少一个实施例中，MMU 1818提供用于实施虚拟地址到存储器中的物理地址的转换的一个或更多个转换后备缓冲区(TLB)。
图19图示出依照至少一个实施例的流式多处理器(“SM”)1900。在至少一个实施例中，SM 1900是图18的SM 1814。在至少一个实施例中，SM 1900包括但不限于指令缓存1902、一个或更多个调度器单元1904、寄存器文件1908、一个或更多个处理核心(“核心”)1910、一个或更多个特殊功能单元(“SFU”)1912、一个或更多个LSU 1914、互连网络1916、共享的存储器/L1缓存1918及其任何适当的组合。在至少一个实施例中，工作分发单元分派用于在并行处理单元(PPU)的GPC上执行的任务，每个任务被分配给GPC内的特定数据处理集群(DPC)，并且如果任务与着色器程序关联，那么该任务被分配给SM 1900之一。在至少一个实施例中，调度器单元1904从工作分发单元接收任务，并且为分配给SM 1900的一个或更多个线程块管理指令调度。在至少一个实施例中，调度器单元1904调度线程块以作为并行线程的线程束而执行，其中每个线程块被分配至少一个线程束。在至少一个实施例中，每个线程束执行线程。在至少一个实施例中，调度器单元1904管理多个不同线程块，将线程束分配给不同线程块并且然后在每个时钟周期期间将来自多个不同合作组的指令分派给各种不同的功能单元(例如处理核心1910、SFU 1912和LSU 1914)。
在至少一个实施例中，“合作组”可以指用于组织通信线程组的编程模型，其允许开发者表达线程通信的粒度，允许实现更丰富、更高效的并行分解的表达。在至少一个实施例中，合作启动API支持线程块之间的同步以用于并行算法的执行。在至少一个实施例中，常规编程模型的API提供用于同步合作线程的单个简单结构体(construct)：跨线程块的所有线程的屏障(例如syncthreads()函数)。然而，在至少一个实施例中，程序员可以限定比线程块更小粒度的线程组，并且在限定的组内同步以便允许实现更大的性能、设计灵活性以及集体组范围的功能接口形式的软件重用。在至少一个实施例中，软件组使得程序员能够在子块和多块粒度下显式地限定线程组，并且在合作组中实施诸如线程同步之类的集体操作。在至少一个实施例中，子块粒度与单线程一样小。在至少一个实施例中，编程模型支持跨软件边界的清洁组成，使得库和效用函数在其本地上下文内可以安全地同步而不必对收敛性作出假设。在至少一个实施例中，合作组图元允许实现合作并行的新模式，包括但不限于生产者-消费者并行、机会主义并行以及跨线程块的整个网格的全局同步。
在至少一个实施例中，分派单元1906被配置为将指令传送至功能单元中的一个或更多个，并且调度器单元1904包括但不限于允许在每个时钟周期期间分派来自相同线程束的两个不同指令的两个分派单元1906。在至少一个实施例中，每个调度器单元1904包括单个分派单元1906或者附加的分派单元1906。
在至少一个实施例中，每个SM 1900在至少一个实施例中包括但不限于为SM 1900的功能单元提供寄存器集合的寄存器文件1908。在至少一个实施例中，在每个所述功能单元之间划分寄存器文件1908，使得每个功能单元被分配寄存器文件1908的专用部分。在至少一个实施例中，在正由SM 1900执行的不同线程束之间划分寄存器文件1908，并且寄存器文件1908为连接到功能单元的数据路径的操作数提供临时存储。在至少一个实施例中，每个SM 1900包括但不限于多个L处理核心1910。在至少一个实施例中，SM 1900包括但不限于大量(例如128个或更多)不同的处理核心1910。在至少一个实施例中，每个处理核心1910包括但不限于完全流水线式、单精度、双精度和/或混合精度处理单元，其包括但不限于浮点算术逻辑单元和整数算术逻辑单元。在至少一个实施例中，浮点算术逻辑单元实现用于浮点算术的IEEE 754-2008标准。在至少一个实施例中，处理核心1910包括但不限于64个单精度(32位)浮点核心、64个整数核心、32个双精度(64位)浮点核心和8个张量核心。
在至少一个实施例中，张量核心被配置为实施矩阵运算。在至少一个实施例中，一个或更多个张量核心包括在处理核心1910中。在至少一个实施例中，张量核心被配置为实施深度学习矩阵算术，例如用于神经网络训练和推理的卷积运算。在至少一个实施例中，每个张量核心在4x4矩阵上操作，并且实施矩阵乘法和累加运算D＝A X B+C，其中A、B、C和D为4x4矩阵。
在至少一个实施例中，矩阵乘法输入A和B为16位浮点矩阵，累加矩阵C和D为16位浮点或32位浮点矩阵。在至少一个实施例中，张量核心利用32位浮点累加对16位浮点输入数据进行操作。在至少一个实施例中，对于4x4x4矩阵乘法而言，16位浮点乘法使用64次运算并且导致全精度乘积，该全精度乘积然后使用32位浮点运算与其他中间乘积累加。在至少一个实施例中，张量核心用来实施由这些较小元素组成的、大得多的二维或更高维矩阵运算。在至少一个实施例中，诸如CUDA-C++API之类的API暴露专用的矩阵加载、矩阵乘法和累加以及矩阵存储操作以便从CUDA-C++程序高效地使用张量核心。在至少一个实施例中，在CUDA层级，线程束级接口假定16x16大小的矩阵跨越线程束的所有32个线程。
在至少一个实施例中，每个SM 1900包括但不限于执行特殊功能(例如属性求值、倒数平方根等等)的M个SFU 1912。在至少一个实施例中，SFU 1912包括但不限于被配置为遍历分层树数据结构的树遍历单元。在至少一个实施例中，SFU 1912包括但不限于被配置为实施纹理贴图滤波操作的纹理单元。在至少一个实施例中，纹理单元被配置为加载来自存储器的纹理贴图(例如2D纹素(texel)阵列)和样本纹理贴图以便产生用在SM 1900执行的着色器程序中的采样纹理值。在至少一个实施例中，纹理贴图存储在共享的存储器/L1缓存1918中。在至少一个实施例中，纹理单元使用mip贴图(例如细节级别变化的纹理贴图)实现诸如滤波操作之类的纹理操作。在至少一个实施例中，每个SM 1900包括但不限于两个纹理单元。
在至少一个实施例中，每个SM 1900包括但不限于在共享的存储器/L1缓存1918与寄存器文件1908之间实现加载和存储操作的N个LSU 1914。在至少一个实施例中，每个SM1900包括但不限于将功能单元中的每一个连接到寄存器文件1908以及将LSU 1914连接到寄存器文件1908和共享的存储器/L1缓存1918的互连网络1916。在至少一个实施例中，互连网络1916是一种交叉开关，其可以被配置为将任何功能单元连接到寄存器文件1908中的任何寄存器并且将LSU 1914连接到寄存器文件1908以及共享的存储器/L1缓存1918中的存储器位置。
在至少一个实施例中，共享的存储器/L1缓存1918是一种片上存储器阵列，其允许SM 1900与图元引擎之间以及SM 1900中的线程之间的数据存储和通信。在至少一个实施例中，共享的存储器/L1缓存1918包括但不限于128KB的存储容量，并且处于从SM 1900到分区单元的路径上。在至少一个实施例中，共享的存储器/L1缓存1918用来缓存读和写。在至少一个实施例中，共享的存储器/L1缓存1918、L2缓存和存储器中的一个或更多个是后备贮存器。
在至少一个实施例中，将数据缓存和共享的存储器功能组合到单个存储器块中为这两种类型的存储器访问提供了改善的性能。在至少一个实施例中，容量被不使用共享的存储器的程序用作或者可通过这些程序用作缓存，例如如果共享的存储器被配置为使用一半的容量，那么纹理和加载/存储操作可以使用剩余的容量。在至少一个实施例中，集成到共享的存储器/L1缓存1918中使得共享的存储器/L1缓存1918能够用作高吞吐量导管以便流送数据，同时提供对于经常重用的数据的高带宽和低延迟访问。在至少一个实施例中，当被配置用于通用并行计算时，可以使用与图形处理相比更简单的配置。在至少一个实施例中，固定功能GPU被绕过，创建简单得多的编程模型。在至少一个实施例中以及在通用并行计算配置中，工作分发单元直接将线程块分配和分发至DPC。在至少一个实施例中，块中的线程执行相同的程序，在计算中使用唯一的线程ID以确保每个线程生成唯一的结果，使用SM 1900执行程序并且实施计算，使用共享的存储器/L1缓存1918在线程之间通信，使用LSU1914通过共享的存储器/L1缓存1918和存储器分区单元读写全局存储器。在至少一个实施例中，当被配置用于通用并行计算时，SM 1900写入调度器单元1904可以用来在DPC上启动新工作的命令。
在至少一个实施例中，PPU包含于或者耦合到台式计算机、膝上型计算机、平板计算机、服务器、超级计算机、智能电话(例如无线手持式设备)、PDA、数字照相机、车辆、头戴式显示器、手持式电子设备等等。在至少一个实施例中，PPU在单个半导体衬底上实施。在至少一个实施例中，PPU与一个或更多个其他设备一起包括在SoC中，所述其他设备例如附加的PPU、存储器、RISC CPU、MMU、数模转换器(“DAC”)等等。
在至少一个实施例中，PPU可以包括在包括一个或更多个存储器设备的图形卡上。在至少一个实施例中，图形卡可以被配置为与台式计算机的主板上的PCIe接口。在至少一个实施例中，PPU可以是包括在主板芯片集中的集成GPU(“iGPU”)。
用于通用计算的软件结构体
下面的图非限制性地阐述了用于实现至少一个实施例的示例性软件结构体。
图20图示出依照至少一个实施例的编程平台的软件堆栈。在至少一个实施例中，编程平台是用于利用计算系统上的硬件加速计算任务的平台。在至少一个实施例中，编程平台可以通过库、编译器指令和/或对编程语言的扩展而为软件开发者访问。在至少一个实施例中，编程平台可以是但不限于CUDA、Radeon开放计算平台(“ROCm”)、OpenCL(OpenCLTM由Khronos集团开发)、SYCL或者英特尔One API。
在至少一个实施例中，编程平台的软件堆栈2000为应用2001提供了一种执行环境。在至少一个实施例中，应用2001可以包括能够在软件堆栈2000上启动的任何计算机软件。在至少一个实施例中，应用2001可以包括但不限于人工智能(“AI”)/机器学习(“ML”)应用、高性能计算(“HPC”)应用、虚拟桌面基础结构(“VDI”)或者数据中心工作负载。
在至少一个实施例中，应用2001和软件堆栈2000运行在硬件2007上。在至少一个实施例中，硬件2007可以包括一个或更多个GPU、CPU、FPGA、AI引擎和/或支持编程平台的其他类型的计算设备。在至少一个实施例中，例如对于CUDA，软件堆栈2000可以是特定于供应商的，并且仅仅与来自特定供应商的设备兼容。在至少一个实施例中，例如对于OpenCL，软件堆栈2000可以与来自不同供应商的设备一起使用。在至少一个实施例中，硬件2007包括连接到可以经由应用编程接口(“API”)调用进行访问以实施计算任务的一个或更多个设备的主机。在至少一个实施例中，硬件2007内的设备可以包括但不限于GPU、FPGA、AI引擎或者其他计算设备(但是也可以包括CPU)及其存储器，这与可以包括但不限于CPU(但是也可以包括计算设备)及其存储器的硬件2007内的主机相反。
在至少一个实施例中，编程平台的软件堆栈2000包括但不限于一定数量的库2003、运行时2005和设备内核驱动程序2006。在至少一个实施例中，库2003中的每一个可以包括数据和编程代码，其可以由计算机程序使用并且在软件开发期间加以利用。在至少一个实施例中，库2003可以包括但不限于预写入代码和子例程、类、值、类型规范、配置数据、文档、帮助数据和/或消息模板。在至少一个实施例中，库2003包括为在一种或多种类型的设备上执行而优化的功能。在至少一个实施例中，库2003可以包括但不限于用于在设备上实施数学、深度学习和/或其他类型的运算的功能。在至少一个实施例中，库2103与相应的API 2102关联，其可以包括暴露在库2103中实现的功能的一个或更多个API。
在至少一个实施例中，如下文结合图25-27更详细地讨论的，应用2001被编写为源代码，该源代码被编译成可执行代码。在至少一个实施例中，应用2001的可执行代码可以至少部分地在软件堆栈2000提供的执行环境上运行。在至少一个实施例中，在应用2001执行期间，与主机相反的是，可以触及需要在设备上运行的代码。在至少一个实施例中，在这样的情况下，可以调用运行时2005以便在设备上加载和启动必要的代码。在至少一个实施例中，运行时2005可以包括能够支持应用S01的执行的任何技术上可行的运行时系统。
在至少一个实施例中，运行时2005被实现为与示为API 2004的相应API关联的一个或更多个运行时库。在至少一个实施例中，这样的运行时库中的一个或更多个可以包括但不限于用于除其他以外的存储器管理、执行控制、设备管理、错误处理和/或同步的功能。在至少一个实施例中，存储器管理功能可以包括但不限于分配、解除分配和拷贝设备存储器以及在主机存储器与设备存储器之间传输数据的功能。在至少一个实施例中，执行控制功能可以包括但不限于在设备上启动功能(当功能为可从主机调用的全局功能时，有时称为“内核”)以及对于要在设备上执行的给定功能设置运行时库维持的缓冲区中的属性值的功能。
在至少一个实施例中，运行时库和相应的API 2004可以以任何技术上可行的方式实现。在至少一个实施例中，一个(或者任意数量的)API可以暴露用于设备的细粒度控制的低层级功能集合，而另一个(或者任意数量的)API可以暴露更高层级的这样的功能集合。在至少一个实施例中，高层级运行时API可以建立在低层级API之上。在至少一个实施例中，运行时API中的一个或更多个可以是特定于语言的API，其分层位于独立于语言的运行时API之上。
在至少一个实施例中，设备内核驱动程序2006被配置为促进与底层设备的通信。在至少一个实施例中，设备内核驱动程序2006可以提供诸如API 2004之类的API和/或其他软件所依赖的低层级功能。在至少一个实施例中，设备内核驱动程序2006可以被配置为在运行时将中间表示(“IR”)代码编译成二进制代码。在至少一个实施例中，对于CUDA而言，设备内核驱动程序2006可以在运行时将不特定于硬件的并行线程执行(“PTX”)IR代码编译成用于特定目标设备的二进制代码(缓存编译的二进制代码)，这有时也称为“定稿”代码。在至少一个实施例中，这样做可以允许定稿的代码在目标设备上运行，其在最初将源代码编译成PTX代码时可能并不存在。可替换地，在至少一个实施例中，可以离线地将设备源代码编译成二进制代码，而无需设备内核驱动程序2006在运行时编译IR代码。
图21图示出依照至少一个实施例的图20的软件堆栈2000的CUDA实现方式。在至少一个实施例中，其上可以启动应用2001的CUDA软件堆栈2100包括CUDA库2103、CUDA运行时2105、CUDA驱动程序2107和设备内核驱动程序2108。在至少一个实施例中，CUDA软件堆栈2100在硬件2109上执行，该硬件可以包括支持CUDA并且由加州圣克拉拉英伟达公司开发的GPU。
在至少一个实施例中，应用2101、CUDA运行时2105和设备内核驱动程序2108可以分别实施与上面结合图20描述的应用2001、运行时2005和设备内核驱动程序2006相似的功能。在至少一个实施例中，CUDA驱动程序2107包括实现CUDA驱动程序API 2106的库(libcuda.so)。在至少一个实施例中，类似于由CUDA运行时库(cudart)实现的CUDA运行时API 2104，CUDA驱动程序API 2106可以非限制性地暴露用于除其他以外的存储器管理、执行控制、设备管理、错误处理、同步和/或图形互操作性的功能。在至少一个实施例中，CUDA驱动程序API 2106有别于CUDA运行时API 2104，因为CUDA运行时API 2104通过提供隐式初始化、上下文(类似于进程)管理和模块(类似于动态加载库)管理而简化了设备代码管理。在至少一个实施例中，与高层级CUDA运行时API 2104形成对照的是，CUDA驱动程序API2106是特别在上下文和模块加载方面提供设备的更细粒度的控制的低层级API。在至少一个实施例中，CUDA驱动程序API 2106可以暴露CUDA运行时API 2104不暴露的用于上下文管理的功能。在至少一个实施例中，CUDA驱动程序API 2106也是独立于语言的，并且除了CUDA运行时API 2104之外也支持例如OpenCL。进一步，在至少一个实施例中，包括CUDA运行时2105在内的开发库可以被认为与包括用户模式CUDA驱动程序2107和内核模式设备驱动程序2108(有时也称为“显示”驱动程序)的驱动程序部件分离。
在至少一个实施例中，CUDA库2103可以包括但不限于诸如应用2101之类的并行计算应用可能利用的数学库、深度学习库、并行算法库和/或信号/图像/视频处理库。在至少一个实施例中，CUDA库2103可以包括数学库，例如除其他以外的其是一种用于实施线性代数运算的基本线性代数子程序(“BLAS”)的实现方式的cuBLAS库，用于计算快速傅立叶变换(“FFT”)的cuFFT库，以及用于生成随机数的cuRAND库。在至少一个实施例中，CUDA库2103可以包括深度学习库，例如除其他以外的用于深度神经网络的图元cuDNN库以及用于高性能深度学习推理的TensorRT平台。
图22图示出依照至少一个实施例的图20的软件堆栈2000的ROCm实现方式。在至少一个实施例中，其上可以启动应用2201的ROCm软件堆栈2200包括语言运行时2203、系统运行时2205、形实转换程序(thunk)2207、ROCm内核驱动程序2208和设备内核驱动程序2209。在至少一个实施例中，ROCm软件堆栈2200在硬件2210上执行，该硬件可以包括支持ROCm并且由加州圣克拉拉AMD公司开发的GPU。
在至少一个实施例中，应用2201可以实施与上面结合图20所讨论的应用2001相似的功能。此外，在至少一个实施例中，语言运行时2203和系统运行时2205可以实施与上面结合图20所讨论的运行时2005相似的功能。在至少一个实施例中，语言运行时2203和系统运行时2205区别在于，系统运行时2205是一种实现ROCr系统运行时API 2204并且利用异构系统架构(Heterogeneous System Architecture，“HAS”)运行时API的独立于语言的运行时。在至少一个实施例中，HAS运行时API是一种瘦的用户模式API，其暴露访问AMD GPU并且与之交互的接口，包括用于除其他以外的存储器管理、经由内核的架构式分派的执行控制、错误处理、系统和代理信息以及运行时初始化和关闭的功能。在至少一个实施例中，与系统运行时2205形成对照的是，语言运行时2203是分层位于ROCr系统运行时API 2204之上的特定于语言的运行时API 2202的一种实现方式。在至少一个实施例中，语言运行时API可以包括但不限于除其他以外的可移植异构计算接口(Heterogeneous compute Interface forPortability，“HIP”)语言运行时API、异构计算编译器(Heterogeneous ComputeCompiler，“HCC”)语言运行时API或者OpenCL API。特别地，HIP语言是C++编程语言的扩展，具有CUDA机制的功能相似版本，并且在至少一个实施例中，HIP语言运行时API包括与上面结合图21所讨论的CUDA运行时API 2104的功能相似的功能，例如用于除其他以外的存储器管理、执行控制、设备管理、错误处理和同步化的功能。
在至少一个实施例中，形实转换程序(ROCt)2207是一种可以用来与底层ROCm驱动程序2208交互的接口。在至少一个实施例中，ROCm驱动程序2208是一种ROCk驱动程序，其是AMDGPU驱动程序和HAS内核驱动程序(amdkfd)的组合。在至少一个实施例中，AMDGPU驱动程序是一种用于AMD开发的GPU的设备内核驱动程序，其实施与上面结合图20所讨论的设备内核驱动程序2006相似的功能。在至少一个实施例中，HAS内核驱动程序是一种经由硬件特征允许不同类型的处理器更有效地共享系统资源的驱动程序。
在至少一个实施例中，各种不同的库(未示出)可以在语言运行时2203之上包括在ROCm软件堆栈2200中，并且提供与上面结合图21所讨论的CUDA库2103相似的功能。在至少一个实施例中，各种不同的库可以包括但不限于数学、深度学习和/或其他库，例如除其他以外的实现与CUDA cuBLAS的功能相似的功能的hipBLAS库、与CUDA cuFFT相似的用于计算FFT的rocFFT库。
图23图示出依照至少一个实施例的图STACK A的软件堆栈2000的OpenCL实现方式。在至少一个实施例中，其上可以启动应用2301的OpenCL软件堆栈2300包括OpenCL框架2305、OpenCL运行时2306和驱动程序2307。在至少一个实施例中，OpenCL软件堆栈2300在不特定于供应商的硬件2109上执行。在至少一个实施例中，由于OpenCL受不同供应商开发的设备支持，可能需要特定的OpenCL驱动程序与来自这样的供应商的硬件互操作。
在至少一个实施例中，应用2301、OpenCL运行时2306、设备内核驱动程序2307和硬件2308可以分别实施与上面结合图20讨论的应用2001、运行时2005、设备内核驱动程序2006和硬件2007相似的功能。在至少一个实施例中，应用2301进一步包括具有要在设备上执行的代码的OpenCL内核2302。
在至少一个实施例中，OpenCL限定了一种允许主机控制连接到该主机的设备的“平台”。在至少一个实施例中，OpenCL框架提供了被示为平台API 2303和运行时API 2305的平台层API和运行时API。在至少一个实施例中，运行时API 2305使用上下文管理设备上的内核的执行。在至少一个实施例中，每个标识的设备可以与各自的上下文关联，运行时API 2305可以使用该上下文为该设备管理除其他以外的命令队列、程序对象和内核对象，共享存储器对象。在至少一个实施例中，平台API 2303暴露允许使用设备上下文——除其他以外——还选择和初始化设备、经由命令队列向设备提交工作以及允许实现去往/来自设备的数据传输的功能。此外，在至少一个实施例中，OpenCL框架提供了各种各样的内置函数(未示出)，除其他以外，还包括数学函数、关系函数和图像处理函数。
在至少一个实施例中，编译器2304也包括在OpenCL框架2305中。在至少一个实施例中，可以在执行应用之前离线地或者在执行应用期间在线地编译源代码。在至少一个实施例中，与CUDA和ROCm形成对照的是，OpenCL应用可以由编译器2304在线编译，该编译器被包括以代表可以用来将诸如标准可移植中间表示(“SPIR-V”)代码之类的源代码和/或IR代码编译成二进制代码的任意数量的编译器。可替换地，在至少一个实施例中，OpenCL应用可以在执行这样的应用之前离线地编译。
图24图示出依照至少一个实施例的受编程平台支持的软件。在至少一个实施例中，编程平台2404被配置为支持各种各样的编程模型2403、中间件和/或库2402以及应用2400可能依赖的框架2401。在至少一个实施例中，应用2400可以是使用例如诸如MXNet、PyTorch或TensorFlow之类的深度学习框架实现的AI/ML应用，其可能依赖诸如cuDNN、NVIDIA集体通信库(“NCCL”)和/或NVIDIA开发者数据加载库(“DALI”)CUDA库之类的库以便在底层硬件上提供加速的计算。
在至少一个实施例中，编程平台2404可以是上面分别结合图21、图22和图23描述的CUDA、ROCm或者OpenCL平台之一。在至少一个实施例中，编程平台2404支持多个编程模块2403，其是允许表达算法和数据结构的底层计算系统的抽象。在至少一个实施例中，编程模块2403可以暴露底层硬件的特征以便改善性能。在至少一个实施例中，编程模块2403可以包括但不限于CUDA、HIP、OpenCL、C++加速大规模并行(“C++AMP”)、开放多处理(“OpenMP”)、开放加速器(“OpenACC”)和/或火神计算。
在至少一个实施例中，库和/或中间件2402提供编程模块2404的抽象的实现方式。在至少一个实施例中，这样的库包括可以由计算机程序使用以及在软件开发期间加以利用的数据和编程代码。在至少一个实施例中，这样的中间件包括向应用提供可从编程平台2404获得的服务之外的服务的软件。在至少一个实施例中，库和/或中间件2402可以包括但不限于cuBLAS、cuFFT、cuRAND和其他CUDA库，或者rocBLAS、rocFFT、rocRAND和其他ROCm库。此外，在至少一个实施例中，库和/或中间件2402可以包括为GPU提供通信例程的NCCL和ROCm通信集体库(“RCCL”)库，用于深度学习加速的MIOpen库，和/或用于线性代数、矩阵和矢量运算、几何变换、数值求解器和相关算法的特征库。
在至少一个实施例中，应用框架2401依赖于库和/或中间件2402。在至少一个实施例中，应用框架2401中的每一个是用来实现应用软件的标准结构的软件框架。返回到上面讨论的AI/ML示例，在至少一个实施例中，AI/ML应用可以使用诸如Caffe、Caffe2、TensorFlow、Keras、PyTorch或者MxNet深度学习框架之类的框架实现。
图25图示出依照至少一个实施例编译在图20-23的编程平台之一上执行的代码。在至少一个实施例中，编译器2501接收包括主机代码以及设备代码二者的源代码2500。在至少一个实施例中，编译器2501被配置为将源代码2500转换成用于在主机上执行的主机可执行代码2502和用于在设备上执行的设备可执行代码2503。在至少一个实施例中，源代码2500可以或者在执行应用之前离线地编译，或者在执行应用期间在线地编译。
在至少一个实施例中，源代码2500可以包括诸如C++、C、Fortran等等之类的受编译器2501支持的任何编程语言下的代码。在至少一个实施例中，源代码2500可以包括在混合了主机代码和设备代码的单一源文件中，在该单一源文件中指示了设备代码的位置。在至少一个实施例中，单一源文件可以是包括CUDA代码的.cu文件或者包括HIP代码的.hip.cpp文件。可替换地，在至少一个实施例中，源代码2500可以包括多个源代码文件而不是单一源文件，主机代码和设备代码分离到其中。
在至少一个实施例中，编译器2501被配置为将源代码2500编译成用于在主机上执行的主机可执行代码2502以及用于在设备上执行的设备可执行代码2503。在至少一个实施例中，编译器2501实施包括将源代码2500解析成抽象系统树(AST)、实施优化以及生成可执行代码在内的操作。在其中源代码2500包括单一源文件的至少一个实施例中，如下文关于图26更详细地讨论的，编译器2501可以在这样的单一源文件中将设备代码和主机代码分离，分别将设备代码和主机代码编译成设备可执行代码2503和主机可执行代码2502，并且在单个文件中将设备可执行代码2503和主机可执行代码2502链接在一起。
在至少一个实施例中，主机可执行代码2502和设备可执行代码2503可以为任何适当的格式，例如二进制代码和/或IR代码。在至少一个实施例中，在CUDA的情况下，主机可执行代码2502可以包括本机目标码(object code)，并且设备可执行代码2503可以用PTX中间表示包括代码。在至少一个实施例中，在ROCm的情况下，主机可执行代码2502和设备可执行代码2503二者都可以包括目标二进制代码。
图26是依照至少一个实施例编译在图20-23的编程平台之一上执行的代码的更详细的图示。在至少一个实施例中，编译器2601被配置为接收源代码2600，编译源代码2600，并且输出可执行文件2608。在至少一个实施例中，源代码2600为诸如.cu文件、.hip.cpp文件之类的单一源文件，或者另一种格式的文件，其包括主机和设备代码二者。在至少一个实施例中，编译器2601可以是，但不限于，用于编译.cu文件中的CUDA代码的英伟达CUDA编译器(“NVCC”)，或者用于编译.hip.cpp文件中的HIP代码的HCC编译器。
在至少一个实施例中，编译器2601包括编译器前端2602、主机编译器2605、设备编译器2606和链接器2609。在至少一个实施例中，编译器前端2602被配置为在源代码2600中将设备代码2604与主机代码2603分离。在至少一个实施例中，设备代码2604由设备编译器2606编译成设备可执行代码2608，其如所描述的可以包括二进制代码或者IR代码。另外，在至少一个实施例中，主机代码2603由主机编译器2605编译成主机可执行代码2607。在至少一个实施例中，对于NVCC而言，主机编译器2605可以是，但不限于，输出本机目标码的通用C/C++编译器，而设备编译器2606可以是，但不限于，派生LLVM编译器基础结构并且输出PTX代码或二进制代码的基于低层级虚拟机(“LLVM”)的编译器。在至少一个实施例中，对于HCC而言，主机编译器2605和设备编译器2606二者都可以是，但不限于，输出目标二进制代码的基于LLVM的编译器。
在至少一个实施例中，在将源代码2600编译成主机可执行代码2607和设备可执行代码2608之后，链接器2609在可执行文件2610中将主机和设备可执行代码2607和2608链接在一起。在至少一个实施例中，用于主机的本机目标码以及用于设备的PTX或二进制代码可以在可执行和可链接格式(“ELF”)文件中链接在一起，所述文件是一种用来存储目标码的容器格式。
图27图示出依照至少一个实施例在编译源代码之前转化源代码。在至少一个实施例中，源代码2700通过转化工具2701，该转化工具将源代码2700转化成转化的源代码2702。在至少一个实施例中，编译器2703用来在一定过程中将转化的源代码2702编译成主机可执行代码2704和设备可执行代码2705，该过程类似于上面结合图25所讨论的编译器2501将源代码2500编译成主机可执行代码2502和设备可执行代码2503。
在至少一个实施例中，转化工具2701实施的转化用来移植源2700以便在与它最初预期运行的环境不同的环境中执行。在至少一个实施例中，转化工具2701可以包括但不限于用来将预期用于CUDA平台的CUDA代码“HIP化(hipify)”为可以在ROCm平台上编译和执行的HIP代码的HIP转化器。在至少一个实施例中，如下文结合图28A-29更详细地讨论的，源代码2700的转化可以包括解析源代码2700并且将对一个编程模型(例如CUDA)提供的API的调用转换成对另一个编程模型(例如HIP)提供的API的调用。返回到HIP化CUDA的示例，在至少一个实施例中，对CUDA运行时API、CUDA驱动程序API和/或CUDA库的调用可以被转换为相应的HIP API调用。在至少一个实施例中，转化工具2701实施的自动转化有时可能是不完整的，需要附加的人工努力以完全移植源代码2700。
配置GPU以用于通用计算
下面的图非限制性地阐述了依照至少一个实施例的用于编译和执行计算源代码的示例性架构。
图28A图示出依照至少一个实施例的被配置为使用不同类型的处理单元编译和执行CUDA源代码2810的系统28A00。在至少一个实施例中，系统28A00包括但不限于CUDA源代码2810、CUDA编译器2850、主机可执行代码2870(1)、主机可执行代码2870(2)、CUDA设备可执行代码2884、CPU 2890、CUDA启用的GPU 2894、GPU 2892、CUDA-HIP转化工具2820、HIP源代码2830、HIP编译器驱动程序2840、HCC 2860以及HCC设备可执行代码2882。
在至少一个实施例中，CUDA源代码2810是CUDA编程语言下的人类可读代码集合。在至少一个实施例中，CUDA代码是CUDA编程语言下的人类可读代码。在至少一个实施例中，CUDA编程语言是C++编程语言的扩展，其包括但不限于限定设备代码以及区分设备代码和主机代码的机制。在至少一个实施例中，设备代码是一种源代码，其在编译之后可在设备上并行地执行。在至少一个实施例中，设备可以是针对并行指令处理优化的处理器，例如CUDA启用的GPU 2890、GPU 28192或者另一个GPGPU等等。在至少一个实施例中，主机代码是源代码，其在编译后可在主机上执行。在至少一个实施例中，主机是一种针对顺序指令处理优化的处理器，例如CPU 2890。
在至少一个实施例中，CUDA源代码2810包括但不限于任意数量(包括零个)的全局功能2812、任意数量(包括零个)的设备功能2814、任意数量(包括零个)的主机功能2816以及任意数量(包括零个)的主机/设备功能2818。在至少一个实施例中，全局功能2812、设备功能2814、主机功能2816和主机/设备功能2818可以在CUDA源代码2810中混合。在至少一个实施例中，全局功能2812中的每一个可在设备上执行并且可从主机调用。在至少一个实施例中，全局功能2812中的一个或更多个因此可以充当设备的入口点。在至少一个实施例中，全局功能2812中的每一个为内核。在至少一个实施例中并且在称为动态并行的技术中，全局功能2812中的一个或更多个限定可在设备上执行并且可从这样的设备调用的内核。在至少一个实施例中，内核在执行期间由设备上的N个不同线程并行地执行N次(其中N为任何正整数)。
在至少一个实施例中，设备功能2814中的每一个在设备上执行并且仅仅可从这样的设备调用。在至少一个实施例中，主机功能2816中的每一个在主机上执行并且仅仅可从这样的主机调用。在至少一个实施例中，主机/设备功能2816中的每一个限定可在主机上执行并且仅仅可从这样的主机调用的功能的主机版本以及可在设备上执行并且仅仅可从这样的设备调用的所述功能的设备版本二者。
在至少一个实施例中，CUDA源代码2810也可以包括但不限于对经由CUDA运行时API 2802限定的任意数量的功能的任意数量的调用。在至少一个实施例中，CUDA运行时API2802可以包括但不限于在主机上执行以分配和解除分配设备存储器、在主机存储器与设备存储器之间传输数据、管理具有多个设备的系统等等的任意数量的功能。在至少一个实施例中，CUDA源代码2810也可以包括对在任意数量的其他CUDA API中指定的任意数量的功能的任意数量的调用。在至少一个实施例中，CUDA API可以是被设计以供CUDA代码使用的任何API。在至少一个实施例中，CUDA API包括但不限于CUDA运行时API 2802、CUDA驱动程序API、用于任意数量的CUDA库的API等等。在至少一个实施例中并且相对于CUDA运行时API2802，CUDA驱动程序API是一种较低层级API，但是提供设备的较细粒度的控制。在至少一个实施例中，CUDA库的示例包括但不限于cuBLAS、cuFFT、cuRAND、cuDNN等等。
在至少一个实施例中，CUDA编译器2850编译输入CUDA代码(例如CUDA源代码2810)以生成主机可执行代码2870(1)和CUDA设备可执行代码2884。在至少一个实施例中，CUDA编译器2850为NVCC。在至少一个实施例中，主机可执行代码2870(1)是可在CPU 2890上执行的、包括在输入源代码中的主机代码的编译版本。在至少一个实施例中，CPU 2890可以是针对顺序指令处理优化的任何处理器。
在至少一个实施例中，CUDA设备可执行代码2884是可在CUDA启用的GPU 2894上执行的、包括在输入源代码中的设备代码的编译版本。在至少一个实施例中，CUDA设备可执行代码2884包括但不限于二进制代码。在至少一个实施例中，CUDA设备可执行代码2884包括但不限于诸如PTX代码之类的IR代码，其在运行时通过设备驱动程序进一步编译成用于特定目标设备(例如CUDA启用的GPU 2894)的二进制代码。在至少一个实施例中，CUDA启用的GPU 2894可以是针对并行指令处理优化的并且支持CUDA的任何处理器。在至少一个实施例中，CUDA启用的GPU 2894由加州圣克拉拉英伟达公司开发。
在至少一个实施例中，CUDA-HIP转化工具2820被配置为将CUDA源代码2810转化为功能上相似的HIP源代码2830。在至少一个实施例中，HIP源代码2830是HIP编程语言下的人类可读代码集合。在至少一个实施例中，HIP代码是HIP编程语言下的人类可读代码。在至少一个实施例中，HIP编程语言是C++编程语言的扩展，其包括但不限于限定设备代码以及区分设备代码和主机代码的CUDA机制的功能上相似的版本。在至少一个实施例中，HIP编程语言可以包括CUDA编程语言的功能子集。在至少一个实施例中，例如，HIP编程语言包括但不限于限定全局功能2812的机制，但是这样的HIP编程语言可能缺乏对于动态并行的支持并且因此在HIP代码中限定的全局功能2812可能仅仅可从主机调用。
在至少一个实施例中，HIP源代码2830包括但不限于任意数量(包括零个)的全局功能2812、任意数量(包括零个)的设备功能2814、任意数量(包括零个)的主机功能2816以及任意数量(包括零个)的主机/设备功能2818。在至少一个实施例中，HIP源代码2830也可以包括对于HIP运行时API 2832中指定的任意数量的功能的任意数量的调用。在至少一个实施例中，HIP运行时API 2832包括但不限于包括在CUDA运行时API 2802中的功能子集的功能上相似的版本。在至少一个实施例中，HIP源代码2830也可以包括对在任意数量的其他HIP API中指定的任意数量的功能的任意数量的调用。在至少一个实施例中，HIP API可以是被设计以供HIP代码和/或ROCm使用的任何API。在至少一个实施例中，HIP API包括但不限于HIP运行时API 2832、HIP驱动程序API、用于任意数量的HIP库的API、用于任意数量的ROCm库的API等等。
在至少一个实施例中，CUDA-HIP转化工具2820将CUDA代码中的每个内核调用从CUDA句法转换成HIP句法，并且将CUDA代码中的任意数量的其他CUDA调用转换成任意数量的其他功能上相似的HIP调用。在至少一个实施例中，CUDA调用是对CUDA API中指定的功能的调用，并且HIP调用是对HIP API中指定的功能的调用。在至少一个实施例中，CUDA-HIP转化工具2820将对于CUDA运行时API 2802中指定的功能的任意数量的调用转换成对于HIP运行时API 2832中指定的功能的任意数量的调用。
在至少一个实施例中，CUDA-HIP转化工具2820是一种执行基于文本的转化过程的称为hipify-perl的工具。在至少一个实施例中，CUDA-HIP转化工具2820是一种相对于hipify-perl执行更加复杂和更加鲁棒的转化过程的称为hipify-clang的工具，该过程涉及使用clang(一种编译器前端)解析CUDA代码并且然后转化得到的符号。在至少一个实施例中，适当地将CUDA代码转换成HIP代码可能需要除了CUDA-HIP转化工具2820实施的修改以外的修改(例如人工编辑)。
在至少一个实施例中，HIP编译器驱动程序2840是一种前端，其确定目标设备2846并且然后配置与目标设备2846兼容的编译器以编译HIP源代码2830。在至少一个实施例中，目标设备2846是一种针对并行指令处理优化的处理器。在至少一个实施例中，HIP编译器驱动程序2840可以以任何技术上可行的方式确定目标设备2846。
在至少一个实施例中，如果目标设备2846与CUDA兼容(例如CUDA启用的GPU2894)，那么HIP编译器驱动程序2840生成HIP/NVCC编译命令2842。在至少一个实施例中并且如结合图28B更详细地描述的，HIP/NVCC编译命令2842将CUDA编译器2850配置为非限制性地使用HIP-CUDA转化标头和CUDA运行时库编译HIP源代码2830。在至少一个实施例中并且响应于HIP/NVCC编译命令2842，CUDA编译器2850生成主机可执行代码2870(1)和CUDA设备可执行代码2884。
在至少一个实施例中，如果目标设备2846与CUDA不兼容，那么HIP编译器驱动程序2840生成HIP/HCC编译命令2844。在至少一个实施例中并且如结合图28C更详细地描述的，HIP/HCC编译命令2844将HCC 2860配置为非限制性地使用HCC标头和HIP/HCC运行时库编译HIP源代码2830。在至少一个实施例中并且响应于HIP/HCC编译命令2844，HCC 2860生成主机可执行代码2870(2)和HCC设备可执行代码2882。在至少一个实施例中，HCC设备可执行代码2882是可在GPU 2892上执行的、包括在HIP源代码2830中的设备代码的编译版本。在至少一个实施例中，GPU 2892可以是针对并行指令处理优化的任何处理器，与CUDA不兼容，并且与HCC兼容。在至少一个实施例中，GPU 2892由加州圣克拉拉AMD公司开发。在至少一个实施例中，GPU 2892是非CUDA启用的GPU 2892。
仅仅出于解释的目的，图28A中描绘了可以在至少一个实施例中实现以编译在CPU2890和不同设备上执行的CUDA源代码2810的三个不同的流。在至少一个实施例中，直接CUDA流在不将CUDA源代码2810转化成HIP源代码2830的情况下编译CUDA源代码2810以便在CPU 2890和CUDA启用的GPU 2894上执行。在至少一个实施例中，间接CUDA流将CUDA源代码2810转化成HIP源代码2830，并且然后编译HIP源代码2830以便在CPU 2890和CUDA启用的GPU 2894上执行。在至少一个实施例中，CUDA/HCC流将CUDA源代码2810转化成HIP源代码2830，并且然后编译HIP源代码2830以便在CPU 2890和GPU 2892上执行。
可以在至少一个实施例中实现的直接CUDA流经由虚线和一系列标注为A1-A3的气泡描绘。在至少一个实施例中并且如利用标注为A1的气泡所描绘的，CUDA编译器2850接收CUDA源代码2810以及将CUDA编译器2850配置为编译CUDA源代码2810的CUDA编译命令2848。在至少一个实施例中，直接CUDA流中使用的CUDA源代码2810用基于不同于C++的编程语言(例如C、Fortran、Python、Java等等)的CUDA编程语言编写。在至少一个实施例中并且响应于CUDA编译命令2848，CUDA编译器2850生成主机可执行代码2870(1)和CUDA设备可执行代码2884(用标注为A2的气泡描绘)。在至少一个实施例中并且如利用标注为A3的气泡所描绘的，主机可执行代码2870(1)和CUDA设备可执行代码2884可以分别在CPU 2890和CUDA启用的GPU 2894上执行。在至少一个实施例中，CUDA设备可执行代码2884可以包括但不限于二进制代码。在至少一个实施例中，CUDA设备可执行代码2884包括但不限于PTX代码，并且在运行时进一步编译成用于特定目标设备的二进制代码。
可以在至少一个实施例中实现的间接CUDA流经由点线和一系列标注为B1-B6的气泡描绘。在至少一个实施例中并且如利用标注为B1的气泡所描绘的，CUDA-HIP转化工具2820接收CUDA源代码2810。在至少一个实施例中并且如利用标注为B2的气泡所描绘的，CUDA-HIP转化工具2820将CUDA源代码2810转化为HIP源代码2830。在至少一个实施例中并且如利用标注为B3的气泡所描绘的，HIP编译器驱动程序2840接收HIP源代码2830并且确定目标设备2846是CUDA启用的。
在至少一个实施例中并且如利用标注为B4的气泡所描绘的，HIP编译器驱动程序2840生成HIP/NVCC编译命令2842并且将HIP/NVCC编译命令2842和HIP源代码2830二者传送至CUDA编译器2850。在至少一个实施例中并且如结合图28B更详细地描述的，HIP/NVCC编译命令2842将CUDA编译器2850配置为非限制性地使用HIP-CUDA转化标头和CUDA运行时库编译HIP源代码2830。在至少一个实施例中并且响应于HIP/NVCC编译命令2842，CUDA编译器2850生成主机可执行代码2870(1)和CUDA设备可执行代码2884(用标注为B5的气泡描绘)。在至少一个实施例中并且如利用标注为B6的气泡所描绘的，主机可执行代码2870(1)和CUDA设备可执行代码2884可以分别在CPU 2890和CUDA启用的GPU 2894上执行。在至少一个实施例中，CUDA设备可执行代码2884包括但不限于二进制代码。在至少一个实施例中，CUDA设备可执行代码2884包括但不限于PTX代码，并且在运行时进一步编译成用于特定目标设备的二进制代码。
可以在至少一个实施例中实现的CUDA/HCC流经由实线和一系列标注为C1-C6的气泡描绘。在至少一个实施例中并且如利用标注为C1的气泡所描绘的，CUDA-HIP转化工具2820接收CUDA源代码2810。在至少一个实施例中并且如利用标注为C2的气泡所描绘的，CUDA-HIP转化工具2820将CUDA源代码2810转化为HIP源代码2830。在至少一个实施例中并且如利用标注为C3的气泡所描绘的，HIP编译器驱动程序2840接收HIP源代码2830并且确定目标设备2846不是CUDA启用的。
在至少一个实施例中，HIP编译器驱动程序2840生成HIP/HCC编译命令2844并且将HIP/HCC编译命令2844和HIP源代码2830二者传送至HCC 2860(用标注为C4的气泡描绘)。在至少一个实施例中并且如结合图28C更详细地描述的，HIP/HCC编译命令2844将HCC 2860配置为非限制性地使用HCC标头和HIP/HCC运行时库编译HIP源代码2830。在至少一个实施例中并且响应于HIP/HCC编译命令2844，HCC 2860生成主机可执行代码2870(2)和HCC设备可执行代码2882(用标注为C5的气泡描绘)。在至少一个实施例中并且如利用标注为C6的气泡所描绘的，主机可执行代码2870(2)和HCC设备可执行代码2882可以分别在CPU 2890和GPU2892上执行。
在至少一个实施例中，在将CUDA源代码2810转化为HIP源代码2830之后，HIP编译器驱动程序2840可以随后用来在不重新执行CUDA-HIP转化工具2820的情况下生成用于CUDA启用的GPU 2894或者GPU 2892的可执行代码。在至少一个实施例中，CUDA-HIP转化工具2820将CUDA源代码2810转化为HIP源代码2830，该HIP源代码然后存储在存储器中。在至少一个实施例中，HIP编译器驱动程序2840然后将HCC 2860配置为基于HIP源代码2830生成主机可执行代码2870(2)和HCC设备可执行代码2882。在至少一个实施例中，HIP编译器驱动程序2840随后将CUDA编译器2850配置为基于存储的HIP源代码2830生成主机可执行代码2870(1)和CUDA设备可执行代码2884。
图28B图示出依照至少一个实施例的被配置为使用CPU 2890和CUDA启用的GPU2894编译和执行图28A的CUDA源代码2810的系统2804。在至少一个实施例中，系统2804包括但不限于CUDA源代码2810、CUDA-HIP转化工具2820、HIP源代码2830、HIP编译器驱动程序2840、CUDA编译器2850、主机可执行代码2870(1)、CUDA设备可执行代码2884、CPU 2890以及CUDA启用的GPU 2894。
在至少一个实施例中并且如本文先前结合图28A所描述的，CUDA源代码2810包括但不限于任意数量(包括零个)的全局功能2812、任意数量(包括零个)的设备功能2814、任意数量(包括零个)的主机功能2816以及任意数量(包括零个)的主机/设备功能2818。在至少一个实施例中，CUDA源代码2810也包括但不限于对任意数量的CUDA API中指定的任意数量的功能的任意数量的调用。
在至少一个实施例中，CUDA-HIP转化工具2820将CUDA源代码2810转化为HIP源代码2830。在至少一个实施例中，CUDA-HIP转化工具2820将CUDA源代码2810中的每个内核调用从CUDA句法转换成HIP句法，并且将CUDA源代码2810中的任意数量的其他CUDA调用转换成任意数量的其他功能上相似的HIP调用。
在至少一个实施例中，HIP编译器驱动程序2840确定目标设备2846是CUDA启用的并且生成HIP/NVCC编译命令2842。在至少一个实施例中，HIP编译器驱动程序2840然后经由HIP/NVCC编译命令2842将CUDA编译器2850配置为编译HIP源代码2830。在至少一个实施例中，作为配置CUDA编译器2850的部分，HIP编译器驱动程序2840提供对于HIP-CUDA转化标头2852的访问。在至少一个实施例中，HIP-CUDA转化标头2852将任意数量的HIP API中指定的任意数量的机制(例如功能)转化为任意数量的CUDA API中指定的任意数量的机制。在至少一个实施例中，CUDA编译器2850使用HIP-CUDA转化标头2852结合与CUDA运行时API 2802相应的CUDA运行时库2854生成主机可执行代码2870(1)和CUDA设备可执行代码2884。在至少一个实施例中，主机可执行代码2870(1)和CUDA设备可执行代码2884然后可以分别在CPU2890和CUDA启用的GPU 2894上执行。在至少一个实施例中，CUDA设备可执行代码2884包括但不限于二进制代码。在至少一个实施例中，CUDA设备可执行代码2884包括但不限于PTX代码，并且在运行时进一步编译成用于特定目标设备的二进制代码。
图28C图示出依照至少一个实施例的被配置为使用CPU 2890和非CUDA启用的GPU2892编译和执行图28A的CUDA源代码2810的系统2806。在至少一个实施例中，系统2806包括但不限于CUDA源代码2810、CUDA-HIP转化工具2820、HIP源代码2830、HIP编译器驱动程序2840、HCC 2860、主机可执行代码2870(2)、HCC设备可执行代码2882、CPU 2890以及GPU2892。
在至少一个实施例中并且如本文先前结合图28A所描述的，CUDA源代码2810包括但不限于任意数量(包括零个)的全局功能2812、任意数量(包括零个)的设备功能2814、任意数量(包括零个)的主机功能2816以及任意数量(包括零个)的主机/设备功能2818。在至少一个实施例中，CUDA源代码2810也包括但不限于对任意数量的CUDA API中指定的任意数量的功能的任意数量的调用。
在至少一个实施例中，CUDA-HIP转化工具2820将CUDA源代码2810转化为HIP源代码2830。在至少一个实施例中，CUDA-HIP转化工具2820将CUDA源代码2810中的每个内核调用从CUDA句法转换成HIP句法，并且将源代码2810中的任意数量的其他CUDA调用转换成任意数量的其他功能上相似的HIP调用。
在至少一个实施例中，HIP编译器驱动程序2840随后确定目标设备2846不是CUDA启用的并且生成HIP/HCC编译命令2844。在至少一个实施例中，HIP编译器驱动程序2840然后将HCC 2860配置为执行HIP/HCC编译命令2844以便编译HIP源代码2830。在至少一个实施例中，HIP/HCC编译命令2844将HCC 2860配置为非限制性地使用HIP/HCC运行时库2858和HCC标头2856生成主机可执行代码2870(2)和HCC设备可执行代码2882。在至少一个实施例中，HIP/HCC运行时库2858与HIP运行时API 2832相应。在至少一个实施例中，HCC标头2856包括但不限于用于HIP和HCC的任意数量和类型的互操作机制。在至少一个实施例中，主机可执行代码2870(2)和HCC设备可执行代码2882可以分别在CPU 2890和GPU 2892上执行。
图29图示出依照至少一个实施例的通过图28C的CUDA-HIP转化工具2820转化的示例性内核。在至少一个实施例中，CUDA源代码2810将给定内核被设计来解决的总体问题分割成可以使用线程块独立解决的相对粗糙的子问题。在至少一个实施例中，每个线程块包括但不限于任意数量的线程。在至少一个实施例中，每个子问题被分割成可以通过线程块内的线程并行合作地解决的相对精细的碎片。在至少一个实施例中，线程块内的线程可以通过经由共享的存储器共享数据并且通过使执行同步以协调存储器访问来合作。
在至少一个实施例中，CUDA源代码2810将与给定内核关联的线程块组织成一维、二维或者三维线程块网格。在至少一个实施例中，每个线程块包括但不限于任意数量的线程，并且网格包括但不限于任意数量的线程块。
在至少一个实施例中，内核是设备代码中的功能，其使用“_global_”声明说明符进行限定。在至少一个实施例中，针对给定内核调用和关联的流执行内核的网格的维度使用CUDA内核启动句法2910指定。在至少一个实施例中，CUDA内核启动句法2910被指定为“KernelName<<<GridSize,BlockSize,SharedMemorySize,Stream>>>(KernelArguments)；”。在至少一个实施例中，执行配置句法为插入到内核名称(“KernelName”)与带括号的内核参量(argument)列表(“KernelArguments”)之间的“<<<...>>>”结构体。在至少一个实施例中，CUDA内核启动句法2910包括但不限于CUDA启动功能句法而不是执行配置句法。
在至少一个实施例中，“GridSize”为dim3类型并且指定了网格的维度和大小。在至少一个实施例中，dim3类型是一种CUDA限定的结构，其包括但不限于无符号整数x、y和z。在至少一个实施例中，如果z未被指定，那么z缺省为1。在至少一个实施例中，如果y未被指定，那么y缺省为1。在至少一个实施例中，网格中线程块的数量等于GridSize.x、GridSize.y和GridSize.z的乘积。在至少一个实施例中，“BlockSize”为dim3类型并且指定了每个线程块的维度和大小。在至少一个实施例中，每线程块的线程的数量等于BlockSize.x、BlockSize.y和BlockSize.z的乘积。在至少一个实施例中，给予执行内核的每个线程唯一的线程ID，该线程ID在该内核内可通过内置变量(例如“threadIdx”)访问。
在至少一个实施例中并且关于CUDA内核启动句法2910，“SharedMemorySize”是一个可选的参量，其指定对于给定内核调用而言除了静态分配的存储器之外每线程块动态分配的共享的存储器中的字节数量。在至少一个实施例中并且关于CUDA内核启动句法2910，SharedMemorySize缺省为零。在至少一个实施例中并且关于CUDA内核启动句法2910，“Stream”是一个可选的参量，其指定关联流并且缺省为零以指定缺省流。在至少一个实施例中，流是一种按顺序执行的命令序列(可能由不同主机线程发出)。在至少一个实施例中，不同的流可以相对于彼此无序地或者并发地执行命令。
在至少一个实施例中，CUDA源代码2810包括但不限于用于示例性内核“MatAdd”的内核定义以及主功能。在至少一个实施例中，主功能为主机代码，其在主机上执行并且包括但不限于使得内核MatAdd在设备上执行的内核调用。在至少一个实施例中并且如图所示，内核MatAdd将大小为NxN的两个矩阵A和B相加，其中N为正整数，并且将结果存储在剧中C中。在至少一个实施例，主功能将threadsPerBlock变量限定为16x16，将numBlocks变量限定为N/16x N/16。在至少一个实施例，主功能然后指定内核调用“MatAdd<<<numBlocks,threadsPerBlock>>>(A,B,C)；”。在至少一个实施例并且根据CUDA内核启动句法2910，内核MatAdd使用具有维度N/16x N/16的线程块网格执行，其中每个线程块具有16x16的维度。在至少一个实施例，每个线程块包括256个线程，网格利用足够的块创建以便每矩阵元素具有一个线程，并且这样的网格中的每个线程执行内核MatAdd以实施一次成对加法。
在至少一个实施例，在将CUDA源代码2810转化为HIP源代码2830时，CUDA-HIP转化工具2820将CUDA源代码2810中的每个内核调用从CUDA内核启动句法2910转化为HIP内核启动句法2920，并且将源代码2810中的任意数量的其他CUDA调用转换成任意数量的其他功能上相似的HIP调用。在至少一个实施例，HIP内核启动句法2920被指定为“hipLaunchKernelGGL(KernelName,GridSize,BlockSize,SharedMemorySize,Stream,KernelArguments)；”。在至少一个实施例，KernelName、GridSize、BlockSize、ShareMemorySize、Stream和KernelArguments中的每一个在HIP内核启动句法2920中具有与在CUDA内核启动句法2910(如本文先前所描述的)中相同的含义。在至少一个实施例，参量SharedMemorySize和Stream在HIP内核启动句法2920中是所需的，并且在CUDA内核启动句法2910中是可选的。
在至少一个实施例，除了使得内核MatAdd在设备上执行的内核调用之外，图29中描绘的HIP源代码2830的部分与图29中描绘的CUDA源代码2810的部分相同。在至少一个实施例，内核MatAdd在HIP源代码2830中用相同的“_global_”声明说明符限定，内核MatAdd在CUDA源代码2810中用该声明说明符进行限定。在至少一个实施例，HIP源代码2830中的内核调用是“hipLaunchKernelGGL(MatAdd,numBlocks,threadsPerBlock,0,0,A,B,C)；”，而CUDA源代码2810中的相应内核调用是“MatAdd<<<numBlocks,threadsPerBlock>>>(A,B,C)；”。
图30更详细地图示出依照至少一个实施例的图28C的非CUDA启用的GPU 2892。在至少一个实施例，GPU 2892由圣克拉拉AMD公司开发。在至少一个实施例，GPU 2892可以被配置为以高度并行的方式实施计算操作。在至少一个实施例，GPU 2892被配置为执行诸如绘制命令、像素操作、几何计算之类的图形管线操作以及与将图像再现到显示器关联的其他操作。在至少一个实施例，GPU 2892被配置为执行与图形无关的操作。在至少一个实施例，GPU 2892被配置为执行与图形有关的操作和与图形无关的操作二者。在至少一个实施例，GPU 2892可以被配置为执行包括在HIP源代码2830中的设备代码。
在至少一个实施例，GPU 2892包括但不限于任意数量的可编程处理单元3020、命令处理器3010、L2缓存3022、存储器控制器3070、DMA引擎3080(1)、系统存储器控制器3082、DMA引擎3080(2)以及GPU控制器3084。在至少一个实施例，每个可编程处理单元3020包括但不限于工作负载管理器3030和任意数量的计算单元3040。在至少一个实施例，命令处理器3010读取来自一个或更多个命令队列(未示出)的命令，并且将命令分发至工作负载管理器3030。在至少一个实施例，对于每个可编程处理单元3020，关联的工作负载管理器3030将工作分发给可编程处理单元3020中包括的计算单元3040。在至少一个实施例，每个计算单元3040可以执行任意数量的线程块，但是每个线程块在单个计算单元3040上执行。在至少一个实施例，工作组是一种线程块。
在至少一个实施例，每个计算单元3040包括但不限于任意数量的SIMD单元3050和共享的存储器3060。在至少一个实施例，每个SIMD单元3050实现一种SIMD架构并且被配置为并行地实施操作。在至少一个实施例，每个SIMD单元3050包括但不限于矢量ALU 3052和矢量寄存器文件3054。在至少一个实施例，每个SIMD单元3050执行不同的线程束。在至少一个实施例，线程束是一组线程(例如16个线程)，其中线程束中的每个线程属于单个线程块并且被配置为基于单个指令集处理不同的数据集合。在至少一个实施例，谓词可以用来禁用线程束中的一个或更多个线程。在至少一个实施例中，巷道为一种线程。在至少一个实施例中，工作项目为一种线程。在至少一个实施例中，波前为一种线程束。在至少一个实施例中，线程块中的不同波前可以一起同步，并且经由共享的存储器3060通信。
在至少一个实施例中，可编程处理单元3020称为“着色器引擎”。在至少一个实施例中，每个可编程处理单元3020包括但不限于除了计算单元3040以外的任意数量的专用图形硬件。在至少一个实施例中，每个可编程处理单元3020包括但不限于任意数量(包括零个)的几何处理器、任意数量(包括零个)的光栅化器、任意数量(包括零个)的再现后端、工作负载管理器3030以及任意数量的计算单元3040。
在至少一个实施例中，计算单元3040共享L2缓存3022。在至少一个实施例中，L2缓存3022被分区。在至少一个实施例中，GPU存储器3090可由GPU 2892中的所有计算单元3040访问。在至少一个实施例中，存储器控制器3070和系统存储器控制器3082促进GPU 2892与主机之间的数据传输，并且DMA引擎3080(1)允许实现GPU 2892与这样的主机之间的异步存储器传输。在至少一个实施例中，存储器控制器3070和GPU控制器3084促进GPU 2892与其他GPU 2892之间的数据传输，并且DMA引擎3080(2)允许实现GPU 2892与其他GPU 2892之间的异步存储器传输。
在至少一个实施例中，GPU 2892包括但不限于促进跨任意数量和类型的、可以在GPU 2892内部或外部的直接或间接链接的部件的数据和控制传送的任意数量和类型的系统互连线。在至少一个实施例中，GPU 2892包括但不限于耦合到任意数量和类型的外围设备的任意数量和类型的I/O接口(例如PCIe)。在至少一个实施例中，GPU 2892可以包括但不限于任意数量(包括零个)的显示引擎和任意数量(包括零个)的多媒体引擎。在至少一个实施例中，GPU 2892实现一种存储器子系统，其包括但不限于可以专用于一个部件或者在多个部件之间共享的任意数量和类型的存储器控制器(例如存储器控制器3070和系统存储器控制器3082)和存储器设备(例如共享的存储器3060)。在至少一个实施例中，GPU 2892实现一种缓存子系统，其包括但不限于一个或更多个缓存存储器(例如L2缓存3022)，所述缓存存储器中的每一个可以专用于任意数量的部件(例如SIMD单元3050、计算单元3040和可编程处理单元3020)或者在这些部件之间共享。
图31图示出依照至少一个实施例如何将示例性CUDA网格3120的线程映射为图30的不同计算单元3040。在至少一个实施例中并且仅仅出于解释的目的，网格3120具有BX xBY x 1的GridSize以及TX x TY x 1的BlockSize。在至少一个实施例中，网格3120因此包括但不限于(BX*BY)个线程块3130，并且每个线程块3130包括但不限于(TX*TY)个线程3140。线程3140在图31中描绘为弯弯曲曲的箭头。
在至少一个实施例中，将网格3120映射到包括但不限于计算单元3040(1)-3040(C)的可编程处理单元3020(1)。在至少一个实施例中并且如图所示，将(BJ*BY)个线程块3130映射到计算单元3040(1)，并且将其余线程块3130映射到计算单元3040(2)。在至少一个实施例中，每个线程块3130可以包括但不限于任意数量的线程束，并且每个线程束被映射到图30的不同SIMD单元3050。
在至少一个实施例中，给定线程块3130中的线程束可以一起同步，并且通过关联的计算单元3040中包括的共享的存储器3060通信。例如并且在至少一个实施例中，线程块3130(BJ,1)中的线程束可以一起同步并且通过共享的存储器3060(1)通信。例如并且在至少一个实施例中，线程块3130(BJ+1,1)中的线程束可以一起同步并且通过共享的存储器3060(2)通信。
任务图和可执行图
在至少一个实施例中，专用计算资源和/或并行计算资源阵列的使用提供了更有效且更快速地执行复杂工作负载的机会。在至少一个实施例中，许多图像处理、机器学习和相关工作负载的执行通常可以通过使用并行操作的一个或更多个专用计算资源和/或两个或更多计算资源而加速。在至少一个实施例中，诸如主机中央处理单元(“CPU”)之类的主机处理器负责准备一个或更多个任务以及将其分发至一个或更多个其他计算资源，并且然后从一个或更多个其他计算资源取回一个或更多个任务的结果。在至少一个实施例中，只要一个或更多个其他计算资源能够比主机CPU可能直接实施一个或更多个任务更快地完成一个或更多个任务，那么使用一个或更多个其他计算资源实施一个或更多个任务通常被认为是有利的。
在至少一个实施例中，将一个或更多个任务卸载到一个或更多个其他计算资源可能是开销密集的。在至少一个实施例中，除其他可能性以外，开销可以包括选择一个或更多个其他计算资源、将可执行指令配置和/或传送到一个或更多个其他计算资源、将数据传送到一个或更多个其他计算资源以及从一个或更多个其他计算资源取回结果。在至少一个实施例中，如果不仔细管理该开销，那么开销的成本可能显著地降低使用一个或更多个其他计算资源实施计算任务的效率。
在至少一个实施例中，一种用于降低开销的方法是将多个任务看作一个集成整体，而不是单独地考虑多个任务中的每一个。在至少一个实施例中，在工作负载中的各种不同的任务之间存在从属性，其中一些任务生成其他任务所需的数据，一些任务需要在其他任务可能开始之前完成，一些任务需要与其他任务共享独占使用资源，和/或一些任务需要与其他任务同步执行。在至少一个实施例中，一些计算资源比其他计算资源更适合某些任务。在至少一个实施例中，相对于孤立地管理多个任务中的每一个，考虑任务之间的从属性和/或考虑计算资源的哪些组合在计算上最高效的技术通常能够降低与跨多个计算资源执行多个任务关联的开销和计算时间。
在至少一个实施例中，可以通过使用任务图将多个任务组织为集成整体。在至少一个实施例中，在任务图中，将包括多个任务的工作负载组织为有向图，其中每个节点与要实施的任务相应，并且两个节点之间的每个有向边与两个节点之间的数据从属性、执行从属性或者某种其他的从属性相应。在至少一个实施例中，从属性可以指示何时一个节点的任务必须在另一个节点的任务可能开始之前完成。在至少一个实施例中，从属性可以指示何时一个节点必须在该节点可能开始和/或继续其任务之前等待来自另一个节点的数据。在至少一个实施例中，一旦准备好任务图，就将任务图转换成任务图的可执行版本(“可执行图”)。在至少一个实施例中，可以使用实例化从任务图生成可执行图。在至少一个实施例中，由于将任务图转换成可执行图可访问整个任务图，因而可以实施各种不同的优化，这可以降低工作负载的总体执行时间。
在至少一个实施例中，可执行图可以多次使用以便在不必从任务图重新生成的情况下让计算资源实施相同的工作负载。在至少一个实施例中，使用可执行图实施工作负载的一个缺点是，可执行图仅限于从其生成可执行图的任务图的相同工作负载。
修改可执行图以实施新任务图的工作负载
图32图示出依照至少一个实施例的流程图。尽管结合图1-31的系统和技术描述了方法3200的方法步骤3210-3260，但是应当理解的是，被配置为实施步骤3210-3260的任何系统与至少一个实施例一致。在至少一个实施例中，方法3200的步骤3210-3260中的一个或更多个可以至少部分地实现为存储在非暂时性有形机器可读介质上的可执行代码，其在由诸如图1-31的任何处理器或计算设备之类的一个或更多个处理器运行时，可以使得一个或更多个处理器实施步骤3210-3260中的一个或更多个。
在至少一个实施例中，方法3200开始于步骤3210，其中构建用于工作负载的任务图。在至少一个实施例中，可以通过进行一个或更多个应用编程接口(API)或者库函数调用显式地构建任务图。API或者库函数调用用来创建任务图的每个节点并且将节点与计算任务关联。在至少一个实施例中，计算任务可以相应于要由计算资源实施的内核任务、与用来启动可执行图的主机CPU上实施的回调函数相应的主机任务、存储器设置任务、存储器拷贝任务等等。在至少一个实施例中，相同和/或不同API或者库函数调用可以用来将有向边添加到任务图以便指定两个任务之间的数据、执行和/或其他从属性。在至少一个实施例中，可以基于其中一个任务生成的数据是另一个任务的源数据的数据从属性间接地将有向边添加到任务图。在至少一个实施例中，可替换地，可以通过记录任务流间接地构建任务图。在至少一个实施例中，通过以下方式记录任务流：创建用于记录任务图的任务流和/或使用一个或更多个API或者库函数调用将现有的任务流转换成用于记录任务图的任务流，进行指定要添加到任务流的任务的一个或更多个附加的API或者库函数调用，并且然后使用一个或更多个第三API或者库函数调用完成任务流到任务图的转换。在至少一个实施例中，任务图是不可执行的任务图。
在至少一个实施例中，当把每个任务作为节点添加到任务图时，每个节点的创建顺序与每个节点关联。在至少一个实施例中，与每个节点关联的创建顺序与序数相应，该序数指示每个节点是否为创建并且添加到任务图的第一、第二、第三等等的节点。在至少一个实施例中，每个节点的创建顺序可以与计数器的计数值相应，该计数器在初始创建任务图时被初始化(例如为0或1)并且每当新节点添加到任务图时递增。
图33A-C图示出依照至少一个实施例的示例性任务图。在至少一个实施例中，任务图3310、3320和/或3330可以与用于指定一个或更多个CUDA任务的CUDA图一致。在至少一个实施例中，任务图3310、3320和/或3330可以与用于指定一个或更多个HIP任务的HIP图一致。如图33所示，任务图3310为包括示为具有标签A-G和结束的圆圈的八个节点/任务的有向图。任务图3300开始于任务A。示为箭头的任务A分别与任务B和F之间的有向边相应于任务A与任务B和F之间的从属性。更具体地，任务A与任务B之间的有向边指示任务A必须在任务B可能开始之前完成，并且任务A与任务F之间的有向边指示任务A必须在任务F可能开始之前完成。由于任务B与F之间不存在有向边，因而任务B和F的实施彼此独立，这意味着任务B和F可以相对于彼此以任何顺序实施，包括：任务B在任务F之前实施，任务F在任务B之前实施，任务B和F至少部分地并发实施，任务B和F以相对于彼此交错的方式实施，和/或其任意组合。任务图3300中的附加任务从属性包括：任务B必须在任务C可能开始之前完成，任务B必须在任务D可能开始之前完成，任务C和D二者必须在任务E可能开始之前完成，以及任务F必须在任务G可能开始之前完成。此外，任务图3300指示任务E和G必须在到达任务结束时完成任务图3300的工作负载之前完成。在至少一个实施例中，通过任务结束将控制返回到主机CPU。在至少一个实施例中，当到达任务结束时，可以开始另一个任务(未示出)。
在至少一个实施例中并且未在图33A中示出，任务A-G中的每一个可以与包括两个或更多任务/节点的子图相应。在至少一个实施例中，子图可以与实现多任务工作负载的API和/或库函数调用创建的一个或更多个节点和/或从属性相应。在至少一个实施例中，图33B图示出与节点D相应的子图3320。子图3320包括标记为D1-D4的四个任务。子图3320进一步示出任务D1必须在任务D2和D3可能开始之前完成，并且任务D2和D3二者必须在任务D4可能开始之前完成。此外，由于任务D2与D3之间没有从属性，因而任务D2与D3可以以相对于彼此的任何顺序实施。子图3320进一步指示任务D4必须在任务D完成之前完成。
向后参照图32，在至少一个实施例中并且在步骤3220，从任务图生成可执行图。在至少一个实施例中，经由实例化实施步骤3220。在至少一个实施例中，经由编译实施步骤3220。在至少一个实施例中，可执行图的生成包括将任务图的每个任务/节点分配给相应计算资源和/或相应计算资源的子集，例如图1-31的任何计算资源、图形处理器、片段处理器、核心、执行单元、集群、多处理器、计算单元和/或流式多处理器等等。在至少一个实施例中，相应计算资源可以基于要实施的任务进行选择，例如任务是否为内核任务、主机任务、存储器设置任务、存储器拷贝任务等等。在至少一个实施例中，相应计算资源可以附加地和/或可替换地基于要由任务实施的计算、任务所需的诸如线程、线程束等等之类的资源的数量和/或其他准则进行选择。在至少一个实施例中，相应计算资源可以附加地和/或可替换地基于任务图中的任务之间的从属性进行选择。
在至少一个实施例中，与任务图中每个任务/节点关联的诸如序数之类的创建顺序信息与可执行图中的相应任务/节点关联。
在至少一个实施例中，当从任务图生成可执行图时，可以实施一个或更多个优化。在至少一个实施例中，优化可以包括重新布置任务以利用任务之间的并行性。在至少一个实施例中，重新布置任务可以包括以下一个或更多个：确定使用多少计算资源，在计算资源之间转移任务，等等。在至少一个实施例中，优化可以包括例如通过将具有对于另一个任务的数据从属性的任务分配给与另一个任务相同的计算资源而将任务重新分发至计算资源以便降低由于数据从属性引起的开销。在至少一个实施例中，重新分发任务可以降低执行时间，因为另一个任务的结果在不必将结果传输至不同计算资源以供某个任务使用的情况下对于某个任务就是可用的。在至少一个实施例中，重新分发任务可以降低执行时间，因为可以更快速地检测另一个任务的完成，允许在另一个任务完成之后更快速地开始某个任务。在至少一个实施例中，可以将任务分配给计算资源以降低计算成本或者将另一个任务的结果传输至分配给前一个任务的计算资源的延迟。在至少一个实施例中，优化可以包括重新布置任务以更好地平衡和/或优化计算资源的能力，例如对每个计算资源可用的线程、线程束、计算线程阵列、核心、流式多处理器等等的数量。在至少一个实施例中，可以通过进行一个或更多个API或者库函数调用而将任务图转换为可执行图。
在至少一个实施例中，在步骤3230，可执行图启动一次或多次以便使得工作负载被实施。在至少一个实施例中，一旦准备好可执行图，就可以启动可执行图以使得一个或更多个计算资源实施任务图的工作负载。在至少一个实施例中，可以使用一个或更多个API或者库函数调用启动可执行图。在至少一个实施例中，对于可执行图中的每个任务/节点，启动可执行图可以包括以下一个或更多个：将与每个任务相应的可执行代码发送至关联的计算资源，基于每个任务的一个或更多个属性将配置发送至每个关联的计算资源，将用于每个任务的数据传输(如果必要的话)至每个关联的计算资源，和/或当每个关联的计算资源完成每个任务时，将每个任务的结果传输回主机CPU或者需要每个任务的结果作为用于要实施的下一个对应任务的数据的一个或更多个计算资源。在至少一个实施例中，当可执行图中的所有任务都被关联的计算资源执行时，任务图的工作负载已经实施。在至少一个实施例中，可执行图可以启动与实施相同工作负载一样多的次数。
在至少一个实施例中，在步骤3240，构建用于另一个工作负载的另一个任务图。在至少一个实施例中，除了构建另一个任务图之外，步骤3240基本上类似于步骤3210。在至少一个实施例中，另一个任务图的工作负载可以不同于过程3210期间构建的任务图的工作负载。在至少一个实施例中，另一个任务图是不可执行的任务图。
在至少一个实施例中，在步骤3250，将另一个任务图应用到可执行图。在至少一个实施例中，步骤3250的目标是修改可执行图，使得可执行图可用来让一个或更多个计算资源实施另一个任务图的工作负载，而不是实施步骤3210期间构建的任务图的工作负载。在至少一个实施例中，当步骤3250成功地将另一个任务图应用到可执行图时，可执行图可以在无需招致例如通过使用步骤3220将另一个任务图转换成新的可执行图的计算成本和延迟的情况下进行修改。在至少一个实施例中，可以使用一个或更多个API和/或库函数调用将另一个任务图应用到可执行图。
在至少一个实施例中，步骤3250在以下时候可能有帮助：从任务流创建任何任务图，和/或用来创建任何任务图的代码包括一个或更多个API和/或库函数调用，使得任务图的拓扑结构和/或任务图的任务的参数通常不为编写代码以构建任务图的开发者所知悉。在至少一个实施例中，开发者可能不知道是否有任何API和/或库函数调用创建了新的任务节点和/或创建了任务节点之间的新的从属性。在至少一个实施例中，开发者可能具有对于任务图的更简单的理解，例如如任务图3310所示，即使任务图的一个或更多个节点可能具有其中包括的嵌入式子图，例如如任务图3330所示，其中节点D已被嵌入式子图3320代替。在至少一个实施例中，当一个或更多个API和/或库函数调用也进行一个或更多个附加的API或库函数调用时，这可能进一步加剧。
在至少一个实施例中，开发者可能不能够知道两个任务图是否具有相同的拓扑结构和/或两个任务图中的任务的参数是否允许将一个任务图应用到从另一个任务图生成的可执行图。在至少一个实施例中，步骤3460通过使用如图34中所示的方法3400而解决了这些问题，图34图示出依照至少一个实施例的流程图。在至少一个实施例中并且在步骤3460的上下文内，新任务图与步骤3440期间构建的任务图相应，并且可执行图与步骤3420期间生成的可执行图相应。尽管结合图1-31的系统和技术描述了方法3400的方法步骤3410-3460，但是应当理解，被配置为实施步骤3410-3460的任何系统都与至少一个实施例一致。在至少一个实施例中，方法3400的步骤3410-3460中的一个或更多个可以至少部分地实现为存储在非暂时性有形机器可读介质上的可执行代码，其在由诸如图1-31的任何处理器或计算设备之类的一个或更多个处理器运行时，可以使得一个或更多个处理器实施步骤3410-3460中的一个或更多个。
在至少一个实施例中，方法3400开始于步骤3410，其中将新任务图的拓扑结构与可执行图的拓扑结构相比较。在至少一个实施例中，新任务图的拓扑结构和可执行图的拓扑结构可以通过确认新任务图和可执行图具有相同的节点/任务布置以及节点/任务之间相同的有向边/从属性布置而进行比较。在至少一个实施例中，这可能是一个复杂的比较，因为即使一个拓扑结构相对于另一个拓扑结构在各种不同的节点之间具有一个或更多个反射、置换等，两个图也可能具有相同的拓扑结构。在与来自图33A的任务图3310一致的至少一个实施例中，节点C和D交换的任务图在拓扑结构上等效于任务图3310。
然而，在至少一个实施例中，拓扑结构的比较可以通过利用在构建新任务图时与每个节点关联的创建顺序信息以及在生成可执行图时与每个节点关联的创建顺序信息而简化。
图35图示出依照至少一个实施例的流程图。在步骤3410的上下文内，新任务图和可执行图与正被比较的两个图相应。尽管结合图1-31的系统和技术描述了方法3500的方法步骤3510-3570，但是应当理解，被配置为实施步骤3510-3570的任何系统都与至少一个实施例一致。在至少一个实施例中，方法3500的步骤3510-3570中的一个或更多个可以至少部分地实现为存储在非暂时性有形机器可读介质上的可执行代码，其在由诸如图1-31的任何处理器或计算设备之类的一个或更多个处理器运行时，可以使得一个或更多个处理器实施步骤3510-3570中的一个或更多个。
在至少一个实施例中，方法3500开始于步骤3510，其中确定正被比较的两个图是否具有相同的节点数量。在至少一个实施例中，当两个图具有相同的节点数量时，两个图的比较利用步骤3520继续。在至少一个实施例中，当两个图没有相同的节点数量时，两个图具有不同的拓扑结构，该结果使用步骤3570返回。
在至少一个实施例中，在步骤3520，比较两个图的每个相应节点的类型。在至少一个实施例中，来自一个图的节点和另一个图的节点都与相同的创建或序数关联时，这两个节点彼此相应。在至少一个实施例中，一旦标识了相应的节点，则将节点的类型与相应节点的类型相比较以确定那些类型是否为相同类型。在至少一个实施例中，当节点为用于实施特定类型的内核任务的内核节点时，相应的节点必须是用于实施相同特定类型的内核任务的内核节点以便那些类型为相同类型。在至少一个实施例中，当节点与用于将特定存储器类型的一维存储器块设置为填充值的存储器设置任务相应时，相应的节点必须与用于将相同特定存储器类型的一维存储器块设置为填充值的存储器设置任务相应以便那些类型为相同类型。在至少一个实施例中，可以对与存储器设置任务、存储器拷贝任务、主机任务等等相应的节点实施类似的测试。在至少一个实施例中，如果图中的节点在另一个图中没有相应的节点，例如因为在另一个图中没有发现具有相同的关联创建顺序的相应节点，那么不存在具有相同类型的相应节点。在至少一个实施例中，一发现任何节点和相应节点没有相同的类型，那么步骤3520就可以在不比较任何另外的节点的类型的情况下结束。
在至少一个实施例中，在步骤3530，确定每个节点和相应节点是否为相同类型。在至少一个实施例中，当一个图中的任何节点和另一个图中的相应节点没有相同类型时，这些图具有不同的拓扑结构，并且该结果使用步骤3570返回。在至少一个实施例中，当一个图中的每个节点在另一个图中具有相同类型的相应节点并且另一个图中的每个节点在图中具有相同类型的相应节点时，这些图以步骤3540开始进一步进行比较。
在至少一个实施例中，在步骤3540，比较这些图的每个相应节点的从属性。在至少一个实施例中，当一个图中的节点的从属性列表包括具有与另一个图中的相应节点的从属性列表相同的关联创建顺序信息的节点时，节点具有与相应节点相同的从属性。在至少一个实施例中，如果与创建顺序序数15关联的节点对于与创建顺序序数6和11关联的节点有从属性，并且具有创建序数15的相应节点对于与创建顺序序数6和11关联的相应节点有从属性，那么节点和相应节点具有相同的从属性。在至少一个实施例中，一发现任何节点和相应节点没有相同的从属性，那么步骤3540就可以在不比较任何另外的节点的从属性的情况下结束。
在至少一个实施例中，在步骤3550，确定一个图中的每个节点和另一个图中的相应节点具有相同的从属性列表。在至少一个实施例中，当一个图中的任何节点和另一个图中的相应节点没有相同的从属性列表时，这些图具有不同的拓扑结构，并且该结果使用步骤3570返回。在至少一个实施例中，当一个图中的每个节点在另一个图中具有带有相同从属性列表的相应节点时，这些图被认为具有相同的拓扑结构，并且该结果使用步骤3560返回。
在至少一个实施例中，在步骤3560，返回指示两个图具有相同拓扑结构的结果。在至少一个实施例中，两个图具有相同拓扑结构的指示可以作为返回值提供给调用来实施方法3500的函数。在至少一个实施例中，步骤3560一完成，方法3500就结束。
在至少一个实施例中，在步骤3570，返回指示两个图具有不同拓扑结构的结果。在至少一个实施例中，两个图具有不同拓扑结构的指示可以作为返回值提供给调用来实施方法3500的函数。在至少一个实施例中，步骤3570一完成，方法3500就结束。
在至少一个实施例中，方法3500的步骤可以以与图35中所示不同的顺序实施。在至少一个实施例中，步骤3540和3550可以在步骤3520和3530之前实施。在至少一个实施例中，步骤3520和3540可以通过一次一个节点地比较每个节点的类型和从属性二者而并发地实施。
向后参照图34，在至少一个实施例中并且在步骤3420，基于步骤3410期间实施的比较确定新任务图和可执行图的拓扑结构是否为相同的拓扑结构。在至少一个实施例中，步骤3560或3570返回的结果可以用来确定新任务图和可执行图的拓扑结构是否为相同或不同的拓扑结构。在至少一个实施例中，当新任务图和可执行图的拓扑结构是相同的拓扑结构时，以步骤3430开始将新任务图的参数应用到可执行图。在至少一个实施例中，当新任务图和可执行图的拓扑结构是不同的拓扑结构时，使用步骤3460返回将新任务图应用到可执行图方面的失败。
在至少一个实施例中，在步骤3430，将新任务图的参数应用到可执行任务。在至少一个实施例中，为了修改可执行图，使得可执行图可以实施新任务图的工作负载，基于新任务的参数修改可执行图。在至少一个实施例中，可以对可执行图做出的修改可能干扰和/或打断步骤3220期间实施的优化中的一个或更多个。在至少一个实施例中，可以对可执行图做出的修改不干扰步骤3220期间实施的优化。在至少一个实施例中，步骤3410的拓扑结构比较已经确认新任务图和可执行图的任务和从属性具有具有相同类型、相同从属性的相应节点。在至少一个实施例中，将新任务图的参数应用到可执行图是受限制的，使得对于某些类型的任务而言，当应用新任务图的参数时对可执行图中的节点参数的任何修改都限制为某些类型的参数。在至少一个实施例中，对于某些类型的任务而言，对可执行图的允许的修改可以限制为某些类型的参数，其不干扰步骤3220期间实施的优化。
在至少一个实施例中，与涉及关联的计算资源触发主机CPU上的回调函数的主机任务相应的节点可以包括可修改参数，所述可修改参数包括主机CPU上的回调函数的诸如指针之类的标识符和/或要作为一个或更多个参量传递至回调函数的可为关联的计算设备访问的数据中的一个或更多个。在至少一个实施例中，对回调函数和/或参量参数的修改是可能的，因为这些修改并未修改正由关联的计算资源实施以实施主机任务的功能或者干扰典型地在生成可执行图期间应用到主机任务的任何优化。
在至少一个实施例中，与涉及关联的计算资源将可为关联的计算资源访问的存储器块设置为指定的填充值的存储器设置任务相应的节点可以具有可修改参数，所述可修改参数包括以下一个或更多个：存储器块的起始位置，存储器块的大小，和/或存储器块中的存储器位置要设置的填充值。在至少一个实施例中，存储器块的大小可以相应于一维存储器块的存储器字的数量、二维存储器块的行列数等等。在至少一个实施例中，对位置、大小和/或填充值参数的修改并未修改正由关联的计算资源实施以实施存储器设置任务的功能和/或干扰典型地在生成可执行图期间应用到存储器设置任务的任何优化。在至少一个实施例中，可能无法例如从一维块到二维块或者反之亦然修改存储器块的块类型，和/或例如在全局、本地、纹理、常数、共享的、寄存器等等的存储器之间修改被设置的存储器的类型，因为这些类型的修改可能涉及存储器设置任务的不同实现方式。
在至少一个实施例中，与涉及关联的计算资源将来自可为关联的计算资源访问的存储器块的内容移动到可为关联的计算资源访问的另一个存储器块的存储器拷贝任务相应的节点可以具有可修改参数，所述可修改参数包括以下一个或更多个：存储器块的源位置，另一个存储器块的目的地位置，和/或要拷贝的存储器量的大小。在至少一个实施例中，要拷贝的存储器量的大小可以相应于一维存储器块的存储器字的数量、二维存储器块的行列数等等。在至少一个实施例中，对源位置、目的地位置和/或大小参数的修改并未修改正由关联的计算资源实施以实施存储器拷贝任务的功能或者干扰典型地在生成可执行图期间应用到存储器拷贝任务的任何优化。在至少一个实施例中，可能无法例如从一维块到二维块或者反之亦然修改正被拷贝的存储器的块类型，和/或例如在全局、本地、纹理、常数、共享的、寄存器等等的存储器之间修改任一存储器块中的存储器的类型，因为这些类型的修改可能涉及存储器拷贝任务的不同实现方式。
在至少一个实施例中，与涉及关联的计算资源实施计算任务的内核任务相应的节点可以具有可修改参数，所述可修改参数包括其是计算任务的参量的指向可为关联的计算资源访问的数据的指针和/或要使用的线程的数量。在至少一个实施例中，对这些参数的修改是可能的，因为它们并未修改正由关联的计算资源实施以实施内核任务的功能或者干扰典型地在生成可执行图期间应用到内核任务的任何优化。在至少一个实施例中，可能无法修改计算任务的起始地址、要实施的计算任务、指向要实施的下一个任务的指针等等，因为这些修改了底层计算任务和/或可能干扰可执行图中的任务之间的从属性。
在至少一个实施例中，在步骤3440，确定是否通过步骤3430成功将新任务图的参数应用到可执行图。在至少一个实施例中，当成功将新任务图的每个参数应用到可执行图时，使用步骤3450返回成功。在至少一个实施例中，在至少一个实施例中，当新任务图的任何参数都不能成功应用到可执行图时，使用步骤3460返回失败。
在至少一个实施例中，在步骤3450，返回指示新任务图已经成功应用到可执行图的结果。在至少一个实施例中，将新任务图应用到可执行图方面的成功可以作为返回值提供给被调用来实施方法3400的函数。在至少一个实施例中，一完成步骤3450，方法3400就结束。
在至少一个实施例中，在步骤3460，返回指示新任务图尚未成功应用到可执行图的结果。在至少一个实施例中，将新任务图应用到可执行图方面的失败可以作为返回值提供给被调用来实施方法3400的函数。在至少一个实施例中，一完成步骤3460，方法3400就结束。
在至少一个实施例中，方法3400的步骤可以以与图34中所示不同的顺序实施。在至少一个实施例中，步骤3410和3430可以通过比较新任务图和可执行图的拓扑结构并且在同一遍中通过新任务图和可执行图的节点将新任务图的参数应用到可执行图而并发地实施。在至少一个实施例中，在比较新任务图和可执行图的拓扑结构期间和/或在将新任务图的参数应用到可执行图方面检测和/或遇到的任何失败都可以报告给日志、用户、用户接口等等。在至少一个实施例中，失败可以经由警报、警告、错误等等报告。在至少一个实施例中，可以带关于失败源的解释地报告失败。在至少一个实施例中，报告失败可以帮助用户改变如何构建任务图以便避免失败的再次发生。
向后参照图32，在至少一个实施例中并且在步骤3260，确定另一个任务图是否成功应用到可执行图。在至少一个实施例中，步骤3250期间做出的将另一个任务图应用到可执行图的每个API和/或库函数调用的成功可以基于每个API和/或库函数调用返回的状态和/或错误代码确定。在至少一个实施例中，状态和/或错误代码可以与步骤3450返回的成功和/或步骤3460返回的失败相应。在至少一个实施例中，当另一个任务图成功应用到可执行图时，可以通过返回到步骤3230启动可执行图一次或多次，其中可执行图用来使得一个或更多个计算资源利用可执行图的每次启动实施另一个任务图的另一个工作负载。在至少一个实施例中，当另一个任务图未成功应用到可执行图时，控制返回到步骤3220，其中从另一个任务图生成新的可执行图。在至少一个实施例中，新的可执行图于是可以使用步骤3230启动以使得一个或更多个计算资源实施另一个任务图的工作负载。
在至少一个实施例中，即使通过使用步骤3250将另一个任务图应用到可执行图获得的可执行图与可执行图使用步骤3220从另一个任务图生成的情况相比可能导致另一个工作负载的性能较慢，通过避免使用步骤3220从另一个任务图生成另一个可执行图的开销的时间节省也可能导致让另一个工作负载被实施的总体处理时间较短。在至少一个实施例中，当使用步骤3250将另一个任务图应用到可执行图干扰了和/或打断了创建可执行图时应用到可执行图的一个或更多个优化时，通过使用步骤3250将另一个任务图应用到可执行图获得的可执行图可能导致另一个工作负载的性能较慢。
总而言之，所公开的技术可以用来允许可执行图实施与新任务图关联的工作负载。在至少一个实施例中，构建描述工作负载的任务图。接着，通过使用诸如实例化之类的过程从任务图生成可执行图。接着，可执行图启动一次或多次，可执行图的每次启动导致任务图的工作负载被实施。接着，构建描述另一个工作负载的另一个任务图。代替从另一个任务图生成新的可执行图的是，接着，一个或更多个API和/或库函数调用用来尝试将另一个任务图应用到已经生成的可执行图。尝试将另一个任务图应用到可执行图包括比较另一个任务图和可执行图的拓扑结构以了解那些拓扑结构是否为相同的拓扑结构。如果拓扑结构是相同的拓扑结构，那么将另一个任务图的参数应用到可执行图。如果另一个任务图的每个参数都成功地应用到可执行图，那么可以启动可执行图以使得一个或更多个计算资源实施另一个任务图的另一个工作负载，而无需招致从另一个任务图生成新的可执行图的开销。如果另一个任务图不能成功应用到可执行图，则从另一个任务图生成新的可执行图。
所公开的实施例的至少一个技术优点在于，所公开的实施例可以用来在不必生成附加的可执行图和/或知道任务图和/或可执行图中任何一个的拓扑结构的情况下将新任务图应用到从另一个任务图生成的可执行图，使得可执行图可以被执行以实施另一个工作负载。因此，利用所公开的实施例，可以使得一个或更多个计算资源实施不同的工作负载，而无需招致每当构建新任务图时与生成新的可执行图关联的经常显著的开销。与现有技术方法形成对照的是，所公开的实施例使得单个可执行图能够用来使得一个或更多个计算资源实施多个不同的工作负载，这相对于现有技术方法降低了总体执行时间。这些技术优点相对于现有技术方法提供了一项或多项技术进步。
1.在至少一个实施例中，一种非暂时性计算机可读介质，包括指令，这些指令在由一个或多个处理器执行时，使得所述一个或多个处理器执行至少一个应用编程接口(API)调用以便通过将第二任务图的不可执行版本应用到第一任务图的可执行版本而修改第一任务图的可执行版本。
2.依照第1条的非暂时性计算机可读介质，其中通过以下方式将第二任务图的不可执行版本应用到第一图的可执行版本：将与第一任务图的可执行版本关联的第一拓扑结构和与第二任务图的不可执行版本关联的第二拓扑结构相比较；确定第一拓扑结构是与第二拓扑结构相同的拓扑结构；以及将第二任务图的不可执行版本的一个或多个参数应用到第一任务图的可执行版本。
3.依照第1条或第2条的非暂时性计算机可读介质，其中在将第二任务图的不可执行版本的所述一个或多个参数应用到第一任务图的可执行版本之后，第一任务图的可执行版本在启动时将一个或多个计算资源配置为实施与第二任务图的不可执行版本关联的工作负载，而不是与第一任务图的不可执行版本关联的工作负载。
4.依照第1-3条中任一项的非暂时性计算机可读介质，其中当第二任务图的不可执行版本的第一参数应用到第一任务图的可执行版本时，该第一参数并不干扰从第一任务图的不可执行版本生成第一任务图的可执行版本时结合到第一任务图的可执行版本中的优化。
5.依照第1-4条中任一项的非暂时性计算机可读介质，其中将第二任务图的不可执行版本的第一参数应用到第一任务图的可执行版本的相应参数改变了该相应参数，该相应参数包括：针对与第一任务图的可执行版本关联的主机任务，指向主机中央处理单元上的回调函数或者该回调函数的参量的指针；针对与第一任务图的可执行版本关联的存储器设置任务，要设置的存储器块的位置、该存储器块的大小或者该存储器块的填充值；针对与第一任务的可执行版本关联的存储器拷贝任务，源存储器块的位置、其中要拷贝该源存储器的内容的目的地位置或者该源存储器块的大小；或者针对与第一图的可执行版本关联的内核任务，一个或多个参量或者线程数量。
6.依照第1-5条中任一项的非暂时性计算机可读介质，其中通过以下方式将第一拓扑结构与第二拓扑结构相比较：确定第一拓扑结构中包括的每个节点是否与第二拓扑结构中包括的具有相同任务类型的节点相应；以及确定第二拓扑结构中包括的每个节点是否与第一拓扑结构中包括的具有相同任务类型的节点相应。
7.依照第1-6条中任一项的非暂时性计算机可读介质，其中当第一拓扑结构中包括的节点和第二拓扑结构中包括的节点与相同的创建顺序值关联时，这两个节点彼此相应。
8.依照第1-7条中任一项的非暂时性计算机可读介质，其中通过确定第一拓扑结构中包括的每个节点和第二拓扑结构中包括的相应节点是否具有相同的从属性列表而进一步将第一拓扑结构与第二拓扑结构相比较。
9.依照第1-8条中任一项的非暂时性计算机可读介质，其中通过确定第一拓扑结构和第二拓扑结构是否具有相同的节点数量而进一步将第一拓扑结构与第二拓扑结构相比较。
10.在至少一个实施例中，一种用于将新任务图应用到可执行图的计算机实现的方法包括：通过将不可执行第二任务图应用到第一任务图的可执行版本而修改第一任务图的可执行版本。
11.依照第10条的计算机实现的方法，进一步包括：在将不可执行第二任务图应用到第一任务图的可执行版本之后，第一任务图的可执行版本在启动时将一个或多个计算资源配置为实施与第二任务图的不可执行版本关联的工作负载，而不是与第一任务图的不可执行版本关联的工作负载。
12.依照第10条或第11条的计算机实现的方法，其中将第二任务图应用到第一图的可执行版本包括：将第一任务图的可执行版本的第一拓扑结构与第二任务图的不可执行版本的第二拓扑结构相比较；以及当第一拓扑结构是与第二拓扑结构相同的拓扑结构时，将第二任务图的参数应用到第一任务图的可执行版本。
13.依照第10-12条中任一项的计算机实现的方法，其中将不可执行第二任务图的第一参数应用到第一任务图的可执行版本的相应参数改变了该相应参数，该相应参数包括：针对与第一任务图的可执行版本关联的主机任务，指向主机中央处理单元上的回调函数或者该回调函数的参量的指针；针对与第一任务图的可执行版本关联的存储器设置任务，要设置的存储器块的位置、该存储器块的大小或者该存储器块的填充值；针对与第一任务的可执行版本关联的存储器拷贝任务，源存储器块的位置、其中要拷贝该源存储器的内容的目的地位置或者该源存储器块的大小；或者针对与第一图的可执行版本关联的内核任务，一个或多个参量或者线程数量。
14.依照第10-13条中任一项的计算机实现的方法，其中将第一拓扑结构与第二拓扑结构相比较包括以下一个或多个：确定第一拓扑结构和第二拓扑结构是否具有相同的节点数量；确定第一拓扑结构中的每个第一节点是否在第二拓扑结构中具有相同任务类型的相应第二节点；确定第二拓扑结构中的每个第二节点是否在第一拓扑结构中具有相同任务类型的相应第一节点；确定第一拓扑结构中的每个第一节点和第二拓扑结构中的相应第二节点是否具有相同的从属性列表；或者确定第二拓扑结构中的每个第二节点和第一拓扑结构中的相应第一节点是否具有相同的从属性列表。
15.依照第10-14条中任一项的计算机实现的方法，其中第一任务图或者第二任务图中的至少一个包括计算统一设备架构(CUDA)图或者可执行的可移植异构计算接口(HIP)图。
16.在至少一个实施例中，一种用于将新任务图应用到可执行图的计算机实现的方法包括：从不可执行第一任务图生成可执行图，其中该可执行图将一个或多个计算资源配置为实施与不可执行第一任务图关联的第一工作负载；以及通过将不可执行第二任务图的一个或多个参数应用到所述可执行图而修改该可执行图，其中该修改的可执行图将所述一个或多个计算资源配置为实施与不可执行第二任务图关联并且不同于第一工作负载的第二工作负载。
17.依照第16条的计算机实现的方法，其中将所述不可执行第二任务图应用到所述可执行图包括：将该可执行图的第一拓扑结构与所述不可执行第二任务图的第二拓扑结构相比较；以及当第一拓扑结构和第二拓扑结构是相同的拓扑结构时，将所述不可执行第二任务图的参数应用到所述可执行图。
18.依照第16条或第17条的计算机实现的方法，其中将第一拓扑结构与第二拓扑结构相比较包括：确定第一拓扑结构中的第一节点和第二拓扑结构中的相应第二节点是否具有相同的类型和相同的从属性列表。
19.依照第16-18条中任一项的计算机实现的方法，其中当第一节点和第二节点与相同的创建顺序值关联时，所述相应第二节点与第一节点相应。
20.依照第16-19条中任一项的计算机实现的方法，其中第一任务图或者第二任务图中的至少一个包括计算统一设备架构(CUDA)图或者可执行的可移植异构计算接口(HIP)图。
任何权利要求中记载的任何权利要求要素和/或本申请中描述的任何要素以任何方式的任意和所有组合都落入本实施例和保护的预期范围内。
其他变型处于本公开的精神之内。因此，虽然所公开的技术易受各种修改和替代构造的影响，但是其某些例示的实施例在附图中示出并且在上文详细地进行了描述。然而，应当理解，不存在将公开内容限制为所公开的一种或多种具体形式的意图，而是相反地，意图在于覆盖落入如所附权利要求书中所限定的公开内容的精神和范围内的所有修改、替代构造和等效物。
在描述所公开实施例的上下文中(特别是在以下权利要求的上下文中)术语“一”、“一个”、“该/这个”和类似指称的使用应当被解释为覆盖单数和复数二者，除非本文另有指明或者与上下文明显矛盾并且不作为术语的定义。术语“包含”、“具有”、“包括”和“含有”应当被解释为开放式术语(表示“包括但不限于”)，除非另有说明。术语“已连接”在未修改并且涉及物理连接时，应当被解释为部分地或者全部地包含在……内、附接到或者接合在一起，即使有什么东西介入。本文中值的范围的叙述仅仅意在用作一种个别地引用落在范围内的每个单独的值的速记方法，除非本文另有指明，并且每个单独的值结合到说明书中，就好像它在本文中被个别地叙述一样。除非另有说明或者与上下文相矛盾，术语“集合”(例如“项目集合”)或者“子集”的使用应当被解释为包括一个或更多个成员的非空集合。此外，除非另有说明或者与上下文相矛盾，相应集合的术语“子集”不一定表示相应集合的适当子集，而是子集和相应集合可以相等。
除非另有明确说明或与上下文明显矛盾，诸如“A、B和C中的至少一个”或者“A、B和C中的至少一个”形式的短语之类的合取语言以其他方式与上下文一起理解，通常用来呈现项目、项等等可以是或者A或者B或者C，或者A和B和C的集合的任何非空子集。例如，在具有三个成员的集合的说明性示例中，合取短语“A、B和C中的至少一个”以及“A、B和C中的至少一个”指的是任何以下集合：{A}、{B}、{C}、{A,B}、{A,C}、{B,C}、{A,B,C}。因此，这样的合取语言通常并非意在暗示某些实施例要求A的至少一个、B的至少一个和C的至少一个中的每一个存在。此外，除非另有说明或者与上下文矛盾，术语“多个”指示为复数的状态(例如，“多个项目”指示多项目)。复数的项目的数量为至少两个，但是当显式地或者通过上下文这样指示时可以为更多个。进一步，除非另有说明或者以其他方式从上下文清楚，短语“基于”表示“至少部分地基于”，并且不是“仅仅基于”。
本文描述的过程的操作可以以任何适当的顺序实施，除非本文另有指明，或者以其他方式与上下文明显矛盾。在至少一个实施例中，诸如本文描述的那些过程(或者其变型和/或组合)之类的过程在配置有可执行指令的一个或更多个计算机系统的控制下实施，并且被实现为通过硬件或者其组合共同地在一个或更多个处理器上执行的代码(例如，可执行指令、一个或更多个计算机程序或者一个或更多个应用)。在至少一个实施例中，代码例如以包括可由一个或更多个处理器执行的多个指令的计算机程序的形式存储在计算机可读存储介质上。在至少一个实施例中，计算机可读存储介质是一种非暂时性计算机可读存储介质，其排除暂时性信号(例如传播瞬态电或电磁传输)，但是包括暂时性信号收发器内的非暂时性数据存储电路系统(例如缓冲区、缓存和队列)。在至少一个实施例中，代码(例如可执行代码或者源代码)存储在其上存储了可执行指令的一个或更多个非暂时性计算机可读存储介质(或者存储可执行指令的其他存储器)的集合上，这些可执行指令在被计算机系统的一个或更多个处理器执行时(即，作为被其执行的结果)，使得计算机系统实施本文描述的操作。在至少一个实施例中，非暂时性计算机可读存储介质集合包括多个非暂时性计算机可读存储介质，并且多个非暂时性计算机可读存储介质的个别非暂时性存储介质中的一个或更多个缺乏所有代码，而多个非暂时性计算机可读存储介质共同存储所有代码。在至少一个实施例中，可执行指令被执行，使得不同的指令由不同的处理器执行——例如，非暂时性计算机可读存储介质存储指令，主要中央处理单元(“CPU”)执行一些指令，而图形处理单元(“GPU”)执行其他指令。在至少一个实施例中，计算机系统的不同部件具有单独的处理器，并且不同的处理器执行不同的指令子集。
因此，在至少一个实施例中，计算机系统被配置为实现一个或更多个服务，其单独地或者共同地实施本文描述的过程的操作，并且这样的计算机系统被配置有允许操作实施的适用的硬件和/或软件。进一步，实现本公开的至少一个实施例的计算机系统为单个设备，并且在另一个实施例中，为包括多个设备的分布式计算机系统，这些设备不同地操作，使得分布式计算机系统实施本文描述的操作并且使得单个设备不实施所有操作。
除非另有声明，本文提供的任何和所有示例或者示例性语言(例如“诸如/例如”)的使用仅仅意在更好地说明本公开的实施例，并不构成对本公开的范围的限制。说明书的语言不应当被解释为表明任何未要求保护的要素对本公开的实践至关重要。
本文引用的包括出版物、专利申请和专利在内的所有参考文献据此通过引用并入本文，其程度就好像每个参考文献都是单独并特别指出要通过引用并入本文并且在此完整阐述一样。
在说明书和权利要求书中，可以与其派生词一起使用术语“耦合”和“连接”。应当理解，这些术语可能并不预期作为彼此的同义词。相反，在特定示例中，“连接”或“耦合”可以用来指示两个或更多元件彼此直接或间接地物理或电气接触。“耦合”也可以表示两个或更多元件彼此不直接接触，但是仍然彼此合作或交互。
除非另有特别说明，可以理解的是，贯穿整个说明书，诸如“处理”、“计算”、“运算”、“确定”等等之类的术语涉及计算机或计算系统或者类似的电子计算设备的动作和/或过程，其将表示为计算系统的寄存器和/或存储器内的诸如电的之类的物理的量的数据操纵和/或变换为类似地表示为计算系统的存储器、寄存器或者其他这样的信息存储、传送或显示设备内的物理量的其他数据。
以类似的方式，术语“处理器”可以指处理来自寄存器和/或存储器的电子数据并且将该电子数据变换为可以存储在寄存器和/或存储器中的其他电子数据的任何设备或者设备的部分。作为非限制性示例，“处理器”可以是CPU或者GPU。“计算平台”可以包括一个或更多个处理器。当在本文使用时，“软件”过程可以包括例如随着时间的推移实施工作的软件和/或硬件实体，例如任务、线程和智能代理。再者，每个过程可以涉及用于连续或间歇地、顺序或并行地执行指令的多个过程。术语“系统”和“方法”在本文中可互换使用，只要系统可以体现一种或多种方法，并且方法可以被认为是系统。
在本文中，可能提及获得、采集、接收模拟或数字数据或者将该数据输入到子系统、计算机系统或者计算机实现的机器中。获得、采集、接收或者输入模拟和数字数据的过程可以以各种各样的方式——例如通过接收数据作为函数调用或者对应用编程接口的调用的参数的数据而实现。在一些实现方式中，获得、采集、接收或者输入模拟或数字数据的过程可以通过经由串行或并行接口传输数据而实现。在另一个实现方式中，获得、采集、接收或者输入模拟或数字数据的过程可以通过经由计算机网络从提供实体到采集实体传输数据而实现。也可能提及提供、输出、传送、发送或者呈现模拟或数字数据。在各种不同的示例中，提供、输出、传送、发送或者呈现模拟或数字数据的过程可以通过传输作为函数调用的输入或输出参数、应用编程接口或者进程间通信机制的参数的数据而实现。
尽管上面的讨论阐述了所描述的技术的示例实现方式，但是其他架构可以用来实现所描述的功能，并且预期处于本公开的范围之内。此外，尽管上面出于讨论的目的限定了具体职责分布，各种不同的功能和职责可能根据情况以不同的方式分布和划分。
此外，尽管以特定于结构特征和/或方法行为的语言描述了主题，但是应当理解的是，所附权利要求书中要求保护的主题不必限于所描述的具体特征或行为。相反，具体特征和行为作为实现权利要求的示例性形式而公开。

Claims (20)

1.一种包括指令的非暂时性计算机可读介质，所述指令由一个或更多个处理器执行时，使得所述一个或更多个处理器执行至少一个应用编程接口(API)调用以便通过将第二任务图的不可执行版本应用到第一任务图的可执行版本来修改所述第一任务图的可执行版本。
2.如权利要求1所述的非暂时性计算机可读介质，其中通过以下方式将所述第二任务图的不可执行版本应用到第一图的可执行版本：
将与所述第一任务图的可执行版本关联的第一拓扑结构和与所述第二任务图的不可执行版本关联的第二拓扑结构相比较；
确定所述第一拓扑结构是与所述第二拓扑结构相同的拓扑结构；以及
将所述第二任务图的不可执行版本的一个或更多个参数应用到所述第一任务图的可执行版本。
3.如权利要求2所述的非暂时性计算机可读介质，其中在将所述第二任务图的不可执行版本的所述一个或更多个参数应用到所述第一任务图的可执行版本之后，所述第一任务图的可执行版本一旦启动，将一个或更多个计算资源配置为实施与所述第二任务图的不可执行版本关联的工作负载，而不是与所述第一任务图的不可执行版本关联的工作负载。
4.如权利要求2所述的非暂时性计算机可读介质，其中当所述第二任务图的不可执行版本的第一参数应用到所述第一任务图的可执行版本时，所述第一参数并不干扰从所述第一任务图的不可执行版本生成所述第一任务图的可执行版本时结合到所述第一任务图的可执行版本中的优化。
5.如权利要求2所述的非暂时性计算机可读介质，其中将所述第二任务图的不可执行版本的第一参数应用到所述第一任务图的可执行版本的相应参数，改变了所述相应参数，所述相应参数包括：
针对与所述第一任务图的可执行版本关联的主机任务，指向主机中央处理单元上的回调函数或者所述回调函数的参量的指针；
针对与所述第一任务图的可执行版本关联的存储器设置任务，要设置的存储器块的位置、所述存储器块的大小或者所述存储器块的填充值；
针对与所述第一任务的可执行版本关联的存储器拷贝任务，源存储器块的位置、所述源存储器的内容要被拷贝到的目的地的位置或者所述源存储器块的大小；或者
针对与所述第一图的可执行版本关联的内核任务，多个线程或者一个或更多个参量。
6.如权利要求2所述的非暂时性计算机可读介质，其中通过以下方式将所述第一拓扑结构与所述第二拓扑结构相比较：
确定包括在所述第一拓扑结构中的每个节点是否与包括在所述第二拓扑结构中的具有相同任务类型的节点对应；以及
确定包括在所述第二拓扑结构中的每个节点是否与包括在所述第一拓扑结构中的具有相同任务类型的节点对应。
7.如权利要求6所述的非暂时性计算机可读介质，其中当包括在所述第一拓扑结构中的节点和包括在所述第二拓扑结构中的节点与相同的创建顺序值关联时，这两个节点彼此对应。
8.如权利要求6所述的非暂时性计算机可读介质，其中通过确定包括在所述第一拓扑结构中的每个节点和包括在所述第二拓扑结构中的对应节点是否具有相同的从属性列表而进一步将所述第一拓扑结构与所述第二拓扑结构相比较。
9.如权利要求6所述的非暂时性计算机可读介质，其中通过确定所述第一拓扑结构和所述第二拓扑结构是否具有相同的节点数量而进一步将所述第一拓扑结构与所述第二拓扑结构相比较。
10.一种用于将新任务图应用到可执行图的计算机实现的方法，所述方法包括：
通过将不可执行的第二任务图应用到第一任务图的可执行版本来修改所述第一任务图的可执行版本。
11.如权利要求10所述的计算机实现的方法，进一步包括：在将所述不可执行的第二任务图应用到所述第一任务图的可执行版本之后，所述第一任务图的可执行版本一旦启动，将一个或更多个计算资源配置为实施与所述第二任务图的不可执行版本关联的工作负载，而不是与所述第一任务图的不可执行版本关联的工作负载。
12.如权利要求10所述的计算机实现的方法，其中将所述第二任务图应用到第一图的可执行版本包括：
将所述第一任务图的可执行版本的第一拓扑结构与所述第二任务图的不可执行版本的第二拓扑结构相比较；以及
当所述第一拓扑结构是与所述第二拓扑结构相同的拓扑结构时，将所述第二任务图的参数应用到所述第一任务图的可执行版本。
13.如权利要求12所述的计算机实现的方法，其中将所述不可执行的第二任务图的第一参数应用到所述第一任务图的可执行版本的相应参数，改变了所述相应参数，所述相应参数包括：
针对与所述第一任务图的可执行版本关联的主机任务，指向主机中央处理单元上的回调函数或者所述回调函数的参量的指针；
针对与所述第一任务图的可执行版本关联的存储器设置任务，要设置的存储器块的位置、所述存储器块的大小或者所述存储器块的填充值；
针对与第一任务的可执行版本关联的存储器拷贝任务，源存储器块的位置、所述源存储器的内容要被拷贝到的目的地位置或者所述源存储器块的大小；或者
针对与所述第一图的可执行版本关联的内核任务，多个线程或者一个或更多个参量。
14.如权利要求12所述的计算机实现的方法，其中将所述第一拓扑结构与所述第二拓扑结构相比较包括以下中的一个或更多个：
确定所述第一拓扑结构和所述第二拓扑结构是否具有相同数量的节点；
确定所述第一拓扑结构中的每个第一节点是否在所述第二拓扑结构中具有相同任务类型的对应的第二节点；
确定所述第二拓扑结构中的每个第二节点是否在所述第一拓扑结构中具有相同任务类型的对应的第一节点；
确定所述第一拓扑结构中的每个第一节点和所述第二拓扑结构中的对应的第二节点是否具有相同的从属性列表；或者
确定所述第二拓扑结构中的每个第二节点和所述第一拓扑结构中的对应的第一节点是否具有相同的从属性列表。
15.如权利要求10所述的计算机实现的方法，其中所述第一任务图或者所述第二任务图中的至少一个包括计算统一设备架构(CUDA)图或者可执行的可移植异构计算接口(HIP)图。
16.一种用于将新任务图应用到可执行图的计算机实现的方法，所述方法包括：
从不可执行的第一任务图生成可执行图，其中所述可执行图将一个或更多个计算资源配置为实施与所述不可执行的第一任务图关联的第一工作负载；以及
通过将不可执行的第二任务图的一个或更多个参数应用到所述可执行图来修改所述可执行图，
其中经修改的可执行图将所述一个或更多个计算资源配置为实施与所述不可执行的第二任务图关联且不同于所述第一工作负载的第二工作负载。
17.如权利要求16所述的计算机实现的方法，其中将所述不可执行的第二任务图应用到所述可执行图包括：
将所述可执行图的第一拓扑结构与所述不可执行的第二任务图的第二拓扑结构相比较；以及
当所述第一拓扑结构和所述第二拓扑结构是相同的拓扑结构时，将所述不可执行的第二任务图的参数应用到所述可执行图。
18.如权利要求16所述的计算机实现的方法，其中将所述第一拓扑结构与所述第二拓扑结构相比较包括：确定所述第一拓扑结构中的第一节点和所述第二拓扑结构中的相应第二节点是否具有相同的类型和相同的从属性列表。
19.如权利要求18所述的计算机实现的方法，其中当所述第一节点和所述第二节点两者与相同的创建顺序值关联时，所述相应第二节点与第一节点相对应。
20.如权利要求16所述的计算机实现的方法，其中所述第一任务图或者所述第二任务图中的至少一个包括计算统一设备架构(CUDA)图或者可执行的可移植异构计算接口(HIP)图。
CN202011280064.3A 2019-11-15 2020-11-16 用于修改可执行图以实施与新任务图关联的工作负载的技术 Pending CN112817738A (zh)

Applications Claiming Priority (4)

Application Number Priority Date Filing Date Title
IN201941046676 2019-11-15
IN201941046676 2019-11-15
US16/732,014 2019-12-31
US16/732,014 US20210149734A1 (en) 2019-11-15 2019-12-31 Techniques for modifying an executable graph to perform a workload associated with a new task graph

Publications (1)

Publication Number Publication Date
CN112817738A true CN112817738A (zh) 2021-05-18

Family

ID=73198150

Family Applications (1)

Application Number Title Priority Date Filing Date
CN202011280064.3A Pending CN112817738A (zh) 2019-11-15 2020-11-16 用于修改可执行图以实施与新任务图关联的工作负载的技术

Country Status (2)

Country Link
EP (1) EP3822770A1 (zh)
CN (1) CN112817738A (zh)

Cited By (2)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN114721844A (zh) * 2022-03-10 2022-07-08 云和恩墨(北京)信息技术有限公司 数据缓存方法和装置、计算机设备、存储介质
CN115759260A (zh) * 2022-11-17 2023-03-07 北京百度网讯科技有限公司 深度学习模型的推理方法、装置、电子设备和存储介质

Families Citing this family (2)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US11586580B2 (en) * 2021-07-08 2023-02-21 Avago Technologies International Sales Pte. Limited Parallel processor optimized for machine learning
CN114035810B (zh) * 2022-01-10 2022-04-15 北京一流科技有限公司 用于多流并行的同步部署系统及其方法

Family Cites Families (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US7164422B1 (en) * 2000-07-28 2007-01-16 Ab Initio Software Corporation Parameterized graphs with conditional components
US10802945B2 (en) * 2016-12-07 2020-10-13 Ab Initio Technology Llc Differencing of executable dataflow graphs
US10990364B2 (en) * 2018-06-06 2021-04-27 Ab Initio Technology Llc Updating executable graphs

Cited By (4)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN114721844A (zh) * 2022-03-10 2022-07-08 云和恩墨(北京)信息技术有限公司 数据缓存方法和装置、计算机设备、存储介质
CN114721844B (zh) * 2022-03-10 2022-11-25 云和恩墨(北京)信息技术有限公司 数据缓存方法和装置、计算机设备、存储介质
CN115759260A (zh) * 2022-11-17 2023-03-07 北京百度网讯科技有限公司 深度学习模型的推理方法、装置、电子设备和存储介质
CN115759260B (zh) * 2022-11-17 2023-10-03 北京百度网讯科技有限公司 深度学习模型的推理方法、装置、电子设备和存储介质

Also Published As

Publication number Publication date
EP3822770A1 (en) 2021-05-19

Similar Documents

Publication Publication Date Title
US20210248115A1 (en) Compute graph optimization
US20210149734A1 (en) Techniques for modifying an executable graph to perform a workload associated with a new task graph
EP3822770A1 (en) Techniques for modifying an executable graph to perform a workload associated with a new task graph
US11294713B2 (en) Asynchronous data movement pipeline
US11080111B1 (en) Technique for sharing context among multiple threads
US20220300578A1 (en) Application programming interface to accelerate matrix operations
US20210279837A1 (en) Cooperative parallel memory allocation
US20210294673A1 (en) Techniques for orchestrating stages of thread synchronization
US11861758B2 (en) Packet processing acceleration using parallel processing
EP3822785A1 (en) Techniques for modifying executable graphs to perform different workloads
US20210149719A1 (en) Techniques for modifying executable graphs to perform different workloads
US20230222069A1 (en) Application programming interface to disassociate a virtual address
US20230185637A1 (en) Application programming interfaces for interoperability
CN116724292A (zh) 线程组的并行处理
CN116802613A (zh) 同步图形执行
US11568523B1 (en) Techniques to perform fast fourier transform
WO2023077436A1 (en) Thread specialization for collaborative data transfer and computation
US20220365829A1 (en) Data compression api
US20230087457A1 (en) Application programming interface to retrieve data
US20220334900A1 (en) Application programming interface to indicate increased resource usage
US20230185635A1 (en) Application programming interfaces for interoperability
US20230185638A1 (en) Application programming interfaces for interoperability
JP2024514370A (ja) リソース使用量を監視するためのアプリケーション・プログラミング・インターフェース
AU2022204612A1 (en) Synchronization barrier
WO2023044408A1 (en) Application programming interface to retrieve data

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