CN108241508B - 处理OpenCL内核的方法及用于该方法的计算设备 - Google Patents

处理OpenCL内核的方法及用于该方法的计算设备 Download PDF

Info

Publication number
CN108241508B
CN108241508B CN201711216401.0A CN201711216401A CN108241508B CN 108241508 B CN108241508 B CN 108241508B CN 201711216401 A CN201711216401 A CN 201711216401A CN 108241508 B CN108241508 B CN 108241508B
Authority
CN
China
Prior art keywords
core
processing
group
control core
work
Prior art date
Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
Active
Application number
CN201711216401.0A
Other languages
English (en)
Other versions
CN108241508A (zh
Inventor
伯恩哈德·爱格
乌洙林
赵永显
刘东勋
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Samsung Electronics Co Ltd
SNU R&DB Foundation
Original Assignee
Samsung Electronics Co Ltd
SNU R&DB Foundation
Priority date (The priority date is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the date listed.)
Filing date
Publication date
Application filed by Samsung Electronics Co Ltd, SNU R&DB Foundation filed Critical Samsung Electronics Co Ltd
Publication of CN108241508A publication Critical patent/CN108241508A/zh
Application granted granted Critical
Publication of CN108241508B publication Critical patent/CN108241508B/zh
Active legal-status Critical Current
Anticipated expiration legal-status Critical

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/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/38—Concurrent instruction execution, e.g. pipeline or look ahead
    • G06F9/3836—Instruction issuing, e.g. dynamic instruction scheduling or out of order instruction execution
    • G06F9/3851—Instruction issuing, e.g. dynamic instruction scheduling or out of order instruction execution from multiple instruction streams, e.g. multistreaming
    • 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
    • G06F9/5044—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 considering hardware capabilities
    • 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
    • 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/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/38—Concurrent instruction execution, e.g. pipeline or look ahead
    • G06F9/3885—Concurrent instruction execution, e.g. pipeline or look ahead using a plurality of independent parallel functional units
    • G06F9/3887—Concurrent instruction execution, e.g. pipeline or look ahead using a plurality of independent parallel functional units controlled by a single instruction for multiple data lanes [SIMD]
    • 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
    • G06T—IMAGE DATA PROCESSING OR GENERATION, IN GENERAL
    • G06T1/00—General purpose image data processing
    • G06T1/20—Processor architectures; Processor configuration, e.g. pipelining

Landscapes

  • Engineering & Computer Science (AREA)
  • Theoretical Computer Science (AREA)
  • Software Systems (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • General Engineering & Computer Science (AREA)
  • Multimedia (AREA)
  • Stored Programmes (AREA)
  • Memory System Of A Hierarchy Structure (AREA)

Abstract

本公开提供了处理OpenCL内核的方法及用于该方法的计算设备。多核计算设备包括具有第一低级别控制核和处理核组的控制核组。控制核组将用于执行开放计算语言(OpenCL)内核的工作组分配至处理核组中的第一低级别控制核和第一处理核。处理核组执行通过控制核组分配的工作组的处理，并且输出处理的结果。控制核按照分级方式分组。

Description

处理OpenCL内核的方法及用于该方法的计算设备
相关申请的交叉引用
本申请要求于2016年12月27日在韩国知识产权局提交的韩国专利申请No.10-2016-0180133的利益和优先权，该韩国专利申请的公开以引用方式全文并入本文中。
技术领域
本公开涉及处理开放计算语言(OpenCL)内核的方法和用于该方法的计算设备。
背景技术
开发的计算设备具有其中在单个集成电路中包括多个核或处理器以满足应用的性能需求的结构。例如，多核处理器是指其中集成了具有算术功能的两个或更多个核的处理器。众核处理器在单个处理器中集成了十六个或更多个核。多核处理器和众核处理器可被包括在具有多媒体芯片的装置中和内置于TV或便携式装置中。
开放计算语言(OpenCL)是用于编写用于跨多个平台执行的程序的框架。换句话说，OpenCL是一种用于跨多个平台执行的程序的诸如通用多核中央处理单元(CPU)、现场可编程门阵列(FPGA)和图形处理单元(GPU)的开放的通用并行计算框架。OpenCL将GPU的能力扩展至除图形处理之外的领域。因此，进行了各种尝试来利用多核计算设备有效地处理OpenCL程序。
发明内容
本发明构思的至少一个实施例提供了一种用于处理开放计算语言(OpenCL)内核的方法和一种用于该方法的计算设备，其中分级控制核组将用于执行OpenCL内核的工作组分配至处理核组。
根据本发明构思的示例性实施例，一种多核计算设备包括：包括第一低级别控制核和处理核组的控制核组。控制核组被构造为将用于执行OpenCL内核的工作组分配至第一低级别控制核和第一处理核。处理核组包括第一处理核中的至少一个。处理核组被构造为处理通过控制核组分配的工作组并且输出工作组的处理结果。包括在控制核组中的多个控制核根据通过控制核向其分配工作组的第一处理核的数量按照分级方式分组。
根据本发明构思的示例性实施例，一种处理OpenCL内核的方法，包括以下步骤：通过控制核组将用于执行OpenCL内核的工作组分配至第一低级别控制核和第一处理核；通过包括第一处理核中的至少一个的处理核组来处理分配的工作组；以及通过处理核组输出工作组的处理的结果。控制核组的多个控制核根据通过控制核向其分配工作组的第一处理核的数量按照分级方式分组。
根据本发明构思的示例性实施例，提供了一种非暂时性计算机可读记录介质，其上记录有用于在计算机上执行上述OpenCL内核处理方法的指令。
根据本发明构思的示例性实施例，一种多核计算设备，包括：处理核组，其包括具有多个第一处理核和多个第二处理核；根控制核，其被构造为从主机处理器接收与OpenCL内核的执行相关的信息，该信息指示用于执行OpenCL内核的工作组；第一叶控制核，其连接在根控制核与第一处理核之间；以及第二叶控制核，其连接在根控制核与第二处理核之间。根控制核将各工作组的第一部分发送至第一叶控制核，并且将各工作组的第二部分发送至第二中间控制核。
附图说明
从下面结合附图的详细描述中，将更加清楚地理解本发明构思的实施例，附图中：
图1是示出OpenCL平台的图；
图2是示出根据本发明构思的示例性实施例的计算设备的框图；
图3是示出根据本发明构思的示例性实施例的计算设备的框图；
图4是示出根据本发明构思的示例性实施例的通过控制核向其分配工作组的处理核的数量的图；
图5是示出根据本发明构思的示例性实施例的动态地分配用于执行OpenCL内核的工作组的方法的图；
图6是示出根据本发明构思的示例性实施例的计算设备的图；
图7是示出根据本发明构思的示例性实施例的叶控制核和处理核的图；
图8是示出根据本发明构思的示例性实施例的在处理核上执行的OpenCL内核代码的图；
图9是示出根据本发明构思的示例性实施例的计算设备的图；
图10是示出根据本发明构思的示例性实施例的缓冲器的图；
图11是示出根据本发明构思的示例性实施例的在增加了缓冲器的情况下在处理核上执行的OpenCL内核代码的图；
图12是示出根据本发明构思的示例性实施例的当增加了缓冲器时用于处理OpenCL工作组的OpenCL内核代码的图；
图13和图14是示出根据本发明构思的示例性实施例的用于评价计算设备的性能的实验条件的表；
图15是示出根据本发明构思的示例性实施例的计算设备的性能的评价结果的图；
图16是示出根据本发明构思的示例性实施例的处理OpenCL内核的方法的流程图；以及
图17是示出根据本发明构思的示例性实施例的处理OpenCL内核的方法的流程图。
具体实施方式
现在将详细说明本发明构思的示例性实施例，其示例示于附图中，图中相同的附图标记始终指代相同的元件。在这方面，当前示例性实施例可具有不同形式并且不应该被理解为限于本文阐述的说明。
在示例性实施例的以下说明中，当一个部分或元件被称作连接至另一部分或元件时，该部分或元件可直接连接至所述另一部分或元件，或者可在它们之间存在中间部分或元件的情况下电连接至所述另一部分或元件。在实施例的说明中，使用诸如单元或模块的术语来指代具有至少一个功能或操作并且通过硬件、软件或硬件与软件的组合实现的单元。
下文中，将参照附图描述本发明构思的示例性实施例。
图1是示出开放计算语言(OpenCL)平台100的图。
OpenCL是用于同时实现异构计算装置的操作的多种编程语言之一。因此，用OpenCL写的程序可跨多个平台执行。
参照图1，OpenCL平台100包括主机处理器110和诸如计算设备120至140的至少一个计算设备。
主机处理器110可为在其上执行主机程序的处理器并且可限定用于执行内核的工作空间。这里，工作空间可被划分为各自包括多个工作项的工作组。各个工作项是作为最小工作单位的工作空间的点。例如，主机处理器110可为通用中央处理单元(CPU)。然而，主机处理器110不限于此。
在实施例中，内核包括诸如多线程程序代码的内核代码。内核代码可包括一种或多种语句。语句是一种构成程序的单独指令或者有含义的表达，并且可包括指令序列。在多线程程序执行模型中，将执行域划分为多个工作项，以使得可相对于各个工作项执行内核代码。在实施例中，在分离的软件线程中执行各个工作项。在实施例中，一些工作项是同步的，以共享数据。可将要同步的工作项归类为一个工作组。可将多线程程序的执行域划分为多个工作组。
计算设备120至140可包括诸如计算装置121和122的至少一个计算装置。另外，计算装置121和122可包括诸如PE 10至40的至少一个处理元件(PE)。计算装置121和122可为用于处理工作组的单元，并且PE 10至40可为用于处理工作项的单元。例如，计算装置121和122可处理从主机处理器110接收到的工作组，并且PE 10至40可处理包括在工作组中的工作项。因此，可通过PE 10至40并行处理工作组的工作项。计算设备120至140中的每一个可为众核处理器或多核处理器。然而，计算设备120至140不限于此。计算设备120至140可为包括至少一个核的任何处理器。在实施例中，各个计算装置是核并且各个PE是能够处理软件线程的硬件线程。
OpenCL程序可包括一个主机程序和至少一个内核。主机程序可通过主机处理器110执行，并且可通过经由主机处理器110发送命令来要求计算设备120至140执行计算，或者可管理计算设备120至140的存储器。这里，内核是指在计算设备120至140上执行的程序，并且可被称作OpenCL内核或者内核代码。
当在包括多个核的计算设备上执行内核时，各个核可对应于诸如计算装置121和122的计算装置。换句话说，计算设备的核可处理被分配至其的工作组，并且包括在核中的诸如PE 10至40的处理元件可处理包括在工作组中的工作项。
计算设备可从主机处理器110接收工作组。例如，主机处理器110可根据核中的工作组的处理结果将工作组动态地分配至核。在实施例中，主机处理器110包括用于动态地分配工作组的工作分布管理器。
然而，如果工作分布管理器将工作组动态地分配至多个核，工作分布管理器执行的分布工作的量(例如，开销(overhead))可根据核的数量而增大，因此，计算设备的性能可降低。另外，如果工作分布管理器将工作组静态地分配至多个核，则工作组可在不同的时间被处理，因此计算设备的操作效率可降低。另外，如果主机处理器110包括工作分布管理器，则主机处理器110可承担很重的负担。
如果计算设备的核用作工作分布管理器，则该核不用于处理工作项。因此，在具有多个核的计算设备上执行OpenCL内核的情况需要有效地分配工作组的方法。
图2是示出根据本发明构思的示例性实施例的计算设备200的框图。
根据实施例的计算设备200可包括多个核，并且多个核可归为控制核组210和处理核组220。计算设备200还可包括除图2所示的元件之外的通用元件。
图2所示的计算设备200可对应于图1所示的计算设备120至140。
控制核组210是将用于执行OpenCL内核的工作组分配至低级别控制核和处理核的一组控制核。控制核的规格可与处理核的规格相同。然而，控制核不限于此。
处理核组220是被构造为处理通过控制核组210分配的工作组并且输出处理的结果的一组处理核。在实施例中，处理核中的每一个包括至少一个处理元件。例如，处理核可对应于众核处理器的核。然而，处理核不限于此。
在实施例中，控制核组210的控制核根据通过控制核向其分配工作组的处理核的数量按照分级方式分组。将参照图3来描述控制核组210的分级分组的示例。
图3是示出根据本发明构思的示例性实施例的计算设备300的框图。
参照图3，计算设备300包括控制核组320和处理核组360。控制核组320包括根控制核330和叶控制核组350。另外，根据控制核组320的层级，计算设备300还可包括中间控制核组340。图3所示的处理核组360对应于图2所示的处理核组220，因此这里不提供处理核组360的详细描述。
在实施例中，当根控制核330接收关于OpenCL内核的执行的信息时，根控制核330将工作组分配至低级别控制核。在图3中，根控制核330下方的控制核被指示为在中间控制核组340中。根据控制核组320的层级，根控制核330下方的控制核可被包括在中间控制核组340和叶控制核组350中的至少一个中。
在实施例中，根控制核330从主机处理器310接收关于OpenCL内核的执行的信息。例如，当执行OpenCL内核时，主机处理器310可产生关于OpenCL内核的执行的信息，并且将该信息发送至根控制核330。关于OpenCL内核的执行的信息可包括用于执行OpenCL内核的工作组的总量和关于包括在OpenCL内核中的应用的信息中的至少一种。
叶控制核组350可包括多个叶控制核，它们被构造为接收关于通过高级别控制核分配的工作组的信息和将工作组分配至处理核。在图3中，叶控制核组350上方的控制核被指示为在中间控制核组340中。然而，根据控制核组320的层级，叶控制核组350上方的各个控制核可为根控制核330和中间控制核组340的控制核中的至少一个。
中间控制核组340可包括多个中间控制核，并且被构造为接收关于通过根控制核330分配的工作组的信息和将工作组分配至叶控制核组350。另外，中间控制核组340可排列为多个级别。参照图3，中间控制核可根据通过中间控制核向其分配工作组的处理核的数量按照分级方式分组。在这种情况下，高级别中间控制核341将工作组分配至低级别中间控制核342。
图4是示出根据本发明构思的示例性实施例的通过控制核向其分配工作组的处理核的数量的图。
根据实施例，控制核组210的控制核根据通过控制核向其分配工作组的处理核的数量按照分级方式分组。
图4是示出包括8×8处理核的处理核组的图。参照图4，通过根控制核330向其分配工作组的处理核410全是处理核组的处理核。
一个中间控制核可将工作组分配至处理核组的处理核420中的一些。例如，在图4所示的情况下，如果中间控制核组340包括四个中间控制核，则一个中间控制核可向其分配工作组的处理核的数量可为十六个。一个中间控制核向其分配工作组的处理核的数量可与另一中间控制核向其分配工作组的处理核的数量不同。例如，当可使用64个处理核和4个中间控制核时，不是各个中间控制核向16个处理核分配工作组，而是一个中间控制核可向15个处理核分配工作组而另一中间控制核向17个处理核分配工作组。
在实施例中，叶控制核接收关于通过高级别控制核分配的工作组的信息，并且可将工作组直接分配至处理核。例如，一个叶控制核可将工作组分配至通过一个中间控制核向其分配工作组的处理核420的处理核430中的一些。也就是说，通过一个叶控制核向其分配工作组的处理核的数量少于通过一个中间控制核向其分配工作组的处理核的数量。例如，如果一个中间控制核管理16个处理核，则一个叶控制核可仅管理4个处理核。
另外，根据本发明构思的实施例，由于叶控制核将工作组分配至与其邻近的处理核，因此可低开销地根据地域性来分配表现出地域性的工作组。
另外，由于实施例的计算设备200除处理核组220之外还包括控制核组210，因此所有的处理核可处理工作组。另外，由于控制核组210按照分级方式组织，因此可有效地执行工作组的动态分配。现在，将参照图5描述利用按照分级方式分组的控制核动态地分配工作组的方法。
图5是示出根据本发明构思的实施例的动态地分配用于执行OpenCL内核的工作组的方法的图。
根据实施例，控制核组根据通过低级别控制核和处理核对工作组的处理结果将工作组再分配至其它低级别控制核和其它处理核。图5所示的主机处理器510对应于图3所示的主机处理器310，因此将省略对其的详细描述。图5所示的计算设备520可对应于图3所示的计算设备300。
参照图5，如果完全地处理了分配至处理核560的工作组，则可通过叶控制核550将被分配至另一处理核的工作组再分配至该处理核560。另外，叶控制核550可将针对额外工作组的请求发送至中间控制核540。在这种情况下，叶控制核550可从处理核560接收针对额外工作组的请求，或者可周期性地在处理核560中检查工作组的处理，以确定分配至处理核560的工作组是否已完全处理。
如果已完全处理分配至叶控制核550的工作组，则可通过中间控制核540将已分配至另一叶控制核的工作组再分配至叶控制核550。另外，中间控制核540可将针对额外工作组的请求发送至根控制核530。
例如，如果第一处理核已完成处理第一工作组而第二处理核仍在处理第二工作组，则即使计划接着在第二处理核上接着处理第三工作组，与第二处理核关联的叶控制核也可将第三工作组发送至第一处理核，而不等待第二处理核完成处理第二工作组。在被中间控制核通知从与第一处理核关联的叶控制核接收请求之后，与第二处理核关联的叶控制核可确定第一处理核可处理这种分配。与第一处理核关联的叶控制核可在完成第一处理之后响应于第一处理核针对额外工作组的请求而发送所述请求。
如果完全处理了分配至中间控制核540的工作组，则可通过根控制核530将已分配至另一中间控制核的工作组再分配至中间控制核540。
另外，当处理了用于执行OpenCL内核的工作组时，如果根控制核530从主机处理器510接收关于新的OpenCL内核的执行的信息，则根控制核530可再分配被分配至处理核组的工作组。
在一个实施例中，省略了图5所示的中间控制核540，因此根控制核530直接与叶控制核560通信。例如，在该实施例中，叶控制核560将处理核560针对额外工作组的请求传送至根控制核530。
图6是示出根据本发明构思的示例性实施例的计算设备630的图。
参照图6，计算设备630连接至主机处理器610和存储器620。另外，计算设备630包括根控制核631、叶控制核633和处理核634。根据控制核组的层级，计算设备630还可包括中间控制核632。
参照图6，根控制核631连接至中间控制核632。因此，根控制核631可将关于分配至处理核634的工作组的信息发送至中间控制核632或者叶控制核633，而不是直接访问处理核634以分配工作组。由于中间控制核632连接至叶控制核633，因此中间控制核632可将关于分配至处理核组中的一些处理核的工作组的信息发送至叶控制核633。由于叶控制核633连接至处理核634，因此叶控制核633可直接将关于分配的工作组的信息发送至处理核634。
由于根控制核631连接至主机处理器610，因此根控制核631可接收关于当执行OpenCL内核时产生的工作组的信息。另外，控制核组和处理核组的核可从存储器620接收用于分配和处理工作组所需的信息。
图7是示出根据本发明构思的示例性实施例的包括叶控制核730和处理核710的计算设备700的示图。
根据实施例，叶控制核730将工作组分配至处理核710。例如，参照图7，叶控制核730可连接至四个处理核710，以将工作组直接分配至处理核710，并且接收工作组处理的结果或者针对额外工作组的请求。
另外，处理核710中的每一个可包括被构造为处理工作组的核711和与外部装置通信的路由器712。在实施例中，处理核710通过路由器712连接至外部路由器720，并且外部路由器720连接至叶控制核730。因此，处理核710可从叶控制核730接收关于工作组的信息。在图7中，处理核710中的每一个包括一个核711。然而，包括在处理核710中的每一个中的核的数量不限于此。
路由器712和外部路由器720可通过片上网络(NoC)通信方法或者总线通信方法彼此通信。然而，路由器712与外部路由器720之间的通信不限于此。
另外，参照图7，叶控制核730可将工作组分配至四个处理核710。然而，叶控制核730向其分配工作组的处理核的数量不限于此。
图8示出了根据本发明构思的示例性实施例的在处理核上执行的OpenCL内核代码。
图8所示的OpenCL内核代码是用于将矩阵(a)的元素与矩阵(b)的元素相加并且将其结果存储在矩阵(c)中。由于矩阵(a)和矩阵(b)是二维矩阵，因此可使用二维工作组空间来执行图8所示的OpenCL内核代码。换句话说，矩阵(a)和矩阵(b)的元素可为作为工作空间的点的工作项。
如果在图6所示的计算设备630上执行图8所示的OpenCL内核代码，则处理核634可访问存储器620来读取矩阵(a)和矩阵(b)的元素。这里，get_global_id()是用于读取内核代码的索引的函数。因此，当处理核634执行图8所示的OpenCL内核代码时，可平行地处理多个工作项。然而，由于有必要重复访问存储器620来读取工作项，因此用于处理OpenCL内核代码所需的时间可增加。
图9是示出根据本发明构思的示例性实施例的计算设备920的图。
实施例的计算设备920还可包括诸如第一缓冲器965和第二缓冲器970的缓冲器，以存储关于工作组的信息。由于计算设备920包括缓冲器，因此计算设备920可从缓冲器读取关于工作组的信息，因此可以以相对较少的次数来访问存储器。图9所示的主机处理器910和根控制核930对应于图5所示的主机处理器510和根控制核530，因此将省略对其的详细描述。
参照图9，叶控制核950和处理核960连接至第二缓冲器970。在实施例中，第二缓冲器970从叶控制核950接收分配至处理核960的工作组的编号和识别号(ID)，并且存储工作组的编号和ID。因此，叶控制核950可通过从第二缓冲器970读取关于分配至处理核960的工作组的信息而不是访问存储器来检查在处理核960中的工作组的处理结果。如果叶控制核950确定分配至处理核960的工作组已被完全处理，则叶控制核950可将已被分配至另一处理核的工作组再分配至处理核960。在实施例中，第二缓冲器970被包括在处理核960中，并且通过处理核960的路由器连接至叶控制核950。然而，第二缓冲器970的位置不限于此。
在实施例中，中间控制核940和叶控制核950连接至第一缓冲器965。在实施例中，第一缓冲器965从中间控制核940接收分配至叶控制核950的工作组的编号和ID，并且存储工作组的编号和ID。因此，中间控制核940可通过从第一缓冲器965中读取关于分配至叶控制核950的工作组的信息而不是访问存储器来检查在叶控制核950中的工作组的处理结果。如果中间控制核940确定分配至叶控制核950的工作组已被完全处理，则中间控制核940可将已被分配至另一叶控制核的工作组再分配至叶控制核950。在实施例中，第一缓冲器965被包括在叶控制核950中。然而，第一缓冲器965的位置不限于此。
在图9中，计算设备920包括两个缓冲器，也就是说，第一缓冲器965和第二缓冲器970。然而，计算设备920的缓冲器的数量不限于此。
图10是示出根据本发明构思的示例性实施例的缓冲器1030的图。图9的第一缓冲器965或第二缓冲器970可通过图10的缓冲器1030来实现。
参照图10，缓冲器1030连接至叶控制核1010和处理核1020。另外，缓冲器1030存储分配至处理核1020的工作组的编号和ID。
在实施例中，叶控制核1010将工作组分配至处理核1020，并且随后将分配至处理核1020的工作组的ID和编号发送至缓冲器1030。在处理工作组之后，处理核1020可访问缓冲器1030以更新分配至处理核1020的工作组的编号。因此，叶控制核1010可通过访问缓冲器1030而不是访问存储器来检查在处理核1020中处理工作组的结果。例如，处理核1030可在完成工作组的处理之后递减分配的工作组的数量。
图11示出了根据本发明构思的示例性实施例的在包括缓冲器的计算设备的处理核上执行的OpenCL内核代码。
图11所示的OpenCL内核代码是利用缓冲器1030来执行与图8所示的OpenCL内核代码的功能相同的功能的代码。换句话说，在增加了缓冲器1030的处理核1020上执行OpenCL内核代码，以将矩阵(a)和矩阵(b)的元素相加，并且将相加的结果存储在矩阵(c)中。
参照图11，read_wg_id()是用于从缓冲器1030读取分配至处理核1020的工作组的ID的函数。因此，由于OpenCL内核代码通过利用函数read_wg_id()来使用存储在缓冲器1030中的工作组的ID作为索引，因此处理核1020可平行地处理工作组，而不访问存储器。
图12示出了根据本发明构思的示例性实施例的当计算设备包括缓冲器时用于执行OpenCL工作组的OpenCL内核代码。
根据实施例，如果处理核1020包括一个或多个处理元件，并且处理元件的数量少于工作组的工作项的数量，则计算设备转换OpenCL内核，从而处理元件将工作项序列化和按次序执行工作项。
例如，如果处理元件的数量少于工作组的工作项的数量，则处理元件中的每一个可处理多个工作项。在这种情况下，每当处理元件中的每一个完全地处理分配至其的工作项时，该处理元件可将关于完全处理的工作项的信息存储在存储器中，并且切换至下一工作项。因此，当处理元件处理多个工作项时，处理元件访问存储器的次数可增加。
计算设备可基于工作项联合方法来转换OpenCL内核代码，以使得各个处理元件能够按次序处理多个工作项，因此减少处理元件访问存储器的次数。在实施例中，在内核代码中包括迭代语句，从而可将内核代码执行等于将被序列化的工作项的数量的次数。迭代语句可被称作工作项联合循环(WCL)或者线程循环。
参照图12，read_num_wgs()是用于从缓冲器1030读取接连被分配至处理核1020的工作组的编号的函数。
因此，由于图12所示的OpenCL内核代码基于用于从缓冲器1030读取信息的函数read_wg_id()和read_num_wgs()来设置索引，因此，处理核1020访问存储器的次数可减少。另外，由于循环围绕OpenCL内核代码，因此不必每当执行工作项时都调用OpenCL内核代码。也就是说，可通过仅调用一次OpenCL内核代码按次序执行多个工作项。因此，与其中未使用工作项联合方法的情况相比，处理核1020访问存储器的次数可减少。
图13和图14是示出根据本发明构思的示例性实施例的用于评价计算设备的性能的实验条件的表。
图13是示出当将不同的调度方法应用于计算设备时根据实施例提供的计算设备分配工作组所需的延迟时间的表。计算设备中的每一个的处理核组包括九十六个处理核，并且计算设备中的每一个的控制核组包括根控制核、四个中间控制核和十六个叶控制核。在实施例中，控制核以50MHz的时钟频率操作，并且处理核以500MHz的时钟频率操作。
在图13中，集中式调度是指其中控制核将工作组动态地分配至所有处理核的调度方法。另外，分布式调度是指其中按照分级方式分组的控制核将工作组动态地分配至处理核的调度方法。另外，硬件支持是指计算设备是否包括缓冲器。
参照图13，应用了分布式调度方法的计算设备的延迟时间比应用了集中式调度方法的计算设备的延迟时间更短。也就是说，可以理解，按照分级方式分组的控制核将工作组动态地分配至九十六个处理核所需的时间比单个控制核将工作组动态地分配至九十六个处理核所需的时间更短。
另外，当应用了相同的调度方法时，包括缓冲器的计算设备的延迟时间比不包括缓冲器的计算设备的延迟时间更短。因此，可以理解，由于包括缓冲器的计算设备访问存储器的次数相对小，因此该计算设备处理工作组所需的延迟时间相对短。
参照图14，执行四种不同的应用来评价计算设备的性能。例如，PI估计应用比其它应用需要更长的时间来处理分配的工作组。小缓存副本的应用和大缓存副本的应用比其它应用更多次地访问存储器。另外，SpMV应用不定期地访问存储器，并且SpMV应用处理分配的工作组所需的时间段不均匀。
图15是示出对根据本发明构思的示例性实施例的计算设备的性能的评价结果的曲线图。
图15示出了相对于在计算设备上的四种应用的执行，计算设备的调度性能的测量结果，基于理想调度性能将所述结果标准化。
参照图15，应用了集中式调度方法的计算设备的调度性能比应用了分布式调度方法的计算设备的调度性能更低。例如，在SpMV应用的情况下，应用了集中式调度方法的计算设备的调度性能与应用了分布式调度方法的计算设备的调度性能明显不同。
另外，增加了缓冲器并且应用了分布式调度方法的计算设备的调度性能比仅应用了分布式调度方法的计算设备的调度性能更高。例如，可通过在相对多次访问存储器的小缓存副本应用和大缓存副本应用的情况下通过增加缓冲器来导致调度性能的大差异。
图16是示出根据本发明构思的示例性实施例的处理OpenCL内核的方法的流程图。
在操作1610中，控制核组320将用于执行OpenCL内核的工作组分配至低级别控制核和处理核。控制核组320的多个控制核可根据通过控制核向其分配工作组的处理核的数量按照分级方式分组。
例如，控制核组320可包括根控制核330和叶控制核组350。
如果根控制核330接收关于OpenCL内核的执行的信息，则根控制核330可将工作组分配至低级别控制核，并且可将关于分配的工作组的信息发送至低级别控制核。叶控制核组350可包括至少一个叶控制核，并且可被构造为接收关于通过高级别控制核分配的工作组的信息，并且将工作组分配至处理核。
另外，控制核组320还可包括中间控制核组340，中间控制核组340包括多个中间控制核，并且被构造为接收关于通过根控制核330分配的工作组的信息并且将工作组分配至叶控制核组350。
根据示例性实施例，如果中间控制核根据通过中间控制核向其分配工作组的处理核的数量按照分级方式分组，则在操作1610中，高级别中间控制核将工作组分配至低级别中间控制核。
在操作1620中，包括至少一个处理核的处理核组360处理通过控制核组320分配的工作组。
在操作1630中，处理核组360输出工作组的处理的结果。
图17是示出根据本发明构思的示例性实施例的处理OpenCL内核的方法的流程图。
图17示出了利用参照图3描述的计算设备300处理OpenCL内核的方法。然而，OpenCL内核处理方法可根据计算设备300的构造而变化。另外，图17所示的操作1720、1730和1740对应于图16所示的操作1610、1620和1630，因此将省略对其的详细描述。
在操作1710中，在执行OpenCL内核时，根控制核330接收通过主机处理器310产生的关于所述执行的信息。例如，根控制核330可在处理用于执行OpenCL内核的工作组的同时，接收关于新的OpenCL内核的执行的信息，并且在这种情况下，根控制核330可再分配被分配至处理核组360的工作组。这里，关于OpenCL内核的执行的信息可包括用于执行OpenCL内核的工作组的总数和关于包括在OpenCL内核中的应用的信息。
另外，根据实施例，在控制核组320将工作组分配至处理核组360之后，处理OpenCL内核的方法还可包括：主机处理器310停止OpenCL内核的执行，并且还收集和输出工作组的分配的结果。
另外，根据实施例，处理OpenCL内核的方法还可包括在缓冲器中存储关于工作组的信息。这里，关于工作组的信息可包括分配至低级别控制核和处理核的工作组的编号和ID。
在操作1750中，控制核组320根据工作组处理的输出结果再分配被分配至其它低级别控制核和其它处理核的工作组。例如，根据通过中间控制核或叶控制核的工作组处理结果，根控制核330可再分配被分配至其它中间控制核或其它叶控制核的工作组。另外，根据通过低级别中间控制核或叶控制核的工作组处理结果，中间控制核可再分配被分配至其它低级别中间控制核或其它叶控制核的工作组。另外，根据通过处理核的工作组处理的结果，叶控制核可再分配被分配至其它处理核的工作组。
如上所述，根据以上实施例中的一个或多个，由于用于将工作组分配至处理核的控制核按照分级方式分组，因此可使工作组有效地分布。
实施例中的至少一个可按照存储指令的非暂时性记录介质的形式实施，可在诸如程序模块的计算机上执行该指令。非暂时性计算机可读介质可为可由计算机访问的任何非暂时性介质，诸如易失性介质、非易失性介质、可分离介质或者不可分离介质。计算机存储介质的示例包括通过用于存储数据(诸如计算机指令、数据结构或程序模块)的任何方法或技术实施的易失性介质、非易失性介质、可分离介质和不可分离介质。
虽然已参照附图描述了一个或多个实施例，但是本领域普通技术人员应该理解，在不脱离本发明构思的精神和范围的情况下，可在其中作出各种形式和细节上的改变。

Claims (18)

1.一种多核计算设备，包括：
控制核组，其包括第一低级别控制核，该控制核组被构造为将用于执行OpenCL内核的工作组分配至第一低级别控制核和第一处理核；以及
处理核组，其包括第一处理核中的至少一个，该处理核组被构造为处理通过控制核组分配的工作组并且输出工作组的处理结果，
其中，包括在控制核组中的多个控制核根据通过控制核向其分配工作组的第一处理核的数量按照分级方式分组，
其中，控制核组包括：
根控制核，其被构造为在接收与OpenCL内核的执行相关的信息时将工作组分配至第一低级别控制核，并且将与分配的工作组相关的信息发送至第一低级别控制核；以及
叶控制核组，其包括至少一个叶控制核，并且被构造为接收与通过控制核组的高级别控制核分配的工作组相关的信息，并且将工作组分配至第一处理核。
2.根据权利要求1所述的多核计算设备，其中，控制核组被构造为根据通过第一低级别控制核和第一处理核对工作组的处理结果再分配被分配至控制核组的第二低级别控制核和第二处理核的工作组。
3.根据权利要求1所述的多核计算设备，其中，控制核组还包括中间控制核组，中间控制核组包括多个中间控制核并且被构造为接收与通过根控制核分配的工作组相关的信息，并且将工作组分配至叶控制核组，
其中，中间控制核根据通过中间控制核向其分配工作组的第一处理核的数量按照分级方式分组，并且控制核组的高级别中间控制核将工作组分配至低级别中间控制核。
4.根据权利要求1所述的多核计算设备，其中，随着OpenCL内核被执行，根控制核接收通过主机处理器产生的与OpenCL内核的执行相关的信息。
5.根据权利要求4所述的多核计算设备，其中，当用于执行OpenCL内核的工作组被处理时，根控制核在从主机处理器接收与新的OpenCL内核的执行相关的信息时再分配被分配至处理核组的工作组。
6.根据权利要求4所述的多核计算设备，其中，在控制核组将工作组分配至处理核组之后，主机处理器停止执行OpenCL内核，并且收集和输出工作组的分配的结果。
7.根据权利要求1所述的多核计算设备，其中，第一处理核当中的给定处理核包括一个或多个处理元件，并且
当处理元件的数量少于指定给所述给定处理核的工作组中的一个工作组的工作项的数量时，所述给定处理核转换OpenCL内核，并且所述处理元件中的一个处理元件使用转换后的OpenCL内核以按次序执行多个工作项。
8.根据权利要求1所述的多核计算设备，还包括被构造为存储关于工作组的信息的缓冲器。
9.根据权利要求8所述的多核计算设备，其中，关于工作组的信息包括分配至第一低级别控制核和第一处理核的工作组的编号和ID中的至少一种。
10.一种处理OpenCL内核的方法，所述方法包括以下步骤：
通过包括第一低级别控制核的控制核组将用于执行OpenCL内核的工作组分配至第一低级别控制核和第一处理核；
通过包括第一处理核中的至少一个处理核的处理核组来处理分配的工作组；以及
通过处理核组输出工作组的处理的结果，
其中，控制核组的多个控制核根据通过控制核向其分配工作组的第一处理核的数量按照分级方式分组，
其中，控制核组包括：
根控制核，其被构造为在接收与OpenCL内核的执行相关的信息时将工作组分配至第一低级别控制核，并且将与分配的工作组相关的信息发送至第一低级别控制核；以及
叶控制核组，其包括至少一个叶控制核，并且被构造为接收与通过控制核组的高级别控制核分配的工作组相关的信息，并且将工作组分配至第一处理核。
11.根据权利要求10所述的方法，还包括，根据工作组的处理的结果，通过控制核组再分配被分配至控制核组的第二低级别控制核和第二处理核的工作组。
12.根据权利要求10所述的方法，其中，控制核组还包括中间控制核组，中间控制核组包括多个中间控制核，并且被构造为接收与通过根控制核分配的工作组相关的信息并且将工作组分配至叶控制核组，
其中，如果中间控制核根据通过中间控制核向其分配工作组的第一处理核的数量按照分级方式分组，则工作组的分配步骤包括通过控制核组的高级别中间控制核将工作组分配至控制核组的低级别中间控制核。
13.根据权利要求10所述的方法，还包括：当执行OpenCL内核时，通过根控制核接收通过主机处理器产生的与OpenCL内核的执行相关的信息。
14.根据权利要求13所述的方法，还包括：当处理用于执行OpenCL内核的工作组时，一旦从主机处理器接收到与新的OpenCL内核的执行相关的信息，就通过根控制核再分配被分配至处理核组的工作组。
15.根据权利要求10所述的方法，其中，第一处理核当中的给定处理核包括一个或多个处理元件，并且当处理元件的数量少于指定给所述给定处理核的工作组中的一个工作组的工作项的数量时，工作组的处理包括所述给定处理核转换OpenCL内核，并且所述处理元件中的一个处理元件使用转换后的OpenCL内核以按次序执行多个工作项。
16.根据权利要求10所述的方法，还包括：将关于工作组的信息存储在缓冲器中。
17.根据权利要求16所述的方法，其中，关于工作组的信息包括：被分配至第一低级别控制核和第一处理核的工作组的编号和ID中的至少一种。
18.一种非暂时性计算机可读记录介质，其上记录有用于在计算机上执行根据权利要求10所述的方法的指令。
CN201711216401.0A 2016-12-27 2017-11-28 处理OpenCL内核的方法及用于该方法的计算设备 Active CN108241508B (zh)

Applications Claiming Priority (2)

Application Number Priority Date Filing Date Title
KR1020160180133A KR102592330B1 (ko) 2016-12-27 2016-12-27 OpenＣＬ 커널을 처리하는 방법과 이를 수행하는 컴퓨팅 장치
KR10-2016-0180133 2016-12-27

Publications (2)

Publication Number Publication Date
CN108241508A CN108241508A (zh) 2018-07-03
CN108241508B true CN108241508B (zh) 2023-06-13

Family

ID=60473303

Family Applications (1)

Application Number Title Priority Date Filing Date
CN201711216401.0A Active CN108241508B (zh) 2016-12-27 2017-11-28 处理OpenCL内核的方法及用于该方法的计算设备

Country Status (5)

Country Link
US (1) US10503557B2 (zh)
EP (1) EP3343370A1 (zh)
JP (1) JP6951962B2 (zh)
KR (1) KR102592330B1 (zh)
CN (1) CN108241508B (zh)

Families Citing this family (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN111490946B (zh) * 2019-01-28 2023-08-11 阿里巴巴集团控股有限公司 基于OpenCL框架的FPGA连接实现方法及装置

Citations (4)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN102099788A (zh) * 2008-06-06 2011-06-15 苹果公司 用于在多处理器上进行数据并行计算的应用编程接口
CN103608777A (zh) * 2011-06-20 2014-02-26 高通股份有限公司 图形处理单元中的存储器共享
CN104011679A (zh) * 2011-12-01 2014-08-27 超威半导体公司 异构并行处理平台的软件库
CN104834630A (zh) * 2014-02-10 2015-08-12 瑞萨电子株式会社 运算控制装置、运算控制方法、存储有程序的非瞬时计算机可读介质以及OpenCL设备

Family Cites Families (19)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JPS595755B2 (ja) 1977-08-02 1984-02-07 日本軽金属株式会社 太陽熱利用暖房方法
US7418470B2 (en) * 2000-06-26 2008-08-26 Massively Parallel Technologies, Inc. Parallel processing systems and method
US7673011B2 (en) 2007-08-10 2010-03-02 International Business Machines Corporation Configuring compute nodes of a parallel computer in an operational group into a plurality of independent non-overlapping collective networks
US8255345B2 (en) 2009-05-15 2012-08-28 The Aerospace Corporation Systems and methods for parallel processing with infeasibility checking mechanism
US9354944B2 (en) * 2009-07-27 2016-05-31 Advanced Micro Devices, Inc. Mapping processing logic having data-parallel threads across processors
US8782469B2 (en) 2009-09-01 2014-07-15 Hitachi, Ltd. Request processing system provided with multi-core processor
KR101640848B1 (ko) 2009-12-28 2016-07-29 삼성전자주식회사 멀티코어 시스템 상에서 단위 작업을 할당하는 방법 및 그 장치
KR101613971B1 (ko) 2009-12-30 2016-04-21 삼성전자주식회사 프로그램 코드의 변환 방법
JP2012094072A (ja) 2010-10-28 2012-05-17 Toyota Motor Corp 情報処理装置
US8683243B2 (en) * 2011-03-11 2014-03-25 Intel Corporation Dynamic core selection for heterogeneous multi-core systems
EP2711839A4 (en) * 2011-05-19 2014-12-03 Nec Corp PARALLEL PROCESSING DEVICE, PARALLEL PROCESSING METHOD, OPTIMIZATION DEVICE, OPTIMIZATION METHOD AND COMPUTER PROGRAM
US20120331278A1 (en) 2011-06-23 2012-12-27 Mauricio Breternitz Branch removal by data shuffling
JP5238876B2 (ja) * 2011-12-27 2013-07-17 株式会社東芝 情報処理装置及び情報処理方法
KR101284195B1 (ko) * 2012-01-09 2013-07-10 서울대학교산학협력단 개방형 범용 병렬 컴퓨팅 프레임워크 동적 작업 분배 장치
KR20130093995A (ko) 2012-02-15 2013-08-23 한국전자통신연구원 계층적 멀티코어 프로세서의 성능 최적화 방법 및 이를 수행하는 멀티코어 프로세서 시스템
KR20140125893A (ko) 2013-01-28 2014-10-30 한국과학기술원 가상화된 매니코어 서버의 작업분배 시스템과 그 방법 및 기록매체
KR102062208B1 (ko) * 2013-05-03 2020-02-11 삼성전자주식회사 멀티스레드 프로그램 코드의 변환 장치 및 방법
CN104035751B (zh) * 2014-06-20 2016-10-12 深圳市腾讯计算机系统有限公司 基于多图形处理器的数据并行处理方法及装置
US10318261B2 (en) * 2014-11-24 2019-06-11 Mentor Graphics Corporation Execution of complex recursive algorithms

Patent Citations (4)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN102099788A (zh) * 2008-06-06 2011-06-15 苹果公司 用于在多处理器上进行数据并行计算的应用编程接口
CN103608777A (zh) * 2011-06-20 2014-02-26 高通股份有限公司 图形处理单元中的存储器共享
CN104011679A (zh) * 2011-12-01 2014-08-27 超威半导体公司 异构并行处理平台的软件库
CN104834630A (zh) * 2014-02-10 2015-08-12 瑞萨电子株式会社 运算控制装置、运算控制方法、存储有程序的非瞬时计算机可读介质以及OpenCL设备

Also Published As

Publication number Publication date
US10503557B2 (en) 2019-12-10
KR102592330B1 (ko) 2023-10-20
KR20180076051A (ko) 2018-07-05
CN108241508A (zh) 2018-07-03
EP3343370A1 (en) 2018-07-04
US20180181443A1 (en) 2018-06-28
JP6951962B2 (ja) 2021-10-20
JP2018106709A (ja) 2018-07-05

Similar Documents

Publication Publication Date Title
US9354944B2 (en) Mapping processing logic having data-parallel threads across processors
US9081618B2 (en) Method and apparatus for the scheduling of computing tasks
US9483319B2 (en) Job scheduling apparatus and method therefor
JP2020537784A (ja) ニューラルネットワークアクセラレーションのための機械学習ランタイムライブラリ
US20150199214A1 (en) System for distributed processing of stream data and method thereof
US9032411B2 (en) Logical extended map to demonstrate core activity including L2 and L3 cache hit and miss ratio
US20150227586A1 (en) Methods and Systems for Dynamically Allocating Resources and Tasks Among Database Work Agents in an SMP Environment
US9501285B2 (en) Register allocation to threads
KR102338849B1 (ko) 실시간 운영 체제에서 스택 메모리 관리를 제공하는 방법 및 시스템
KR20110075297A (ko) 병렬도를 고려한 병렬 처리 장치 및 방법
US20170344398A1 (en) Accelerator control device, accelerator control method, and program storage medium
US8214585B2 (en) Enabling parallel access volumes in virtual machine environments
CN103425534A (zh) 在许多应用之间共享的图形处理单元
US9792209B2 (en) Method and apparatus for cache memory data processing
Georgiou et al. Topology-aware job mapping
Chisholm et al. Supporting mode changes while providing hardware isolation in mixed-criticality multicore systems
US10360079B2 (en) Architecture and services supporting reconfigurable synchronization in a multiprocessing system
CN108241508B (zh) 处理OpenCL内核的方法及用于该方法的计算设备
US10437736B2 (en) Single instruction multiple data page table walk scheduling at input output memory management unit
CN116010093A (zh) 数据处理方法、装置、计算机设备和可读存储介质
KR101755154B1 (ko) 이종 연산 처리 장치에 대한 동적 작업 할당 방법 및 장치
WO2014027444A1 (ja) スケジューリング装置、及び、スケジューリング方法
US9547522B2 (en) Method and system for reconfigurable virtual single processor programming model
US11755300B1 (en) Systems and methods for array structure processing
US9251100B2 (en) Bitmap locking using a nodal lock

Legal Events

Date Code Title Description
PB01 Publication
PB01 Publication
SE01 Entry into force of request for substantive examination
SE01 Entry into force of request for substantive examination
GR01 Patent grant
GR01 Patent grant