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 PDFInfo
- 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
Links
- 238000012545 processing Methods 0.000 claims description 233
- 238000000034 method Methods 0.000 claims description 61
- 239000000872 buffer Substances 0.000 claims description 45
- 238000010586 diagram Methods 0.000 description 16
- 230000006870 function Effects 0.000 description 14
- 239000011159 matrix material Substances 0.000 description 12
- 238000004891 communication Methods 0.000 description 3
- 241001409283 Spartina mottle virus Species 0.000 description 2
- 238000012986 modification Methods 0.000 description 2
- 230000004048 modification Effects 0.000 description 2
- 238000012795 verification Methods 0.000 description 2
- 239000000470 constituent Substances 0.000 description 1
- 238000005516 engineering process Methods 0.000 description 1
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
-
- 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/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)
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.
本実施形態で使用される用語は、本実施形態での機能を考慮しながら、可能な限り、現在汎用される一般的な用語を選択したが、それは、当技術分野の当業者の意図、または判例、新技術の出現などによっても異なる。また、特定の場合、任意に選定された用語もあり、その場合、当該実施形態の説明部分において、詳細にその意味を記載する。従って、本実施形態で使用される用語は、単純な用語の名称ではない、その用語が有する意味と、本実施形態の全般にわたった内容とを基に定義されなければならない。 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
ホストプロセッサ110は、ホストプログラムが実行されるプロセッサであり、カーネル(kernel)を実行するための作業空間(work space)を定義することができる。ここで、該作業空間はまた、作業グループ(work group)に分割され、該作業グループは、複数個の作業アイテム(work item)を含み得る。ここで、該作業アイテムは、作業空間の各ポイントに該当し、作業の最小単位になる。例えば、ホストプロセッサ110は、汎用CPU(central processing unit)でもあるが、それに限定されるものではない。
The
コンピューティングデバイス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
一方、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
もし複数個のコアが含まれたコンピューティング装置においてカーネルが実行される場合、該コンピューティング装置は、コンピューティングデバイスと対応し、各コアは、コンピュートユニット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
一方、コンピューティング装置100は、ホストプロセッサ110から作業グループを受信することができる。
On the other hand, the
具体的には、ホストプロセッサ110は、各コアによって処理される作業グループの処理結果により、作業グループを各コアに動的に割り当てることができる。かような場合、ホストプロセッサ110は、作業グループを動的に割り当てる作業分配マネージャを含んでもよい。
Specifically, the
しかし、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
一方、コンピューティング装置に含まれた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
図2は、一実施形態による、コンピューティング装置の構成を示したブロック図である。 FIG. 2 is a block diagram showing a configuration of a computing device according to an embodiment.
一実施形態によるコンピューティング装置200は、複数個のコアを含み、複数個のコアは、コントロールコアグループ210とプロセシングコアグループ220とに分類される。図2に図示された構成要素以外に、他の汎用的な構成要素がさらに含まれてもよいということは、関連技術分野で当業者であるならば、理解することができるであろう。
The
一方、図1の構成と比較すれば、図2のコンピューティング装置200は、図1のコンピューティングデバイス120ないし140と対応する。
On the other hand, as compared with the configuration of FIG. 1, the
コントロールコアグループ210は、OpenCLカーネルを実行するための作業グループを、下位階層のコントロールコア及びプロセシングコアに割り当てる、コントロールコアの集合を意味する。ここで、該コントロールコアの仕様は、プロセシングコアの仕様とも同一であるが、作業グループを割り当てる機能を遂行することができる仕様であるならば、それに限定されるものではない。
The
プロセシングコアグループ220は、コントロールコアグループ210によって割り当てられた作業グループを処理し、作業グループの処理結果を出力する、プロセシングコアの集合を意味する。ここで、各プロセシングコアは、少なくとも1つのプロセシングエレメントを含み得る。例えば、該プロセシングコアは、メニーコアプロセッサの各コアとも対応するが、それに限定されるものではない。
The
ここで、コントロールコアグループ210に含まれた複数個のコントロールコアはまた、コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類される。コントロールコアグループ210が階層的に分類される例示は、図3を介して説明する。
Here, the plurality of control cores included in the
図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
最上位コントロールコア330は、OpenCLカーネルの実行情報が受信されることにより、作業グループを、下位階層のコントロールコアに割り当てることができる。図3としては、最上位コントロールコア330の下位階層のコントロールコアを、中間コントロールコアグループ340で表示したが、最上位コントロールコア330の下位階層のコントロールコアは、コントロールコアグループ320の階層構造により、中間コントロールコアグループ340及び最下位コントロールコアグループ350のうち少なくとも一つにもなるということは、当該技術分野の通常の技術者であるならば、理解することができるであろう。
The
一方、最上位コントロールコア330は、ホストプロセッサ310から、OpenCLカーネルの実行情報を受信することができる。具体的には、ホストプロセッサ310は、OpenCLカーネルが実行されることにより、OpenCLカーネルの実行情報を生成し、最上位コントロールコア330に送信することができる。ここで、OpenCLカーネルの実行情報は、OpenCLカーネルが実行されるための全体作業グループの数及びOpenCLカーネルが含まれたアプリケーション情報のうち少なくとも一つを含み得る。
On the other hand, the top-
最下位コントロールコアグループ350は、上位階層のコントロールコアによって割り当てられた作業グループについての情報を受信し、プロセシングコアに作業グループを割り当てる、複数個の最下位コントロールコアを含んでもよい。図3としては、最下位コントロールコアグループ350の上位階層のコントロールコアを、中間コントロールコアグループ340と表示したが、最下位コントロールコアの上位階層のコントロールコアは、コントロールコアグループ320の階層構造により、最上位コントロールコア330及び中間コントロールコアグループ340のうち少なくとも一つにもなるということは、当該技術分野の当業者であるならば、理解することができるであろう。
The lowest
中間コントロールコアグループ340は、最上位コントロールコア330によって割り当てられた作業グループに係わる情報を受信し、最下位コントロールコアグループ350に作業グループを割り当てる、複数個の中間コントロールコアを含んでもよい。また、中間コントロールコアグループ340は、さまざまな階層に分類される。図3を参照すれば、中間コントロールコアもまた、中間コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類される。かような場合、上位階層の中間コントロールコア341は、下位階層の中間コントロールコア342に作業グループを割り当てることができる。
The intermediate
図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
図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
一方、中間コントロールコアは、プロセシングコアグループに含まれたプロセシングコアの一部420に作業グループを割り当てることができる。例えば、図4において、中間コントロールコアグループ340が総4個の中間コントロールコアを含むのであるならば、1つの中間コントロールコアが作業グループを割り当てることができるプロセシングコアの数は、16個でもある。一方、中間コントロールコアが作業グループを割り当てることができるプロセシングコアの数は、各中間コントロールコアごとに異なってもよいということは、当該技術分野の当業者であるならば、理解することができるであろう。
On the other hand, the intermediate control core can assign a working group to a
一方、最下位コントロールコアは、上位階層のコントロールコアから割り当てられた作業グループについての情報を受信し、プロセシングコアに、直接作業グループを割り当てることができる。例えば、該最下位コントロールコアは、中間コントロールコアによって作業グループが割り当てられたプロセシングコアのうち一部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
また、一実施形態による最下位コントロールコアは、近接した複数個のプロセシングコアに作業グループを割り当てるために、ローカル性(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
図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
図5を参照すれば、最下位コントロールコア550は、プロセシングコア560に割り当てられた作業グループの処理が完了すると、他の最下位プロセシングコアに割り当てられた作業グループをプロセシングコア560にさらに割り当てることができる。また、最下位コントロールコア550は、中間コントロールコア540に作業グループの追加要請を送信することができる。ここで、最下位コントロールコア550は、プロセシングコア560から作業グループの追加要請を受信するか、あるいはプロセシングコア560の作業グループの進行状況を周期的に確認することにより、プロセシングコア560に割り当てられた作業グループの処理が完了したかいなかということを判断することができる。
Referring to FIG. 5, when the processing of the work group assigned to the
中間コントロールコア540は、最下位コントロールコア550に割り当てられた作業グループの処理が完了すると、他の最下位コントロールコアに割り当てられた作業グループを、最下位コントロールコア550にさらに割り当てる。また、中間コントロールコア540は、最上位コントロールコア530に作業グループの追加要請を送信することができる。
When the processing of the work group assigned to the
最上位コントロールコア530は、中間コントロールコア540に割り当てられた作業グループの処理が完了すると、他の中間コントロールコアに割り当てられた作業グループを中間コントロールコア540にさらに割り当てる。
When the processing of the work group assigned to the
また、最上位コントロールコア530は、OpenCLカーネルを実行するための作業グループが処理中であるとき、ホストプロセッサ510から新たなOpenCLカーネルの実行情報が受信されることにより、プロセシングコアグループに割り当てられた作業グループを再割り当てすることができる。
Further, the top-
図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
図6を参照すれば、最上位コントロールコア631は、中間コントロールコア632と結合されている。従って、最上位コントロールコア631は、作業グループを割り当てるために、プロセシングコア634に直接アクセスする代わりに、中間コントロールコア632または最下位コントロールコア633に、プロセシングコア634に割り当てられた作業グループについての情報を送信することができる。中間コントロールコア632は、最下位コントロールコア633と結合されているので、プロセシングコアグループの一部に割り当てられた作業グループについての情報を、最下位コントロールコア633に送信することができる。最下位コントロールコア633の場合、プロセシングコア634と結合されており、プロセシングコア634に割り当てられた作業グループに係わる情報を直接伝えることができる。
With reference to FIG. 6, the top-
一方、最上位コントロールコア631は、ホストプロセッサ610と結合され、OpenCLカーネルが実行されるとき、生成された作業グループについての情報を受信することができる。また、コントロールコアグループ及びプロセシングコアグループに含まれたコアは、メモリ620から作業グループを割り当てるのに必要な情報と、作業グループの処理に必要な情報とを受信することができる。
On the other hand, the top-
図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
また、プロセシングコア710は、作業グループを処理するコア711以外に、外部装置と通信するためのルータ712をさらに含んでもよい。プロセシングコア710は、ルータ712を介して、外部ルータ720と結合され、外部ルータ720は、最下位コントロールコア730とも結合される。従って、プロセシングコア710は、最下位コントロールコア730から、作業グループについての情報を受信することができる。図7は、プロセシングコア710に、単一コア711が含まれているように図示されているが、プロセシングコア710内のコア711の数は、それに限定されるものではない。
In addition to the
一方、ルータ712及び外部ルータ720は、NoC(network on chip)通信方式またはバス通信方式により、通信を行うことができるが、ルータ712と外部ルータ720との通信方式は、それらに限定されるものではない。
On the other hand, the
また、図7には、最下位コントロールコア730が4個のプロセシングコア710に作業グループを割り当てることができるように図示されているが、最下位コントロールコア730が作業グループを割り当てることができるプロセシングコア710の数は、それに限定されるものではない。
Further, in FIG. 7, the
図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
図9は、他の一実施形態によるコンピューティング装置について説明するための図である。 FIG. 9 is a diagram for explaining a computing device according to another embodiment.
一実施形態によるコンピューティング装置920は、作業グループについての情報を保存するバッファをさらに含んでもよい。該バッファが含まれたコンピューティング装置920は、作業グループについての情報をバッファから読み取るために、メモリにアクセスする回数を減らすことができる。一方、図9のホストプロセッサ910及び最上位コントロールコア930は、図5のホストプロセッサ510及び最上位コントロールコア530とそれぞれ対応するので、詳細な説明は省略する。
The
図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
また、中間コントロールコア940と最下位コントロールコア950は、第1バッファ965とも結合される。第1バッファ965は、中間コントロールコア940から伝達した、最下位コントロールコア950に割り当てられた作業グループの数と、当該作業グループのIDと、を保存することができる。従って、中間コントロールコア940は、メモリにアクセスせず、第1バッファ965に保存された最下位コントロールコア950に割り当てられた作業グループについての情報を読み取ることにより、最下位コントロールコア950の作業グループの処理結果を確認することができる。中間コントロールコア940は、最下位コントロールコア950に割り当てられた作業グループの処理が完了したと判断されると、他の最下位コントロールコアに割り当てられた作業グループを、最下位コントロールコア950にさらに割り当てることができる。一方、第1バッファ965は、最下位コントロールコア950に含まれてもよいが、第1バッファ965の位置は、それに限定されるものではない。
The
また、図9のコンピューティング装置920は、2個のバッファ965,970を含んでいるが、該バッファの数は、それに限定されるものではない。
Further, the
図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
例えば、最下位コントロールコア1010は、プロセシングコア1020に作業グループを割り当てた後、当該作業グループのIDと、プロセシングコア1020に割り当てられた作業グループの数と、をバッファ1030に伝達することができる。プロセシングコア1020は、作業グループを処理した後、バッファ1030に保存されたプロセシングコア1020に割り当てられた作業グループの数をアップデートすることができる。従って、最下位コントロールコア1010は、メモリにアクセスせず、バッファ1030を確認することにより、プロセシングコア1020の作業グループ処理結果を確認することができる。
For example, the
図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
図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
図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
具体的には、作業グループに含まれた作業アイテムの数より、プロセシングエレメントの数が少ない場合、プロセシングエレメントは、複数個の作業アイテムを処理することができる。かような場合、該プロセシングエレメントは、割り当てられた作業アイテムの処理が完了するたびに、完了した作業アイテムについての情報をメモリに保存した後、新たな作業アイテムにスイッチングすることができる。従って、該プロセシングエレメントが複数個の作業アイテムを処理するとき、該プロセシングエレメントがメモリにアクセスする回数が増加する。 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
従って、図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
図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
具体的には、コントロールコアグループ320は、最上位コントロールコア330及び最下位コントロールコアグループ350を含み得る。
Specifically, the
最上位コントロールコア330は、OpenCLカーネルの実行情報が受信されることにより、作業グループを、下位階層のコントロールコアに割り当て、割り当てられた作業グループについての情報を、下位階層のコントロールコアに送信することができる。最下位コントロールコアグループ350は、上位階層のコントロールコアによって割り当てられた作業グループについての情報を受信し、プロセシングコアに、作業グループを割り当てる、少なくとも1つの最下位コントロールコアを含み得る。
Upon receiving the execution information of the OpenCL kernel, the
また、コントロールコアグループ320は、最上位コントロールコア330によって割り当てられた作業グループに係わる情報を受信し、最下位コントロールコアグループ350に作業グループを割り当てる、複数個の中間コントロールコアが含まれた中間コントロールコアグループ340をさらに含んでもよい。
Further, the
一実施形態により、中間コントロールコアが、中間コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類されるとき、段階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
段階1620において、少なくとも1つのプロセシングコアが含まれたプロセシングコアグループ360は、コントロールコアグループ320によって割り当てられた作業グループを処理することができる。
In
段階1630において、プロセシングコアグループ360は、作業グループの処理結果を出力することができる。
At
図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
段階1710において、最上位コントロールコア330は、OpenCLカーネルが実行されることにより、ホストプロセッサ310によって生成されたOpenCLカーネルの実行情報を受信することができる。かような場合、最上位コントロールコア330は、OpenCLカーネルを実行するための作業グループが処理中であるとき、ホストプロセッサ310から新たなOpenCLカーネルの実行情報が受信されることにより、プロセシングコアグループ360に割り当てられた作業グループを再割り当てすることができる。ここで、OpenCLカーネルの実行情報は、OpenCLカーネルが実行されるための全体作業グループの数、及びOpenCLカーネルが含まれたアプリケーション情報を含んでもよい。
In
また、一実施形態によるOpenCLカーネルを処理する方法は、ホストプロセッサ310が、コントロールコアグループ320によるプロセシングコアグループ360に対する作業グループ割り当てが完了すると、OpenCLカーネルの実行を中断し、割り当て結果を収集して出力する段階をさらに含んでもよい。
Further, in the method of processing the OpenCL kernel according to one embodiment, when the
また、一実施形態による、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
本実施形態は、コンピュータによって実行されるプログラムモジュールのような、コンピュータによって実行可能な命令語を含む記録媒体の形態でも具現される。コンピュータ可読媒体は、コンピュータによってアクセスされる任意の可用媒体でもあり、揮発性及び不揮発性の媒体、分離型及び非分離型の媒体をいずれも含む。コンピュータ記録媒体は、コンピュータ可読命令語、データ構造、プログラムモジュール、またはその他データのような情報保存のための任意の方法または技術によって具現された揮発性及び不揮発性、分離型及び非分離型の媒体をいずれも含む。 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
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.
前記下位階層のコントロールコア及び前記プロセシングコアによる前記作業グループの処理結果に従って、前記コントロールコアグループの他の下位階層のコントロールコア及び他のプロセシングコアに割り当てられた作業グループを再割り当てすることを特徴とする請求項1に記載のコンピューティング装置。 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.
前記最上位コントロールコアによって割り当てられた作業グループに係わる情報を受信し、前記最下位コントロールコアグループに前記作業グループを割り当てる、複数個の中間コントロールコアを含んだ中間コントロールコアグループをさらに含み、
前記中間コントロールコアは、前記中間コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数によって階層的に分類され、前記コントロールコアグループの上位階層の中間コントロールコアは、下位階層の中間コントロールコアに前記作業グループを割り当てることを特徴とする請求項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カーネルの実行情報が受信されることにより、前記プロセシングコアグループに割り当てられた前記作業グループを再割り当てすることを特徴とする請求項4に記載のコンピューティング装置。 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カーネルの実行を中断し、割り当て結果を収集して出力することを特徴とする請求項4に記載のコンピューティング装置。 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.
前記下位階層のコントロールコア及び前記プロセシングコアに割り当てられた作業グループの数と、前記割り当てられた作業グループのIDとのうち少なくとも一つを含むことを特徴とする請求項8に記載のコンピューティング装置。 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. ..
前記プロセシングコアのうちの少なくとも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 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. ..
前記下位階層のコントロールコア及び前記プロセシングコアに割り当てられた作業グループの数と、前記割り当てられた作業グループの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.
A computer-readable recording medium on which a program for executing the method according to any one of claims 10 to 17 is recorded.
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)
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)
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 |
-
2016
- 2016-12-27 KR KR1020160180133A patent/KR102592330B1/en active IP Right Grant
-
2017
- 2017-10-18 US US15/787,219 patent/US10503557B2/en active Active
- 2017-11-21 EP EP17202801.1A patent/EP3343370A1/en not_active Ceased
- 2017-11-28 CN CN201711216401.0A patent/CN108241508B/en active Active
- 2017-12-13 JP JP2017238434A patent/JP6951962B2/en active Active
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 |