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

Текущая версия на 15:19, 18 ноября 2018

[править] 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 — за па­рал­лель­ное ис­пол­не­ние на уров­не кла­сте­ра, ме­ж­ду не­сколь­ки­ми уз­ла­ми. |