[go: up one dir, main page]
More Web Proxy on the site http://driver.im/

JP2018106709A - OpenCLカーネルを処理する方法、及びそれを遂行するコンピューティング装置 - Google Patents

OpenCLカーネルを処理する方法、及びそれを遂行するコンピューティング装置 Download PDF

Info

Publication number
JP2018106709A
JP2018106709A JP2017238434A JP2017238434A JP2018106709A JP 2018106709 A JP2018106709 A JP 2018106709A JP 2017238434 A JP2017238434 A JP 2017238434A JP 2017238434 A JP2017238434 A JP 2017238434A JP 2018106709 A JP2018106709 A JP 2018106709A
Authority
JP
Japan
Prior art keywords
group
processing
control core
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.)
Granted
Application number
JP2017238434A
Other languages
English (en)
Other versions
JP6951962B2 (ja
JP2018106709A5 (ja
Inventor
エガー ベーンハート
Egger Bernhard
エガー ベーンハート
洙 林 烏
Su-Rim Oh
洙 林 烏
永 顯 趙
Young-Hyun Cho
永 顯 趙
劉 東 勳
Dong Hun Yoo
東 勳 劉
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
Seoul National University Industry Foundation
SNU R&DB Foundation
Original Assignee
Samsung Electronics Co Ltd
Seoul National University Industry Foundation
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 Industry Foundation, Seoul National University R&DB Foundation filed Critical Samsung Electronics Co Ltd
Publication of JP2018106709A publication Critical patent/JP2018106709A/ja
Publication of JP2018106709A5 publication Critical patent/JP2018106709A5/ja
Application granted granted Critical
Publication of JP6951962B2 publication Critical patent/JP6951962B2/ja
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

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/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/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/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/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/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)

Abstract

【課題】OpenCLカーネルを処理する方法を提供する。
【解決手段】方法は、OpenCL(open computing language)カーネルを実行するための作業グループを、下位階層のコントロールコア(control core)及びプロセシングコア(processing core)に割り当てるコントロールコアグループ及び少なくとも1つのプロセシングコアを含み、コントロールコアグループによって割り当てられた作業グループを処理し、作業グループの処理結果を出力する、プロセシングコアグループを含む。コントロールコアグループに含まれた複数個のコントロールコアは、コントロールコアにより作業グループが割り当てられるプロセシングコアの数により、階層的に分類される。
【選択図】図3

Description

本開示は、OpenCLカーネルを処理する方法、及びそれを遂行するコンピューティング装置に関する。
コンピューティング装置は、アプリケーションの性能条件を満足するために、単一集積回路に、さまざまなコアまたはプロセッサを統合する形態で発展している。例えば、マルチコアプロセッサ(multi core processor)は、演算機能を有した2個以上のコアが、1つのプロセッサに集積されたものを意味する。また、さらに多数(一般的に、16個以上)のコアが1つのプロセッサに集積されたメニーコアプロセッサ(many core processor)も開発されている。マルチコアプロセッサとメニーコアプロセッサは、タブレットPC(personal computer)、携帯電話、PDA(personal digital assistant)、ラップトップ、メディアプレーヤ、GPS(global position system)装置、電子書籍端末機、MP3プレーヤ、デジタルカメラを含んだ携帯機器やTV(television)に搭載されるマルチメディアチップを含んだ埋め込み装置に搭載される。
一方、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プログラムを効率的に処理するための多様な試みが進められている。
本発明が解決しようとする課題は、階層的に分類されたコントロールコアグループが、OpenCLカーネル(kernel)を実行するための作業グループをプロセシングコアグループに割り当てる、OpenCLカーネルを処理する方法、及びそれを遂行するコンピューティング装置を提供することである。本実施形態がなそうとする技術的課題は、前述のような技術的課題に限定されるものではなく、以下の実施形態から、他の技術的課題が類推されもする。
一側面によれば、OpenCLカーネルを実行するための作業グループを、下位階層のコントロールコア(control core)及びプロセシングコア(processing core)に割り当てるコントロールコアグループ;及び少なくとも1つの前記プロセシングコアを含み、前記コントロールコアグループによって割り当てられた前記作業グループを処理し、前記作業グループの処理結果を出力する、プロセシングコアグループ;を含み、前記コントロールコアグループに含まれた複数個のコントロールコアは、前記コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数により、階層的に分類される。
他の側面によれば、OpenCLカーネルを処理する方法は、コントロールコアグループにより、OpenCLカーネルを実行するための作業グループを、下位階層のコントロールコア及びプロセシングコアに割り当てる段階;少なくとも1つの前記プロセシングコアが含まれたプロセシングコアグループにより、前記割り当てられた前記作業グループを処理する段階;及び前記プロセシングコアグループにより、前記作業グループの処理結果を出力する段階;を含み、前記コントロールコアグループに含まれた複数個のコントロールコアは、前記コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数により、階層的に分類される。
さらに他の側面によれば、前記方法をコンピュータで実行させるためのプログラムを記録したコンピュータ読み取り可能な記録媒体を提供する。
本発明によれば、プロセシングコアに作業グループを割り当てるコントロールコアが階層的に分類されているために、作業グループを効率的に分配することができる。
OpenCLプラットフォームについて説明するための図である。 一実施形態による、コンピューティング装置の構成を示したブロック図である。 他の一実施形態による、コンピューティング装置の構成を示したブロック図である。 一実施形態によって、コントロールコアによって作業グループが割り当てられたプロセシングコアの数について説明するための図である。 一実施形態による、OpenCLカーネルを実行するための作業グループを動的に割り当てる方法について説明するための図である。 一実施形態による、コンピューティング装置の構造について説明するための図である。 他の一実施形態による、最下位コントロールコア及びプロセシングコアについて説明するための図である。 一実施形態による、プロセシングコアで実行されるOpenCLカーネルコードである。 他の一実施形態によるコンピューティング装置について説明するための図である。 一実施形態によるバッファについて説明するための図である。 一実施形態による、バッファが追加された場合、プロセシングコアで実行されるOpenCLカーネルコードである。 他の一実施形態による、バッファが追加された場合、OpenCL作業グループを実行するOpenCLカーネルコードである。 一実施形態によるコンピューティング装置の性能を検証するための実験環境を示した図である。 一実施形態によるコンピューティング装置の性能を検証するための実験環境を示した図である。 一実施形態によるコンピューティング装置の性能の検証結果について説明するための図である。 一実施形態により、OpenCLカーネルを処理する方法を図示したフローチャートである。 一実施形態により、OpenCLカーネルを処理する方法を図示した詳細フローチャートである。
本実施形態で使用される用語は、本実施形態での機能を考慮しながら、可能な限り、現在汎用される一般的な用語を選択したが、それは、当技術分野の当業者の意図、または判例、新技術の出現などによっても異なる。また、特定の場合、任意に選定された用語もあり、その場合、当該実施形態の説明部分において、詳細にその意味を記載する。従って、本実施形態で使用される用語は、単純な用語の名称ではない、その用語が有する意味と、本実施形態の全般にわたった内容とを基に定義されなければならない。
実施形態に係わる説明において、ある部分が他の部分と結合されているとするとき、それは、直接に結合されている場合だけではなく、その中間に、他の構成要素を挟んで電気的に結合されている場合も含む。また、ある部分がある構成要素を含むとするとき、それは、特別に反対となる記載がない限り、他の構成要素を除くものではなく、他の構成要素をさらに含むということを意味する。また、実施形態に記載された「…部」、「…モジュール」の用語は、少なくとも1つの機能や動作を処理する単位を意味し、それは、ハードウェアまたはソフトウェアによっても具現され、ハードウェアとソフトウェアとの結合によっても具現される。
本実施形態で使用される「構成される」または「含む」というような用語は、必ずしも、明細書上に記載されたさまざまな構成要素、またはさまざまな段階をいずれも含むと解釈されるものではなく、そのうち一部構成要素または一部段階が含まれないこともあり、または、追加的な構成要素または段階をさらに含んでもよいと解釈されなければならない。
下記実施形態に係わる説明は、権利範囲を限定すると解釈されるものではなく、当該技術分野の当業者が容易に類推することができることは、実施形態の権利範囲に属すると解釈されなければならないのである。以下、添付された図面を参照しつつ、ただ例示のための実施形態について詳細に説明する。
図1は、OpenCL(open computing language)プラットフォームについて説明するための図である。
OpenCLは、互いに異なる形態のコンピューティングデバイスが同時に動作するように具現することができるプログラミング言語中の一つである。従って、OpenCLによって作成されたプログラムは、複数のプラットフォームで同時に動作することができる。
図1を参照すれば、OpenCLプラットフォーム100は、ホストプロセッサ(host processor)110と、少なくとも1つのコンピューティングデバイス(CD:computing device)120ないし140を含んでもよい。
ホストプロセッサ110は、ホストプログラムが実行されるプロセッサであり、カーネル(kernel)を実行するための作業空間(work space)を定義することができる。ここで、該作業空間はまた、作業グループ(work group)に分割され、該作業グループは、複数個の作業アイテム(work item)を含み得る。ここで、該作業アイテムは、作業空間の各ポイントに該当し、作業の最小単位になる。例えば、ホストプロセッサ110は、汎用CPU(central processing unit)でもあるが、それに限定されるものではない。
コンピューティングデバイス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つのコアを有するプロセッサであるならば、それらに限定されるものではない。
一方、OpenCLプログラムは、1つのホストプログラムと、少なくとも1つのカーネルとを含み得る。該ホストプログラムは、ホストプロセッサ110によって実行され、ホストプロセッサ110から、コマンド(command)を介して、コンピューティングデバイス120ないし140に計算を指示したり、コンピューティングデバイス120ないし140に装着されたメモリを管理したりすることができる。ここで、該カーネルは、コンピューティングデバイス120ないし140で実行されるプログラムを意味し、OpenCLカーネルまたはカーネルコードによっても表現される。
もし複数個のコアが含まれたコンピューティング装置においてカーネルが実行される場合、該コンピューティング装置は、コンピューティングデバイスと対応し、各コアは、コンピュートユニット121ないし122と対応する。言い替えれば、該コンピューティング装置のコアは、作業グループを割り当てられて処理することができ、コアに含まれたプロセシングエレメント10ないし40は、作業グループに含まれた各作業アイテムを処理することができる。
一方、コンピューティング装置100は、ホストプロセッサ110から作業グループを受信することができる。
具体的には、ホストプロセッサ110は、各コアによって処理される作業グループの処理結果により、作業グループを各コアに動的に割り当てることができる。かような場合、ホストプロセッサ110は、作業グループを動的に割り当てる作業分配マネージャを含んでもよい。
しかし、1つの作業分配マネージャが、複数個のコアに作業グループを動的に割り当てる場合、コアの数により、作業分配マネージャが遂行しなければならない分配作業の量が増加し、全体コンピューティング装置の性能が低下してしまう。また、作業分配マネージャが、複数個のコアに作業グループを静的に割り当てる場合、各作業グループの処理時間が異なり、全体コンピューティング装置の作業効率が低下してしまう。併せて、ホストプロセッサ110に作業分配マネージャが含まれていると、ホストプロセッサ110の負担が増大してしまう。
一方、コンピューティング装置に含まれた1つのコアが作業分配マネージャである場合、当該コアは、コンピュートユニット121ないし122に活用されない。従って、複数個のコアが含まれたコンピューティング装置においてOpenCLカーネルが実行される場合に作業グループを効率的に割り当てることができる方法が要求される。
図2は、一実施形態による、コンピューティング装置の構成を示したブロック図である。
一実施形態によるコンピューティング装置200は、複数個のコアを含み、複数個のコアは、コントロールコアグループ210とプロセシングコアグループ220とに分類される。図2に図示された構成要素以外に、他の汎用的な構成要素がさらに含まれてもよいということは、関連技術分野で当業者であるならば、理解することができるであろう。
一方、図1の構成と比較すれば、図2のコンピューティング装置200は、図1のコンピューティングデバイス120ないし140と対応する。
コントロールコアグループ210は、OpenCLカーネルを実行するための作業グループを、下位階層のコントロールコア及びプロセシングコアに割り当てる、コントロールコアの集合を意味する。ここで、該コントロールコアの仕様は、プロセシングコアの仕様とも同一であるが、作業グループを割り当てる機能を遂行することができる仕様であるならば、それに限定されるものではない。
プロセシングコアグループ220は、コントロールコアグループ210によって割り当てられた作業グループを処理し、作業グループの処理結果を出力する、プロセシングコアの集合を意味する。ここで、各プロセシングコアは、少なくとも1つのプロセシングエレメントを含み得る。例えば、該プロセシングコアは、メニーコアプロセッサの各コアとも対応するが、それに限定されるものではない。
ここで、コントロールコアグループ210に含まれた複数個のコントロールコアはまた、コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類される。コントロールコアグループ210が階層的に分類される例示は、図3を介して説明する。
図3は、他の一実施形態による、コンピューティング装置の構成を示したブロック図である。図3を参照すれば、コンピューティング装置300は、コントロールコアグループ320とプロセシングコアグループ360とを含み得る。ここで、コントロールコアグループ320は、1つの最上位(root;ルート)コントロールコア330と最下位(leaf;リーフ)コントロールコアグループ350とを含み得る。また、コンピューティング装置300は、コントロールコアグループ320の階層構造により、中間コントロールコアグループ340をさらに含んでもよい。一方、図3のプロセシングコアグループ360は、図2のプロセシングコアグループ220と対応するので、詳細な説明は省略する。
最上位コントロールコア330は、OpenCLカーネルの実行情報が受信されることにより、作業グループを、下位階層のコントロールコアに割り当てることができる。図3としては、最上位コントロールコア330の下位階層のコントロールコアを、中間コントロールコアグループ340で表示したが、最上位コントロールコア330の下位階層のコントロールコアは、コントロールコアグループ320の階層構造により、中間コントロールコアグループ340及び最下位コントロールコアグループ350のうち少なくとも一つにもなるということは、当該技術分野の通常の技術者であるならば、理解することができるであろう。
一方、最上位コントロールコア330は、ホストプロセッサ310から、OpenCLカーネルの実行情報を受信することができる。具体的には、ホストプロセッサ310は、OpenCLカーネルが実行されることにより、OpenCLカーネルの実行情報を生成し、最上位コントロールコア330に送信することができる。ここで、OpenCLカーネルの実行情報は、OpenCLカーネルが実行されるための全体作業グループの数及びOpenCLカーネルが含まれたアプリケーション情報のうち少なくとも一つを含み得る。
最下位コントロールコアグループ350は、上位階層のコントロールコアによって割り当てられた作業グループについての情報を受信し、プロセシングコアに作業グループを割り当てる、複数個の最下位コントロールコアを含んでもよい。図3としては、最下位コントロールコアグループ350の上位階層のコントロールコアを、中間コントロールコアグループ340と表示したが、最下位コントロールコアの上位階層のコントロールコアは、コントロールコアグループ320の階層構造により、最上位コントロールコア330及び中間コントロールコアグループ340のうち少なくとも一つにもなるということは、当該技術分野の当業者であるならば、理解することができるであろう。
中間コントロールコアグループ340は、最上位コントロールコア330によって割り当てられた作業グループに係わる情報を受信し、最下位コントロールコアグループ350に作業グループを割り当てる、複数個の中間コントロールコアを含んでもよい。また、中間コントロールコアグループ340は、さまざまな階層に分類される。図3を参照すれば、中間コントロールコアもまた、中間コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類される。かような場合、上位階層の中間コントロールコア341は、下位階層の中間コントロールコア342に作業グループを割り当てることができる。
図4は、一実施形態によって、コントロールコアによって作業グループが割り当てられるプロセシングコアの数について説明するための図である。
一実施形態によれば、コントロールコアグループ210に含まれたコントロールコアは、コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類される。
図4は、8x8プロセシングコアが含まれたプロセシングコアグループを示した図である。図4を参照すれば、最上位コントロールコア330によって作業グループが割り当てられるプロセシングコア410は、全体プロセシングコアグループでもある。
一方、中間コントロールコアは、プロセシングコアグループに含まれたプロセシングコアの一部420に作業グループを割り当てることができる。例えば、図4において、中間コントロールコアグループ340が総4個の中間コントロールコアを含むのであるならば、1つの中間コントロールコアが作業グループを割り当てることができるプロセシングコアの数は、16個でもある。一方、中間コントロールコアが作業グループを割り当てることができるプロセシングコアの数は、各中間コントロールコアごとに異なってもよいということは、当該技術分野の当業者であるならば、理解することができるであろう。
一方、最下位コントロールコアは、上位階層のコントロールコアから割り当てられた作業グループについての情報を受信し、プロセシングコアに、直接作業グループを割り当てることができる。例えば、該最下位コントロールコアは、中間コントロールコアによって作業グループが割り当てられたプロセシングコアのうち一部430に作業グループを割り当てることができる。従って、最下位コントロールコアによって作業グループが割り当てられるプロセシングコアの数は、中間コントロールコアによって作業グループが割り当てられるプロセシングコアの数より少ない。
また、一実施形態による最下位コントロールコアは、近接した複数個のプロセシングコアに作業グループを割り当てるために、ローカル性(locality)を有する作業グループを割り当てる場合、少ないオーバーヘッドでもって、作業グループのローカル性を提供することができる。
また、一実施形態によるコンピューティング装置200は、プロセシングコアグループ220以外に、コントロールコアグループ210を追加して含んでいるために、全てのプロセシングコアが作業グループを処理することができる。また、コントロールコアグループ210が階層的に分類されているために、作業グループを動的に割り当てるのに効率的である。階層的に分類されたコントロールコアが作業グループを動的に割り当てる方法は、図5を介して説明する。
図5は、一実施形態による、OpenCLカーネルを実行するための作業グループを動的に割り当てる方法について説明するための図である。
一実施形態によるコントロールコアグループは、下位階層のコントロールコア及びプロセシングコアの作業グループの処理結果により、他の下位階層のコントロールコア及び他のプロセシングコアに割り当てられた作業グループを再割り当てすることができる。一方、図5のホストプロセッサ510は、図3のホストプロセッサ310と対応するので、詳細な説明は省略する。
図5を参照すれば、最下位コントロールコア550は、プロセシングコア560に割り当てられた作業グループの処理が完了すると、他の最下位プロセシングコアに割り当てられた作業グループをプロセシングコア560にさらに割り当てることができる。また、最下位コントロールコア550は、中間コントロールコア540に作業グループの追加要請を送信することができる。ここで、最下位コントロールコア550は、プロセシングコア560から作業グループの追加要請を受信するか、あるいはプロセシングコア560の作業グループの進行状況を周期的に確認することにより、プロセシングコア560に割り当てられた作業グループの処理が完了したかいなかということを判断することができる。
中間コントロールコア540は、最下位コントロールコア550に割り当てられた作業グループの処理が完了すると、他の最下位コントロールコアに割り当てられた作業グループを、最下位コントロールコア550にさらに割り当てる。また、中間コントロールコア540は、最上位コントロールコア530に作業グループの追加要請を送信することができる。
最上位コントロールコア530は、中間コントロールコア540に割り当てられた作業グループの処理が完了すると、他の中間コントロールコアに割り当てられた作業グループを中間コントロールコア540にさらに割り当てる。
また、最上位コントロールコア530は、OpenCLカーネルを実行するための作業グループが処理中であるとき、ホストプロセッサ510から新たなOpenCLカーネルの実行情報が受信されることにより、プロセシングコアグループに割り当てられた作業グループを再割り当てすることができる。
図6は、一実施形態による、コンピューティング装置の構造について説明するための図である。図6を参照すれば、コンピューティング装置630は、ホストプロセッサ610及びメモリ620と結合される。また、コンピューティング装置630は、最上位コントロールコア631、最下位コントロールコア633及びプロセシングコア634を含むことができ、コントロールコアグループの階層構造により、中間コントロールコア632をさらに含んでもよい。
図6を参照すれば、最上位コントロールコア631は、中間コントロールコア632と結合されている。従って、最上位コントロールコア631は、作業グループを割り当てるために、プロセシングコア634に直接アクセスする代わりに、中間コントロールコア632または最下位コントロールコア633に、プロセシングコア634に割り当てられた作業グループについての情報を送信することができる。中間コントロールコア632は、最下位コントロールコア633と結合されているので、プロセシングコアグループの一部に割り当てられた作業グループについての情報を、最下位コントロールコア633に送信することができる。最下位コントロールコア633の場合、プロセシングコア634と結合されており、プロセシングコア634に割り当てられた作業グループに係わる情報を直接伝えることができる。
一方、最上位コントロールコア631は、ホストプロセッサ610と結合され、OpenCLカーネルが実行されるとき、生成された作業グループについての情報を受信することができる。また、コントロールコアグループ及びプロセシングコアグループに含まれたコアは、メモリ620から作業グループを割り当てるのに必要な情報と、作業グループの処理に必要な情報とを受信することができる。
図7は、他の一実施形態による、最下位コントロールコア及びプロセシングコアについて説明するための図である。
一実施形態による1つの最下位コントロールコア730は、複数個のプロセシングコア710に作業グループを割り当てることができる。例えば、図7を参照すれば、最下位コントロールコア730は、4個のプロセシングコア710に結合され、直接作業グループを割り当て、作業グループの処理結果、または作業グループの追加要請を受信することができる。
また、プロセシングコア710は、作業グループを処理するコア711以外に、外部装置と通信するためのルータ712をさらに含んでもよい。プロセシングコア710は、ルータ712を介して、外部ルータ720と結合され、外部ルータ720は、最下位コントロールコア730とも結合される。従って、プロセシングコア710は、最下位コントロールコア730から、作業グループについての情報を受信することができる。図7は、プロセシングコア710に、単一コア711が含まれているように図示されているが、プロセシングコア710内のコア711の数は、それに限定されるものではない。
一方、ルータ712及び外部ルータ720は、NoC(network on chip)通信方式またはバス通信方式により、通信を行うことができるが、ルータ712と外部ルータ720との通信方式は、それらに限定されるものではない。
また、図7には、最下位コントロールコア730が4個のプロセシングコア710に作業グループを割り当てることができるように図示されているが、最下位コントロールコア730が作業グループを割り当てることができるプロセシングコア710の数は、それに限定されるものではない。
図8は、一実施形態による、プロセシングコアで実行されるOpenCLカーネルコードである。図8は、行列aの要素と、行列bの要素とを加え、その結果を行列cに保存するカーネルコードである。行列aと行列bは、二次元マトリックスであるので、図8のカーネルコードを実行するためには、二次元作業グループ空間が必要である。言い替えれば、行列aと行列bとの各要素は、作業空間の各ポイントである作業アイテムにもなる。
一方、図6のコンピューティング装置630において、図8のカーネルコードが実行される場合、プロセシングコア634は、行列aの要素と、行列bの要素とを読み出すために、メモリ620にアクセスすることができる。ここで、関数get_global_id()は、メモリ620からカーネルコードのインデックスを読み出す関数である。従って、プロセシングコア634が、図8のカーネルコードを実行するとき、複数個の作業アイテムが並列的に処理されるが、作業アイテムを読み出すために、メモリ620に何回もアクセスしなければならないので、カーネルコードを処理する時間が長くなる。
図9は、他の一実施形態によるコンピューティング装置について説明するための図である。
一実施形態によるコンピューティング装置920は、作業グループについての情報を保存するバッファをさらに含んでもよい。該バッファが含まれたコンピューティング装置920は、作業グループについての情報をバッファから読み取るために、メモリにアクセスする回数を減らすことができる。一方、図9のホストプロセッサ910及び最上位コントロールコア930は、図5のホストプロセッサ510及び最上位コントロールコア530とそれぞれ対応するので、詳細な説明は省略する。
図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の位置は、それに限定されるものではない。
また、中間コントロールコア940と最下位コントロールコア950は、第1バッファ965とも結合される。第1バッファ965は、中間コントロールコア940から伝達した、最下位コントロールコア950に割り当てられた作業グループの数と、当該作業グループのIDと、を保存することができる。従って、中間コントロールコア940は、メモリにアクセスせず、第1バッファ965に保存された最下位コントロールコア950に割り当てられた作業グループについての情報を読み取ることにより、最下位コントロールコア950の作業グループの処理結果を確認することができる。中間コントロールコア940は、最下位コントロールコア950に割り当てられた作業グループの処理が完了したと判断されると、他の最下位コントロールコアに割り当てられた作業グループを、最下位コントロールコア950にさらに割り当てることができる。一方、第1バッファ965は、最下位コントロールコア950に含まれてもよいが、第1バッファ965の位置は、それに限定されるものではない。
また、図9のコンピューティング装置920は、2個のバッファ965,970を含んでいるが、該バッファの数は、それに限定されるものではない。
図10は、一実施形態によるバッファについて説明するための図である。図10を参照すれば、バッファ1030は、最下位コントロールコア1010とプロセシングコア1020とに結合される。また、バッファ1030は、プロセシングコア1020に割り当てられた作業グループのIDと、プロセシングコア1020に割り当てられた作業グループの数と、を保存することができる。
例えば、最下位コントロールコア1010は、プロセシングコア1020に作業グループを割り当てた後、当該作業グループのIDと、プロセシングコア1020に割り当てられた作業グループの数と、をバッファ1030に伝達することができる。プロセシングコア1020は、作業グループを処理した後、バッファ1030に保存されたプロセシングコア1020に割り当てられた作業グループの数をアップデートすることができる。従って、最下位コントロールコア1010は、メモリにアクセスせず、バッファ1030を確認することにより、プロセシングコア1020の作業グループ処理結果を確認することができる。
図11は、一実施形態による、バッファが追加された場合、プロセシングコアで実行されるOpenCLカーネルコードである。図11のカーネルコードは、バッファ1030を利用して、図8のカーネルコードと同一機能を遂行するコードである。言い替えれば、図11のカーネルコードは、バッファ1030が追加されたプロセシングコア1020によって実行される、行列aの要素と、行列bの要素とを加え、その結果を行列cに保存するカーネルコードである。
図11を参照すれば、関数read_wg_id()は、プロセシングコア1020に割り当てられた作業グループのIDをバッファ1030から読み取る関数である。従って、図11のカーネルコードは、関数read_wg_id()を介して、バッファ1030に保存された作業グループのIDを、カーネルコードのインデックスとして活用するために、メモリにアクセスせずとも、各プロセシングコア1020が並列的に作業グループを処理することができる。
図12は、他の一実施形態による、バッファが追加された場合、OpenCL作業グループを実行するOpenCLカーネルコードである。
一実施形態によって、プロセシングコア1020が少なくとも1つのプロセシングエレメントを含み、作業グループに含まれた作業アイテムの数より、プロセシングエレメントの数が少なければ、コンピューティング装置は、該プロセシングエレメントが該作業アイテムを直列化し、順次に実行するように、OpenCLカーネルを変換することができる。
具体的には、作業グループに含まれた作業アイテムの数より、プロセシングエレメントの数が少ない場合、プロセシングエレメントは、複数個の作業アイテムを処理することができる。かような場合、該プロセシングエレメントは、割り当てられた作業アイテムの処理が完了するたびに、完了した作業アイテムについての情報をメモリに保存した後、新たな作業アイテムにスイッチングすることができる。従って、該プロセシングエレメントが複数個の作業アイテムを処理するとき、該プロセシングエレメントがメモリにアクセスする回数が増加する。
プロセシングエレメントによるメモリのアクセス回数を減らすために、コンピューティング装置は、作業アイテム合体(work item coalescing)方法に基づいて、プロセシングエレメントが、複数個の作業アイテムを直列化して処理するように、カーネルコードを変換することができる。
図12を参照すれば、関数read_num_wgs()は、プロセシングコア1020に連続的に割り当てられた作業グループの数を、バッファ1030から読み取る関数である。
従って、図12のカーネルコードは、バッファ1030に保存された情報を読み取る関数read_wg_id()及び関数read_num_wgs()に基づいて、カーネルコードのインデックスを設定するために、プロセシングコア1020によるメモリアクセス回数を減少させることができる。また、カーネルコードを取り囲むループが適用されており、実行されなければならないカーネルコードを、1つの作業アイテムを実行するたびに呼び出す必要なしに、1回のカーネルコードの呼び出しにより、複数個の作業アイテムを順次に実行することができる。従って、作業アイテム合体方法を適用しない場合に比べ、プロセシングコア1020のメモリアクセス回数を減少させることができる。
図13及び図14は、一実施形態によるコンピューティング装置の性能を検証するための実験環境を示した図である。
図13は、一実施形態によるコンピューティング装置に、異なるスケジューリング技法を適用した場合、作業グループを割り当てるのに必要な遅延時間を示した図である。ここでは、コンピューティング装置に含まれたプロセシンググループは、96個のプロセシングコアを含み、該コンピューティング装置に含まれたコントロールコアグループは、1個の最上位コントロールコア、4個の中間コントロールコア、及び16個の最下位コントロールコアを含むと仮定する。また、コントロールコアとプロセシングコアは、それぞれ50MHz及び500MHzのクロック周波数(clock frequency)によって動作すると仮定する。
図13の集中的スケジューリング方法は、単一コントロールコアが、全てのプロセシングコアに、作業グループを動的に割り当てるスケジューリング方法を意味する。分散スケジューリング方法は、階層的に分類されたコントロールコアが、プロセシングコアに作業グループを動的に割り当てるスケジューリング方法を意味する。また、ハードウェアサポート項目は、コンピューティング装置がバッファを含むか否かということを示す項目である。
図13を参照すれば、分散スケジューリング方法が適用されたコンピューティング装置の遅延時間は、集中的スケジューリング方法が適用されたコンピューティング装置の遅延時間より短い。従って、階層的に分類されたさまざまなコントロールコアが、96個のプロセシングコアに作業グループを動的に割り当てる時間が、単一コントロールコアが作業グループを動的に割り当てる時間より短くなるということが分かる。
一方、同一スケジューリング方法が適用された場合、バッファが追加されたコンピューティング装置の遅延時間は、バッファなしのコンピューティング装置の遅延時間より短い。従って、バッファが追加されたコンピューティング装置は、メモリにアクセスする回数が減少されるために、作業グループを処理する遅延時間が短くなるということが分かる。
図14を参照すれば、コンピューティング装置の性能を検証するために、異なる4個のアプリケーションが実行されると仮定する。具体的には、PI estimationアプリケーションは、他のアプリケーションに比べ、割り当てられる作業グループの処理時間が長いという特性を有する。Small buffer copy及びBig buffer copyアプリケーションは、他のアプリケーションに比べ、メモリアクセス回数が多いという特性を有する。また、SpMVアプリケーションは、他のアプリケーションに比べ、メモリアクセスが一定ではなく、割り当てられた作業グループを処理する時間が非均一であるという特性を有する。
図15は、一実施形態によるコンピューティング装置の性能検証結果について説明するための図である。図15は、理想的なスケジューリングの性能に基づいて、コンピューティング装置によって、4個のアプリケーションが実行されるとき、スケジューリング性能を正規化した結果である。
図15を参照すれば、集中的スケジューリング方法が適用されたコンピューティング装置のスケジューリング性能は、分散スケジューリング方法が適用されたコンピューティング装置のスケジューリング性能より全体的に低いということが分かる。具体的には、SpMVアプリケーションの場合、集中的スケジューリング方法が適用されたコンピューティング装置と、分散スケジューリング方法が適用されたコンピューティング装置とのスケジューリング性能差が大きいということが分かる。
また、バッファが追加されたコンピューティング装置が、分散スケジューリングに基づいて、作業グループのスケジューリングを行う場合、分散スケジューリングだけ適用されたコンピューティング装置に比べ、スケジューリング性能がさらに良好であるということが分かる。具体的には、メモリアクセス回数が多いSmall buffer copyアプリケーション及びBig buffer copyアプリケーションが、バッファ有無によるスケジューリング性能差が大きいということが分かる。
図16は、一実施形態による、OpenCLカーネルを処理する方法を図示したフローチャートである。
段階1610において、コントロールコアグループ320は、OpenCLカーネルを実行するための作業グループを、下位階層のコントロールコア及びプロセシングコアに割り当てることができる。ここで、コントロールコアグループ320に含まれた複数個のコントロールコアは、コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類される。
具体的には、コントロールコアグループ320は、最上位コントロールコア330及び最下位コントロールコアグループ350を含み得る。
最上位コントロールコア330は、OpenCLカーネルの実行情報が受信されることにより、作業グループを、下位階層のコントロールコアに割り当て、割り当てられた作業グループについての情報を、下位階層のコントロールコアに送信することができる。最下位コントロールコアグループ350は、上位階層のコントロールコアによって割り当てられた作業グループについての情報を受信し、プロセシングコアに、作業グループを割り当てる、少なくとも1つの最下位コントロールコアを含み得る。
また、コントロールコアグループ320は、最上位コントロールコア330によって割り当てられた作業グループに係わる情報を受信し、最下位コントロールコアグループ350に作業グループを割り当てる、複数個の中間コントロールコアが含まれた中間コントロールコアグループ340をさらに含んでもよい。
一実施形態により、中間コントロールコアが、中間コントロールコアによって作業グループが割り当てられるプロセシングコアの数により、階層的に分類されるとき、段階1610において、上位階層の中間コントロールコアは、下位階層の中間コントロールコアに作業グループを割り当てることができる。
段階1620において、少なくとも1つのプロセシングコアが含まれたプロセシングコアグループ360は、コントロールコアグループ320によって割り当てられた作業グループを処理することができる。
段階1630において、プロセシングコアグループ360は、作業グループの処理結果を出力することができる。
図17は、一実施形態による、OpenCLカーネルを処理する方法を図示した詳細フローチャートである。図17は、図16によるコンピューティング装置300が、OpenCLカーネルを処理する方法について詳細に説明するための図であり、コンピューティング装置300の構成によって、OpenCLカーネルを処理する方法が異なるということは、当該技術分野の当業者であるならば、理解することができるであろう。また、図17の段階1720、段階1730及び段階1740は、図16の段階1610、段階1620及び段階1630とそれぞれ対応するので、詳細な説明は省略する。
段階1710において、最上位コントロールコア330は、OpenCLカーネルが実行されることにより、ホストプロセッサ310によって生成されたOpenCLカーネルの実行情報を受信することができる。かような場合、最上位コントロールコア330は、OpenCLカーネルを実行するための作業グループが処理中であるとき、ホストプロセッサ310から新たなOpenCLカーネルの実行情報が受信されることにより、プロセシングコアグループ360に割り当てられた作業グループを再割り当てすることができる。ここで、OpenCLカーネルの実行情報は、OpenCLカーネルが実行されるための全体作業グループの数、及びOpenCLカーネルが含まれたアプリケーション情報を含んでもよい。
また、一実施形態によるOpenCLカーネルを処理する方法は、ホストプロセッサ310が、コントロールコアグループ320によるプロセシングコアグループ360に対する作業グループ割り当てが完了すると、OpenCLカーネルの実行を中断し、割り当て結果を収集して出力する段階をさらに含んでもよい。
また、一実施形態による、OpenCLカーネルを処理する方法は、バッファに作業グループについての情報を保存する段階をさらに含んでもよい。ここで、該作業グループについての情報は、下位階層のコントロールコア及びプロセシングコアに割り当てられた作業グループのID、及び割り当てられた作業グループの数を含んでもよい。
段階1750において、コントロールコアグループ320は、出力された作業グループの処理結果に従って、下位階層の他のコントロールコア、及び他のプロセシングコアに割り当てられた作業グループを再割り当てすることができる。例えば、最上位コントロールコア330は、中間コントロールコアまたは最下位コントロールコアの作業グループの処理結果に従って、他の中間コントロールコア、または他の最下位コントロールコアに割り当てられた作業グループを再割り当てすることができる。また、中間コントロールコアは、下位階層の中間コントロールコアまたは最下位コントロールコアの作業グループの処理結果に従って、下位階層の他の中間コントロールコア、または他の最下位コントロールコアに割り当てられた作業グループを再割り当てすることができる。また、最下位コントロールコアは、プロセシングコアの作業グループの処理結果に従って、他のプロセシングコアに割り当てられた作業グループを再割り当てすることができる。
本実施形態は、コンピュータによって実行されるプログラムモジュールのような、コンピュータによって実行可能な命令語を含む記録媒体の形態でも具現される。コンピュータ可読媒体は、コンピュータによってアクセスされる任意の可用媒体でもあり、揮発性及び不揮発性の媒体、分離型及び非分離型の媒体をいずれも含む。コンピュータ記録媒体は、コンピュータ可読命令語、データ構造、プログラムモジュール、またはその他データのような情報保存のための任意の方法または技術によって具現された揮発性及び不揮発性、分離型及び非分離型の媒体をいずれも含む。
本実施形態の範囲は、前述の説明よりは、特許請求の範囲によって示され、特許請求の範囲の意味及び範囲、そしてその均等概念から導き出される全ての変更、または変形された形態が含まれるものであると解釈されなければならない。
本発明の、OpenCLカーネルを処理する方法、及びそれを遂行するコンピューティング装置は、例えば、電子装置関連の技術分野に効果的に適用可能である。
10,20,30,40 プロセシングエレメント
100 OpenCLプラットフォーム
110,310 ホストプロセッサ
120,130,140 コンピューティングデバイス
121,122 コンピューティングユニット
200,300 コンピューティング装置
210,320 コントロールコアグループ
220,360 プロセシングコアグループ
330 最上位コントロールコア
340 中間コントロールコアグループ
350 最下位コントロールコア
360 プロセシングコアグループ

Claims (20)

  1. 複数個のコアを含むコンピューティング装置において、
    下位階層のコントロールコアを含み、OpenCL(open computing language)カーネルを実行するための作業グループを、前記下位階層のコントロールコア及びプロセシングコアに割り当てるコントロールコアグループと、
    少なくとも1つの前記プロセシングコアを含み、前記コントロールコアグループによって割り当てられた前記作業グループを処理し、前記作業グループの処理結果を出力する、プロセシングコアグループと、を含み、
    前記コントロールコアグループに含まれた複数個のコントロールコアは、前記コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数によって階層的に分類される、
    コンピューティング装置。
  2. 前記コントロールコアグループは、
    前記OpenCLカーネルの実行情報が受信されることにより、前記作業グループを前記下位階層のコントロールコアに割り当て、前記割り当てられた作業グループについての情報を前記下位階層のコントロールコアに送信する最上位コントロールコアと、
    前記コントロールコアグループの上位階層のコントロールコアによって割り当てられた作業グループについての情報を受信し、前記プロセシングコアに前記作業グループを割り当てる、少なくとも1つの最下位コントロールコアを含んだ最下位コントロールコアグループと、を含むことを特徴とする請求項1に記載のコンピューティング装置。
  3. 前記コントロールコアグループは、
    前記下位階層のコントロールコア及び前記プロセシングコアの前記作業グループの処理結果に従って、前記コントロールコアグループの他の下位階層のコントロールコア、及び他のプロセシングコアに割り当てられた作業グループを再割り当てすることを特徴とする請求項2に記載のコンピューティング装置。
  4. 前記コントロールコアグループは、
    前記最上位コントロールコアによって割り当てられた作業グループに係わる情報を受信し、前記最下位コントロールコアグループに前記作業グループを割り当てる、複数個の中間コントロールコアを含んだ中間コントロールコアグループをさらに含み、
    前記中間コントロールコアは、前記中間コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数によって階層的に分類され、前記コントロールコアグループの上位階層の中間コントロールコアは、下位階層の中間コントロールコアに前記作業グループを割り当てることを特徴とする請求項2に記載のコンピューティング装置。
  5. 前記最上位コントロールコアは、前記OpenCLカーネルが実行されることにより、ホストプロセッサによって生成された前記OpenCLカーネルの実行情報を受信することを特徴とする請求項2に記載のコンピューティング装置。
  6. 前記最上位コントロールコアは、
    前記OpenCLカーネルを実行するための作業グループを処理しているとき、前記ホストプロセッサから新たなOpenCLカーネルの実行情報が受信されることにより、前記プロセシングコアグループに割り当てられた前記作業グループを再割り当てすることを特徴とする請求項5に記載のコンピューティング装置。
  7. 前記ホストプロセッサは、
    前記コントロールコアグループによる前記プロセシングコアグループへの前記作業グループの割り当てが完了すると、前記OpenCLカーネルの実行を中断し、割り当て結果を収集して出力することを特徴とする請求項5に記載のコンピューティング装置。
  8. 前記プロセシングコアは、
    少なくとも1つのプロセシングエレメントを含み、
    前記作業グループに含まれた作業アイテムの数より、前記プロセシングエレメントの数が少ないとき、前記プロセシングエレメントが前記作業アイテムを直列化し、順次に実行するように、前記OpenCLカーネルを変換することを特徴とする請求項1に記載のコンピューティング装置。
  9. 前記作業グループについての情報を保存するバッファをさらに含むことを特徴とする請求項1に記載のコンピューティング装置。
  10. 前記作業グループについての情報は、
    前記下位階層のコントロールコア及び前記プロセシングコアに割り当てられた作業グループの数と、前記割り当てられた作業グループのIDとのうち少なくとも一つを含むことを特徴とする請求項9に記載のコンピューティング装置。
  11. 下位階層のコントロールコアを含むコントロールコアグループにより、OpenCL(open computing language)カーネルを実行するための作業グループを、前記下位階層のコントロールコア及びプロセシングコアに割り当てる段階と、
    少なくとも1つの前記プロセシングコアが含まれたプロセシングコアグループにより、前記割り当てられた前記作業グループを処理する段階と、
    前記プロセシングコアグループにより、前記作業グループの処理結果を出力する段階と、を含み、
    前記コントロールコアグループに含まれた複数個のコントロールコアは、前記コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数によって階層的に分類される、
    OpenCLカーネルを処理する方法。
  12. 前記コントロールコアグループは、
    前記OpenCLカーネルの実行情報が受信されることにより、前記作業グループを前記下位階層のコントロールコアに割り当て、前記割り当てられた作業グループについての情報を前記下位階層のコントロールコアに送信する最上位コントロールコアと、
    前記コントロールコアグループの上位階層のコントロールコアによって割り当てられた作業グループについての情報を受信し、前記プロセシングコアに前記作業グループを割り当てる、少なくとも1つの最下位コントロールコアを含んだ最下位コントロールコアグループと、を含むことを特徴とする請求項11に記載のOpenCLカーネルを処理する方法。
  13. 前記コントロールコアグループにより、前記出力された前記作業グループの処理結果に従って、前記コントロールコアグループの下位階層の他のコントロールコア、及び他のプロセシングコアに割り当てられた作業グループを再割り当てする段階をさらに含むことを特徴とする請求項12に記載のOpenCLカーネルを処理する方法。
  14. 前記コントロールコアグループは、
    前記最上位コントロールコアによって割り当てられた作業グループに係わる情報を受信し、前記最下位コントロールコアグループに前記作業グループを割り当てる、複数個の中間コントロールコアを含んだ中間コントロールコアグループをさらに含み、
    前記中間コントロールコアが、前記中間コントロールコアにより前記作業グループが割り当てられる前記プロセシングコアの数によって階層的に分類されるとき、
    前記作業グループを割り当てる段階は、
    前記コントロールコアグループの上位階層の中間コントロールコアにより、前記コントロールコアグループの下位階層の中間コントロールコアに前記作業グループを割り当てる段階を含むことを特徴とする請求項12に記載のOpenCLカーネルを処理する方法。
  15. 前記OpenCLカーネルが実行されることにより、ホストプロセッサによって生成された前記OpenCLカーネルの実行情報を、前記最上位コントロールコアにより受信する段階をさらに含むことを特徴とする請求項12に記載のOpenCLカーネルを処理する方法。
  16. 前記OpenCLカーネルを実行するための作業グループを処理しているとき、前記ホストプロセッサから新たなOpenCLカーネルの実行情報が受信されることにより、前記最上位コントロールコアにより、前記プロセシングコアグループに割り当てられた前記作業グループを再割り当てする段階をさらに含むことを特徴とする請求項15に記載のOpenCLカーネルを処理する方法。
  17. 前記作業グループを処理する段階は、
    前記プロセシングコアが少なくとも1つのプロセシングエレメントを含み、前記作業グループに含まれた作業アイテムの数より前記プロセシングエレメントの数が少ないとき、前記プロセシングエレメントが前記作業アイテムを直列化し、順次に実行するように前記OpenCLカーネルを変換する段階をさらに含むことを特徴とする請求項11に記載のOpenCLカーネルを処理する方法。
  18. 前記作業グループについての情報をバッファに保存する段階を含むことを特徴とする請求項11に記載のOpenCLカーネルを処理する方法。
  19. 前記作業グループについての情報は、
    前記下位階層のコントロールコア及び前記プロセシングコアに割り当てられた作業グループの数と、前記割り当てられた作業グループのIDとのうち少なくとも一つを含むことを特徴とする請求項18に記載のOpenCLカーネルを処理する方法。
  20. 請求項11ないし19のうちいずれか1項に記載の方法をコンピュータで実行させるためのプログラムを記録したコンピュータ読み取り可能な記録媒体。
JP2017238434A 2016-12-27 2017-12-13 OpenCLカーネルを処理する方法、及びそれを遂行するコンピューティング装置 Active JP6951962B2 (ja)

Applications Claiming Priority (2)

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

Publications (3)

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

Family

ID=60473303

Family Applications (1)

Application Number Title Priority Date Filing Date
JP2017238434A Active JP6951962B2 (ja) 2016-12-27 2017-12-13 OpenCLカーネルを処理する方法、及びそれを遂行するコンピューティング装置

Country Status (5)

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

Families Citing this family (2)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN111490946B (zh) * 2019-01-28 2023-08-11 阿里巴巴集团控股有限公司 基于OpenCL框架的FPGA连接实现方法及装置
KR20240136005A (ko) 2023-03-06 2024-09-13 한국전자통신연구원 병렬연산작업 오프로딩 장치 및 방법

Citations (5)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JP2013500543A (ja) * 2009-07-27 2013-01-07 アドバンスト・マイクロ・ディバイシズ・インコーポレイテッド データ並列スレッドを有する処理論理の複数のプロセッサにわたるマッピング
JP2013134670A (ja) * 2011-12-27 2013-07-08 Toshiba Corp 情報処理装置及び情報処理方法
JP2014507737A (ja) * 2011-03-11 2014-03-27 インテル・コーポレーション 異種マルチコアシステム用のダイナミックコア選択
JPWO2012157786A1 (ja) * 2011-05-19 2014-07-31 日本電気株式会社 並列処理装置、並列処理方法、最適化装置、最適化方法、および、コンピュータ・プログラム
US20160321777A1 (en) * 2014-06-20 2016-11-03 Tencent Technology (Shenzhen) Company Limited Data parallel processing method and apparatus based on multiple graphic processing units

Family Cites Families (18)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JPS595755B2 (ja) 1977-08-02 1984-02-07 日本軽金属株式会社 太陽熱利用暖房方法
US7418470B2 (en) * 2000-06-26 2008-08-26 Massively Parallel Technologies, Inc. Parallel processing systems and method
US7673011B2 (en) 2007-08-10 2010-03-02 International Business Machines Corporation Configuring compute nodes of a parallel computer in an operational group into a plurality of independent non-overlapping collective networks
US8286198B2 (en) * 2008-06-06 2012-10-09 Apple Inc. Application programming interfaces for data parallel computing on multiple processors
US8255344B2 (en) 2009-05-15 2012-08-28 The Aerospace Corporation Systems and methods for parallel processing optimization for an evolutionary algorithm
US8782469B2 (en) 2009-09-01 2014-07-15 Hitachi, Ltd. Request processing system provided with multi-core processor
KR101640848B1 (ko) 2009-12-28 2016-07-29 삼성전자주식회사 멀티코어 시스템 상에서 단위 작업을 할당하는 방법 및 그 장치
KR101613971B1 (ko) 2009-12-30 2016-04-21 삼성전자주식회사 프로그램 코드의 변환 방법
JP2012094072A (ja) 2010-10-28 2012-05-17 Toyota Motor Corp 情報処理装置
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
KR101284195B1 (ko) * 2012-01-09 2013-07-10 서울대학교산학협력단 개방형 범용 병렬 컴퓨팅 프레임워크 동적 작업 분배 장치
KR20130093995A (ko) 2012-02-15 2013-08-23 한국전자통신연구원 계층적 멀티코어 프로세서의 성능 최적화 방법 및 이를 수행하는 멀티코어 프로세서 시스템
KR20140125893A (ko) 2013-01-28 2014-10-30 한국과학기술원 가상화된 매니코어 서버의 작업분배 시스템과 그 방법 및 기록매체
KR102062208B1 (ko) * 2013-05-03 2020-02-11 삼성전자주식회사 멀티스레드 프로그램 코드의 변환 장치 및 방법
JP6200824B2 (ja) * 2014-02-10 2017-09-20 ルネサスエレクトロニクス株式会社 演算制御装置及び演算制御方法並びにプログラム、OpenCLデバイス
US10318261B2 (en) * 2014-11-24 2019-06-11 Mentor Graphics Corporation Execution of complex recursive algorithms

Patent Citations (5)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JP2013500543A (ja) * 2009-07-27 2013-01-07 アドバンスト・マイクロ・ディバイシズ・インコーポレイテッド データ並列スレッドを有する処理論理の複数のプロセッサにわたるマッピング
JP2014507737A (ja) * 2011-03-11 2014-03-27 インテル・コーポレーション 異種マルチコアシステム用のダイナミックコア選択
JPWO2012157786A1 (ja) * 2011-05-19 2014-07-31 日本電気株式会社 並列処理装置、並列処理方法、最適化装置、最適化方法、および、コンピュータ・プログラム
JP2013134670A (ja) * 2011-12-27 2013-07-08 Toshiba Corp 情報処理装置及び情報処理方法
US20160321777A1 (en) * 2014-06-20 2016-11-03 Tencent Technology (Shenzhen) Company Limited Data parallel processing method and apparatus based on multiple graphic processing units

Also Published As

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

Similar Documents

Publication Publication Date Title
JP7382925B2 (ja) ニューラルネットワークアクセラレーションのための機械学習ランタイムライブラリ
US10896064B2 (en) Coordinated, topology-aware CPU-GPU-memory scheduling for containerized workloads
KR20210067415A (ko) 클라우드 컴퓨팅 기반의 플랫폼 제공 시스템 및 이를 이용한 플랫폼 제공 방법
US9063918B2 (en) Determining a virtual interrupt source number from a physical interrupt source number
JP7486575B2 (ja) コンテナに対する外部操作とミューテーション・イベントとの対応関係
US20210158131A1 (en) Hierarchical partitioning of operators
US20180314550A1 (en) Cluster topology aware container scheduling for efficient data transfer
US10572421B2 (en) Topology-aware parallel reduction in an accelerator
US10635492B2 (en) Leveraging shared work to enhance job performance across analytics platforms
US11409564B2 (en) Resource allocation for tuning hyperparameters of large-scale deep learning workloads
JP6951962B2 (ja) OpenCLカーネルを処理する方法、及びそれを遂行するコンピューティング装置
US9317328B2 (en) Strategic placement of jobs for spatial elasticity in a high-performance computing environment
US20060020701A1 (en) Thread transfer between processors
JP2015022763A (ja) オペレ−ティングシステム構成装置及び方法
US12131188B1 (en) Scheduling for locality of reference to memory
US10956512B2 (en) Document link migration
US11372677B1 (en) Efficient scheduling of load instructions
US11651235B2 (en) Generating a candidate set of entities from a training set
US20210073033A1 (en) Memory management using coherent accelerator functionality
US11182314B1 (en) Low latency neural network model loading
US11620120B1 (en) Configuration of secondary processors
US9251101B2 (en) Bitmap locking using a nodal lock
US20240256360A1 (en) Numa awareness architecture for vm-based container in kubernetes environment
US11080258B2 (en) Table generation based on scripts for existing tables
CN116594773A (zh) 向Ray集群提交计算任务的方法和装置

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

R250 Receipt of annual fees

Free format text: JAPANESE INTERMEDIATE CODE: R250