JP6951962B2 - How to handle the OpenCL kernel and the computing equipment that carries it out - Google Patents

How to handle the OpenCL kernel and the computing equipment that carries it out Download PDF

Info

Publication number
JP6951962B2
JP6951962B2 JP2017238434A JP2017238434A JP6951962B2 JP 6951962 B2 JP6951962 B2 JP 6951962B2 JP 2017238434 A JP2017238434 A JP 2017238434A JP 2017238434 A JP2017238434 A JP 2017238434A JP 6951962 B2 JP6951962 B2 JP 6951962B2
Authority
JP
Japan
Prior art keywords
group
processing
core
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
JP2017238434A
Other languages
Japanese (ja)
Other versions
JP2018106709A5 (en
JP2018106709A (en
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
Seoul National University 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, Seoul National University R&DB Foundation filed Critical Samsung Electronics Co Ltd
Publication of JP2018106709A publication Critical patent/JP2018106709A/en
Publication of JP2018106709A5 publication Critical patent/JP2018106709A5/ja
Application granted granted Critical
Publication of JP6951962B2 publication Critical patent/JP6951962B2/en
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Classifications

    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00Arrangements for program control, e.g. control units
    • G06F9/06Arrangements 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/30Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/38Concurrent instruction execution, e.g. pipeline or look ahead
    • G06F9/3836Instruction issuing, e.g. dynamic instruction scheduling or out of order instruction execution
    • G06F9/3851Instruction issuing, e.g. dynamic instruction scheduling or out of order instruction execution from multiple instruction streams, e.g. multistreaming
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00Arrangements for program control, e.g. control units
    • G06F9/06Arrangements 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/46Multiprogramming arrangements
    • G06F9/50Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5005Allocation of resources, e.g. of the central processing unit [CPU] to service a request
    • G06F9/5027Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resource being a machine, e.g. CPUs, Servers, Terminals
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00Arrangements for program control, e.g. control units
    • G06F9/06Arrangements 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/46Multiprogramming arrangements
    • G06F9/50Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5005Allocation of resources, e.g. of the central processing unit [CPU] to service a request
    • G06F9/5027Allocation 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/5044Allocation 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
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00Arrangements for program control, e.g. control units
    • G06F9/06Arrangements 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/30Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/38Concurrent instruction execution, e.g. pipeline or look ahead
    • G06F9/3885Concurrent instruction execution, e.g. pipeline or look ahead using a plurality of independent parallel functional units
    • G06F9/3887Concurrent 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]
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00Arrangements for program control, e.g. control units
    • G06F9/06Arrangements 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/46Multiprogramming arrangements
    • G06F9/50Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5061Partitioning or combining of resources
    • G06F9/5066Algorithms for mapping a plurality of inter-dependent sub-tasks onto a plurality of physical CPUs
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06TIMAGE DATA PROCESSING OR GENERATION, IN GENERAL
    • G06T1/00General purpose image data processing
    • G06T1/20Processor 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)

Description

本開示は、OpenCLカーネルを処理する方法、及びそれを遂行するコンピューティング装置に関する。 The present disclosure relates to a method of processing an OpenCL kernel and a computing device carrying it out.

コンピューティング装置は、アプリケーションの性能条件を満足するために、単一集積回路に、さまざまなコアまたはプロセッサを統合する形態で発展している。例えば、マルチコアプロセッサ(multi core processor)は、演算機能を有した2個以上のコアが、1つのプロセッサに集積されたものを意味する。また、さらに多数(一般的に、16個以上)のコアが1つのプロセッサに集積されたメニーコアプロセッサ(many core processor)も開発されている。マルチコアプロセッサとメニーコアプロセッサは、タブレットPC(personal computer)、携帯電話、PDA(personal digital assistant)、ラップトップ、メディアプレーヤ、GPS(global position system)装置、電子書籍端末機、MP3プレーヤ、デジタルカメラを含んだ携帯機器やTV(television)に搭載されるマルチメディアチップを含んだ埋め込み装置に搭載される。 Computing equipment has evolved in the form of integrating various cores or processors into a single integrated circuit to meet the performance requirements of an application. For example, a multi-core processor means that two or more cores having an arithmetic function are integrated in one processor. In addition, many core processors (many core processors) in which a large number (generally 16 or more) of cores are integrated in one processor have also been developed. Multi-core processors and many-core processors include tablet PCs (personal computers), mobile phones, personal digital assistants (PDAs), laptops, media players, GPS (global position system) devices, electronic book terminals, MP3 players, and digital cameras. However, it is installed in embedded devices including multimedia chips installed in mobile devices and TVs (televisions).

一方、OpenCL(open computing language)は、複数のプラットフォームで動作するプログラムを作成可能にする。言い替えれば、OpenCLは、GPU(graphics processing unit)だけではなく、汎用マルチコアCPU(central processing unit)、FPGA(field-programmable gate array)などの複数のプラットフォームにおいて、プログラム動作を可能にする開放型汎用並列コンピューティングフレームワークであり、グラフィック処理装置(GPU)の力量を、グラフィック処理以外の領域(例えば、GPGPU(general-purpose computing on graphics processing unit))に拡張させる。それにより、複数個のコアが含まれたコンピューティング装置がOpenCLプログラムを効率的に処理するための多様な試みが進められている。 On the other hand, OpenCL (open computing language) makes it possible to create programs that run on multiple platforms. In other words, OpenCL is an open general-purpose parallel that enables program operation not only on GPU (graphics processing unit) but also on multiple platforms such as general-purpose multi-core CPU (central processing unit) and FPGA (field-programmable gate array). It is a computing framework that extends the capabilities of a graphics processing unit (GPU) to areas other than graphics processing (eg, GPGPU (general-purpose computing on graphics processing unit)). As a result, various attempts are being made for a computing device containing a plurality of cores to efficiently process an OpenCL program.

本発明が解決しようとする課題は、階層的に分類されたコントロールコアグループが、OpenCLカーネル(kernel)を実行するための作業グループをプロセシングコアグループに割り当てる、OpenCLカーネルを処理する方法、及びそれを遂行するコンピューティング装置を提供することである。本実施形態がなそうとする技術的課題は、前述のような技術的課題に限定されるものではなく、以下の実施形態から、他の技術的課題が類推されもする。 The problem to be solved by the present invention is how a hierarchically classified control core group allocates a working group for executing an OpenCL kernel (kernel) to a processing core group, how to process an OpenCL kernel, and how to handle it. It is to provide a computing device to carry out. The technical problem to be solved by this embodiment is not limited to the above-mentioned technical problem, and other technical problems can be inferred from the following embodiments.

一側面によれば、OpenCLカーネルを実行するための作業グループを、下位階層のコントロールコア(control core)及びプロセシングコア(processing core)に割り当てるコントロールコアグループ;及び少なくとも1つの前記プロセシングコアを含み、前記コントロールコアグループによって割り当てられた前記作業グループを処理し、前記作業グループの処理結果を出力する、プロセシングコアグループ;を含み、前記コントロールコアグループに含まれた複数個のコントロールコアは、前記コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数により、階層的に分類される。 According to one aspect, a control core group that assigns a working group for running an OpenCL kernel to a lower level control core and a processing core; and at least one said processing core, said. A plurality of control cores included in the control core group include a processing core group that processes the work group assigned by the control core group and outputs the processing result of the work group; The work groups are classified hierarchically according to the number of the processing cores to which they are assigned.

他の側面によれば、OpenCLカーネルを処理する方法は、コントロールコアグループにより、OpenCLカーネルを実行するための作業グループを、下位階層のコントロールコア及びプロセシングコアに割り当てる段階;少なくとも1つの前記プロセシングコアが含まれたプロセシングコアグループにより、前記割り当てられた前記作業グループを処理する段階;及び前記プロセシングコアグループにより、前記作業グループの処理結果を出力する段階;を含み、前記コントロールコアグループに含まれた複数個のコントロールコアは、前記コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数により、階層的に分類される。 According to another aspect, the method of processing the OpenCL kernel is the step of assigning a working group for executing the OpenCL kernel to the lower level control core and processing core by the control core group; at least one said processing core. A plurality of steps included in the control core group, including a step of processing the assigned work group by the included processing core group; and a step of outputting the processing result of the work group by the processing core group. The control cores are hierarchically classified according to the number of the processing cores to which the work group is assigned by the control cores.

さらに他の側面によれば、前記方法をコンピュータで実行させるためのプログラムを記録したコンピュータ読み取り可能な記録媒体を提供する。 According to yet another aspect, there is provided a computer-readable recording medium on which a program for executing the method on a computer is recorded.

本発明によれば、プロセシングコアに作業グループを割り当てるコントロールコアが階層的に分類されているために、作業グループを効率的に分配することができる。 According to the present invention, since the control cores that assign work groups to the processing cores are hierarchically classified, the work groups can be efficiently distributed.

OpenCLプラットフォームについて説明するための図である。It is a figure for demonstrating the OpenCL platform. 一実施形態による、コンピューティング装置の構成を示したブロック図である。It is a block diagram which showed the structure of the computing apparatus by one Embodiment. 他の一実施形態による、コンピューティング装置の構成を示したブロック図である。It is a block diagram which showed the structure of the computing apparatus by another embodiment. 一実施形態によって、コントロールコアによって作業グループが割り当てられたプロセシングコアの数について説明するための図である。It is a figure for demonstrating the number of processing cores to which the work group was assigned by the control core by one Embodiment. 一実施形態による、OpenCLカーネルを実行するための作業グループを動的に割り当てる方法について説明するための図である。It is a figure for demonstrating the method of dynamically allocating the work group for executing an OpenCL kernel according to one Embodiment. 一実施形態による、コンピューティング装置の構造について説明するための図である。It is a figure for demonstrating the structure of the computing apparatus by one Embodiment. 他の一実施形態による、最下位コントロールコア及びプロセシングコアについて説明するための図である。It is a figure for demonstrating the lowest control core and the processing core by another embodiment. 一実施形態による、プロセシングコアで実行されるOpenCLカーネルコードである。An OpenCL kernel code that is executed on the processing core according to one embodiment. 他の一実施形態によるコンピューティング装置について説明するための図である。It is a figure for demonstrating the computing apparatus by another embodiment. 一実施形態によるバッファについて説明するための図である。It is a figure for demonstrating the buffer by one Embodiment. 一実施形態による、バッファが追加された場合、プロセシングコアで実行されるOpenCLカーネルコードである。According to one embodiment, OpenCL kernel code that is executed on the processing core when a buffer is added. 他の一実施形態による、バッファが追加された場合、OpenCL作業グループを実行するOpenCLカーネルコードである。According to another embodiment, OpenCL kernel code that executes an OpenCL working group when a buffer is added. 一実施形態によるコンピューティング装置の性能を検証するための実験環境を示した図である。It is a figure which showed the experimental environment for verifying the performance of the computing apparatus by one Embodiment. 一実施形態によるコンピューティング装置の性能を検証するための実験環境を示した図である。It is a figure which showed the experimental environment for verifying the performance of the computing apparatus by one Embodiment. 一実施形態によるコンピューティング装置の性能の検証結果について説明するための図である。It is a figure for demonstrating the verification result of the performance of the computing apparatus by one Embodiment. 一実施形態により、OpenCLカーネルを処理する方法を図示したフローチャートである。It is a flowchart which illustrated the method of processing the OpenCL kernel by one Embodiment. 一実施形態により、OpenCLカーネルを処理する方法を図示した詳細フローチャートである。It is a detailed flowchart which illustrated the method of processing the OpenCL kernel by one Embodiment.

本実施形態で使用される用語は、本実施形態での機能を考慮しながら、可能な限り、現在汎用される一般的な用語を選択したが、それは、当技術分野の当業者の意図、または判例、新技術の出現などによっても異なる。また、特定の場合、任意に選定された用語もあり、その場合、当該実施形態の説明部分において、詳細にその意味を記載する。従って、本実施形態で使用される用語は、単純な用語の名称ではない、その用語が有する意味と、本実施形態の全般にわたった内容とを基に定義されなければならない。 As for the terms used in the present embodiment, the general terms currently commonly used are selected as much as possible in consideration of the functions in the present embodiment, but it is the intention of a person skilled in the art, or it is the intention of a person skilled in the art. It also depends on precedents and the emergence of new technologies. In addition, in a specific case, there is a term arbitrarily selected, and in that case, the meaning is described in detail in the explanation part of the embodiment. Therefore, the term used in the present embodiment must be defined based on the meaning of the term and the general content of the present embodiment, which is not a simple name of the term.

実施形態に係わる説明において、ある部分が他の部分と結合されているとするとき、それは、直接に結合されている場合だけではなく、その中間に、他の構成要素を挟んで電気的に結合されている場合も含む。また、ある部分がある構成要素を含むとするとき、それは、特別に反対となる記載がない限り、他の構成要素を除くものではなく、他の構成要素をさらに含むということを意味する。また、実施形態に記載された「…部」、「…モジュール」の用語は、少なくとも1つの機能や動作を処理する単位を意味し、それは、ハードウェアまたはソフトウェアによっても具現され、ハードウェアとソフトウェアとの結合によっても具現される。 In the description of the embodiment, when one part is connected to another part, it is not only when it is directly connected, but also electrically connected with another component in the middle. Including the case where it is. Also, when a part contains a component, it does not exclude the other component, but further includes the other component, unless otherwise stated to be the opposite. Further, the terms "... part" and "... module" described in the embodiment mean a unit for processing at least one function or operation, which is also embodied by hardware or software, and hardware and software. It is also embodied by the combination with.

本実施形態で使用される「構成される」または「含む」というような用語は、必ずしも、明細書上に記載されたさまざまな構成要素、またはさまざまな段階をいずれも含むと解釈されるものではなく、そのうち一部構成要素または一部段階が含まれないこともあり、または、追加的な構成要素または段階をさらに含んでもよいと解釈されなければならない。 The terms "constituent" or "contains" as used in this embodiment are not necessarily construed to include any of the various components or stages described herein. It must be construed that some of the components or stages may not be included, or that additional components or stages may be included.

下記実施形態に係わる説明は、権利範囲を限定すると解釈されるものではなく、当該技術分野の当業者が容易に類推することができることは、実施形態の権利範囲に属すると解釈されなければならないのである。以下、添付された図面を参照しつつ、ただ例示のための実施形態について詳細に説明する。 The following description of the embodiment is not construed as limiting the scope of rights, and what can be easily inferred by a person skilled in the art of the art must be construed as belonging to the scope of rights of the embodiment. be. Hereinafter, embodiments for illustration purposes will be described in detail with reference to the accompanying drawings.

図1は、OpenCL(open computing language)プラットフォームについて説明するための図である。 FIG. 1 is a diagram for explaining an OpenCL (open computing language) platform.

OpenCLは、互いに異なる形態のコンピューティングデバイスが同時に動作するように具現することができるプログラミング言語中の一つである。従って、OpenCLによって作成されたプログラムは、複数のプラットフォームで同時に動作することができる。 OpenCL is one of the programming languages that can be embodied so that different forms of computing devices operate at the same time. Therefore, programs created by OpenCL can run on multiple platforms at the same time.

図1を参照すれば、OpenCLプラットフォーム100は、ホストプロセッサ(host processor)110と、少なくとも1つのコンピューティングデバイス(CD:computing device)120ないし140を含んでもよい。 With reference to FIG. 1, the OpenCL platform 100 may include a host processor 110 and at least one computing device (CD) 120-140.

ホストプロセッサ110は、ホストプログラムが実行されるプロセッサであり、カーネル(kernel)を実行するための作業空間(work space)を定義することができる。ここで、該作業空間はまた、作業グループ(work group)に分割され、該作業グループは、複数個の作業アイテム(work item)を含み得る。ここで、該作業アイテムは、作業空間の各ポイントに該当し、作業の最小単位になる。例えば、ホストプロセッサ110は、汎用CPU(central processing unit)でもあるが、それに限定されるものではない。 The host processor 110 is a processor on which a host program is executed, and a workspace for executing a kernel can be defined. Here, the work space is also divided into work groups, which may include a plurality of work items. Here, the work item corresponds to each point in the work space and becomes the minimum unit of work. For example, the host processor 110 is also a general-purpose CPU (central processing unit), but is not limited thereto.

コンピューティングデバイス120ないし140は、少なくとも1つのコンピュートユニット(CU:compute unit)121ないし122を含み得る。また、コンピュートユニット121ないし122は、少なくとも1つのプロセシングエレメント(PE:processing element)10ないし40を含み得る。ここで、コンピュートユニット121ないし122は、作業グループを処理する単位にもなり、プロセシングエレメント10ないし40は、作業アイテムを処理する単位にもなる。例えば、コンピュートユニット121ないし122は、ホストプロセッサ110から受信された1つの作業グループを処理することができ、プロセシングエレメント10ないし40は、作業グループに含まれた作業アイテムそれぞれを処理することができる。従って、該作業グループに含まれた作業アイテムは、プロセシングエレメント10ないし40によって並列的に処理される。一方、コンピューティングデバイス120ないし140は、メニーコアプロセッサ、マルチコアプロセッサでもあるが、少なくとも1つのコアを有するプロセッサであるならば、それらに限定されるものではない。 The computing devices 120-140 may include at least one compute unit (CU) 121-122. Further, the compute units 121 to 122 may include at least one processing element (PE) 10 to 40. Here, the compute units 121 to 122 also serve as units for processing work groups, and the processing elements 10 to 40 also serve as units for processing work items. For example, compute units 121 to 122 can process one work group received from the host processor 110, and processing elements 10 to 40 can process each work item contained in the work group. Therefore, the work items included in the work group are processed in parallel by the processing elements 10 to 40. On the other hand, the computing devices 120 to 140 are also many-core processors and multi-core processors, but are not limited to those as long as they are processors having at least one core.

一方、OpenCLプログラムは、1つのホストプログラムと、少なくとも1つのカーネルとを含み得る。該ホストプログラムは、ホストプロセッサ110によって実行され、ホストプロセッサ110から、コマンド(command)を介して、コンピューティングデバイス120ないし140に計算を指示したり、コンピューティングデバイス120ないし140に装着されたメモリを管理したりすることができる。ここで、該カーネルは、コンピューティングデバイス120ないし140で実行されるプログラムを意味し、OpenCLカーネルまたはカーネルコードによっても表現される。 An OpenCL program, on the other hand, may include one host program and at least one kernel. The host program is executed by the host processor 110, and the host processor 110 instructs the computing devices 120 to 140 to perform calculations or uses the memory installed in the computing devices 120 to 140 via a command. Can be managed. Here, the kernel means a program executed by computing devices 120 to 140, and is also represented by an OpenCL kernel or kernel code.

もし複数個のコアが含まれたコンピューティング装置においてカーネルが実行される場合、該コンピューティング装置は、コンピューティングデバイスと対応し、各コアは、コンピュートユニット121ないし122と対応する。言い替えれば、該コンピューティング装置のコアは、作業グループを割り当てられて処理することができ、コアに含まれたプロセシングエレメント10ないし40は、作業グループに含まれた各作業アイテムを処理することができる。 If the kernel is run on a computing device that contains multiple cores, the computing device corresponds to the computing device, and each core corresponds to a compute unit 121-122. In other words, the core of the computing device can be assigned a work group to process, and the processing elements 10 to 40 contained in the core can process each work item contained in the work group. ..

一方、コンピューティング装置100は、ホストプロセッサ110から作業グループを受信することができる。 On the other hand, the computing device 100 can receive a working group from the host processor 110.

具体的には、ホストプロセッサ110は、各コアによって処理される作業グループの処理結果により、作業グループを各コアに動的に割り当てることができる。かような場合、ホストプロセッサ110は、作業グループを動的に割り当てる作業分配マネージャを含んでもよい。 Specifically, the host processor 110 can dynamically assign a work group to each core according to the processing result of the work group processed by each core. In such cases, the host processor 110 may include a work distribution manager that dynamically allocates work groups.

しかし、1つの作業分配マネージャが、複数個のコアに作業グループを動的に割り当てる場合、コアの数により、作業分配マネージャが遂行しなければならない分配作業の量が増加し、全体コンピューティング装置の性能が低下してしまう。また、作業分配マネージャが、複数個のコアに作業グループを静的に割り当てる場合、各作業グループの処理時間が異なり、全体コンピューティング装置の作業効率が低下してしまう。併せて、ホストプロセッサ110に作業分配マネージャが含まれていると、ホストプロセッサ110の負担が増大してしまう。 However, if one work distribution manager dynamically assigns work groups to multiple cores, the number of cores increases the amount of distribution work that the work distribution manager must perform, and the overall computing device. Performance will be reduced. Further, when the work distribution manager statically assigns work groups to a plurality of cores, the processing time of each work group is different, and the work efficiency of the entire computing device is lowered. At the same time, if the host processor 110 includes a work distribution manager, the load on the host processor 110 increases.

一方、コンピューティング装置に含まれた1つのコアが作業分配マネージャである場合、当該コアは、コンピュートユニット121ないし122に活用されない。従って、複数個のコアが含まれたコンピューティング装置においてOpenCLカーネルが実行される場合に作業グループを効率的に割り当てることができる方法が要求される。 On the other hand, when one core included in the computing device is a work distribution manager, the core is not utilized by the computing units 121 to 122. Therefore, there is a need for a method that can efficiently allocate work groups when the OpenCL kernel is executed in a computing device that includes a plurality of cores.

図2は、一実施形態による、コンピューティング装置の構成を示したブロック図である。 FIG. 2 is a block diagram showing a configuration of a computing device according to an embodiment.

一実施形態によるコンピューティング装置200は、複数個のコアを含み、複数個のコアは、コントロールコアグループ210とプロセシングコアグループ220とに分類される。図2に図示された構成要素以外に、他の汎用的な構成要素がさらに含まれてもよいということは、関連技術分野で当業者であるならば、理解することができるであろう。 The computing device 200 according to one embodiment includes a plurality of cores, and the plurality of cores are classified into a control core group 210 and a processing core group 220. Those skilled in the art will appreciate that other general purpose components may be further included in addition to the components illustrated in FIG.

一方、図1の構成と比較すれば、図2のコンピューティング装置200は、図1のコンピューティングデバイス120ないし140と対応する。 On the other hand, as compared with the configuration of FIG. 1, the computing device 200 of FIG. 2 corresponds to the computing devices 120 to 140 of FIG.

コントロールコアグループ210は、OpenCLカーネルを実行するための作業グループを、下位階層のコントロールコア及びプロセシングコアに割り当てる、コントロールコアの集合を意味する。ここで、該コントロールコアの仕様は、プロセシングコアの仕様とも同一であるが、作業グループを割り当てる機能を遂行することができる仕様であるならば、それに限定されるものではない。 The control core group 210 means a set of control cores that allocate working groups for executing the OpenCL kernel to the control cores and processing cores in the lower hierarchy. Here, the specifications of the control core are the same as the specifications of the processing core, but the specifications are not limited to those that can perform the function of assigning a work group.

プロセシングコアグループ220は、コントロールコアグループ210によって割り当てられた作業グループを処理し、作業グループの処理結果を出力する、プロセシングコアの集合を意味する。ここで、各プロセシングコアは、少なくとも1つのプロセシングエレメントを含み得る。例えば、該プロセシングコアは、メニーコアプロセッサの各コアとも対応するが、それに限定されるものではない。 The processing core group 220 means a set of processing cores that processes the work group assigned by the control core group 210 and outputs the processing result of the work group. Here, each processing core may include at least one processing element. For example, the processing core corresponds to, but is not limited to, each core of the manycore processor.

ここで、コントロールコアグループ210に含まれた複数個のコントロールコアはまた、コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類される。コントロールコアグループ210が階層的に分類される例示は、図3を介して説明する。 Here, the plurality of control cores included in the control core group 210 are also hierarchically classified according to the number of processing cores to which the work group is assigned by the control cores. An example in which the control core group 210 is hierarchically classified will be described with reference to FIG.

図3は、他の一実施形態による、コンピューティング装置の構成を示したブロック図である。図3を参照すれば、コンピューティング装置300は、コントロールコアグループ320とプロセシングコアグループ360とを含み得る。ここで、コントロールコアグループ320は、1つの最上位(root;ルート)コントロールコア330と最下位(leaf;リーフ)コントロールコアグループ350とを含み得る。また、コンピューティング装置300は、コントロールコアグループ320の階層構造により、中間コントロールコアグループ340をさらに含んでもよい。一方、図3のプロセシングコアグループ360は、図2のプロセシングコアグループ220と対応するので、詳細な説明は省略する。 FIG. 3 is a block diagram showing a configuration of a computing device according to another embodiment. With reference to FIG. 3, the computing device 300 may include a control core group 320 and a processing core group 360. Here, the control core group 320 may include one top (root) control core 330 and one bottom (leaf) control core group 350. Further, the computing device 300 may further include an intermediate control core group 340 due to the hierarchical structure of the control core group 320. On the other hand, the processing core group 360 of FIG. 3 corresponds to the processing core group 220 of FIG. 2, and therefore detailed description thereof will be omitted.

最上位コントロールコア330は、OpenCLカーネルの実行情報が受信されることにより、作業グループを、下位階層のコントロールコアに割り当てることができる。図3としては、最上位コントロールコア330の下位階層のコントロールコアを、中間コントロールコアグループ340で表示したが、最上位コントロールコア330の下位階層のコントロールコアは、コントロールコアグループ320の階層構造により、中間コントロールコアグループ340及び最下位コントロールコアグループ350のうち少なくとも一つにもなるということは、当該技術分野の通常の技術者であるならば、理解することができるであろう。 The uppermost control core 330 can assign a work group to a lower layer control core by receiving the execution information of the OpenCL kernel. In FIG. 3, the control cores in the lower hierarchy of the uppermost control core 330 are displayed in the intermediate control core group 340, but the control cores in the lower hierarchy of the uppermost control core 330 are based on the hierarchical structure of the control core group 320. It will be understood by ordinary engineers in the art that it will also be at least one of the intermediate control core group 340 and the lowest control core group 350.

一方、最上位コントロールコア330は、ホストプロセッサ310から、OpenCLカーネルの実行情報を受信することができる。具体的には、ホストプロセッサ310は、OpenCLカーネルが実行されることにより、OpenCLカーネルの実行情報を生成し、最上位コントロールコア330に送信することができる。ここで、OpenCLカーネルの実行情報は、OpenCLカーネルが実行されるための全体作業グループの数及びOpenCLカーネルが含まれたアプリケーション情報のうち少なくとも一つを含み得る。 On the other hand, the top-level control core 330 can receive the execution information of the OpenCL kernel from the host processor 310. Specifically, the host processor 310 can generate execution information of the OpenCL kernel by executing the OpenCL kernel and send it to the uppermost control core 330. Here, the execution information of the OpenCL kernel may include at least one of the total number of working groups for the OpenCL kernel to be executed and the application information including the OpenCL kernel.

最下位コントロールコアグループ350は、上位階層のコントロールコアによって割り当てられた作業グループについての情報を受信し、プロセシングコアに作業グループを割り当てる、複数個の最下位コントロールコアを含んでもよい。図3としては、最下位コントロールコアグループ350の上位階層のコントロールコアを、中間コントロールコアグループ340と表示したが、最下位コントロールコアの上位階層のコントロールコアは、コントロールコアグループ320の階層構造により、最上位コントロールコア330及び中間コントロールコアグループ340のうち少なくとも一つにもなるということは、当該技術分野の当業者であるならば、理解することができるであろう。 The lowest control core group 350 may include a plurality of lowest control cores that receive information about work groups assigned by higher level control cores and assign work groups to processing cores. In FIG. 3, the control cores in the upper hierarchy of the lowest control core group 350 are displayed as the intermediate control core group 340, but the control cores in the upper hierarchy of the lowest control core are due to the hierarchical structure of the control core group 320. Those skilled in the art will appreciate that it will also be at least one of the top control core 330 and the intermediate control core group 340.

中間コントロールコアグループ340は、最上位コントロールコア330によって割り当てられた作業グループに係わる情報を受信し、最下位コントロールコアグループ350に作業グループを割り当てる、複数個の中間コントロールコアを含んでもよい。また、中間コントロールコアグループ340は、さまざまな階層に分類される。図3を参照すれば、中間コントロールコアもまた、中間コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類される。かような場合、上位階層の中間コントロールコア341は、下位階層の中間コントロールコア342に作業グループを割り当てることができる。 The intermediate control core group 340 may include a plurality of intermediate control cores that receive information relating to the work group assigned by the top control core 330 and assign the work group to the lowest control core group 350. Further, the intermediate control core group 340 is classified into various hierarchies. With reference to FIG. 3, intermediate control cores are also hierarchically classified according to the number of processing cores to which work groups are assigned by the intermediate control cores. In such a case, the intermediate control core 341 of the upper hierarchy can assign a work group to the intermediate control core 342 of the lower hierarchy.

図4は、一実施形態によって、コントロールコアによって作業グループが割り当てられるプロセシングコアの数について説明するための図である。 FIG. 4 is a diagram for explaining the number of processing cores to which a work group is assigned by the control cores according to one embodiment.

一実施形態によれば、コントロールコアグループ210に含まれたコントロールコアは、コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類される。 According to one embodiment, the control cores included in the control core group 210 are hierarchically classified according to the number of processing cores to which the work group is assigned by the control cores.

図4は、8x8プロセシングコアが含まれたプロセシングコアグループを示した図である。図4を参照すれば、最上位コントロールコア330によって作業グループが割り当てられるプロセシングコア410は、全体プロセシングコアグループでもある。 FIG. 4 is a diagram showing a processing core group including an 8x8 processing core. Referring to FIG. 4, the processing core 410 to which the working group is assigned by the top-level control core 330 is also the overall processing core group.

一方、中間コントロールコアは、プロセシングコアグループに含まれたプロセシングコアの一部420に作業グループを割り当てることができる。例えば、図4において、中間コントロールコアグループ340が総4個の中間コントロールコアを含むのであるならば、1つの中間コントロールコアが作業グループを割り当てることができるプロセシングコアの数は、16個でもある。一方、中間コントロールコアが作業グループを割り当てることができるプロセシングコアの数は、各中間コントロールコアごとに異なってもよいということは、当該技術分野の当業者であるならば、理解することができるであろう。 On the other hand, the intermediate control core can assign a working group to a part 420 of the processing cores included in the processing core group. For example, in FIG. 4, if the intermediate control core group 340 includes a total of 4 intermediate control cores, the number of processing cores to which one intermediate control core can be assigned a working group is also 16. On the other hand, it can be understood by those skilled in the art that the number of processing cores to which an intermediate control core can be assigned a work group may be different for each intermediate control core. There will be.

一方、最下位コントロールコアは、上位階層のコントロールコアから割り当てられた作業グループについての情報を受信し、プロセシングコアに、直接作業グループを割り当てることができる。例えば、該最下位コントロールコアは、中間コントロールコアによって作業グループが割り当てられたプロセシングコアのうち一部430に作業グループを割り当てることができる。従って、最下位コントロールコアによって作業グループが割り当てられるプロセシングコアの数は、中間コントロールコアによって作業グループが割り当てられるプロセシングコアの数より少ない。 On the other hand, the lowest control core can receive information about the work group assigned from the control core of the upper hierarchy, and can directly assign the work group to the processing core. For example, the lowest control core can assign a work group to a part 430 of the processing cores to which the work group is assigned by the intermediate control core. Therefore, the number of processing cores to which work groups are assigned by the lowest control core is less than the number of processing cores to which work groups are assigned by intermediate control cores.

また、一実施形態による最下位コントロールコアは、近接した複数個のプロセシングコアに作業グループを割り当てるために、ローカル性(locality)を有する作業グループを割り当てる場合、少ないオーバーヘッドでもって、作業グループのローカル性を提供することができる。 Further, when the lowest control core according to one embodiment allocates a work group having locality in order to allocate a work group to a plurality of adjacent processing cores, the locality of the work group is reduced with a small overhead. Can be provided.

また、一実施形態によるコンピューティング装置200は、プロセシングコアグループ220以外に、コントロールコアグループ210を追加して含んでいるために、全てのプロセシングコアが作業グループを処理することができる。また、コントロールコアグループ210が階層的に分類されているために、作業グループを動的に割り当てるのに効率的である。階層的に分類されたコントロールコアが作業グループを動的に割り当てる方法は、図5を介して説明する。 Further, since the computing device 200 according to the embodiment includes the control core group 210 in addition to the processing core group 220, all the processing cores can process the working group. Moreover, since the control core group 210 is classified hierarchically, it is efficient to dynamically assign the work group. A method in which the hierarchically classified control cores dynamically assign work groups will be described with reference to FIG.

図5は、一実施形態による、OpenCLカーネルを実行するための作業グループを動的に割り当てる方法について説明するための図である。 FIG. 5 is a diagram for explaining a method of dynamically allocating a work group for executing the OpenCL kernel according to the embodiment.

一実施形態によるコントロールコアグループは、下位階層のコントロールコア及びプロセシングコアの作業グループの処理結果により、他の下位階層のコントロールコア及び他のプロセシングコアに割り当てられた作業グループを再割り当てすることができる。一方、図5のホストプロセッサ510は、図3のホストプロセッサ310と対応するので、詳細な説明は省略する。 The control core group according to one embodiment can reassign work groups assigned to other lower layer control cores and other processing cores depending on the processing results of lower layer control cores and processing core work groups. .. On the other hand, since the host processor 510 of FIG. 5 corresponds to the host processor 310 of FIG. 3, detailed description thereof will be omitted.

図5を参照すれば、最下位コントロールコア550は、プロセシングコア560に割り当てられた作業グループの処理が完了すると、他の最下位プロセシングコアに割り当てられた作業グループをプロセシングコア560にさらに割り当てることができる。また、最下位コントロールコア550は、中間コントロールコア540に作業グループの追加要請を送信することができる。ここで、最下位コントロールコア550は、プロセシングコア560から作業グループの追加要請を受信するか、あるいはプロセシングコア560の作業グループの進行状況を周期的に確認することにより、プロセシングコア560に割り当てられた作業グループの処理が完了したかいなかということを判断することができる。 Referring to FIG. 5, when the processing of the work group assigned to the processing core 560 is completed, the lowest control core 550 may further allocate the work group assigned to the other lowest processing core to the processing core 560. can. Further, the lowest control core 550 can send a request for adding a working group to the intermediate control core 540. Here, the lowest control core 550 is assigned to the processing core 560 by receiving a request for adding a work group from the processing core 560 or periodically checking the progress of the work group of the processing core 560. It can be determined whether the processing of the work group has been completed.

中間コントロールコア540は、最下位コントロールコア550に割り当てられた作業グループの処理が完了すると、他の最下位コントロールコアに割り当てられた作業グループを、最下位コントロールコア550にさらに割り当てる。また、中間コントロールコア540は、最上位コントロールコア530に作業グループの追加要請を送信することができる。 When the processing of the work group assigned to the lowest control core 550 is completed, the intermediate control core 540 further allocates the work group assigned to the other lowest control core to the lowest control core 550. In addition, the intermediate control core 540 can send a request for adding a working group to the uppermost control core 530.

最上位コントロールコア530は、中間コントロールコア540に割り当てられた作業グループの処理が完了すると、他の中間コントロールコアに割り当てられた作業グループを中間コントロールコア540にさらに割り当てる。 When the processing of the work group assigned to the intermediate control core 540 is completed, the uppermost control core 530 further allocates the work group assigned to the other intermediate control core to the intermediate control core 540.

また、最上位コントロールコア530は、OpenCLカーネルを実行するための作業グループが処理中であるとき、ホストプロセッサ510から新たなOpenCLカーネルの実行情報が受信されることにより、プロセシングコアグループに割り当てられた作業グループを再割り当てすることができる。 Further, the top-level control core 530 is assigned to the processing core group by receiving new OpenCL kernel execution information from the host processor 510 when the work group for executing the OpenCL kernel is being processed. Working groups can be reassigned.

図6は、一実施形態による、コンピューティング装置の構造について説明するための図である。図6を参照すれば、コンピューティング装置630は、ホストプロセッサ610及びメモリ620と結合される。また、コンピューティング装置630は、最上位コントロールコア631、最下位コントロールコア633及びプロセシングコア634を含むことができ、コントロールコアグループの階層構造により、中間コントロールコア632をさらに含んでもよい。 FIG. 6 is a diagram for explaining the structure of the computing device according to the embodiment. With reference to FIG. 6, the computing device 630 is coupled with the host processor 610 and the memory 620. The computing device 630 may also include a top control core 631, a bottom control core 633, and a processing core 634, and may further include an intermediate control core 632 due to the hierarchical structure of the control core groups.

図6を参照すれば、最上位コントロールコア631は、中間コントロールコア632と結合されている。従って、最上位コントロールコア631は、作業グループを割り当てるために、プロセシングコア634に直接アクセスする代わりに、中間コントロールコア632または最下位コントロールコア633に、プロセシングコア634に割り当てられた作業グループについての情報を送信することができる。中間コントロールコア632は、最下位コントロールコア633と結合されているので、プロセシングコアグループの一部に割り当てられた作業グループについての情報を、最下位コントロールコア633に送信することができる。最下位コントロールコア633の場合、プロセシングコア634と結合されており、プロセシングコア634に割り当てられた作業グループに係わる情報を直接伝えることができる。 With reference to FIG. 6, the top-level control core 631 is coupled to the intermediate control core 632. Therefore, instead of directly accessing the processing core 634 for the top-level control core 631 to assign a work group, the intermediate control core 632 or the bottom-level control core 633 is informed about the work group assigned to the processing core 634. Can be sent. Since the intermediate control core 632 is coupled to the lowest control core 633, information about the work group assigned to a part of the processing core group can be transmitted to the lowest control core 633. In the case of the lowest control core 633, it is combined with the processing core 634 and can directly convey information related to the work group assigned to the processing core 634.

一方、最上位コントロールコア631は、ホストプロセッサ610と結合され、OpenCLカーネルが実行されるとき、生成された作業グループについての情報を受信することができる。また、コントロールコアグループ及びプロセシングコアグループに含まれたコアは、メモリ620から作業グループを割り当てるのに必要な情報と、作業グループの処理に必要な情報とを受信することができる。 On the other hand, the top-level control core 631 can be coupled with the host processor 610 to receive information about the generated workgroup when the OpenCL kernel is executed. Further, the cores included in the control core group and the processing core group can receive the information necessary for allocating the work group from the memory 620 and the information necessary for the processing of the work group.

図7は、他の一実施形態による、最下位コントロールコア及びプロセシングコアについて説明するための図である。 FIG. 7 is a diagram for explaining the lowest control core and the processing core according to another embodiment.

一実施形態による1つの最下位コントロールコア730は、複数個のプロセシングコア710に作業グループを割り当てることができる。例えば、図7を参照すれば、最下位コントロールコア730は、4個のプロセシングコア710に結合され、直接作業グループを割り当て、作業グループの処理結果、または作業グループの追加要請を受信することができる。 One lowest control core 730 according to one embodiment can assign work groups to a plurality of processing cores 710. For example, referring to FIG. 7, the lowest control core 730 can be coupled to four processing cores 710, directly assign work groups, and receive work group processing results or work group addition requests. ..

また、プロセシングコア710は、作業グループを処理するコア711以外に、外部装置と通信するためのルータ712をさらに含んでもよい。プロセシングコア710は、ルータ712を介して、外部ルータ720と結合され、外部ルータ720は、最下位コントロールコア730とも結合される。従って、プロセシングコア710は、最下位コントロールコア730から、作業グループについての情報を受信することができる。図7は、プロセシングコア710に、単一コア711が含まれているように図示されているが、プロセシングコア710内のコア711の数は、それに限定されるものではない。 In addition to the core 711 that processes the work group, the processing core 710 may further include a router 712 for communicating with an external device. The processing core 710 is coupled to the external router 720 via the router 712, and the external router 720 is also coupled to the lowest control core 730. Therefore, the processing core 710 can receive information about the working group from the lowest control core 730. FIG. 7 shows that the processing core 710 includes a single core 711, but the number of cores 711 in the processing core 710 is not limited thereto.

一方、ルータ712及び外部ルータ720は、NoC(network on chip)通信方式またはバス通信方式により、通信を行うことができるが、ルータ712と外部ルータ720との通信方式は、それらに限定されるものではない。 On the other hand, the router 712 and the external router 720 can communicate by the NoC (network on chip) communication method or the bus communication method, but the communication method between the router 712 and the external router 720 is limited to them. is not it.

また、図7には、最下位コントロールコア730が4個のプロセシングコア710に作業グループを割り当てることができるように図示されているが、最下位コントロールコア730が作業グループを割り当てることができるプロセシングコア710の数は、それに限定されるものではない。 Further, in FIG. 7, the lowest control core 730 is shown so that a work group can be assigned to the four processing cores 710, but the lowest control core 730 can assign a work group to the processing core. The number of 710 is not limited to that.

図8は、一実施形態による、プロセシングコアで実行されるOpenCLカーネルコードである。図8は、行列aの要素と、行列bの要素とを加え、その結果を行列cに保存するカーネルコードである。行列aと行列bは、二次元マトリックスであるので、図8のカーネルコードを実行するためには、二次元作業グループ空間が必要である。言い替えれば、行列aと行列bとの各要素は、作業空間の各ポイントである作業アイテムにもなる。 FIG. 8 is the OpenCL kernel code executed on the processing core according to one embodiment. FIG. 8 is a kernel code that adds the elements of the matrix a and the elements of the matrix b and stores the result in the matrix c. Since the matrix a and the matrix b are two-dimensional matrices, a two-dimensional working group space is required to execute the kernel code of FIG. In other words, each element of the matrix a and the matrix b is also a work item which is a point of the workspace.

一方、図6のコンピューティング装置630において、図8のカーネルコードが実行される場合、プロセシングコア634は、行列aの要素と、行列bの要素とを読み出すために、メモリ620にアクセスすることができる。ここで、関数get_global_id()は、メモリ620からカーネルコードのインデックスを読み出す関数である。従って、プロセシングコア634が、図8のカーネルコードを実行するとき、複数個の作業アイテムが並列的に処理されるが、作業アイテムを読み出すために、メモリ620に何回もアクセスしなければならないので、カーネルコードを処理する時間が長くなる。 On the other hand, when the kernel code of FIG. 8 is executed in the computing device 630 of FIG. 6, the processing core 634 may access the memory 620 to read the elements of the matrix a and the elements of the matrix b. can. Here, the function get_global_id () is a function that reads the kernel code index from the memory 620. Therefore, when the processing core 634 executes the kernel code shown in FIG. 8, a plurality of work items are processed in parallel, but the memory 620 must be accessed many times in order to read the work items. , Takes longer to process kernel code.

図9は、他の一実施形態によるコンピューティング装置について説明するための図である。 FIG. 9 is a diagram for explaining a computing device according to another embodiment.

一実施形態によるコンピューティング装置920は、作業グループについての情報を保存するバッファをさらに含んでもよい。該バッファが含まれたコンピューティング装置920は、作業グループについての情報をバッファから読み取るために、メモリにアクセスする回数を減らすことができる。一方、図9のホストプロセッサ910及び最上位コントロールコア930は、図5のホストプロセッサ510及び最上位コントロールコア530とそれぞれ対応するので、詳細な説明は省略する。 The computing device 920 according to one embodiment may further include a buffer for storing information about the working group. The computing device 920 including the buffer can reduce the number of times the memory is accessed in order to read information about the working group from the buffer. On the other hand, since the host processor 910 and the top-level control core 930 of FIG. 9 correspond to the host processor 510 and the top-level control core 530 of FIG. 5, detailed description thereof will be omitted.

図9を参照すれば、最下位コントロールコア950とプロセシングコア960は、第2バッファ970と結合される。第2バッファ970は、最下位コントロールコア950から伝達された、プロセシングコア960に割り当てられた作業グループの数と、当該作業グループのIDと、を保存することができる。従って、最下位コントロールコア950は、メモリにアクセスせず、第2バッファ970に保存されたプロセシングコア960に割り当てられた作業グループについての情報を読み取ることにより、プロセシングコア960の作業グループの処理結果を確認することができる。最下位コントロールコア950は、プロセシングコア960に割り当てられた作業グループの処理が完了したと判断されると、他のプロセシングコアに割り当てられた作業グループを、プロセシングコア960にさらに割り当てることができる。一方、第2バッファ970は、プロセシングコア960に含まれ、プロセシングコア960のルータ712を介して、最下位コントロールコア950にも結合されるが、第2バッファ970の位置は、それに限定されるものではない。 Referring to FIG. 9, the lowest control core 950 and the processing core 960 are coupled to the second buffer 970. The second buffer 970 can store the number of work groups assigned to the processing core 960 transmitted from the lowest control core 950 and the ID of the work group. Therefore, the lowest control core 950 does not access the memory, and reads the information about the work group assigned to the processing core 960 stored in the second buffer 970 to obtain the processing result of the work group of the processing core 960. You can check. When the lowest control core 950 determines that the processing of the work group assigned to the processing core 960 is completed, the work group assigned to the other processing core can be further assigned to the processing core 960. On the other hand, the second buffer 970 is included in the processing core 960 and is also coupled to the lowest control core 950 via the router 712 of the processing core 960, but the position of the second buffer 970 is limited to that. is not it.

また、中間コントロールコア940と最下位コントロールコア950は、第1バッファ965とも結合される。第1バッファ965は、中間コントロールコア940から伝達した、最下位コントロールコア950に割り当てられた作業グループの数と、当該作業グループのIDと、を保存することができる。従って、中間コントロールコア940は、メモリにアクセスせず、第1バッファ965に保存された最下位コントロールコア950に割り当てられた作業グループについての情報を読み取ることにより、最下位コントロールコア950の作業グループの処理結果を確認することができる。中間コントロールコア940は、最下位コントロールコア950に割り当てられた作業グループの処理が完了したと判断されると、他の最下位コントロールコアに割り当てられた作業グループを、最下位コントロールコア950にさらに割り当てることができる。一方、第1バッファ965は、最下位コントロールコア950に含まれてもよいが、第1バッファ965の位置は、それに限定されるものではない。 The intermediate control core 940 and the lowest control core 950 are also coupled to the first buffer 965. The first buffer 965 can store the number of work groups assigned to the lowest control core 950 transmitted from the intermediate control core 940 and the ID of the work group. Therefore, the intermediate control core 940 does not access the memory and reads the information about the work group assigned to the lowest control core 950 stored in the first buffer 965 to read the information about the work group of the lowest control core 950. You can check the processing result. When the intermediate control core 940 determines that the processing of the work group assigned to the lowest control core 950 is completed, the intermediate control core 940 further allocates the work group assigned to the other lowest control core to the lowest control core 950. be able to. On the other hand, the first buffer 965 may be included in the lowest control core 950, but the position of the first buffer 965 is not limited thereto.

また、図9のコンピューティング装置920は、2個のバッファ965,970を含んでいるが、該バッファの数は、それに限定されるものではない。 Further, the computing device 920 of FIG. 9 includes two buffers 965 and 970, but the number of the buffers is not limited thereto.

図10は、一実施形態によるバッファについて説明するための図である。図10を参照すれば、バッファ1030は、最下位コントロールコア1010とプロセシングコア1020とに結合される。また、バッファ1030は、プロセシングコア1020に割り当てられた作業グループのIDと、プロセシングコア1020に割り当てられた作業グループの数と、を保存することができる。 FIG. 10 is a diagram for explaining a buffer according to one embodiment. With reference to FIG. 10, the buffer 1030 is coupled to the lowest control core 1010 and the processing core 1020. Further, the buffer 1030 can store the ID of the work group assigned to the processing core 1020 and the number of work groups assigned to the processing core 1020.

例えば、最下位コントロールコア1010は、プロセシングコア1020に作業グループを割り当てた後、当該作業グループのIDと、プロセシングコア1020に割り当てられた作業グループの数と、をバッファ1030に伝達することができる。プロセシングコア1020は、作業グループを処理した後、バッファ1030に保存されたプロセシングコア1020に割り当てられた作業グループの数をアップデートすることができる。従って、最下位コントロールコア1010は、メモリにアクセスせず、バッファ1030を確認することにより、プロセシングコア1020の作業グループ処理結果を確認することができる。 For example, the lowest control core 1010 can assign a work group to the processing core 1020, and then transmit the ID of the work group and the number of work groups assigned to the processing core 1020 to the buffer 1030. After processing the work groups, the processing core 1020 can update the number of work groups assigned to the processing core 1020 stored in the buffer 1030. Therefore, the lowest control core 1010 can confirm the work group processing result of the processing core 1020 by checking the buffer 1030 without accessing the memory.

図11は、一実施形態による、バッファが追加された場合、プロセシングコアで実行されるOpenCLカーネルコードである。図11のカーネルコードは、バッファ1030を利用して、図8のカーネルコードと同一機能を遂行するコードである。言い替えれば、図11のカーネルコードは、バッファ1030が追加されたプロセシングコア1020によって実行される、行列aの要素と、行列bの要素とを加え、その結果を行列cに保存するカーネルコードである。 FIG. 11 is an OpenCL kernel code according to one embodiment that is executed on the processing core when a buffer is added. The kernel code of FIG. 11 is a code that uses the buffer 1030 to perform the same function as the kernel code of FIG. In other words, the kernel code of FIG. 11 is the kernel code executed by the processing core 1020 to which the buffer 1030 is added, which adds the elements of the matrix a and the elements of the matrix b and stores the result in the matrix c. ..

図11を参照すれば、関数read_wg_id()は、プロセシングコア1020に割り当てられた作業グループのIDをバッファ1030から読み取る関数である。従って、図11のカーネルコードは、関数read_wg_id()を介して、バッファ1030に保存された作業グループのIDを、カーネルコードのインデックスとして活用するために、メモリにアクセスせずとも、各プロセシングコア1020が並列的に作業グループを処理することができる。 Referring to FIG. 11, the function read_wg_id () is a function that reads the ID of the work group assigned to the processing core 1020 from the buffer 1030. Therefore, in the kernel code of FIG. 11, each processing core 1020 is used without accessing the memory in order to utilize the ID of the working group stored in the buffer 1030 as an index of the kernel code via the function read_wg_id (). Can process working groups in parallel.

図12は、他の一実施形態による、バッファが追加された場合、OpenCL作業グループを実行するOpenCLカーネルコードである。 FIG. 12 is OpenCL kernel code that executes an OpenCL working group when a buffer is added, according to another embodiment.

一実施形態によって、プロセシングコア1020が少なくとも1つのプロセシングエレメントを含み、作業グループに含まれた作業アイテムの数より、プロセシングエレメントの数が少なければ、コンピューティング装置は、該プロセシングエレメントが該作業アイテムを直列化し、順次に実行するように、OpenCLカーネルを変換することができる。 In one embodiment, if the processing core 1020 contains at least one processing element and the number of processing elements is less than the number of work items contained in the work group, the computing device will perform the processing element with the work item. You can convert the OpenCL kernel to serialize and run sequentially.

具体的には、作業グループに含まれた作業アイテムの数より、プロセシングエレメントの数が少ない場合、プロセシングエレメントは、複数個の作業アイテムを処理することができる。かような場合、該プロセシングエレメントは、割り当てられた作業アイテムの処理が完了するたびに、完了した作業アイテムについての情報をメモリに保存した後、新たな作業アイテムにスイッチングすることができる。従って、該プロセシングエレメントが複数個の作業アイテムを処理するとき、該プロセシングエレメントがメモリにアクセスする回数が増加する。 Specifically, when the number of processing elements is smaller than the number of work items included in the work group, the processing element can process a plurality of work items. In such a case, the processing element can switch to a new work item after storing the information about the completed work item in the memory each time the processing of the assigned work item is completed. Therefore, when the processing element processes a plurality of work items, the number of times the processing element accesses the memory increases.

プロセシングエレメントによるメモリのアクセス回数を減らすために、コンピューティング装置は、作業アイテム合体(work item coalescing)方法に基づいて、プロセシングエレメントが、複数個の作業アイテムを直列化して処理するように、カーネルコードを変換することができる。 To reduce the number of times memory is accessed by a processing element, the computing device is based on a work item coalescing method, so that the processing element processes multiple work items in series in the kernel code. Can be converted.

図12を参照すれば、関数read_num_wgs()は、プロセシングコア1020に連続的に割り当てられた作業グループの数を、バッファ1030から読み取る関数である。 Referring to FIG. 12, the function read_num_wgs () is a function that reads the number of working groups continuously assigned to the processing core 1020 from the buffer 1030.

従って、図12のカーネルコードは、バッファ1030に保存された情報を読み取る関数read_wg_id()及び関数read_num_wgs()に基づいて、カーネルコードのインデックスを設定するために、プロセシングコア1020によるメモリアクセス回数を減少させることができる。また、カーネルコードを取り囲むループが適用されており、実行されなければならないカーネルコードを、1つの作業アイテムを実行するたびに呼び出す必要なしに、1回のカーネルコードの呼び出しにより、複数個の作業アイテムを順次に実行することができる。従って、作業アイテム合体方法を適用しない場合に比べ、プロセシングコア1020のメモリアクセス回数を減少させることができる。 Therefore, the kernel code of FIG. 12 reduces the number of memory accesses by the processing core 1020 to index the kernel code based on the function read_wg_id () and the function read_num_wgs () that read the information stored in buffer 1030. Can be made to. In addition, a loop surrounding the kernel code is applied, and multiple work items can be called by one kernel code call without having to call the kernel code that must be executed each time one work item is executed. Can be executed sequentially. Therefore, the number of memory accesses of the processing core 1020 can be reduced as compared with the case where the work item combination method is not applied.

図13及び図14は、一実施形態によるコンピューティング装置の性能を検証するための実験環境を示した図である。 13 and 14 are diagrams showing an experimental environment for verifying the performance of the computing device according to the embodiment.

図13は、一実施形態によるコンピューティング装置に、異なるスケジューリング技法を適用した場合、作業グループを割り当てるのに必要な遅延時間を示した図である。ここでは、コンピューティング装置に含まれたプロセシンググループは、96個のプロセシングコアを含み、該コンピューティング装置に含まれたコントロールコアグループは、1個の最上位コントロールコア、4個の中間コントロールコア、及び16個の最下位コントロールコアを含むと仮定する。また、コントロールコアとプロセシングコアは、それぞれ50MHz及び500MHzのクロック周波数(clock frequency)によって動作すると仮定する。 FIG. 13 is a diagram showing the delay time required to allocate a work group when different scheduling techniques are applied to the computing device according to one embodiment. Here, the processing group included in the computing device includes 96 processing cores, and the control core group included in the computing device is one top control core, four intermediate control cores, and so on. And 16 lowest control cores are included. It is also assumed that the control core and the processing core operate at clock frequencies of 50 MHz and 500 MHz, respectively.

図13の集中的スケジューリング方法は、単一コントロールコアが、全てのプロセシングコアに、作業グループを動的に割り当てるスケジューリング方法を意味する。分散スケジューリング方法は、階層的に分類されたコントロールコアが、プロセシングコアに作業グループを動的に割り当てるスケジューリング方法を意味する。また、ハードウェアサポート項目は、コンピューティング装置がバッファを含むか否かということを示す項目である。 The centralized scheduling method of FIG. 13 means a scheduling method in which a single control core dynamically allocates work groups to all processing cores. The distributed scheduling method means a scheduling method in which hierarchically classified control cores dynamically assign work groups to processing cores. The hardware support item is an item indicating whether or not the computing device includes a buffer.

図13を参照すれば、分散スケジューリング方法が適用されたコンピューティング装置の遅延時間は、集中的スケジューリング方法が適用されたコンピューティング装置の遅延時間より短い。従って、階層的に分類されたさまざまなコントロールコアが、96個のプロセシングコアに作業グループを動的に割り当てる時間が、単一コントロールコアが作業グループを動的に割り当てる時間より短くなるということが分かる。 Referring to FIG. 13, the delay time of the computing device to which the distributed scheduling method is applied is shorter than the delay time of the computing device to which the centralized scheduling method is applied. Therefore, it can be seen that the time for various hierarchically classified control cores to dynamically allocate work groups to 96 processing cores is shorter than the time for a single control core to dynamically allocate work groups. ..

一方、同一スケジューリング方法が適用された場合、バッファが追加されたコンピューティング装置の遅延時間は、バッファなしのコンピューティング装置の遅延時間より短い。従って、バッファが追加されたコンピューティング装置は、メモリにアクセスする回数が減少されるために、作業グループを処理する遅延時間が短くなるということが分かる。 On the other hand, when the same scheduling method is applied, the delay time of the computing device to which the buffer is added is shorter than the delay time of the computing device without the buffer. Therefore, it can be seen that the computing device to which the buffer is added has a shorter delay time for processing the work group because the number of times to access the memory is reduced.

図14を参照すれば、コンピューティング装置の性能を検証するために、異なる4個のアプリケーションが実行されると仮定する。具体的には、PI estimationアプリケーションは、他のアプリケーションに比べ、割り当てられる作業グループの処理時間が長いという特性を有する。Small buffer copy及びBig buffer copyアプリケーションは、他のアプリケーションに比べ、メモリアクセス回数が多いという特性を有する。また、SpMVアプリケーションは、他のアプリケーションに比べ、メモリアクセスが一定ではなく、割り当てられた作業グループを処理する時間が非均一であるという特性を有する。 With reference to FIG. 14, it is assumed that four different applications are executed to verify the performance of the computing device. Specifically, the PI estimation application has a characteristic that the processing time of the assigned work group is longer than that of other applications. The Small buffer copy and Big buffer copy applications have the characteristic that the number of memory accesses is larger than that of other applications. Further, the SpMV application has a characteristic that the memory access is not constant and the time for processing the assigned work group is non-uniform as compared with other applications.

図15は、一実施形態によるコンピューティング装置の性能検証結果について説明するための図である。図15は、理想的なスケジューリングの性能に基づいて、コンピューティング装置によって、4個のアプリケーションが実行されるとき、スケジューリング性能を正規化した結果である。 FIG. 15 is a diagram for explaining the performance verification result of the computing device according to the embodiment. FIG. 15 shows the result of normalizing the scheduling performance when four applications are executed by the computing device based on the ideal scheduling performance.

図15を参照すれば、集中的スケジューリング方法が適用されたコンピューティング装置のスケジューリング性能は、分散スケジューリング方法が適用されたコンピューティング装置のスケジューリング性能より全体的に低いということが分かる。具体的には、SpMVアプリケーションの場合、集中的スケジューリング方法が適用されたコンピューティング装置と、分散スケジューリング方法が適用されたコンピューティング装置とのスケジューリング性能差が大きいということが分かる。 With reference to FIG. 15, it can be seen that the scheduling performance of the computing device to which the centralized scheduling method is applied is generally lower than the scheduling performance of the computing device to which the distributed scheduling method is applied. Specifically, in the case of the SpMV application, it can be seen that the scheduling performance difference between the computing device to which the centralized scheduling method is applied and the computing device to which the distributed scheduling method is applied is large.

また、バッファが追加されたコンピューティング装置が、分散スケジューリングに基づいて、作業グループのスケジューリングを行う場合、分散スケジューリングだけ適用されたコンピューティング装置に比べ、スケジューリング性能がさらに良好であるということが分かる。具体的には、メモリアクセス回数が多いSmall buffer copyアプリケーション及びBig buffer copyアプリケーションが、バッファ有無によるスケジューリング性能差が大きいということが分かる。 Further, it can be seen that when the computing device to which the buffer is added performs the scheduling of the work group based on the distributed scheduling, the scheduling performance is further better than that of the computing device to which only the distributed scheduling is applied. Specifically, it can be seen that the Small buffer copy application and the Big buffer copy application, which have a large number of memory accesses, have a large difference in scheduling performance depending on the presence or absence of a buffer.

図16は、一実施形態による、OpenCLカーネルを処理する方法を図示したフローチャートである。 FIG. 16 is a flowchart illustrating a method of processing an OpenCL kernel according to an embodiment.

段階1610において、コントロールコアグループ320は、OpenCLカーネルを実行するための作業グループを、下位階層のコントロールコア及びプロセシングコアに割り当てることができる。ここで、コントロールコアグループ320に含まれた複数個のコントロールコアは、コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類される。 At stage 1610, the control core group 320 can allocate working groups for running the OpenCL kernel to lower hierarchy control cores and processing cores. Here, the plurality of control cores included in the control core group 320 are hierarchically classified according to the number of processing cores to which the work group is assigned by the control cores.

具体的には、コントロールコアグループ320は、最上位コントロールコア330及び最下位コントロールコアグループ350を含み得る。 Specifically, the control core group 320 may include a top control core 330 and a bottom control core group 350.

最上位コントロールコア330は、OpenCLカーネルの実行情報が受信されることにより、作業グループを、下位階層のコントロールコアに割り当て、割り当てられた作業グループについての情報を、下位階層のコントロールコアに送信することができる。最下位コントロールコアグループ350は、上位階層のコントロールコアによって割り当てられた作業グループについての情報を受信し、プロセシングコアに、作業グループを割り当てる、少なくとも1つの最下位コントロールコアを含み得る。 Upon receiving the execution information of the OpenCL kernel, the uppermost control core 330 allocates a work group to the lower layer control core, and transmits information about the assigned work group to the lower layer control core. Can be done. The lowest control core group 350 may include at least one lowest control core that receives information about work groups assigned by higher level control cores and assigns work groups to processing cores.

また、コントロールコアグループ320は、最上位コントロールコア330によって割り当てられた作業グループに係わる情報を受信し、最下位コントロールコアグループ350に作業グループを割り当てる、複数個の中間コントロールコアが含まれた中間コントロールコアグループ340をさらに含んでもよい。 Further, the control core group 320 receives information related to the work group assigned by the uppermost control core 330, and assigns the work group to the lowermost control core group 350, which is an intermediate control including a plurality of intermediate control cores. The core group 340 may be further included.

一実施形態により、中間コントロールコアが、中間コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類されるとき、段階1610において、上位階層の中間コントロールコアは、下位階層の中間コントロールコアに作業グループを割り当てることができる。 In one embodiment, when the intermediate control cores are hierarchically classified by the number of processing cores to which work groups are assigned by the intermediate control cores, in stage 1610, the upper layer intermediate control cores are the lower layer intermediate controls. You can assign work groups to the core.

段階1620において、少なくとも1つのプロセシングコアが含まれたプロセシングコアグループ360は、コントロールコアグループ320によって割り当てられた作業グループを処理することができる。 In step 1620, the processing core group 360, which includes at least one processing core, can process the work group assigned by the control core group 320.

段階1630において、プロセシングコアグループ360は、作業グループの処理結果を出力することができる。 At step 1630, the processing core group 360 can output the processing results of the working group.

図17は、一実施形態による、OpenCLカーネルを処理する方法を図示した詳細フローチャートである。図17は、図16によるコンピューティング装置300が、OpenCLカーネルを処理する方法について詳細に説明するための図であり、コンピューティング装置300の構成によって、OpenCLカーネルを処理する方法が異なるということは、当該技術分野の当業者であるならば、理解することができるであろう。また、図17の段階1720、段階1730及び段階1740は、図16の段階1610、段階1620及び段階1630とそれぞれ対応するので、詳細な説明は省略する。 FIG. 17 is a detailed flowchart illustrating a method of processing an OpenCL kernel according to one embodiment. FIG. 17 is a diagram for explaining in detail the method of processing the OpenCL kernel by the computing device 300 according to FIG. 16, and the method of processing the OpenCL kernel differs depending on the configuration of the computing device 300. Anyone skilled in the art will understand. Further, since the stages 1720, 1730 and 1740 in FIG. 17 correspond to the stages 1610, 1620 and 1630 in FIG. 16, detailed description thereof will be omitted.

段階1710において、最上位コントロールコア330は、OpenCLカーネルが実行されることにより、ホストプロセッサ310によって生成されたOpenCLカーネルの実行情報を受信することができる。かような場合、最上位コントロールコア330は、OpenCLカーネルを実行するための作業グループが処理中であるとき、ホストプロセッサ310から新たなOpenCLカーネルの実行情報が受信されることにより、プロセシングコアグループ360に割り当てられた作業グループを再割り当てすることができる。ここで、OpenCLカーネルの実行情報は、OpenCLカーネルが実行されるための全体作業グループの数、及びOpenCLカーネルが含まれたアプリケーション情報を含んでもよい。 In step 1710, the top-level control core 330 can receive the execution information of the OpenCL kernel generated by the host processor 310 by executing the OpenCL kernel. In such a case, the top-level control core 330 receives the execution information of the new OpenCL kernel from the host processor 310 while the working group for executing the OpenCL kernel is being processed, so that the processing core group 360 You can reassign the workgroups assigned to. Here, the execution information of the OpenCL kernel may include the number of the entire working group for the OpenCL kernel to be executed, and the application information including the OpenCL kernel.

また、一実施形態によるOpenCLカーネルを処理する方法は、ホストプロセッサ310が、コントロールコアグループ320によるプロセシングコアグループ360に対する作業グループ割り当てが完了すると、OpenCLカーネルの実行を中断し、割り当て結果を収集して出力する段階をさらに含んでもよい。 Further, in the method of processing the OpenCL kernel according to one embodiment, when the host processor 310 completes the work group allocation to the processing core group 360 by the control core group 320, the execution of the OpenCL kernel is interrupted and the allocation result is collected. It may further include a step of output.

また、一実施形態による、OpenCLカーネルを処理する方法は、バッファに作業グループについての情報を保存する段階をさらに含んでもよい。ここで、該作業グループについての情報は、下位階層のコントロールコア及びプロセシングコアに割り当てられた作業グループのID、及び割り当てられた作業グループの数を含んでもよい。 Also, the method of processing the OpenCL kernel according to one embodiment may further include the step of storing information about the working group in a buffer. Here, the information about the work group may include the ID of the work group assigned to the control core and the processing core in the lower hierarchy, and the number of the assigned work groups.

段階1750において、コントロールコアグループ320は、出力された作業グループの処理結果に従って、下位階層の他のコントロールコア、及び他のプロセシングコアに割り当てられた作業グループを再割り当てすることができる。例えば、最上位コントロールコア330は、中間コントロールコアまたは最下位コントロールコアの作業グループの処理結果に従って、他の中間コントロールコア、または他の最下位コントロールコアに割り当てられた作業グループを再割り当てすることができる。また、中間コントロールコアは、下位階層の中間コントロールコアまたは最下位コントロールコアの作業グループの処理結果に従って、下位階層の他の中間コントロールコア、または他の最下位コントロールコアに割り当てられた作業グループを再割り当てすることができる。また、最下位コントロールコアは、プロセシングコアの作業グループの処理結果に従って、他のプロセシングコアに割り当てられた作業グループを再割り当てすることができる。 In step 1750, the control core group 320 can reassign the work groups assigned to the other control cores in the lower hierarchy and the other processing cores according to the processing result of the output work group. For example, the top-level control core 330 may reassign work groups assigned to other intermediate control cores or other lowest-level control cores according to the processing results of the work groups of the intermediate control core or the lowest control core. can. In addition, the intermediate control core reassigns the work groups assigned to other intermediate control cores or other lowest control cores in the lower hierarchy according to the processing results of the work groups of the intermediate control cores or the lowest control cores in the lower hierarchy. Can be assigned. Further, the lowest control core can reassign the work group assigned to another processing core according to the processing result of the work group of the processing core.

本実施形態は、コンピュータによって実行されるプログラムモジュールのような、コンピュータによって実行可能な命令語を含む記録媒体の形態でも具現される。コンピュータ可読媒体は、コンピュータによってアクセスされる任意の可用媒体でもあり、揮発性及び不揮発性の媒体、分離型及び非分離型の媒体をいずれも含む。コンピュータ記録媒体は、コンピュータ可読命令語、データ構造、プログラムモジュール、またはその他データのような情報保存のための任意の方法または技術によって具現された揮発性及び不揮発性、分離型及び非分離型の媒体をいずれも含む。 The present embodiment is also embodied in the form of a recording medium containing a computer-executable instruction word, such as a computer-executed program module. A computer-readable medium is also any usable medium accessed by a computer, including both volatile and non-volatile media, separable and non-separable media. Computer recording media are volatile and non-volatile, separable and non-separable media embodied by any method or technique for storing information such as computer-readable instructions, data structures, program modules, or other data. Including any.

本実施形態の範囲は、前述の説明よりは、特許請求の範囲によって示され、特許請求の範囲の意味及び範囲、そしてその均等概念から導き出される全ての変更、または変形された形態が含まれるものであると解釈されなければならない。 The scope of the present embodiment is indicated by the scope of claims, rather than the above description, and includes the meaning and scope of the scope of claims, and all modifications or modifications derived from the concept of equality thereof. Must be interpreted as.

本発明の、OpenCLカーネルを処理する方法、及びそれを遂行するコンピューティング装置は、例えば、電子装置関連の技術分野に効果的に適用可能である。 The method of processing the OpenCL kernel and the computing device for carrying out the method of the present invention are effectively applicable to, for example, electronic device-related technical fields.

10,20,30,40 プロセシングエレメント
100 OpenCLプラットフォーム
110,310 ホストプロセッサ
120,130,140 コンピューティングデバイス
121,122 コンピューティングユニット
200,300 コンピューティング装置
210,320 コントロールコアグループ
220,360 プロセシングコアグループ
330 最上位コントロールコア
340 中間コントロールコアグループ
350 最下位コントロールコア
360 プロセシングコアグループ
10, 20, 30, 40 Processing element 100 OpenCL platform 110, 310 Host processor 120, 130, 140 Computing device 121, 122 Computing unit 200, 300 Computing device 210, 320 Control core group 220, 360 Processing core group 330 Top Control Core 340 Middle Control Core Group 350 Bottom Control Core 360 Processing Core Group

Claims (18)

複数個のコアを含むコンピューティング装置において、
下位階層のコントロールコアを含み、OpenCL(open computing language)カーネルを実行するための作業グループを、前記下位階層のコントロールコア及びプロセシングコアに割り当てるコントロールコアグループと、
記プロセシングコアのうちの少なくとも1つを含み、前記コントロールコアグループによって割り当てられた前記作業グループを処理し、前記作業グループの処理結果を出力する、プロセシングコアグループと、を含み、
前記コントロールコアグループに含まれた複数個のコントロールコアは、前記コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数によって階層的に分類され、
前記コントロールコアグループは、
前記OpenCLカーネルの実行情報が受信されることにより、前記作業グループを前記下位階層のコントロールコアに割り当て、前記割り当てられた作業グループについての情報を前記下位階層のコントロールコアに送信する最上位コントロールコアと、
前記コントロールコアグループの上位階層のコントロールコアによって割り当てられた作業グループについての情報を受信し、前記プロセシングコアに前記作業グループを割り当てる、少なくとも1つの最下位コントロールコアを含んだ最下位コントロールコアグループと、を含む、
コンピューティング装置。
In a computing device that contains multiple cores
A control core group that includes a lower-level control core and allocates a working group for executing an OpenCL (open computing language) kernel to the lower-level control core and processing core, and a control core group.
Wherein at least one of the previous SL processing cores to process the working group assigned by the control core group, and outputs the processing result of the task group includes a processing core group, and
The plurality of control cores included in the control core group are hierarchically classified according to the number of the processing cores to which the working group is assigned by the control cores .
The control core group
Upon receiving the execution information of the OpenCL kernel, the work group is assigned to the control core of the lower layer, and the information about the assigned work group is transmitted to the control core of the lower layer. ,
A lowest control core group containing at least one lowest control core that receives information about a work group assigned by a higher level control core of the control core group and assigns the work group to the processing core. including,
Computing equipment.
前記コントロールコアグループは、
前記下位階層のコントロールコア及び前記プロセシングコアによる前記作業グループの処理結果に従って、前記コントロールコアグループの他の下位階層のコントロールコア及び他のプロセシングコアに割り当てられた作業グループを再割り当てすることを特徴とする請求項に記載のコンピューティング装置。
The control core group
It is characterized in that the work groups assigned to other lower layer control cores and other processing cores of the control core group are reassigned according to the processing result of the work group by the lower layer control core and the processing core. The computing device according to claim 1.
前記コントロールコアグループは、
前記最上位コントロールコアによって割り当てられた作業グループに係わる情報を受信し、前記最下位コントロールコアグループに前記作業グループを割り当てる、複数個の中間コントロールコアを含んだ中間コントロールコアグループをさらに含み、
前記中間コントロールコアは、前記中間コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数によって階層的に分類され、前記コントロールコアグループの上位階層の中間コントロールコアは、下位階層の中間コントロールコアに前記作業グループを割り当てることを特徴とする請求項に記載のコンピューティング装置。
The control core group
Further including an intermediate control core group including a plurality of intermediate control cores, which receives information relating to the work group assigned by the highest control core and assigns the work group to the lowest control core group.
The intermediate control cores are hierarchically classified according to the number of the processing cores to which the working group is assigned by the intermediate control cores, and the intermediate control cores in the upper layer of the control core group are assigned to the intermediate control cores in the lower layer. The computing device according to claim 1 , wherein a work group is assigned.
前記最上位コントロールコアは、前記OpenCLカーネルが実行されることにより、ホストプロセッサによって生成された前記OpenCLカーネルの実行情報を受信することを特徴とする請求項に記載のコンピューティング装置。 The computing device according to claim 1 , wherein the top-level control core receives execution information of the OpenCL kernel generated by a host processor when the OpenCL kernel is executed. 前記最上位コントロールコアは、
前記OpenCLカーネルを実行するための作業グループを処理しているとき、前記ホストプロセッサから新たなOpenCLカーネルの実行情報が受信されることにより、前記プロセシングコアグループに割り当てられた前記作業グループを再割り当てすることを特徴とする請求項に記載のコンピューティング装置。
The top control core is
While processing the work group for executing the OpenCL kernel, the work group assigned to the processing core group is reassigned by receiving the execution information of the new OpenCL kernel from the host processor. The computing device according to claim 4 , wherein the computing device is characterized by the above.
前記ホストプロセッサは、
前記コントロールコアグループによる前記プロセシングコアグループへの前記作業グループの割り当てが完了すると、前記OpenCLカーネルの実行を中断し、割り当て結果を収集して出力することを特徴とする請求項に記載のコンピューティング装置。
The host processor
The computing according to claim 4 , wherein when the allocation of the working group to the processing core group by the control core group is completed, the execution of the OpenCL kernel is interrupted, and the allocation result is collected and output. Device.
前記プロセシングコアの中の所与のプロセシングコアが、少なくとも1つのプロセシングエレメントを含み、前記所与のプロセシングコアに割り当てられた前記作業グループのうちの1つに含まれた作業アイテムの数より、前記プロセシングエレメントの数が少ないとき、前記所与のプロセシングコアは、前記OpenCLカーネルを変換し、前記プロセシングエレメントのうちの1つが、前記変換されたOpenCLカーネルを用いて、複数個の前記作業アイテムを順次に実行することを特徴とする請求項1に記載のコンピューティング装置。 Given processing core in said processing core includes at least one processing element, than the number of work items included in one of the given processing cores assigned to the work group, the When the number of processing elements is small, the given processing core transforms the OpenCL kernel, and one of the processing elements uses the transformed OpenCL kernel to sequentially process a plurality of said work items. The computing device according to claim 1, wherein the computing device is executed. 前記作業グループについての情報を保存するバッファをさらに含むことを特徴とする請求項1に記載のコンピューティング装置。 The computing device according to claim 1, further comprising a buffer for storing information about the working group. 前記作業グループについての情報は、
前記下位階層のコントロールコア及び前記プロセシングコアに割り当てられた作業グループの数と、前記割り当てられた作業グループのIDとのうち少なくとも一つを含むことを特徴とする請求項に記載のコンピューティング装置。
Information about the working group
The computing device according to claim 8 , further comprising at least one of the number of work groups assigned to the lower layer control core and the processing core, and the ID of the assigned work group. ..
下位階層のコントロールコアを含むコントロールコアグループにより、OpenCL(open computing language)カーネルを実行するための作業グループを、前記下位階層のコントロールコア及びプロセシングコアに割り当てる段階と、
記プロセシングコアのうちの少なくとも1つが含まれたプロセシングコアグループにより、前記割り当てられた前記作業グループを処理する段階と、
前記プロセシングコアグループにより、前記作業グループの処理結果を出力する段階と、を含み、
前記コントロールコアグループに含まれた複数個のコントロールコアは、前記コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数によって階層的に分類され、
前記コントロールコアグループは、
前記OpenCLカーネルの実行情報が受信されることにより、前記作業グループを前記下位階層のコントロールコアに割り当て、前記割り当てられた作業グループについての情報を前記下位階層のコントロールコアに送信する最上位コントロールコアと、
前記コントロールコアグループの上位階層のコントロールコアによって割り当てられた作業グループについての情報を受信し、前記プロセシングコアに前記作業グループを割り当てる、少なくとも1つの最下位コントロールコアを含んだ最下位コントロールコアグループと、を含む、
OpenCLカーネルを処理する方法。
A stage in which a work group for executing an OpenCL (open computing language) kernel is assigned to the lower layer control core and processing core by a control core group including a lower layer control core, and
The at least one processing core group that contains one of the previous SL processing core, a step of treating the said working group allocated,
The processing core group includes a step of outputting the processing result of the working group.
The plurality of control cores included in the control core group are hierarchically classified according to the number of the processing cores to which the working group is assigned by the control cores .
The control core group
Upon receiving the execution information of the OpenCL kernel, the work group is assigned to the control core of the lower layer, and the information about the assigned work group is transmitted to the control core of the lower layer. ,
A lowest control core group containing at least one lowest control core that receives information about a work group assigned by a higher level control core of the control core group and assigns the work group to the processing core. including,
How to handle the OpenCL kernel.
前記コントロールコアグループにより、前記出力された前記作業グループの処理結果に従って、前記コントロールコアグループの下位階層の他のコントロールコア及び他のプロセシングコアに割り当てられた作業グループを再割り当てする段階をさらに含むことを特徴とする請求項10に記載のOpenCLカーネルを処理する方法。 The control core group further includes a step of reallocating work groups assigned to other control cores and other processing cores in the lower hierarchy of the control core group according to the output processing result of the work group. 10. The method of processing the OpenCL kernel according to claim 10. 前記コントロールコアグループは、
前記最上位コントロールコアによって割り当てられた作業グループに係わる情報を受信し、前記最下位コントロールコアグループに前記作業グループを割り当てる、複数個の中間コントロールコアを含んだ中間コントロールコアグループをさらに含み、
前記中間コントロールコアが、前記中間コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数によって階層的に分類されるとき、
前記作業グループを割り当てる段階は、
前記コントロールコアグループの上位階層の中間コントロールコアにより、前記コントロールコアグループの下位階層の中間コントロールコアに前記作業グループを割り当てる段階を含むことを特徴とする請求項10に記載のOpenCLカーネルを処理する方法。
The control core group
Further including an intermediate control core group including a plurality of intermediate control cores, which receives information relating to the work group assigned by the highest control core and assigns the work group to the lowest control core group.
When the intermediate control cores are hierarchically classified by the number of the processing cores to which the working group is assigned by the intermediate control cores.
The stage of assigning the work group is
The method of processing the OpenCL kernel according to claim 10 , wherein the intermediate control core in the upper layer of the control core group includes a step of allocating the work group to the intermediate control core in the lower layer of the control core group. ..
前記OpenCLカーネルが実行されることにより、ホストプロセッサによって生成された前記OpenCLカーネルの実行情報を、前記最上位コントロールコアにより受信する段階をさらに含むことを特徴とする請求項10に記載のOpenCLカーネルを処理する方法。 The OpenCL kernel according to claim 10 , further comprising a step of receiving the execution information of the OpenCL kernel generated by the host processor by the top-level control core by executing the OpenCL kernel. How to handle. 前記OpenCLカーネルを実行するための作業グループを処理しているとき、前記ホストプロセッサから新たなOpenCLカーネルの実行情報が受信されることにより、前記最上位コントロールコアにより、前記プロセシングコアグループに割り当てられた前記作業グループを再割り当てする段階をさらに含むことを特徴とする請求項13に記載のOpenCLカーネルを処理する方法。 While processing the work group for executing the OpenCL kernel, the top-level control core assigned the processing core group to the processing core group by receiving the execution information of the new OpenCL kernel from the host processor. The method of processing an OpenCL kernel according to claim 13 , further comprising the step of reallocating the working group. 記プロセシングコアの中の所与のプロセシングコアが、少なくとも1つのプロセシングエレメントを含み、前記所与のプロセシングコアに割り当てられた前記作業グループに含まれた作業アイテムの数より前記プロセシングエレメントの数が少ないとき、前記作業グループを処理する段階は、前記所与のプロセシングコアが前記OpenCLカーネルを変換することと、前記プロセシングエレメントのうちの1つが、前記変換されたOpenCLカーネルを用いて、複数個の前記作業アイテムを順次に実行することとを含むことを特徴とする請求項10に記載のOpenCLカーネルを処理する方法。 Given processing core in the prior SL processing core includes at least one processing element, the number of the processing elements than the number of work items included in the working group assigned to the given processing core When less, the stage of processing the working group is that the given processing core transforms the OpenCL kernel and one of the processing elements uses the transformed OpenCL kernel to perform a plurality of steps. method of processing OpenCL kernel of claim 10, characterized in that it comprises the method comprising sequentially performing the work item. 前記作業グループについての情報をバッファに保存する段階を含むことを特徴とする請求項10に記載のOpenCLカーネルを処理する方法。 The method of processing an OpenCL kernel according to claim 10 , further comprising storing information about the working group in a buffer. 前記作業グループについての情報は、
前記下位階層のコントロールコア及び前記プロセシングコアに割り当てられた作業グループの数と、前記割り当てられた作業グループのIDとのうち少なくとも一つを含むことを特徴とする請求項16に記載のOpenCLカーネルを処理する方法。
Information about the working group
The OpenCL kernel according to claim 16 , wherein the OpenCL kernel includes at least one of the number of work groups assigned to the lower layer control core and the processing core, and the ID of the assigned work group. How to handle.
請求項10ないし17のうちいずれか1項に記載の方法をコンピュータで実行させるためのプログラムを記録したコンピュータ読み取り可能な記録媒体。
A computer-readable recording medium on which a program for executing the method according to any one of claims 10 to 17 is recorded.
JP2017238434A 2016-12-27 2017-12-13 How to handle the OpenCL kernel and the computing equipment that carries it out Active JP6951962B2 (en)

Applications Claiming Priority (2)

Application Number Priority Date Filing Date Title
KR10-2016-0180133 2016-12-27
KR1020160180133A KR102592330B1 (en) 2016-12-27 2016-12-27 Method for processing OpenCL kernel and computing device thereof

Publications (3)

Publication Number Publication Date
JP2018106709A JP2018106709A (en) 2018-07-05
JP2018106709A5 JP2018106709A5 (en) 2021-01-07
JP6951962B2 true JP6951962B2 (en) 2021-10-20

Family

ID=60473303

Family Applications (1)

Application Number Title Priority Date Filing Date
JP2017238434A Active JP6951962B2 (en) 2016-12-27 2017-12-13 How to handle the OpenCL kernel and the computing equipment that carries it out

Country Status (5)

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

Families Citing this family (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN111490946B (en) * 2019-01-28 2023-08-11 阿里巴巴集团控股有限公司 FPGA connection realization method and device based on OpenCL framework

Family Cites Families (23)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JPS595755B2 (en) 1977-08-02 1984-02-07 日本軽金属株式会社 Solar heating method
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
US8286198B2 (en) * 2008-06-06 2012-10-09 Apple Inc. Application programming interfaces for data parallel computing on multiple processors
US8285653B2 (en) 2009-05-15 2012-10-09 The Aerospace Corporation Systems and methods for generating random feasible solutions for an evolutionary process
US9354944B2 (en) * 2009-07-27 2016-05-31 Advanced Micro Devices, Inc. Mapping processing logic having data-parallel threads across processors
WO2011027382A1 (en) 2009-09-01 2011-03-10 Hitachi, Ltd. Request processing system provided with multi-core processor
KR101640848B1 (en) 2009-12-28 2016-07-29 삼성전자주식회사 Job Allocation Method on Multi-core System and Apparatus thereof
KR101613971B1 (en) 2009-12-30 2016-04-21 삼성전자주식회사 Method for transforming program code
JP2012094072A (en) 2010-10-28 2012-05-17 Toyota Motor Corp Information processing apparatus
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
US9092267B2 (en) * 2011-06-20 2015-07-28 Qualcomm Incorporated Memory sharing in graphics processing unit
US20120331278A1 (en) 2011-06-23 2012-12-27 Mauricio Breternitz Branch removal by data shuffling
US20130141443A1 (en) * 2011-12-01 2013-06-06 Michael L. Schmit Software libraries for heterogeneous parallel processing platforms
JP5238876B2 (en) * 2011-12-27 2013-07-17 株式会社東芝 Information processing apparatus and information processing method
KR101284195B1 (en) * 2012-01-09 2013-07-10 서울대학교산학협력단 Dynamic workload distribution apparatus for opencl
KR20130093995A (en) 2012-02-15 2013-08-23 한국전자통신연구원 Method for performance optimization of hierarchical multi-core processor and the multi-core processor system of performing the method
KR20140125893A (en) 2013-01-28 2014-10-30 한국과학기술원 Virtualized many-core environment job distribution system and method, and recording medium thereof
KR102062208B1 (en) * 2013-05-03 2020-02-11 삼성전자주식회사 Apparatus and Method for translating multithreaded program code
JP6200824B2 (en) * 2014-02-10 2017-09-20 ルネサスエレクトロニクス株式会社 Arithmetic control apparatus, arithmetic control method, program, and OpenCL device
CN104035751B (en) * 2014-06-20 2016-10-12 深圳市腾讯计算机***有限公司 Data parallel processing method based on multi-graphics processor and device
US10318261B2 (en) * 2014-11-24 2019-06-11 Mentor Graphics Corporation Execution of complex recursive algorithms

Also Published As

Publication number Publication date
US20180181443A1 (en) 2018-06-28
KR20180076051A (en) 2018-07-05
US10503557B2 (en) 2019-12-10
JP2018106709A (en) 2018-07-05
CN108241508B (en) 2023-06-13
CN108241508A (en) 2018-07-03
EP3343370A1 (en) 2018-07-04
KR102592330B1 (en) 2023-10-20

Similar Documents

Publication Publication Date Title
US10248444B2 (en) Method of migrating virtual machines between non-uniform memory access nodes within an information handling system
US8661435B2 (en) System and method for affinity dispatching for task management in an emulated multiprocessor environment
CN103797462B (en) A kind of method and apparatus creating virtual machine
US9922045B2 (en) Data management in a multi-tenant distributive environment
US9244629B2 (en) Method and system for asymmetrical processing with managed data affinity
US8650296B1 (en) Workload reallocation involving inter-server transfer of software license rights and intra-server transfer of hardware resources
US8413158B2 (en) Processor thread load balancing manager
US10013264B2 (en) Affinity of virtual processor dispatching
US10467052B2 (en) Cluster topology aware container scheduling for efficient data transfer
US20060020701A1 (en) Thread transfer between processors
US8352702B2 (en) Data processing system memory allocation
CN112424765A (en) Container framework for user-defined functions
CN114138422A (en) Extensible NVMe storage virtualization method and system
US20060212642A1 (en) Partition allocation method and computer system
JP6951962B2 (en) How to handle the OpenCL kernel and the computing equipment that carries it out
KR101535792B1 (en) Apparatus for configuring operating system and method thereof
CN105677481A (en) Method and system for processing data and electronic equipment
US9176910B2 (en) Sending a next request to a resource before a completion interrupt for a previous request
Kim et al. Sophy+: Programming model and software platform for hybrid resource management of many-core accelerators
EP4323868A1 (en) System and method of dynamically partitioning computers at runtime
CN116258622A (en) GPU distribution method and device based on container, electronic equipment and medium
JP2016139171A (en) Resource allocation device, resource allocation system and resource allocation method
JP2022539291A (en) Dynamic allocation of computing resources
Kim et al. SoPHy: A Software Platform for Hybrid Resource Management of Homogeneous Many-core Accelerators
CN117441161A (en) Software optimization method and equipment of NUMA architecture

Legal Events

Date Code Title Description
A521 Request for written amendment filed

Free format text: JAPANESE INTERMEDIATE CODE: A523

Effective date: 20201120

A621 Written request for application examination

Free format text: JAPANESE INTERMEDIATE CODE: A621

Effective date: 20201120

TRDD Decision of grant or rejection written
A977 Report on retrieval

Free format text: JAPANESE INTERMEDIATE CODE: A971007

Effective date: 20210908

A01 Written decision to grant a patent or to grant a registration (utility model)

Free format text: JAPANESE INTERMEDIATE CODE: A01

Effective date: 20210914

A61 First payment of annual fees (during grant procedure)

Free format text: JAPANESE INTERMEDIATE CODE: A61

Effective date: 20210927

R150 Certificate of patent or registration of utility model

Ref document number: 6951962

Country of ref document: JP

Free format text: JAPANESE INTERMEDIATE CODE: R150