LXF171:Вы­чис­лять па­рал­лель­но

Подписка на печатную версию Весь 2015 год (12 номеров) Первое полугодие (6 номеров) Второе полугодие (6 номеров)

Подписка на электронную версию Весь 2015 год (12 номеров) Первое полугодие (6 номеров) Второе полугодие (6 номеров)

Подшивки старых номеров журнала (печатные версии) Весь 2014 год (12 номеров) Первое полугодие (6 номеров) Второе полугодие (6 номеров)

OpenCL: Стан­дарт для бу­ду­ще­го

Ми­ха­ил Ос­тап­ке­вич и Ев­ге­ний Бал­дин оты­ска­ли спо­соб при­звать не­сколь­ко раз­но­тип­ных вы­чис­ли­те­лей под зна­ме­на од­ной про­грам­мы.

Наш эксперт Михаил Остапкевич Романтик, очарованный компьютерами и создаваемыми в них идеальными мирами; верит, что сложнейшие новые технологии могут и должны служить во благо человечеству.

Наш эксперт Ев­ге­ний Бал­дин Физик, который действительно знает, что такое нехватка вычислительных ресурсов.

Вы пла­ни­руе­те за­гру­зить на­гло про­стаи­ваю­щие про­цес­со­ры, но еще не оп­ре­де­ли­лись с лю­би­мой ар­хи­тек­ту­рой? Вы не про­тив за­гру­зить ви­део­кар­ту, при­чем не обя­за­тель­но ви­део­кар­ту фир­мы N, а в пер­спек­ти­ве пе­ре­не­сти про­грам­му на эн­тер­прайз­ный Intel Xeon Phi и ин­ди-про­цес­сор Epiphany, или про­шить ее код в про­грам­ми­руе­мую ло­ги­че­­скую ин­те­граль­ную схе­му — ПЛИС? То­гда OpenCL — ваш вы­бор, а кро­ме то­го, свет­лое бу­ду­щее па­рал­лель­но­го про­грам­ми­ро­ва­ния!

От­ку­да и за­чем все это?

OpenCL (Open Computing Language) — это стан­дарт, пред­ло­жен­ный Apple в 2008 го­ду. Как и MPI, он по­зво­ля­ет соз­да­вать плат­фор­мен­но-не­за­ви­си­мые па­рал­лель­ные про­грам­мы. Но ес­ли MPI ори­ен­ти­ро­ван на муль­ти­ком­пь­ю­те­ры, то OpenCL рас­счи­тан на муль­ти­про­цес­со­ры (на­при­мер, SMP-сис­те­мы или сис­те­мы с мно­го­ядер­ны­ми про­цес­со­ра­ми), GPU и все­воз­мож­ные ус­ко­ри­те­ли. Идея Apple за­клю­ча­лась в том, что­бы ра­зо­гнать свой гра­фи­че­­ский ин­тер­фейс с по­мо­щью про­стаи­ваю­щих спе­циа­ли­зи­ро­ван­ных мощ­но­стей во вре­мя не­иг­ро­вой дея­тель­но­сти. Чуть по­го­дя к Apple ре­ши­ла примк­нуть фир­ма AMD, у ко­то­рой на тот мо­мент уже был свой соб­ст­вен­ный ве­ло­си­пед Close To Metal, ко­то­рый она уби­ла в поль­зу от­кры­то­го стан­дар­та. Те­перь OpenCL офи­ци­аль­но под­дер­жи­ва­ют фак­ти­че­­ски все ос­нов­ные про­из­во­ди­те­ли вы­чис­ли­тель­ной тех­ни­ки, вклю­чая IBM, Intel, AMD и NVIDIA. Раз­ра­бот­кой OpenCL за­ни­ма­ет­ся не­ком­мер­че­­ская не­за­ви­си­мая ор­га­ни­за­ция Khronos Group (http://www.khronos.org/).

Есть не­сколь­ко реа­ли­за­ций OpenCL для раз­лич­ных вы­чис­ли­те­лей: CPU (IA-32, AMD64, ARM, STL Cell, Intel Xeon Phi), GPU (NVIDIA, AMD, S3 Imagination Technologies PowerVR) и да­же ПЛИС (Xilinx, Altera). На­род­ная су­пер­ком­пь­ю­тер­ная ар­хи­тек­ту­ра Epiphany от Adapteva так­же име­ет SDK OpenCL. В пер­спек­ти­ве поя­вят­ся реа­ли­за­ции для про­цес­со­ров ZiiLabs ZMS-40 и циф­ро­вых сиг­на­ль­ных про­цес­со­ров TI. OpenCL по­зво­ля­ет ра­бо­тать как с па­рал­ле­лиз­мом по дан­ным, ко­гда од­на и та же опе­ра­ция од­но­вре­мен­но при­ме­ня­ет­ся к раз­ным ком­плек­там дан­ных, так и с па­рал­ле­лиз­мом на уров­не за­дач, ко­гда не­за­ви­си­мые по­то­ки ко­манд ра­бо­та­ют од­но­вре­мен­но и вы­пол­ня­ют раз­ные функ­ции.

В ми­ре со­вре­мен­ных вы­чис­ли­тель­ных сис­тем есть весь­ма раз­но­об­раз­ный и не­од­но­род­ный на­бор вы­чис­ли­те­лей. Во-пер­вых, есть мощ­ные цен­траль­ные про­цес­со­ры, спо­соб­ные вы­пол­нять од­но­вре­мен­но не­боль­шое чис­ло по­то­ков ко­манд (ни­тей, threads). Так­же лег­ко дос­туп­ны гра­фи­че­­ские про­цес­со­ры, ко­то­рые мож­но ис­поль­зо­вать для сче­та. Счет на чис­ло ни­тей в со­вре­мен­ных гра­фи­че­­ских про­цес­со­рах идет уже на ты­ся­чи, хоть они и не та­кие «тол­стые», как ни­ти цен­траль­но­го про­цес­со­ра. Есть и дру­гие разно­об­раз­ные спе­циа­ли­зи­ро­ван­ные вы­чис­ли­тель­ные ус­ко­ри­тели (ак­се­ле­ра­то­ры) — на­при­мер, ПЛИС’ы (FPGA) или сиг­наль­ные про­цес­со­ры (DSP).

Да­же в со­вер­шен­но обыч­ном ком­пь­ю­те­ре уже до­воль­но дав­но со­жи­тель­ст­ву­ют не­сколь­ко раз­но­тип­ных вы­чис­ли­те­лей — как ми­ни­мум один про­цес­сор и гра­фи­че­­ская кар­та. А ес­ли по­смот­реть на са­мые вы­со­ко­про­из­во­ди­тель­ные кла­сте­ры, то раз­ме­ще­ние в их уз­лах не­сколь­ких мно­го­ядер­ных цен­траль­ных про­цес­со­ров и не­сколь­ких GPU ста­ло те­перь уже прак­ти­че­­ски стан­дарт­ным. Но на этом эво­лю­ция не за­кон­чи­лась. Сле­дую­щий шаг на пу­ти ус­ко­ре­ния ви­дит­ся в до­бав­ле­нии в вы­чис­ли­тель­ный узел ПЛИС’ов, ло­ги­ка ра­бо­ты ко­то­рых мо­жет мно­го­крат­но пе­ре­про­грам­ми­ро­вать­ся, а дос­ти­жи­мая про­из­во­ди­тель­ность вы­чис­ле­ний в ря­де по­лез­ных слу­ча­ев на по­ря­док вы­ше да­же чем у вы­со­ко­про­из­во­ди­тель­ных GPU.

Ра­нее в рам­ках дан­ной се­рии бы­ло рас­ска­за­но о тех­но­ло­гии CUDA. Эта тех­но­ло­гия по­зво­ля­ет эф­фек­тив­но ис­поль­зо­вать для па­рал­лель­но­го сче­та ши­ро­кий спектр уст­ройств GPU от NVIDIA пря­мо здесь и сей­час. Но как быть, ес­ли у нас есть гра­фи­че­­ский про­цес­сор от дру­го­го про­из­во­ди­те­ля? Или что де­лать, ес­ли мы хо­тим, что­бы на­ша про­грам­ма без пе­ре­дел­ки мог­ла счи­тать не толь­ко на GPU, но и на яд­рах цен­траль­но­го про­цес­со­ра?

Как-то со­вер­шен­но не хо­чет­ся пе­ре­пи­сы­вать про­грам­му за­но­во ка­ж­дый раз, ко­гда объ­яв­ля­ет­ся но­вый вид вы­чис­ли­те­лей. OpenCL в ка­кой-то ме­ре по­зво­ля­ет из­бе­жать это­го. Цель соз­да­ния OpenCL — дать еди­ную тех­но­ло­гию по­строе­ния про­грамм на раз­но­род­ных вы­чис­ли­те­лях. На­пи­сав толь­ко од­ну про­грам­му, мож­но за­пус­тить ее фак­ти­че­­ски на всех па­рал­лель­ных уст­рой­ст­вах. А ес­ли сей­час это­го сде­лать нель­зя — то это не на­дол­го. Бо­лее то­го, OpenCL по­зво­лит за­пус­тить про­грам­му, на­пи­сан­ную се­го­дня, на вы­чис­ли­те­ле, ко­то­рый соз­дадут в бу­ду­щем, да­же без необ­хо­ди­мо­сти пе­ре­ком­пи­ли­ро­вать ее за­но­во! Сле­ду­ет, од­на­ко, осоз­на­вать, что хоть пра­виль­но «твик­ну­тая» про­грам­ма под со­от­вет­ст­вую­щую ар­хи­тек­ту­ру не ус­ту­па­ет по про­из­во­ди­тель­но­сти той же CUDA, но при пе­ре­но­се на прин­ци­пи­аль­но дру­гой вы­чис­ли­тель с другой про­из­во­ди­тель­но­стью без­ус­лов­но мо­гут воз­ник­нуть серь­ез­ные про­бле­мы.

Сло­ва­рик

В OpenCL ап­па­рат­ное обес­пе­че­ние под­раз­де­ля­ет­ся на две час­ти: од­на «хост-ма­ши­на» и од­но или не­сколь­ко уст­ройств OpenCL — при­мер­но так, как это по­ка­за­но на ри­сун­ке. Хост-ма­ши­на — это обыч­ный ком­пь­ю­тер, на цен­траль­ном про­цес­со­ре ко­то­ро­го вы­пол­ня­ет­ся про­грам­ма, ко­ор­ди­ни­рую­щая ход вы­чис­ле­ний. Уст­рой­ст­во OpenCL — это уст­рой­ст­во, ис­поль­зуе­мое для вы­пол­не­ния вы­чис­ле­ний. В ка­че­­ст­ве уст­рой­ст­ва OpenCL мо­гут, на­при­мер, ис­поль­зо­вать­ся GPU или сам про­цес­сор хост-ма­ши­ны.

Уст­рой­ст­во OpenCL со­сто­ит из вы­чис­ли­тель­ных мо­ду­лей [compute unit], ка­ж­дый из ко­то­рых мо­жет вы­пол­нять свой по­ток ко­манд. Вы­чис­ли­тель­ный мо­дуль со­дер­жит об­ра­ба­ты­ваю­щие эле­мен­ты (ОЭ), ка­ж­дый из ко­то­рых од­но­вре­мен­но с ос­таль­ны­ми об­ра­ба­ты­ваю­щим эле­мен­та­ми об­ра­ба­ты­ва­ет свой ком­плект дан­ных.

Ес­ли в ка­че­­ст­ве OpenCL-уст­рой­ст­ва ис­поль­зу­ет­ся GPU, то в нем мо­жет быть 8 – 16 вы­чис­ли­тель­ных мо­ду­лей, ка­ж­дый из ко­то­рых со­дер­жит до не­сколь­ких де­сят­ков об­ра­ба­ты­ваю­щих эле­мен­тов. Ес­ли в ка­че­­ст­ве OpenCL-уст­рой­ст­ва ис­поль­зу­ет­ся CPU, то чис­ло вы­чис­ли­тель­ных мо­ду­лей в нем сов­па­да­ет с чис­лом ядер, или в два раза боль­ше не­го, ес­ли в яд­ре под­дер­жи­ва­ет­ся тех­но­ло­гия HyperThreading, а в ка­ж­дом та­ком вы­чис­ли­тель­ном мо­ду­ле толь­ко один об­ра­ба­ты­ваю­щий эле­мент.

OpenCL-при­ло­же­ние со­сто­ит из двух час­тей: хост-про­грам­мы и счет­ной про­грам­мы. Хост-про­грам­ма иг­ра­ет цен­траль­ную роль. Она соз­да­ет оче­ре­ди для вы­пол­не­ния вы­чис­ле­ний счет­ной про­грам­мой, ини­ции­ру­ет пе­ре­ме­ще­ния дан­ных ме­ж­ду хост-ма­ши­ной и уст­рой­ст­ва­ми OpenCL, а так­же за­пус­ка­ет на ис­пол­не­ние функ­ции счет­ной про­грам­мы. Счет­ная про­грам­ма иг­ра­ет под­чи­нен­ную роль. Она со­сто­ит из на­бо­ра функ­ций, про­из­во­дя­щих вы­чис­ле­ния. Те ее функ­ции, ко­то­рые мож­но вы­зы­вать из хост-про­грам­мы, на­зы­ва­ют яд­ра­ми. Ини­циа­то­ром за­пус­ка яд­ра на ис­пол­не­ние обыч­но вы­сту­па­ет хост-про­грам­ма. При этом за­прос на его за­пуск раз­ме­ща­ет­ся в оче­ре­ди. Ти­пич­ный по­ря­док ис­пол­не­ния за­про­сов в оче­ре­ди — в по­ряд­ке по­сту­п­ле­ния за­про­сов в нее, хо­тя OpenCL по­зво­ля­ет отой­ти от это­го ес­те­ст­вен­но­го пра­ви­ла. За­пуск яд­ра за­клю­ча­ет­ся в соз­да­нии боль­шо­го чис­ла про­цес­сов для его ис­пол­не­ния. Все эти про­цес­сы бу­дут вы­пол­нять од­ну и ту же функ­цию, ко­то­рая за­про­грам­ми­ро­ва­на в яд­ре, но ка­ж­дый из них бу­дет об­ра­ба­ты­вать свой фраг­мент дан­ных. Рус­скоя­зыч­ный тер­мин для этих про­цес­сов еще не ус­то­ял­ся, по­это­му ис­поль­зу­ем тер­мин «ра­бо­чий эле­мент» (от англ. Work-item).

Оп­ре­де­лить, ка­кой фраг­мент дан­ных нуж­но об­ра­ба­ты­вать в дан­ном ра­бо­чем эле­мен­те, мож­но по его уни­каль­но­му иден­ти­фи­ка­то­ру. В за­ви­си­мо­сти от ха­рак­те­ра об­ра­ба­ты­вае­мых дан­ных в ка­че­­ст­ве иден­ти­фи­ка­то­ра мож­но ис­поль­зо­вать один, два или три це­ло­чис­лен­ных ин­дек­са.

Ра­бо­чий эле­мент мо­жет ис­пол­нять­ся на од­ном или не­сколь­ких об­ра­ба­ты­ваю­щих эле­мен­тах. Ра­бо­чие эле­мен­ты, ис­пол­няе­мые на од­ном вы­чис­ли­тель­ном мо­ду­ле, об­ра­зу­ют ра­бо­чую груп­пу [work-group]. Ес­ли чис­ло об­ра­ба­ты­ваю­щих эле­мен­тов боль­ше чис­ла ра­бо­чих эле­мен­тов, то воз­мож­но их од­но­вре­мен­ное ис­пол­не­ние. Од­на­ко го­раз­до ча­ще воз­ни­ка­ет си­туа­ция, ко­гда ра­бо­чих эле­мен­тов боль­ше, чем об­ра­ба­ты­ваю­щих. То­гда ис­пол­нить сра­зу все ра­бо­чие эле­мен­ты не­воз­мож­но, и ор­га­ни­зу­ет­ся их по­сле­до­ва­тель­ное ис­пол­не­ние на об­ра­ба­ты­ваю­щих эле­мен­тах.

Ре­аль­на си­туа­ция, ко­гда в од­ном ра­бо­чем эле­мен­те тре­бу­ют­ся дан­ные, вы­чис­ляе­мые в дру­гом, ко­то­рый, мо­жет быть, еще не на­чи­нал ис­пол­нять­ся. В этом слу­чае для кор­рект­ной об­ра­бот­ки тре­бу­ет­ся ис­поль­зо­ва­ние син­хро­ни­за­ции, при ко­то­рой ра­бо­чий эле­мент при­ос­та­нав­ли­ва­ет­ся до то­го мо­мен­та, ко­гда тре­буе­мые ему дан­ные бу­дут по­счи­та­ны. Рас­ста­нов­ка то­чек син­хро­ни­за­ции — это про­сто не­ис­чер­пае­мый ис­точ­ник раз­ло­жен­ных на до­ро­ге раз­но­го ро­да «граб­лей». Без хо­ро­ше­го ал­го­рит­ма вы­чис­ле­ний лег­ко за­гу­бить про­из­во­ди­тель­ность па­рал­лель­ной про­грам­мы, по­это­му кро­ме про­грам­ми­ст­ских на­вы­ков, спе­циа­ли­стам по па­рал­лель­ным вы­чис­ле­ни­ям не­об­хо­ди­ма фун­да­мен­таль­ная ма­те­ма­ти­че­­ская под­го­тов­ка.

Раз­ные реа­ли­за­ции OpenCL мо­гут со­су­ще­ст­во­вать на од­ной хост-ма­ши­не. Бо­лее то­го, их мож­но од­но­вре­мен­но ис­поль­зо­вать в од­ной хост-про­грам­ме. Фор­маль­но они пред­став­ля­ют­ся как от­дель­ные «OpenCL-плат­фор­мы», ка­ж­дая из ко­то­рых со­дер­жит на­бор под­дер­жи­вае­мых кон­крет­но ими OpenCL-уст­ройств. Од­но и то же уст­рой­ст­во мо­жет при­сут­ст­во­вать в бо­лее чем од­ной плат­фор­ме. На­при­мер, цен­траль­ный про­цес­сор Intel Core2 Quad Q8300 мо­жет быть пе­ре­чис­лен как в реа­ли­за­ции Intel OpenCL, так и в AMD OpenCL. При этом на хост-ма­ши­не мо­жет быть ус­та­нов­лен па­кет CUDA, ко­то­рый, кро­ме все­го про­че­го, так­же со­дер­жит реа­ли­за­цию OpenCL. В плат­фор­ме, со­от­вет­ст­вую­щей ему, бу­дет пе­ре­чис­ле­но од­но уст­рой­ст­во GPU, на­при­мер, NVIDIA GTS250. Со­всем аб­ст­ра­ги­ро­вать­ся, ес­те­ст­вен­но, не уда­ст­ся, так как не­об­хо­ди­мо хо­тя бы при­мер­но пред­став­лять, где про­грам­ма бу­дет счи­тать­ся, но стан­дар­ти­за­ция пра­виль­ная. [[Файл: |right | thumb|100px|]] Па­мять OpenCL-уст­рой­ст­ва от­де­ле­на от па­мя­ти хост-ма­ши­ны. Пе­ред на­ча­лом вы­чис­ле­ний не­об­хо­ди­мо ско­пи­ро­вать ис­ход­ные дан­ные из ос­нов­ной па­мя­ти хост-ма­ши­ны в гло­баль­ную па­мять OpenCL-уст­рой­ст­ва, а по­сле за­вер­ше­ния вы­чис­ле­ний пе­ре­мес­тить ре­зуль­та­ты об­рат­но. Гло­баль­ная па­мять дос­туп­на всем ра­бо­чим эле­мен­там счет­ной про­грам­мы. Кро­ме нее, OpenCL-уст­рой­ст­во име­ет кон­стант­ную, ло­каль­ную и соб­ст­вен­ную па­мять [private memory]. Кон­стант­ная па­мять дос­туп­на всем ра­бо­чим эле­мен­там по чте­нию. За­пи­сы­вать в нее мож­но из хост-про­грам­мы. Ло­каль­ная па­мять мо­жет ис­поль­зо­вать­ся толь­ко ра­бо­чи­ми эле­мен­та­ми од­ной ра­бо­чей груп­пы. Соб­ст­вен­ная па­мять дос­туп­на толь­ко од­но­му ра­бо­че­му эле­мен­ту. Та­кое раз­де­ле­ние на ти­пы па­мя­ти по­зво­ля­ет бо­лее пол­но ис­поль­зо­вать осо­бен­но­сти гра­фи­че­­ских ус­ко­ри­те­лей. В ча­ст­но­сти, в слу­чае GPU ло­каль­ная па­мять про­еци­ру­ет­ся на раз­де­ляе­мую, а соб­ст­вен­ная — на ре­ги­ст­ро­вую; обе они су­ще­ст­вен­но бы­ст­рее гло­баль­ной. При ис­поль­зо­ва­нии же цен­траль­но­го про­цес­со­ра в ка­че­­ст­ве OpenCL-уст­рой­ст­ва гло­баль­ная, ло­каль­ная и соб­ст­вен­ная па­мя­ти про­еци­ру­ют­ся на ос­нов­ную па­мять, и ско­рость их ра­бо­ты со­от­вет­ст­вен­но оди­на­ко­ва.

В стан­дар­те OpenCL опи­сан функ­цио­наль­ный ин­тер­фейс для хост-ма­ши­ны, язык “OpenCL C” для реа­ли­за­ции счет­ной про­грам­мы и функ­цио­наль­ный ин­тер­фейс для OpenCL-уст­рой­ст­ва.

Ин­тер­фейс для хост-ма­ши­ны — это тот на­бор функ­ций, ко­то­рый мож­но вы­зы­вать из хост-про­грам­мы, на­пи­сан­ной, на­при­мер, на C. Эти функ­ции по­зво­ля­ют по­лу­чать спи­сок OpenCL-уст­ройств, оп­ре­де­лять их па­ра­мет­ры (clGetPlatformIDs, clGetPlatformInfo), ком­пи­ли­ро­вать счет­ную про­грам­му (clCreateProgramWithSource), вы­де­лять и ос­во­бо­ж­дать па­мять на OpenCL-уст­рой­ст­ве (clCreateBuffer, clReleaseMemObject), пе­ре­сы­лать дан­ные (clEnqueueReadBuffer, clEnqueueWriteBuffer), за­пус­кать вы­чис­ле­ния (clEnqueueNDRangeKernel) и осу­ще­ст­в­лять син­хро­ни­за­цию на OpenCL-уст­рой­ст­вах.

Язык OpenCL C ос­но­ван на стан­дарт­ном C, од­на­ко в нем нет стан­дарт­ной биб­лио­те­ки функ­ций ANSI C и ее за­го­ло­воч­ных фай­лов, ука­за­те­лей на функ­ции, мас­си­вов пе­ре­мен­ной дли­ны и би­то­вых по­лей, а так­же в нем за­пре­ще­на ре­кур­сия. Но, с дру­гой сто­ро­ны, в не­го до­бав­ле­ны ра­бо­чие эле­мен­ты, ра­бо­чие груп­пы, век­тор­ные ти­пы, син­хро­ни­за­ция и ква­ли­фи­ка­то­ры ад­рес­но­го про­стран­ст­ва, ко­то­рые оп­ре­де­ля­ют, в гло­баль­ной, ло­каль­ной или соб­ст­вен­ной па­мя­ти рас­по­ла­га­ют­ся дан­ные.

Сбор­ка про­грам­мы

OpenCL C вклю­чен в биб­лио­те­ку для хост-про­грам­мы и вы­зы­ва­ется че­рез ин­тер­фейс для хост-ма­ши­ны. Ком­пи­ля­ция счет­ной про­грам­мы про­из­во­дит­ся в про­цес­се ра­бо­ты хост-про­грам­мы. Это по­зво­ля­ет на­страи­вать­ся на ис­поль­зо­ва­ние то­го или ино­го уст­рой­ст­ва OpenCL в про­цес­се вы­пол­не­ния про­грам­мы. Так­же это даст воз­мож­ность ис­пол­нять OpenCL-про­грам­му без пе­ре­ком­пи­ля­ции хост-про­грам­мы да­же на тех уст­рой­ст­вах, ко­то­рые еще не раз­ра­бо­та­ны. В от­ли­чие от CUDA или MPI, ком­пи­ля­ция OpenCL-про­грам­мы не тре­бу­ет за­пус­ка пре­про­цес­со­ра, хо­тя и не ис­клю­ча­ет его:

gcc helloworld.c -o helloworld -lOpenCL

Воз­мож­но, при ком­пи­ля­ции при­дет­ся ука­зать ме­сто­по­ло­же­ние за­го­ло­воч­ных фай­лов и биб­лио­те­ки libOpenCL. Пред­по­ла­га­ет­ся, что со­от­вет­ст­вую­щие драй­ве­ра и SDK для имею­щих­ся вы­чис­ли­тель­ных уст­ройств уже ус­та­нов­ле­ны. Ес­ли вы не в кур­се, о чем тут го­во­рит­ся, то для оку­чи­ва­ния GPU от Nvidia име­ет смысл за­гля­нуть по ад­ре­су https://developer.nvidia.com/opencl, для Intel есть со­от­вет­ст­вую­щий ре­сурс http://software.intel.com/en-us/vcsource/tools/opencl, а в слу­чае ви­део­кар­ты от AMD мож­но по­рыть­ся на сай­те http://developer.amd.com.

Ин­тер­фейс для вы­чис­ли­тель­но­го уст­рой­ст­ва по­зво­ля­ет ра­бо­тать с ра­бо­чи­ми эле­мен­та­ми и век­тор­ны­ми дан­ны­ми и про­из­во­дить син­хро­ни­за­цию по­то­ков, и со­дер­жит боль­шое чис­ло ма­те­ма­ти­че­­ских и гео­мет­ри­че­­ских функ­ций, функ­ций пе­ча­ти, от­лад­ки, асин­хрон­но­го ко­пи­ро­ва­ния ме­ж­ду раз­ны­ми ви­да­ми па­мя­ти.

При­мер хост-про­грам­мы

В ти­пич­ной хост-про­грам­ме мож­но вы­де­лить сле­дую­щие эта­пы ра­бо­ты:

1 Оп­ре­де­ле­ние дос­туп­ных уст­ройств OpenCL и их ха­рак­те­ри­стик.

2 Фор­ми­ро­ва­ние кон­тек­ста, ко­то­рый вклю­ча­ет вы­бран­ные для ис­поль­зо­ва­ния уст­рой­ст­ва и на­строй­ки.

3 Соз­да­ние оче­ре­ди ко­манд.

4 Соз­да­ние дво­ич­но­го ис­пол­няе­мо­го ко­да счет­ной про­грам­мы по ее ис­ход­но­му тек­сту.

5 Дек­ла­ра­ция яд­ра в дво­ич­ном ко­де счет­ной про­грам­мы.

6 Ре­зер­ви­ро­ва­ние па­мя­ти в OpenCL-уст­рой­ст­вах и ко­пи­ро­ва­ние ту­да дан­ных для счет­ной про­грам­мы.

7 Ини­циа­ли­за­ция за­пус­ка ядер на счет.

8 Ко­пи­ро­ва­ние ре­зуль­та­тов вы­чис­ле­ния из па­мя­ти OpenCL-уст­рой­ст­ва в ос­нов­ную па­мять хост-ма­ши­ны и ос­во­бо­ж­де­ние бло­ков па­мя­ти, за­ре­зер­ви­ро­ван­ных ра­нее на OpenCL-уст­рой­ст­вах.

9 Ос­во­бо­ж­де­ние ре­сур­сов и уда­ле­ние соз­дан­ные ра­нее объ­ек­тов в хост-про­грам­ме по ра­бо­те с OpenCL.

Рас­смот­рим при­мер про­стей­шей хост-про­грам­мы по­эле­мент­но­го сло­же­ния двух век­то­ров.

Для на­ча­ла нуж­но под­клю­чить OpenCL-ин­тер­фейс для хост-про­грам­мы (воз­мож­но, для ва­шей реа­ли­за­ции OpenCL путь до за­го­ло­воч­но­го фай­ла бу­дет не­мно­го дру­гой):

#include <CL/opencl.h>

Под­клю­ча­ем стан­дарт­ные за­го­ло­воч­ные фай­лы:

#include <stdio.h>

#include <stdlib.h>

Для удоб­ст­ва оп­ре­де­ля­ем кон­стан­ты:

#define TEST1_OPENCL1_BUF1_SIZE1 4096

#define SIZE 5 // За­да­ние раз­ме­ра век­то­ров

До­бав­ля­ем ис­ход­ник счет­ной про­грам­мы, со­дер­жа­щей яд­ро vadd для сло­же­ния двух чи­сел:

const char *vaddsrc =

“__kernel void vadd(“

“ __global float *d_A,”

“ __global float *d_B,”

“ __global float *d_C”

“){“

“ unsigned int n;”

“”

// По­лу­че­ние уни­каль­но­го иден­ти­фи­ка­то­ра ра­бо­че­го эле­мен­та

“ n = get_global_id(0);”

“ d_C[n] = d_A[n] + d_B[n];”

”}”;

Здесь d_A и d_B — ука­за­те­ли на складываемые век­то­ры в гло­баль­ной па­мя­ти, d_A — ука­за­тель на ре­зуль­ти­рую­щий век­тор.

Оп­ре­де­ля­ем век­то­ры в ос­нов­ной па­мя­ти хост-ком­пь­ю­те­ра:

float hostdata_A[SIZE] = {10., 20., 30., 40., 50.};

float hostdata_B[SIZE] = {1., 2., 3., 4., 5.};

float hostdata_C[SIZE];

char clcompileflags[TEST1_OPENCL1_BUF1_SIZE1];

char buf[TEST1_OPENCL1_BUF1_SIZE1];

// Од­но­мер­ное про­стран­ст­во уни­каль­ных иден­ти­фи­ка­то­ров ра­бо­чих эле­мен­тов

size_t sizeWork[1] = {SIZE};

На­чи­на­ем про­грам­му с оп­ре­де­ле­ния вспо­мо­га­тель­ных пе­ре­мен­ных:

int main(){

// Код ошиб­ки, ко­то­рый воз­вра­ща­ет­ся OpenCL-функ­ция­ми

cl_int clerr;

// Ко­ли­че­ст­во дос­туп­ных OpenCL-плат­форм

cl_uint qty_platforms = 0;

// Спи­сок иден­ти­фи­ка­то­ров OpenCL-плат­форм

cl_platform_id* platforms;

cl_uint ui;

// Ко­ли­че­ст­во дос­туп­ных OpenCL-уст­ройств для ка­ж­дой плат­фор­мы

cl_uint *qty_devices;

cl_device_id **devices;

cl_uint i;

// Де­ск­рип­тор кон­тек­ста

cl_context context1;

size_t parmsz;

// Де­ск­рип­тор оче­ре­ди ко­манд

cl_command_queue queue1;

cl_program program1;

// Де­ск­рип­тор объ­ек­та яд­ра

cl_kernel kernel1;

Оп­ре­де­ля­ем ука­за­те­ли на век­то­ры A и B с ис­ход­ны­ми дан­ны­ми и век­тор C с ре­зуль­ти­рую­щи­ми дан­ны­ми в гло­баль­ной па­мя­ти:

cl_mem remotedata_A;

cl_mem remotedata_B;

cl_mem remotedata_C;

Этап 1 Оп­ре­де­ле­ние дос­туп­ных OpenCL-уст­ройств.

// Оп­ре­де­ле­ние чис­ла OpenCL-плат­форм

clerr = clGetPlatformIDs(0, NULL, &qty_platforms);

// Об­ра­бот­ка оши­бок

if(clerr != CL_SUCCESS){

fprintf(stderr, “Ошиб­ка, код = %d.\n”, clerr);

// Да­лее ана­ло­гич­ные бло­ки опу­ще­ны

return 1;

}

// Вы­де­ле­ние па­мя­ти для хра­не­ния ин­фор­ма­ции о кон­фи­гу­ра­ции и па­ра­мет­ров уст­ройств

platforms = (cl_platform_id*)

malloc(sizeof(cl_platform_id)*qty_platforms);

devices = (cl_device_id**)

malloc(sizeof(cl_device_id*)*qty_platforms);

qty_devices = (cl_uint*)malloc(sizeof(cl_uint)*qty_platforms);

// Об­ра­бот­ка оши­бок вы­де­ле­ния па­мя­ти здесь и да­лее опу­ще­на

// По­лу­че­ние спи­ска иден­ти­фи­ка­то­ров плат­форм

clerr = clGetPlatformIDs(qty_platforms, platforms, NULL);

for (ui=0; ui < qty_platforms; ui++){

// По­лу­че­ние ко­ли­че­ст­ва имею­щих­ся OpenCL-уст­ройств для ка­ж­дой плат­фор­мы

clerr = clGetDeviceIDs(platforms[ui], CL_DEVICE_TYPE_ALL, 0, NULL, &qty_devices[ui]);

// По­лу­че­ние спи­ска иден­ти­фи­ка­то­ров OpenCL-уст­ройств

if(qty_devices[ui]){

devices[ui] = (cl_device_id*)

malloc(qty_devices[ui] * sizeof(cl_device_id));

clerr = clGetDeviceIDs(platforms[ui], CL_DEVICE_TYPE_ALL, qty_devices[ui], devices[ui], NULL);

}

}

Этап 2 Фор­ми­ро­ва­ние кон­тек­ста. Соз­да­ет­ся кон­текст, в ко­то­рый вклю­че­ны OpenCL-уст­рой­ст­ва плат­фор­мы 0 (уп­ро­ще­ние):

clerr = CL_SUCCESS;

context1 = clCreateContext(0, qty_devices[0], devices[0], NULL, NULL, &clerr);

Этап 3 Соз­да­ние оче­ре­ди ко­манд. Оче­редь соз­да­ет­ся для ну­ле­во­го уст­рой­ст­ва ну­ле­вой плат­фор­мы:

queue1 = clCreateCommandQueue(context1, devices[0][0], 0, &clerr);

Этап 4 Ком­пи­ля­ция счет­ной про­грам­мы.

program1 = clCreateProgramWithSource(context1, 1, &vaddsrc, NULL, &clerr);

snprintf(clcompileflags, TEST1_OPENCL1_BUF1_SIZE1, “-cl-mad-enable”);

clerr = clBuildProgram(program1, 0, NULL, clcompileflags, NULL, NULL);

Этап 5 Дек­ла­ра­ция яд­ра в счет­ной про­грам­ме.

kernel1 = clCreateKernel(program1, “vadd”, &clerr);

Этап 6 Ре­зер­ви­ро­ва­ние па­мя­ти в OpenCL-уст­рой­ст­ве. При этом век­то­ры A и B с исход­ными дан­ными ко­пи­ру­ют­ся из ос­нов­ной па­мя­ти хост-ма­ши­ны в па­мять OpenCL-уст­рой­ст­ва.

remotedata_A = clCreateBuffer( context1, CL_MEM_READ_ONLY |CL_MEM_COPY_HOST_PTR, sizeof(int) * SIZE, hostdata_A, NULL);

remotedata_B = clCreateBuffer(context1, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * SIZE, hostdata_B, NULL);

remotedata_C = clCreateBuffer(context1, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * SIZE, hostdata_C, NULL);

Этап 7 За­пуск яд­ра на вы­чис­ле­ния.

clSetKernelArg(kernel1, 0, sizeof(cl_mem), (void*)&remotedata_A);

clSetKernelArg(kernel1, 1, sizeof(cl_mem), (void*)&remotedata_B);

clSetKernelArg(kernel1, 2, sizeof(cl_mem), (void*)&remotedata_C);

clerr = clEnqueueNDRangeKernel(queue1, kernel1, 1, NULL, sizeWork, NULL, 0, NULL, NULL);

Этап 8 Ко­пи­ро­ва­ние ре­зуль­та­тов в ос­нов­ную па­мять хост-ма­ши­ны и ос­во­бо­ж­де­ние па­мя­ти в OpenCL-уст­рой­ст­ве.

// Ко­пи­ро­ва­ние ре­зуль­ти­рую­ще­го век­то­ра

clEnqueueReadBuffer(queue1, remotedata_C, CL_TRUE, 0, SIZE * sizeof(int), hostdata_C, 0, NULL, NULL);

// Ос­во­бо­ж­де­ние па­мя­ти на OpenCL-уст­рой­ст­ве

clReleaseMemObject(remotedata_A);

clReleaseMemObject(remotedata_B);

clReleaseMemObject(remotedata_C);

Этап 9 Ос­во­бо­ж­де­ние ре­сур­сов.

clReleaseKernel(kernel1);

clReleaseProgram(program1);

clReleaseCommandQueue(queue1);

clReleaseContext(context1);

free(qty_devices);

free(devices);

free(platforms);

А те­перь на­ко­нец-то мож­но вы­дох­нуть и рас­пе­ча­тать ре­зуль­ти­рую­щий век­тор:

for(i = 0; i < SIZE; i++)

printf(“ %f\n”, hostdata_C[i]);

}

• Метамодернизм в позднем творчестве В.Г. Сорокина
• ЛитРПГ - последняя отрыжка постмодерна
• "Ричард III и семиотика"
• 3D-визуализация обложки Ridero создаем обложку книги при работе над самиздатом.
• Архитектура метамодерна - говоря о современном искусстве, невозможно не поговорить об архитектуре. В данной статье будет отмечено несколько интересных принципов, характерных для построек "новой волны", столь притягательных и скандальных.
• Литература
• Метамодерн
• Рокер-Прометей против изначального зла в «Песне про советскую милицию» Вени Дркина, Автор: Нина Ищенко, к.ф.н, член Союза Писателей ЛНР - перепубликация из журнала "Топос".
• Как избавиться от комаров? Лучшие типы ловушек.
• Что делать если роблокс вылетает на windows
• Что делать, если ребенок смотрит порно?
• Почему собака прыгает на людей при встрече?
• Какое масло лить в Задний дифференциал (мост) Visco diff 38434AA050
• О чем может рассказать хвост вашей кошки?
• Верветки
• Отчетность бюджетных учреждений при закупках по Закону № 223-ФЗ
• Срок исковой давности как правильно рассчитать
• Дмитрий Патрушев минсельхоз будет ли преемником Путина
• Кто такой Владислав Поздняков? Что такое "Мужское Государство" и почему его признали экстремистским в России?
• Как правильно выбрать машинное масло в Димитровграде?
• Как стать богатым и знаменитым в России?
• Почему фильм "Пипец" (Kick-Ass) стал популярен по всему миру?
• Как стать мудрецом?
• Как правильно установить FreeBSD
• Как стать таким как Путин?
• Где лучше жить - в Димитровграде или в Ульяновске?
• Почему город Димитровград так называется?
• Что такое метамодерн?
• ВАЖНО! Временное ограничение движения автотранспортных средств в Димитровграде
• Тарифы на электроэнергию для майнеров предложено повысить

Не­смот­ря на ви­ди­мую слож­ность, рас­смот­рен­ный при­мер мак­си­маль­но уп­ро­щен. В ча­ст­но­сти, мы про­из­во­ди­ли счет на фик­си­ро­ван­ном ну­ле­вом уст­рой­ст­ве ну­ле­вой плат­фор­мы. В ре­аль­но­сти мы пред­по­чли бы бо­лее осоз­нан­но вы­брать уст­рой­ст­во для сче­та — на­при­мер, воз­мож­но, за­хо­те­ли бы по­ис­кать сна­ча­ла GPU от NVIDIA. Для это­го в функ­ции clGetDeviceIDs не­об­хо­ди­мо за­ме­нить CL_DEVICE_TYPE_ALL на CL_DEVICE_TYPE_GPU.

В бо­лее слож­ных реа­ли­за­ци­ях счет мо­жет про­из­во­дить­ся на не­сколь­ких раз­но­тип­ных OpenCL-уст­рой­ст­вах. Ти­пич­ная схе­ма ор­га­ни­за­ции та­ко­го сче­та за­клю­ча­ет­ся в соз­да­нии не­сколь­ких кон­тек­стов или не­сколь­ких оче­ре­дей в од­ном кон­тек­сте, ка­ж­дая из ко­то­рых ас­со­ции­ро­ва­на с не­ко­то­рым уст­рой­ст­вом.

Дру­гое уп­ро­ще­ние свя­за­но с тем, что текст счет­ной про­грам­мы с ядром мы хра­ним как кон­стан­ту в хост-про­грам­ме. Счет­ных про­грамм и их ядер мо­жет быть дос­та­точ­но мно­го, и в пол­но­цен­ной реа­ли­за­ции име­ет смысл хра­нить их в от­дель­ных фай­лах и за­гру­жать при не­об­хо­ди­мо­сти.

В за­клю­че­ние не­об­хо­ди­мо от­ме­тить, что ис­поль­зо­ва­ние OpenCL не ис­клю­ча­ет воз­мож­ность при­ме­не­ния дру­гих тех­но­ло­гий по­строе­ния па­рал­лель­ных про­грамм в од­ной про­грам­ме. На­при­мер, на кла­сте­рах с мно­го­про­цес­сор­ны­ми уз­ла­ми или уз­ла­ми с GPU це­ле­со­об­раз­но стро­ить гиб­рид­ные про­грам­мы на ба­зе OpenCL и MPI. OpenCL в них бу­дет от­ве­чать за ор­га­ни­за­цию па­рал­лель­но­го сче­та внут­ри од­но­го уз­ла, а MPI — за па­рал­лель­ное ис­пол­не­ние на уров­не кла­сте­ра, ме­ж­ду не­сколь­ки­ми уз­ла­ми. |