Жаңа технология жаңадан пайда болған эволюциялық түрге ұқсайды. Көптеген ескі адамдарға ұқсамайтын біртүрлі жаратылыс. Кейде ыңғайсыз, кейде күлкілі. Алғашында оның жаңа қасиеттері осы өмір сүруге қолайлы және тұрақты әлемге мүлдем сәйкес келмейді.

Дегенмен, аз уақыт өтеді, және жаңадан бастаушы жылдамырақ жүгіреді, жоғары секіреді және жалпы күштірек болады. Және ол ретроградтық көршілеріне қарағанда шыбындарды көбірек жейді. Содан кейін бұл көршілер бұрынғы ебедейсізімен ұрысудың қажеті жоқ екенін түсіне бастайды. Онымен дос болған дұрыс, тіпті симбиозды ұйымдастырған дұрыс. Қарайсың, сонда шыбын көбейеді.

GPGPU (General-Purpose Graphics Processing Units) технологиясы ежелден ақылды академиктердің теориялық есептеулерінде ғана бар. Басқа қалай? Ондаған жылдардағы түбегейлі өзгерістерді ұсыныңыз есептеу процесі, оның параллель тармақтарын есептеуді видеокартаға тапсыру - тек теоретиктер бұған қабілетті.

CUDA технологиясының логотипі оның тереңдікте өскенін еске салады
3D графикасы.

Бірақ GPGPU технологиясы университет журналдарының беттеріне ұзақ уақыт шаң жинамады. Өзінің ең жақсы қасиеттерінің қауырсындарын жұлып алған ол продюсерлердің назарын аударды. NVIDIA GeForce графикалық процессорларында GPGPU жүзеге асырылатын CUDA осылай дүниеге келді.

CUDA арқасында GPGPU технологиялары негізгі ағымға айналды. Ал енді бағдарламалау жүйелерінің ең шолақ және ең жалқау әзірлеушісі ғана өз өнімімен CUDA-ға қолдау көрсетпейді. IT-басылымдары технологияның егжей-тегжейлерін көптеген ғылыми-танымал мақалаларда ұсыну құрметіне ие болды, ал бәсекелестер тез арада ұқсас нәрсені жасау үшін шаблондар мен кросс-компиляторларға отырды.

Жұртшылықтың мойындалуы тек қана жұлдыздардың арманы ғана емес, сонымен қатар жаңадан дамып келе жатқан технологиялардың да арманы. Ал CUDA жолы болды. Ол естіледі, олар ол туралы айтады және ол туралы жазады.

Олар қалың ғылыми журналдарда GPGPU-ны талқылауды жалғастыра бергендей жазады. Олар оқырманды «тор», «SIMD», «warp», «хост», «текстік және тұрақты жады» сияқты көптеген терминдермен бомбалайды. Олар оны nVidia GPU ұйымдастыру схемаларына ең жоғарғы деңгейге дейін түсіреді, оны параллель алгоритмдердің орама жолдары арқылы жүргізеді және (ең күшті қозғалыс) C тіліндегі кодтың ұзын тізімін көрсетеді. Нәтижесінде, мақаланың кірісінде бізде CUDA оқырманын, ал шығысында - дәл сол оқырманды, бірақ фактілерге, диаграммаларға, кодтарға толы басы ісіп жатқанын түсінуге деген жаңа және жалындаған тілек бар екені белгілі болды. , алгоритмдер мен терминдер.

Сонымен қатар, кез келген технологияның мақсаты - өмірімізді жеңілдету. Ал CUDA онымен тамаша жұмыс істейді. Оның жұмысының нәтижелері кез келген скептикті жүздеген схемалар мен алгоритмдерден жақсырақ сендіреді.

Барлық жерден алыс

CUDA өнімділігі жоғары суперкомпьютерлермен қолдау көрсетеді
NVIDIA Tesla.

Дегенмен, қарапайым пайдаланушының өмірін жеңілдету саласындағы CUDA жұмысының нәтижелерін қарастырмас бұрын, оның барлық шектеулерін түсіну керек. Жын сияқты: кез келген тілек, бірақ бір. CUDA-ның сонымен қатар Ахиллес өкшесі бар. Олардың бірі - ол жұмыс істей алатын платформалардың шектеулері.

CUDA қолдайтын nVidia шығарған бейне карталардың тізімі CUDA қосылған өнімдер деп аталатын арнайы тізімде берілген. Тізім өте әсерлі, бірақ оңай жіктеледі. CUDA қолдауынан бас тартылмайды:

Бортта кемінде 256 Мбайт бейне жады бар nVidia GeForce 8-ші, 9-шы, 100-ші, 200-ші және 400-ші сериялы үлгілері. Қолдау жұмыс үстеліне де, мобильді карталарға да таралады.

Жұмыс үстелі және мобильді графикалық карталардың басым көпшілігі nVidia Quadro болып табылады.

nvidia ION нетбук сериясының барлық шешімдері.

Жоғары өнімді HPC (High Performance Computing) және nVidia Tesla суперкомпьютерлік шешімдері жеке есептеулер үшін де, масштабталатын кластерлік жүйелерді ұйымдастыру үшін де қолданылады.

Сондықтан, CUDA негізіндегі бағдарламалық өнімдерді қолданбас бұрын, таңдаулылардың осы тізімін тексерген жөн.

Графикалық картаның өзінен басқа, CUDA қолдауы тиісті драйверді қажет етеді. Ол GPU көп ядролы қазынасына кодтар мен бағдарламалық деректерге қол жеткізу үшін бағдарламалық интерфейстің бір түрі ретінде әрекет ететін орталық және графикалық процессор арасындағы байланыстырушы болып табылады. Қателік жасамау үшін nVidia драйверлер бетіне кіріп, ең соңғы нұсқасын алуды ұсынады.

...бірақ процесс

CUDA қалай жұмыс істейді? Оқырманды белгілі бір терминдердің тұңғиығына түсірмей, нақты GPU аппараттық архитектурасында параллельді есептеудің күрделі процесін қалай түсіндіруге болады?

Мұны орталық процессордың графикалық процессормен симбиозда бағдарламаны қалай орындайтынын елестету арқылы жасауға болады.

Архитектуралық тұрғыдан орталық процессор (CPU) және оның графикалық аналогы (GPU) басқаша орналастырылған. Егер біз автомобиль өнеркәсібі әлеміне ұқсастық жасасақ, онда орталық процессор - бұл «сарай» деп аталатындардың бірі. Бұл жеңіл көлікке ұқсайды, бірақ сонымен бірге (әзірлеушілердің көзқарасы бойынша) «және швейцариялық, және орақшы және құбырдағы ойыншы». Бір уақытта шағын жүк көлігі, автобус және гипертрофияланған хэтчбек рөлін орындайды. Әмбебап, қысқаша. Оның цилиндрлік өзектері аз, бірақ олар кез келген тапсырманы дерлік «тартады» және әсерлі кэш жады көптеген деректерді сақтай алады.

Бірақ GPU - бұл спорттық автомобиль. Бір ғана функция бар: ұшқышты мәреге мүмкіндігінше тез жеткізу. Сондықтан, үлкен жады жоқ, қосымша орын жоқ. Бірақ ядролық цилиндрлер орталық процессорға қарағанда жүздеген есе көп.

CUDA көмегімен GPGPU әзірлеушілеріне бағдарламалаудың күрделілігін түсінудің қажеті жоқ.
DirectX және OpenGL сияқты графикалық қозғалтқыштар үшін

Кез келген тапсырманы, соның ішінде графиканы шешуге қабілетті, бірақ орташа өнімділігі бар орталық процессордан айырмашылығы, графикалық процессор бір тапсырманың жоғары жылдамдықты шешіміне бейімделген: кірістегі көпбұрыштардың үйінділерін нүктедегі пикселдер тобына айналдыру. шығару. Сонымен қатар, бұл тапсырманы GPU бөлігі ретінде жүздеген салыстырмалы қарапайым есептеуіш ядроларда параллель шешуге болады.

Сонымен, вагон мен спорттық машинаның тандемі қандай болуы мүмкін? CUDA жұмысы келесідей болады: бағдарлама процессорда параллель орындалатын код бөлігі болғанша орындалады. Содан кейін ол ең керемет процессордың екі (тіпті сегіз) ядросында баяу орындалмай, жүздеген GPU өзектеріне ауыстырылады. Сонымен бірге бұл бөлімнің орындалу уақыты бірнеше есе қысқарады, яғни бүкіл бағдарламаның орындалу уақыты да қысқарады.

Технологиялық тұрғыдан бағдарламашы үшін ештеңе өзгермейді. CUDA бағдарламаларының коды Си тілінде жазылған. Дәлірек айтсақ, оның арнайы диалектісінде «С с ағынмен» (С с с ағынды). Стэнфордта әзірленген Си тілінің бұл кеңейтімі Брук деп аталады. Brook кодын GPU-ға жіберетін интерфейс CUDA-ны қолдайтын видеокартаның драйвері болып табылады. Ол бағдарламаның осы бөлімін өңдеудің бүкіл процесін GPU бағдарламалаушы үшін процессордың сопроцессоры сияқты көрінетіндей етіп ұйымдастырады. Дербес компьютерлердің алғашқы күндерінде математикалық сопроцессорды пайдалану сияқты. Brook, CUDA қолдауы бар видеокарталар және оларға арналған драйверлердің пайда болуымен кез келген бағдарламашы өз бағдарламаларында GPU-ға қол жеткізе алды. Бірақ бұған дейін шаманизм бірнеше жылдар бойы DirectX немесе OpenGL графикалық қозғалтқыштары үшін бағдарламалау әдістерін жетілдіретін элитаның тар шеңберіне ие болды.

Бұл сыпайы балдың бөшкесінде - CUDA мақтайды - бұл жақпаға шыбын салуға тұрарлық, яғни шектеулер. Бағдарламалауды қажет ететін әрбір тапсырма CUDA көмегімен шешуге жарамайды. Кәдімгі кеңсе тапсырмаларын шешуде жеделдетуге қол жеткізу мүмкін болмайды, бірақ World of Warcraft ойынындағы мыңдаған бірдей типтегі жауынгерлердің мінез-құлқын есептейтін CUDA-ға сенуге болады. Бірақ бұл саусақтан сорылған міндет. CUDA қазірдің өзінде өте тиімді шешетін мысалдарды қарастырайық.

Әділ істер

CUDA - өте прагматикалық технология. Бейне карталарында қолдауды жүзеге асырған nVidia CUDA баннерін университеттік ортада да, саудада да көптеген энтузиастар алады деп күткен. Осылайша болды. CUDA негізіндегі жобалар тірі және жақсы.

NVIDIA PhysX

Келесі ойын шедеврін жарнамалай отырып, өндірушілер оның 3D реализмін жиі атап өтеді. Бірақ 3D ойын әлемі қаншалықты шынайы болғанымен, ауырлық күші, үйкеліс, гидродинамика сияқты физиканың қарапайым заңдары қате орындалса, жалғандық бірден сезіледі.

NVIDIA PhysX физика қозғалтқышының ерекшеліктерінің бірі - тіндермен шынайы жұмыс.

Негізгі физикалық заңдылықтарды компьютерлік модельдеу алгоритмдерін енгізу өте еңбекті қажет ететін жұмыс. Бұл саладағы ең танымал компаниялар - Havok Physics кросс-платформалық физикасы бар ирландиялық Havok компаниясы және Калифорниялық Ageia - әлемдегі алғашқы физикалық процессордың (PPU - Physics Processing Unit) және сәйкес PhysX физика қозғалтқышының бастаушысы. Олардың біріншісін Intel сатып алғанымен, қазір ATI видеокарталары мен AMD процессорлары үшін Havok қозғалтқышын оңтайландыру саласында белсенді жұмыс істейді. Бірақ Ageia өзінің PhysX қозғалтқышымен nVidia бөлігі болды. Сонымен қатар, nVidia PhysX-ті CUDA технологиясына бейімдеу бойынша өте күрделі тапсырманы шешті.

Бұл статистиканың арқасында мүмкін болды. GPU қаншалықты күрделі рендеринг орындаса да, оның кейбір ядролары әлі де жұмыс істемейтіні статистикалық түрде дәлелденді. Дәл осы ядроларда PhysX қозғалтқышы жұмыс істейді.

CUDA арқасында ойын әлемінің физикасына қатысты есептеулердің негізгі бөлігі бейне картада орындала бастады. Орталық процессордың босатылған қуаты геймплейдің басқа тапсырмаларына лақтырылды. Нәтижесі көп күттірмеді. Сарапшылардың пікірінше, CUDA жүйесінде жұмыс істейтін PhysX геймплейінің өнімділігі кем дегенде бір ретке өсті. Физикалық заңдарды жүзеге асырудың орындылығы да өсті. CUDA көп өлшемді объектілер үшін үйкеліс, гравитация және бізге таныс басқа нәрселерді іске асырудың күнделікті есебімен айналысады. Енді кейіпкерлер мен олардың жабдықтары бізге таныс физикалық әлемнің заңдарына тамаша сәйкес келеді, сонымен қатар шаң, тұман, жарылыс толқыны, жалын және су.

NVIDIA Texture Tools 2 текстураны қысу пакетінің CUDA нұсқасы

Сізге заманауи ойындардағы шынайы нысандар ұнай ма? Текстура жасаушыларға алғыс айта кеткен жөн. Бірақ текстурадағы шындық неғұрлым көп болса, оның көлемі соғұрлым үлкен болады. Ол қымбат жадты көбірек алады. Бұған жол бермеу үшін текстуралар алдын ала қысылады және қажет болған жағдайда динамикалық түрде ашылады. Ал қысу және декомпрессия үздіксіз есептеулер. Текстуралармен жұмыс істеу үшін nVidia NVIDIA Texture Tools бумасын шығарды. Ол DirectX текстураларының (HFC форматы деп аталатын) тиімді қысылуын және декомпрессиясын қолдайды. Бұл пакеттің екінші нұсқасы DirectX 11 технологиясында енгізілген BC4 және BC5 қысу алгоритмдерін қолдауымен мақтана алады.Бірақ ең бастысы NVIDIA Texture Tools 2 CUDA-ны қолдайды. nVidia мәліметтері бойынша, бұл текстураларды қысу және декомпрессиялау тапсырмаларында өнімділікті 12 есе арттыруға мүмкіндік береді. Бұл геймплейдің кадрлары тезірек жүктелетінін және ойыншыны шынайылығымен қуантады дегенді білдіреді.

NVIDIA Texture Tools 2 пакеті CUDA бағдарламасымен жұмыс істеуге бейімделген. Текстураларды қысу және ашу кезінде өнімділіктің жоғарылауы анық.

CUDA пайдалану бейнебақылаудың тиімділігін айтарлықтай жақсарта алады.

Нақты уақыттағы бейне ағынын өңдеу

Өзіңізге ұнайтын нәрсені айтыңыз, бірақ қазіргі әлем, тыңшылық тұрғысынан, Оруэллиан Үлкен ағасының әлеміне көрінгеннен әлдеқайда жақын. Бейнекамералардың көзқарасын көлік жүргізушілері де, қоғамдық орындарға келушілер де сезеді.

Бейне ақпараттың толып жатқан өзендері оны өңдеу орталықтарына құйылады және ... тығырыққа тіреледі - адамға. Ол көп жағдайда бейне әлемін бақылайтын соңғы құрал болып табылады. Оның үстіне, агенттік ең тиімді емес. Жыпылықтап, көңілі ауып, ұйықтауға ұмтылады.

CUDA арқасында бейне ағынында бірнеше нысандарды бір уақытта қадағалау алгоритмдерін енгізу мүмкін болды. Бұл жағдайда процесс нақты уақытта орын алады, ал бейне толық 30 кадр/сек. Қазіргі көп ядролы процессорларда мұндай алгоритмді жүзеге асырумен салыстырғанда, CUDA өнімділікті екі-үш есе арттырады, және бұл, көріп тұрсыз, көп нәрсе.

Бейнені түрлендіру, аудио сүзгілеу

Badaboom Video Converter - түрлендіруді жылдамдату үшін CUDA қолданбасын бірінші болып қолданды.

Жаңа бейне прокатын FullHD сапасында және үлкен экранда көру өте жақсы. Бірақ сіз жолда үлкен экранды ала алмайсыз, ал FullHD бейне кодегі мобильді гаджеттің қуаттылығы аз процессорында ықырады. Конверсия көмекке келеді. Бірақ іс жүзінде оны кездестіргендердің көпшілігі конверсияның ұзақ уақытына шағымданады. Түсінікті, процесс әдеттегі, параллелизацияға қолайлы және оның орталық процессорда орындалуы өте оңтайлы емес.

Бірақ CUDA оны жарылыспен жеңеді. Бірінші белгі - Elevental ұсынған Badaboom түрлендіргіші. Badaboom әзірлеушілері CUDA таңдау арқылы қате есептеген жоқ. Тесттер стандартты бір жарым сағаттық фильм iPhone/iPod Touch пішіміне жиырма минуттан аз уақыт ішінде түрленетінін көрсетеді. Және бұл тек процессорды пайдаланған кезде бұл процесс бір сағаттан астам уақытты қажет ететініне қарамастан.

CUDA және кәсіби музыка әуесқойларына көмектеседі. Олардың кез келгені тиімді FIR-кроссовер үшін патшалықтың жартысын береді - дыбыс спектрін бірнеше жолақтарға бөлетін сүзгілер жиынтығы. Бұл процесс өте көп уақытты қажет етеді және дыбыстық материалдың үлкен көлемімен дыбыс инженерін бірнеше сағат бойы «темекілеуге» мәжбүр етеді. CUDA негізіндегі FIR кроссоверін енгізу оның жұмысын жүздеген есе жылдамдатады.

CUDA болашағы

GPGPU технологиясын шындыққа айналдыра отырып, CUDA жетістіктерімен тоқтап қалмайды. Барлық жерде болатындай, CUDA-да рефлексия принципі жұмыс істейді: енді CUDA SDK нұсқаларының дамуына nVidia бейне процессорларының архитектурасы ғана әсер етіп қоймайды, сонымен қатар CUDA технологиясының өзі nVidia-ны чиптерінің архитектурасын қайта қарауға мәжбүр етеді. Мұндай көріністің мысалы nVidia ION платформасы болып табылады. Оның екінші нұсқасы CUDA тапсырмаларын шешу үшін арнайы оңтайландырылған. Бұл тіпті салыстырмалы түрде арзан аппараттық шешімдерде тұтынушылар CUDA-ның барлық күші мен тамаша мүмкіндіктерін алатынын білдіреді.

– төмен деңгейлі бағдарламалау интерфейстерінің жиынтығы ( API) ойындарды және басқа да өнімділігі жоғары мультимедиялық қолданбаларды жасауға арналған. Жоғары өнімділікті қолдауды қамтиды 2D- және 3D-графикалық, дыбыстық және енгізу құрылғылары.

Direct3D (D3D) – 3D шығыс интерфейсі примитивтер(геометриялық денелер). Құрамына кіреді .

OpenGL(ағылшын тілінен. Графикалық кітапхананы ашыңыз, сөзбе-сөз – ашық графикалық кітапхана) – екі өлшемді және үш өлшемді компьютерлік графиканы пайдаланатын қолданбаларды жазу үшін программалау тілінен тәуелсіз кросс-платформалық бағдарламалау интерфейсін анықтайтын спецификация. Қарапайым примитивтерден күрделі 3D көріністерін салуға арналған 250-ден астам функцияларды қамтиды. Ол ғылыми зерттеулерде бейне ойындар, виртуалды шындық, визуализация жасауда қолданылады. Платформада Windows-мен жарысады .

OpenCL(ағылшын тілінен. Есептеу тілін ашыңыз, сөзбе-сөз - ашық есептеу тілі) - шеңбер(бағдарламалық жүйенің негізі) әртүрлі графикада параллельді есептеулерге байланысты компьютерлік бағдарламаларды жазуға арналған ( GPU) және ( ). Рамкаға OpenCLпрограммалау тілін және қолданбалы бағдарламалау интерфейсін қамтиды ( API). OpenCLнұсқау деңгейінде және деректер деңгейінде параллелизмді қамтамасыз етеді және әдістемені жүзеге асыру болып табылады GPGPU.

GPGPU(ағылшын тілінен қысқартылған. Жалпы мақсаттағы графикалық өңдеу қондырғылары, сөзбе-сөз - GPUжалпы мақсат) - әдетте орындалатын жалпы есептеулер үшін видеокартаның графикалық процессорын пайдалану әдістемесі.

шейдер(ағылшын) шейдер) — объектінің немесе кескіннің соңғы параметрлерін анықтау үшін үш өлшемді графикада қолданылатын синтезделген кескіндерге көлеңкелер салуға арналған бағдарлама. Әдетте жарықтың жұтылуы мен шашырауы, текстураның кескіні, шағылысу және сыну, көлеңкелеу, беттік орын ауыстыру және ерікті күрделіліктің өңдеуден кейінгі әсерлерінің сипаттамаларын қамтиды. Күрделі беттерді қарапайым геометриялық фигуралар арқылы бейнелеуге болады.

көрсету(ағылшын) көрсету) – визуализация, компьютерлік графикада программалық құралдарды қолдану арқылы модельден кескін алу процесі.

SDK(ағылшын тілінен қысқартылған. Бағдарламалық қамтамасыз етуді әзірлеу жинағы) бағдарламалық қамтамасыз етуді әзірлеу құралдарының жиынтығы болып табылады.

Орталық Есептеуіш Бөлім(ағылшын тілінен қысқартылған. Орталық процессор, сөзбе-сөз - орталық / негізгі / негізгі есептеу құрылғысы) - орталық (микро);машина нұсқауларын орындайтын құрылғы; есептеу операцияларын орындауға (операциялық жүйе мен қолданбалы бағдарламалық қамтамасыз етумен берілген) және барлық құрылғылардың жұмысын үйлестіруге жауапты аппараттық құрал.

GPU(ағылшын тілінен қысқартылған. Графикалық өңдеу блогы, сөзбе-сөз – графикалық есептеуіш құрылғы) – графикалық процессор; графиканы көрсетуді (визуализацияны) орындайтын бөлек құрылғы немесе ойын консолі. Қазіргі графикалық процессорлар компьютерлік графиканы нақты өңдеуде және көрсетуде өте тиімді. Қазіргі заманғы бейне адаптерлерде графикалық процессор 3D графикалық үдеткіш ретінде пайдаланылады, бірақ кейбір жағдайларда оны есептеулер үшін де пайдалануға болады ( GPGPU).

Мәселелер Орталық Есептеуіш Бөлім

Ұзақ уақыт бойы дәстүрлі өнімділіктің артуы, негізінен, тактілік жиіліктің дәйекті ұлғаюымен (өнімділіктің шамамен 80% тактілік жиілікпен анықталды) бір мезгілде транзисторлар санының бір мезгілде ұлғаюымен байланысты болды. чип. Дегенмен, тактілік жиіліктің одан әрі ұлғаюы (3,8 ГГц-тен астам тактілік жиілікте чиптер жай қызып кетеді!) Бірқатар іргелі физикалық кедергілерге төтеп береді (себебі технологиялық процесс атом өлшеміне дерлік жақындады: , және кремний атомының өлшемі шамамен 0,543 нм):

Біріншіден, кристалл өлшемдерінің кішіреюімен және тактілік жиілігінің жоғарылауымен транзисторлардың ағып кету тогы артады. Бұл электр қуатын тұтынудың ұлғаюына және жылу шығарудың ұлғаюына әкеледі;

Екіншіден, жоғары тактілік жылдамдықтардың артықшылығы ішінара жадқа қол жеткізу кідірістерімен өтеледі, өйткені жадқа қол жеткізу уақыттары сағаттық жылдамдықтардың артуына сәйкес келмейді;

Үшіншіден, кейбір қолданбалар үшін дәстүрлі сериялық архитектуралар тиімсіз болады, өйткені есептеулердің дәйекті ағынынан туындайтын өнімділік кедергісі «Вон Нейманның тығырыққа тірелуі» деп аталатындықтан, сағат жылдамдығы артады. Бұл ретте резистивті-сыйымдылық сигналды беру кідірістері артады, бұл тактілік жиіліктің ұлғаюымен байланысты қосымша кедергі болып табылады.

Даму GPU

дамуымен қатар GPU:

…

2008 жылдың қарашасы - Intel 4 ядролы желіні енгізді Intel Core i7келесі ұрпақ микроархитектурасына негізделген Нехалем. Процессорлар 2,6-3,2 ГГц тактілік жиілікте жұмыс істейді. 45 нм технологиялық технологияда жасалған.

2008 жылдың желтоқсаны – төрт ядролы жөнелтілімдер басталды AMD Phenom II 940(код атауы - Денеб). 3 ГГц жиілікте жұмыс істейді, 45 нм технологиялық технология бойынша шығарылады.

…

2009 жылдың мамыры – Компания AMD GPU нұсқасын ұсынды ATI Radeon HD 4890негізгі тактілік жиілігімен 850 МГц-тен 1 ГГц-ке дейін өсті. Бұл бірінші графика 1 ГГц жиілікте жұмыс істейтін процессор. Жиіліктің артуына байланысты чиптің өңдеу қуаты 1,36-дан 1,6 терафлопқа дейін өсті. Процессордың құрамында 800 (!) ядро ​​бар, бейне жадыны қолдайды GDDR5, DirectX 10.1, ATI CrossFireXжәне қазіргі заманғы бейне карта үлгілеріне тән барлық басқа технологиялар. Чип 55 нм технологиясы негізінде жасалған.

Негізгі айырмашылықтар GPU

Айырықша белгілері GPU(салыстырғанда ) мыналар:

– текстуралар мен күрделі графикалық объектілерді есептеу жылдамдығын барынша арттыруға бағытталған архитектура;

типтік қуаттың шыңы болып табылады GPUқарағанда әлдеқайда жоғары ;

– арнаулы құбыр сәулетінің арқасында, GPUқарағанда графикалық ақпаратты өңдеуде әлдеқайда тиімді.

«Жанр дағдарысы»

үшін «Жанр дағдарысы». 2005 жылға қарай жетілді - олар сол кезде пайда болды. Бірақ, технологияның дамуына қарамастан, кәдімгі өнімділіктің өсуі айтарлықтай төмендеді. Сонымен қатар өнімділік GPUөсуін жалғастыруда. Осылайша, 2003 жылға қарай бұл революциялық идея кристалданды - графиканың есептеу қуатын пайдалану. Графиктік процессорлар «графикалық емес» есептеулер үшін белсенді түрде қолданыла бастады (физикалық модельдеу, сигналдарды өңдеу, есептеу математикасы/геометрия, деректер қорымен операциялар, есептеу биологиясы, есептеу экономикасы, компьютерлік көру және т.б.).

Негізгі мәселе бағдарламалау үшін стандартты интерфейстің жоқтығы болды GPU. Әзірлеушілер пайдаланды OpenGLнемесе Direct3Dбірақ бұл өте ыңғайлы болды. корпорациясы NVIDIA(Графика, медиа және коммуникациялық процессорлардың, сондай-ақ сымсыз медиа процессорларының ірі өндірушілерінің бірі; 1993 жылы негізі қаланған) қандай да бір бірыңғай және ыңғайлы стандартты әзірлеумен айналысты - және технологияны енгізді. CUDA.

Қалай басталды

2006 - NVIDIAкөрсетеді CUDA™; есептеу техникасындағы революцияның басталуы GPU.

2007 - NVIDIAархитектурасын шығарады CUDA(түпнұсқа нұсқасы CUDA SDK 2007 жылы 15 ақпанда ұсынылды); номинация» Ең жақсы жаңа» журналынан Танымал ғылымжәне басылымнан «Оқырмандар таңдауы». HPCWire.

2008 жыл – технология NVIDIA CUDAбастап «Техникалық үздігі» номинациясында жеңіске жетті ДК журналы.

Не болды CUDA

CUDA(ағылшын тілінен қысқартылған. Бірыңғай құрылғы архитектурасын есептеу, сөзбе-сөз - құрылғылардың біртұтас есептеу архитектурасы) - сәулет (бағдарламалық және аппараттық құралдар жиынтығы) бойынша өндіруге мүмкіндік береді. GPUжалпы мақсаттағы есептеулер GPUшын мәнінде қуатты сопроцессор ретінде әрекет етеді.

Технология NVIDIA CUDA™программалау тіліндегі жалғыз өңдеу ортасы болып табылады C, бұл әзірлеушілерге графикалық процессорлардың өңдеу қуатының арқасында аз уақыт ішінде күрделі есептеу мәселелерін шешу үшін бағдарламалық жасақтаманы жасауға мүмкіндік береді. Қазірдің өзінде әлемде миллиондаған адамдар жұмыс істейді GPUқолдауымен CUDA, және мыңдаған бағдарламашылар қазірдің өзінде (тегін!) құралдарды пайдалануда CUDAқолданбаларды жеделдету және ең күрделі ресурсты қажет ететін тапсырмаларды шешу үшін - бейне және аудио кодтаудан мұнай мен газды барлауға, өнімді модельдеуге, медициналық бейнелеуге және ғылыми зерттеулерге дейін.

CUDAәзірлеушіге өз қалауы бойынша графикалық үдеткіштің нұсқаулар жинағына қол жеткізуді ұйымдастыруға және оның жадын басқаруға, онда күрделі параллельді есептеулерді ұйымдастыруға мүмкіндік береді. Қолдау көрсетілетін графикалық үдеткіш CUDAқазіргі сияқты қуатты бағдарламаланатын ашық архитектураға айналады. Мұның бәрі әзірлеушіге жабдыққа төмен деңгейлі, бөлінген және жоғары жылдамдықты қол жеткізуді қамтамасыз етеді CUDAкомпиляторлар, жөндеушілер, математикалық кітапханалар, бағдарламалық платформалар сияқты жоғары деңгейлі күрделі құралдарды құру үшін қажетті негіз.

Уральский, жетекші технология маманы NVIDIA, салыстыру GPUжәне , былай дейді: - Бұл жол талғамайтын көлік. Ол әрқашан және барлық жерде саяхаттайды, бірақ өте жылдам емес. А GPUспорттық көлік болып табылады. Нашар жолда ол жай ғана ешқайда кетпейді, бірақ жақсы қамтуды береді - және ол жол талғамайтын көлік ешқашан армандамаған барлық жылдамдығын көрсетеді! ..».

Технологиялық мүмкіндіктер CUDA

Дербес компьютерлерді шағын суперкомпьютерлерге айналдыруға арналған құрылғылар бұрыннан белгілі. Өткен ғасырдың 80-жылдарында нарықта сол кездегі жалпы ISA кеңейту слоттарына енгізілген транспьютерлер ұсынылды. Алғашында олардың сәйкес тапсырмалардағы көрсеткіштері әсерлі болды, бірақ кейін әмбебап процессорлардың өнімділігінің өсуі жеделдеді, олар параллельді есептеулерде өз позицияларын нығайтты, ал трансформаторларда ешқандай мән болмады. Мұндай құрылғылар әлі де бар болса да, олар әртүрлі мамандандырылған үдеткіштер болып табылады. Бірақ көбінесе олардың қолдану аясы тар және мұндай үдеткіштер кеңінен қолданылмайды.

Бірақ жақында параллельді есептеулер эстафетасы үш өлшемді ойындармен байланысты жаппай нарыққа көшті. 3D графикасында қолданылатын параллельді векторлық есептеулерге арналған көп ядролы процессорлары бар жалпы мақсаттағы құрылғылар жалпы мақсаттағы процессорлар жасай алмайтын ең жоғары өнімділікке жетеді. Әрине, максималды жылдамдыққа тек бірқатар ыңғайлы тапсырмаларда қол жеткізіледі және кейбір шектеулер бар, бірақ мұндай құрылғылар бастапқыда арналмаған салаларда кеңінен қолданыла бастады. тамаша үлгіМұндай параллельді процессор Sony-Toshiba-IBM альянсы әзірлеген және Sony PlayStation 3 ойын консолінде қолданылатын Cell процессоры, сондай-ақ нарық көшбасшыларының - Nvidia және AMD барлық заманауи бейне карталары.

Біз бүгін ұяшыққа қол тигізбейміз, бірақ ол бұрын пайда болған және қосымша векторлық мүмкіндіктері бар әмбебап процессор болса да, біз бүгін бұл туралы айтпаймыз. 3D бейне үдеткіштері үшін алғашқы жалпы мақсаттағы графикалық емес есептеу технологиялары GPGPU (General-Purpose computation on GPUs) бірнеше жыл бұрын пайда болды. Өйткені, заманауи бейне чиптер жүздеген математикалық орындау бірліктерін қамтиды және бұл қуат көптеген есептеуді қажет ететін қосымшаларды айтарлықтай жылдамдату үшін пайдаланылуы мүмкін. Ал графикалық процессорлардың қазіргі ұрпақтары жоғары деңгейлі бағдарламалау тілдерімен және осы мақалада қарастырылғандай аппараттық-бағдарламалық қамтамасыз ету архитектураларымен бірге осы мүмкіндіктерді ашады және оларды әлдеқайда қолжетімді етеді.

GPCPU құру қазіргі заманғы бейне чиптерді орындауға қабілетті жеткілікті жылдам және икемді шейдер бағдарламаларының пайда болуымен түрткі болды. Әзірлеушілер GPU-ны 3D қолданбаларындағы кескінді ғана емес, сонымен қатар басқа параллельді есептеулерде де қолдануға шешім қабылдады. GPGPU-де бұл үшін графикалық API пайдаланылды: OpenGL және Direct3D, деректер текстуралар түрінде бейне чипке жіберілген кезде және есептеу бағдарламаларышейдерлер ретінде жүктеледі. Бұл әдістің кемшіліктері бағдарламалаудың салыстырмалы түрде жоғары күрделілігі, төмен жылдамдықпроцессор мен графикалық процессор арасындағы деректер алмасу және басқа шектеулер, олар туралы кейінірек талқылаймыз.

GPU есептеулері дамыды және өте жылдам дамып келеді. Әрі қарай, екі ірі бейне чип өндірушілері, Nvidia және AMD сәйкесінше CUDA (Compute Unified Device Architecture) және CTM (Close To Metal немесе AMD Stream Computing) деп аталатын тиісті платформаларды әзірледі және жариялады. Алдыңғы GPU бағдарламалау үлгілерінен айырмашылығы, олар графикалық карталардың аппараттық мүмкіндіктеріне тікелей қол жеткізу арқылы жасалды. Платформалар бір-бірімен үйлеспейді, CUDA - C бағдарламалау тілінің кеңейтімі, ал CTM виртуалды машина, құрастыру кодын орындау. Бірақ екі платформа да дәстүрлі графикалық құбырды және сәйкес Direct3D немесе OpenGL интерфейстерін пайдалана отырып, алдыңғы GPGPU үлгілерінің кейбір маңызды шектеулерін жойды.

Әрине, OpenGL қолданатын ашық стандарттар ең портативті және әмбебап болып көрінеді, олар әртүрлі өндірушілердің бейне чиптері үшін бірдей кодты пайдалануға мүмкіндік береді. Бірақ мұндай әдістердің көптеген кемшіліктері бар, олар әлдеқайда аз икемді және қолдануға ыңғайлы емес. Сонымен қатар, олар қазіргі есептеуіш процессорларда бар жылдам ортақ (ортақ) жады сияқты белгілі бір бейне карталардың ерекше мүмкіндіктерін пайдалануға жол бермейді.

Сондықтан Nvidia өзінің компиляторы мен GPU есептеулеріне арналған кітапханалары бар C-тәрізді бағдарламалау тілі CUDA платформасын шығарды. Әрине, бейне чиптері үшін оңтайлы кодты жазу оңай емес және бұл тапсырма ұзақ қолмен жұмыс істеуді қажет етеді, бірақ CUDA жай ғана барлық мүмкіндіктерді ашады және бағдарламашыға GPU аппараттық мүмкіндіктерін көбірек басқаруға мүмкіндік береді. Nvidia CUDA қолдауы өте кең таралған Geforce 8, 9 және 200 сериялы бейне карталарында қолданылатын G8x, G9x және GT2xx чиптері үшін қол жетімді болуы маңызды. CUDA 2.0 соңғы нұсқасы енді шығарылды, оның бірнеше жаңа мүмкіндіктері бар, мысалы, қос дәлдіктегі есептеулерді қолдау. CUDA 32 биттік және 64 биттік Linux, Windows және MacOS X операциялық жүйелерінде қол жетімді.

Параллельді есептеудегі процессор мен графикалық процессор арасындағы айырмашылық

Әмбебап процессорлардың жиіліктерінің өсуі физикалық шектеулерге және жоғары қуат тұтынуына әкелді және олардың өнімділігі бірнеше ядролардың бір чипте орналасуына байланысты барған сайын артып келеді. Қазір сатылған процессорларда тек төрт ядроға дейін бар (бұдан әрі өсу жылдам болмайды) және олар жалпы қолданбаларға арналған, MIMD пайдаланыңыз - бірнеше нұсқаулар мен деректер ағыны. Әрбір ядро ​​басқаларынан бөлек жұмыс істейді, әртүрлі процестер үшін әртүрлі нұсқауларды орындайды.

4 компонентті (бір дәлдіктегі өзгермелі нүкте) және екі компонентті (қос дәлдік) векторларға арналған мамандандырылған векторлық мүмкіндіктер (SSE2 және SSE3) бірінші кезекте графикалық қолданбаларға қойылатын талаптардың артуына байланысты жалпы мақсаттағы процессорларда пайда болды. Сондықтан белгілі бір тапсырмалар үшін графикалық процессорларды пайдалану тиімдірек, өйткені олар бастапқыда олар үшін жасалған.

Мысалы, Nvidia бейне чиптерінде негізгі блок сегізден онға дейінгі ядролары және барлығы жүздеген ALU, бірнеше мың регистрлері және ортақ жадының шағын көлемі бар мультипроцессор болып табылады. Сонымен қатар, видеокартада оған барлық мультипроцессорлардың қол жеткізе алатын жылдам ғаламдық жады, әрбір мультипроцессордағы жергілікті жады және тұрақтыларға арналған арнайы жады бар.

Ең бастысы, GPU-дағы бұл бірнеше мультипроцессорлық ядролар SIMD (бір нұсқаулық ағыны, бірнеше деректер ағыны) ядролары болып табылады. Және бұл ядролар бір уақытта бірдей нұсқауларды орындайды, бағдарламалаудың бұл стилі графикалық алгоритмдерде және көптеген ғылыми тапсырмаларда жиі кездеседі, бірақ нақты бағдарламалауды қажет етеді. Бірақ бұл тәсіл оларды жеңілдету есебінен орындау бірліктерінің санын көбейтуге мүмкіндік береді.

Сонымен, CPU және GPU архитектураларының негізгі айырмашылықтарын тізіп көрейік. Орталық процессордың өзектері ең жоғары өнімділікпен дәйекті нұсқаулардың бір ағынын орындауға арналған, ал графикалық процессорлар параллель нұсқаулар ағындарының көп санын жылдам орындауға арналған. Жалпы мақсаттағы процессорлар бүтін сандарды да, өзгермелі нүкте сандарын да өңдейтін бір нұсқаулар ағынында жоғары өнімділікке қол жеткізу үшін оңтайландырылған. Жадқа кіру кездейсоқ.

CPU дизайнерлері өнімділікті арттыру үшін параллельді орындау үшін мүмкіндігінше көп нұсқаулар алуға тырысады. Ол үшін Intel Pentium процессорларынан бастап сағатына екі команданың орындалуын қамтамасыз ететін суперскалярлық орындау пайда болды, ал Pentium Pro инструкциялардың ретсіз орындалуымен ерекшеленді. Бірақ командалардың дәйекті ағынының параллель орындалуы белгілі бір негізгі шектеулерге ие және орындау бірліктерінің санын көбейту арқылы жылдамдықты бірнеше есе арттыруға қол жеткізу мүмкін емес.

Бейне чиптер ең басынан бастап қарапайым және параллельді жұмыс істейді. Бейне чип кірісте көпбұрыштар тобын алады, барлық қажетті операцияларды орындайды және шығыста пикселдер шығарады. Көпбұрыштар мен пикселдерді өңдеу тәуелсіз, оларды бір-бірінен бөлек, параллельді өңдеуге болады. Сондықтан, графикалық процессордағы жұмыстың параллельді ұйымдастырылуына байланысты, орталық процессорға арналған нұсқаулардың дәйекті ағынынан айырмашылығы оңай жүктелетін орындау бірліктерінің үлкен саны қолданылады. Сонымен қатар, қазіргі заманғы графикалық процессорлар сағатына бірден көп нұсқауларды орындай алады (қос шығарылым). Осылайша, Tesla архитектурасы белгілі бір жағдайларда MAD+MUL немесе MAD+SFU операцияларын бір уақытта іске қосады.

GPU процессордан жадқа қол жеткізу принциптері бойынша да ерекшеленеді. GPU-да ол қосылған және оңай болжауға болады - егер текстуралық текст жадтан оқылса, біраз уақыттан кейін көрші текстер үшін уақыт келеді. Иә, және дәл солай жазу кезінде - пиксель кадр буферіне жазылады және бірнеше циклден кейін оның жанында орналасқаны жазылады. Сондықтан жадты ұйымдастыру орталық процессорда қолданылатыннан өзгеше. Ал бейне чип, әмбебап процессорлардан айырмашылығы, жай ғана үлкен кэш жадын қажет етпейді, ал текстуралар тек бірнеше (қазіргі GPU-да 128-256 дейін) килобайтты қажет етеді.

Және GPU және CPU үшін жадпен жұмыс біршама ерекшеленеді. Сонымен, барлық процессорларда кірістірілген жад контроллері бола бермейді және барлық графикалық процессорларда әдетте Nvidia GT200 чипіндегі сегіз 64 биттік арнаға дейін бірнеше контроллер болады. Сонымен қатар, видеокарталар жылдамырақ жадты пайдаланады және нәтижесінде бейне микросхемаларда жад өткізу қабілеттілігі бірнеше есе көп, бұл үлкен деректер ағындарымен жұмыс істейтін параллельді есептеулер үшін де өте маңызды.

Жалпы мақсаттағы процессорларда транзисторлардың және микросхемалардың үлкен саны нұсқау буферлеріне, аппараттық құрал тармақтарын болжауға және микросхемадағы кэш жадының үлкен көлеміне барады. Барлық осы аппараттық блоктар бірнеше нұсқаулар ағынының орындалуын жылдамдату үшін қажет. Бейне чиптер транзисторларды орындау блоктарының массивтеріне, ағынды басқару блоктарына, шағын ортақ жадыға және көп арналы жады контроллерлеріне жұмсайды. Жоғарыда айтылғандар жеке ағындардың орындалуын тездетпейді, ол микросхемаға бір уақытта микросхемада орындалатын және жад өткізу қабілетінің жоғары болуын қажет ететін бірнеше мың ағындарды өңдеуге мүмкіндік береді.

Кэштеудегі айырмашылықтар туралы. Жалпы мақсаттағы процессорлар жадқа кіру кідірісін азайту арқылы өнімділікті арттыру үшін кэшті пайдаланады, ал GPU өткізу қабілеттілігін арттыру үшін кэшті немесе ортақ жадты пайдаланады. Орталық процессорлар үлкен кэштермен және код тармағын болжаумен жадқа кіру кідірістерін азайтады. Бұл аппараттық бөліктер чип аймағының көп бөлігін алады және көп қуатты тұтынады. Бейне чиптері бір уақытта мыңдаған ағындарды орындау арқылы жадқа қол жеткізудің кідірістері мәселесін айналып өтеді - ағындардың бірі жадтан деректерді күтіп тұрғанда, бейне чип күтусіз және кідіріссіз басқа ағынның есептеулерін орындай алады.

Көп ағынды қолдауда да көптеген айырмашылықтар бар. Орталық процессор бір процессор өзегіне 1-2 есептеу ағынын орындайды, ал бейне чиптер бір мультипроцессорға 1024 ағынға дейін қолдау көрсете алады, оның ішінде чипте бірнеше бар. Ал егер процессор үшін бір ағыннан екіншісіне ауысу жүздеген циклді қажет ететін болса, онда GPU бір циклде бірнеше ағындарды ауыстырады.

Сонымен қатар, процессорлар векторлық есептеулер үшін SIMD (бір нұсқау, бірнеше деректер) блоктарын, ал графикалық процессорлар скалярлық ағынды өңдеу үшін SIMT (бір нұсқау, бірнеше ағындар) пайдаланады. SIMT әзірлеушіден деректерді векторларға түрлендіруді талап етпейді және ағындарда еркін тармақталуға мүмкіндік береді.

Қысқаша айтқанда, қазіргі әмбебап процессорлардан айырмашылығы, бейне чиптер арифметикалық операциялардың көп санымен параллельді есептеулерге арналған деп айта аламыз. Ал GPU транзисторларының анағұрлым көп саны өз мақсатына - деректер массивтерін өңдеуге жұмыс істейді және бірнеше реттік есептеу ағындарының орындалуын (ағынын басқару) бақыламайды. Бұл процессор мен графикалық процессорда қанша орын алатынын көрсететін диаграмма әртүрлі логика:

Нәтижесінде, графикалық процессордың қуатын ғылыми және басқа графикалық емес есептеулерде тиімді пайдаланудың негізі алгоритмдерді бейне чиптерде қол жетімді жүздеген орындау бірліктеріне параллельизациялау болып табылады. Мысалы, молекулярлық модельдеудің көптеген қосымшалары бейне микросхемалардағы есептеулер үшін өте қолайлы, олар үлкен есептеу қуатын қажет етеді, сондықтан параллельді есептеулерге ыңғайлы. Бірнеше GPU пайдалану мұндай мәселелерді шешу үшін одан да көп есептеу қуатын береді.

GPU-да есептеулерді орындау деректерді параллель өңдеуді пайдаланатын алгоритмдерде тамаша нәтижелерді көрсетеді. Яғни, деректердің үлкен көлеміне бірдей математикалық операциялар тізбегі қолданылғанда. Бұл жағдайда ең жақсы нәтижеге арифметикалық нұсқаулар санының жадыға кіру санына қатынасы жеткілікті үлкен болған жағдайда қол жеткізіледі. Бұл ағынды басқаруға аз талаптар қояды, ал математиканың жоғары тығыздығы мен деректердің үлкен көлемі орталық процессордағы сияқты үлкен кэштерді қажет етпейді.

Жоғарыда сипатталған барлық айырмашылықтардың нәтижесінде бейне чиптердің теориялық өнімділігі орталық процессордың өнімділігінен айтарлықтай асып түседі. Nvidia соңғы бірнеше жылдағы CPU және GPU өнімділігінің өсу графигін ұсынады:

Әрине, бұл деректерде қулық жоқ емес. Шынында да, процессорда теориялық көрсеткіштерге практикада қол жеткізу оңайырақ, ал цифрлар графикалық процессор жағдайында бір дәлдік үшін, ал орталық процессор жағдайында қос дәлдік үшін берілген. Кез келген жағдайда, кейбір параллельді тапсырмалар үшін бір дәлдік жеткілікті және әмбебап және графикалық процессорлар арасындағы жылдамдық айырмашылығы өте үлкен, сондықтан ойын шамға тұрарлық.

GPU-да есептеулерді қолданудың алғашқы әрекеттері

Бейне чиптер ұзақ уақыт бойы параллельді математикалық есептеулерде қолданылған. Мұндай қолданбаның алғашқы әрекеттері өте қарапайым болды және растризация және Z-буферлеу сияқты кейбір аппараттық мүмкіндіктерді пайдаланумен шектелді. Бірақ қазіргі ғасырда шейдерлердің пайда болуымен олар матрицаларды есептеуді тездете бастады. 2003 жылы SIGGRAPH-те GPU есептеулері үшін жеке бөлім бөлінді және ол GPGPU (GPU бойынша жалпы мақсаттағы есептеу) - әмбебап GPU есептеуі) деп аталды.

Ең танымал BrookGPU графикалық процессорында графикалық емес есептеулерді орындауға арналған Brook stream бағдарламалау тілінің компиляторы болып табылады. Оның пайда болуына дейін есептеулер үшін бейне чиптердің мүмкіндіктерін пайдаланатын әзірлеушілер екі жалпы API интерфейсінің бірін таңдады: Direct3D немесе OpenGL. Бұл GPU пайдалануды айтарлықтай шектеді, өйткені 3D графикасы параллельді бағдарламашылар білуі талап етілмейтін шейдерлер мен текстураларды пайдаланады, олар ағындар мен өзектерді пайдаланады. Брук олардың жұмысын жеңілдетуге көмектесе алды. Стэнфорд университетінде әзірленген Си тіліне арналған бұл ағындық кеңейтімдер 3D API-ді бағдарламашылардан жасырды және бейне чипті параллельді сопроцессор ретінде ұсынды. Компилятор DirectX, OpenGL немесе x86 қосылған кітапханамен байланыстырылған кодты шығара отырып, C++ коды және кеңейтімдері бар .br файлын талдады.

Әрине, Бруктың көптеген кемшіліктері болды, біз оларға тоқталамыз және оларды кейінірек толығырақ қарастырамыз. Бірақ оның пайда болуының өзі Nvidia мен ATI-дің GPU есептеу бастамасына назарының айтарлықтай артуына әкелді, өйткені бұл мүмкіндіктердің дамуы болашақта нарықты түбегейлі өзгертіп, оның жаңа секторын - параллельді есептеулерді ашады. бейне чиптері.

Әрі қарай, Brook жобасының кейбір зерттеушілері Nvidia әзірлеу тобына жаңа нарық үлесін ашып, аппараттық-бағдарламалық қамтамасыз ету параллельді есептеу стратегиясын енгізу үшін қосылды. Бұл Nvidia бастамасының басты артықшылығы - әзірлеушілер өздерінің GPU-ларының барлық мүмкіндіктерін ең кішкентай бөлшектерге дейін жақсы біледі және графикалық API пайдаланудың қажеті жоқ, сонымен қатар драйверді пайдаланып жабдықпен тікелей жұмыс істей аласыз. Осы топтың күш-жігерінің нәтижесі - Nvidia CUDA (Есептеу біртұтас құрылғы архитектурасы), осы мақаланың тақырыбы болып табылатын Nvidia GPU-да параллельді есептеулерге арналған жаңа аппараттық және бағдарламалық қамтамасыз ету архитектурасы.

GPU-да параллельді есептеулерді қолдану аймақтары

Есептеулерді бейне чиптерге көшіру қандай артықшылықтар әкелетінін түсіну үшін біз бүкіл әлем бойынша зерттеушілер алған орташа көрсеткіштерді береміз. Орташа алғанда, есептеулерді GPU-ге тасымалдау кезінде, көптеген тапсырмаларда жылдам әмбебап процессорлармен салыстырғанда жеделдету 5-30 есеге жетеді. Ең үлкен сандар (100 есе жылдамдату тәртібі және одан да көп!) SSE блоктарын қолданатын есептеулер үшін өте қолайлы емес, бірақ GPU үшін өте ыңғайлы кодта қол жеткізіледі.

Бұл процессордағы SSE векторланған кодқа қарсы GPU-дағы синтетикалық кодтың жылдамдығын арттырудың кейбір мысалдары (Nvidia сәйкес):

• Флуоресцентті микроскопия: 12x;
• Молекулалық динамика (байланыссыз күш есептеу): 8-16x;
• Электростатика (тікелей және көп деңгейлі кулон қосындысы): 40-120x және 7x.

Бұл Nvidia өте жақсы көретін тақтайша, оны барлық презентацияларда көрсетеді, біз оны мақаланың екінші бөлігінде толығырақ қарастырамыз, CUDA есептеулерін практикалық қолданудың нақты мысалдарына арналған:

Көріп отырғаныңыздай, сандар өте тартымды, әсіресе 100-150 еселік табыстар әсерлі. Келесі CUDA мақаласында біз осы сандардың кейбірін егжей-тегжейлі қарастырамыз. Енді біз GPU есептеулері қолданылатын негізгі қолданбаларды тізімдейміз: кескіндер мен сигналдарды талдау және өңдеу, физика модельдеу, есептеу математикасы, есептеу биологиясы, қаржылық есептеулер, мәліметтер базасы, газдар мен сұйықтықтардың динамикасы, криптография, адаптивті сәулелік терапия, астрономия , дыбысты өңдеу, биоинформатика, биологиялық модельдеу, компьютерлік көру, деректерді өндіру, цифрлық кино және теледидар, электромагниттік модельдеу, географиялық ақпараттық жүйелер, әскери қолданбалар, тау-кен жұмыстарын жоспарлау, молекулалық динамика, магниттік-резонанстық бейнелеу (МРТ), нейрондық желілер, океанографиялық зерттеулер, бөлшектер физика, ақуыздың қатпарлануын модельдеу, кванттық химия, сәулелік бақылау, бейнелеу, радар, резервуарды модельдеу, жасанды интеллект, спутниктік деректерді талдау, сейсмикалық барлау, хирургия, ультрадыбыстық, бейнеконференция.

Көптеген қолданбалар туралы мәліметтерді Nvidia веб-сайтында бөлімде табуға болады. Көріп отырғаныңыздай, тізім өте үлкен, бірақ бұл бәрі емес! Оны жалғастыруға болады және болашақта бейне микросхемаларда параллельді есептеулерді қолданудың басқа салалары табылатыны сөзсіз, бұл туралы бізде әлі де түсінік жоқ.

Nvidia CUDA мүмкіндіктері

CUDA технологиясы - бұл графикалық үдеткіштің нұсқаулар жинағына қол жеткізуге және оның жадын параллельді есептеуде басқаруға мүмкіндік беретін Си тілінің кеңейтіміне негізделген Nvidia бағдарламалық және аппараттық есептеулер архитектурасы. CUDA сегізінші және одан жоғары буын Geforce бейне үдеткіштерінің (Geforce 8, Geforce 9, Geforce 200 сериялары), сондай-ақ Quadro және Tesla графикалық процессорларында іске асырылуы мүмкін алгоритмдерді жүзеге асыруға көмектеседі.

CUDA көмегімен GPU бағдарламалаудың күрделілігі айтарлықтай жоғары болғанымен, ол GPGPU ерте шешімдеріне қарағанда төмен. Мұндай бағдарламалар қолданбаны MPI бағдарламалауға ұқсас, бірақ ортақ бейне жадында сақталған деректерді ортақ пайдаланбай бірнеше мультипроцессорлар арасында бөлуді талап етеді. Әр мультипроцессорға арналған CUDA бағдарламалау OpenMP бағдарламалауға ұқсас болғандықтан, ол жадты ұйымдастыруды жақсы түсінуді талап етеді. Бірақ, әрине, әзірлеу және CUDA-ға көшіру күрделілігі қолданбаға өте тәуелді.

Әзірлеушілер жинағы көптеген код мысалдарын қамтиды және жақсы құжатталған. OpenMP және MPI-мен бұрыннан таныс адамдар үшін оқу процесі шамамен екі-төрт аптаға созылады. API кеңейтілген C тіліне негізделген және осы тілден кодты аудару үшін CUDA SDK ашық Open64 компиляторына негізделген nvcc пәрмен жолы компиляторын қамтиды.

Біз CUDA негізгі сипаттамаларын тізімдейміз:

• Nvidia бейне чиптерінде параллельді есептеулер үшін бірыңғай бағдарламалық және аппараттық шешім;
• мобильдіден мультичипке дейін қолдау көрсетілетін шешімдердің кең ауқымы
• стандартты C бағдарламалау тілі;
• FFT (Fast Furier Transform) және BLAS (Сызықтық алгебра) сандық талдауға арналған стандартты кітапханалар;
• CPU және GPU арасындағы оңтайландырылған деректер алмасу;
• API OpenGL және DirectX графикасымен өзара әрекеттесу;
• 32 және 64 биттік операциялық жүйелерді қолдау: Windows XP, Windows Vista, Linux және MacOS X;
• төмен деңгейде даму мүмкіндігі.

Операциялық жүйелерді қолдауға келетін болсақ, барлық негізгілерді қосу керек Linux дистрибутивтері(Red Hat Enterprise Linux 3.x/4.x/5.x, SUSE Linux 10.x), бірақ энтузиастардың пікірінше, CUDA басқа құрылымдарда жақсы жұмыс істейді: Fedora Core, Ubuntu, Gentoo және т.б.

CUDA әзірлеу ортасы (CUDA құралдар жинағы) мыналарды қамтиды:

• nvcc компиляторы;
• FFT және BLAS кітапханалары;
• профиль жасаушы;
• GPU үшін gdb жөндеу құралы;
• CUDA жұмыс уақыты драйвері стандартты Nvidia драйверлеріне кіреді
• бағдарламалау бойынша нұсқаулық;
• CUDA Developer SDK (бастапқы код, утилиталар және құжаттама).

Мысалдарда бастапқы код: параллельдік битоникалық сұрыптау (битоникалық сұрыптау), матрицалық транспозиция, үлкен массивтердің параллель префикстерінің қосындысы, кескін конволюциясы, дискретті толқындық түрлендіру, OpenGL және Direct3D-мен өзара әрекеттесу мысалы, CUBLAS және CUFFT кітапханаларын пайдалану, опцион бағасын есептеу (Блэк-Скоулз формуласы, биномдық модель, Монте-Карло әдісі), Mersenne Twister параллель кездейсоқ сандар генераторы, үлкен массив гистограммасын есептеу, шуды азайту, Собель сүзгісі (шеттерді табу).

CUDA артықшылықтары мен шектеулері

Программист көзқарасы бойынша графикалық конвейер өңдеу кезеңдерінің жиынтығы болып табылады. Геометриялық блок үшбұрыштарды, ал растризация блогы мониторда көрсетілетін пикселдерді жасайды. Дәстүрлі GPGPU бағдарламалау үлгісі келесідей:

Осындай модель шеңберіндегі есептеулерді GPU-ға беру үшін арнайы тәсіл қажет. Тіпті екі векторды элемент бойынша қосу үшін кескінді экранға немесе экраннан тыс буферге салу қажет болады. Сурет растерленген, әрбір пиксельдің түсі берілген бағдарлама бойынша есептеледі (пиксель шейдер). Бағдарлама әрбір пиксель үшін текстуралардан енгізілген деректерді оқиды, оларды қосады және шығыс буферіне жазады. Және бұл көптеген операциялардың барлығы кәдімгі бағдарламалау тілінде бір операторда жазылған нәрсе үшін қажет!

Сондықтан GPGPU-ны жалпы мақсаттағы есептеулер үшін пайдалану әзірлеушілерге үйрену үшін тым көп күрделілік түріндегі шектеулерге ие. Және басқа да шектеулер жеткілікті, өйткені пиксельдік шейдер - бұл пикселдің соңғы түсінің оның координаталарына тәуелділігінің формуласы ғана, ал пиксель шейдерінің тілі - бұл формулаларды С тәрізді синтаксиспен жазуға арналған тіл. Алғашқы GPGPU әдістері GPU қуатын пайдаланудың ақылды айласы болып табылады, бірақ ешқандай ыңғайлылықсыз. Ондағы деректер кескіндермен (текстуралармен), ал алгоритм растризация процесімен көрсетіледі. Айта кету керек, жад пен орындаудың өте нақты үлгісі.

Nvidia компаниясының GPU құрылғыларында есептеуге арналған аппараттық және бағдарламалық жасақтама архитектурасы Nvidia компаниясының GPGPU үлгілерінен ерекшеленеді, ол стандартты синтаксиспен, көрсеткіштермен және бейне чиптерінің есептеу ресурстарына қол жеткізу үшін ең аз кеңейтімдерді қажет ететін нақты C тілінде GPU үшін бағдарламаларды жазуға мүмкіндік береді. CUDA графикалық API интерфейстеріне тәуелді емес және жалпы мақсаттағы есептеулер үшін арнайы жасалған кейбір мүмкіндіктерге ие.

GPGPU есептеуіндегі дәстүрлі тәсілмен салыстырғанда CUDA артықшылықтары:

• CUDA қолданбалы бағдарламалау интерфейсі кеңейтімдері бар стандартты Си бағдарламалау тіліне негізделген, бұл CUDA архитектурасын үйрену және енгізу процесін жеңілдетеді;
• CUDA бір мультипроцессорға 16 Кбайт ортақ жадқа қолжетімділікті қамтамасыз етеді, ол текстураны алудан жоғары өткізу қабілеттілігі бар кэшті ұйымдастыру үшін пайдаланылуы мүмкін;
• жүйе мен бейне жады арасында деректерді тиімдірек тасымалдау
• артық және үстеме шығындары бар графикалық API қажет емес;
• желілік жадты адрестеу, жинақтау және шашырату, ерікті адрестерге жазу мүмкіндігі;
• бүтін және разрядтық операцияларды аппараттық қамтамасыз ету.

CUDA негізгі шектеулері:

• орындалатын функциялар үшін рекурсиялық қолдаудың болмауы;
• блоктың ең төменгі ені - 32 жіп;
• Nvidia иелігіндегі меншікті CUDA архитектурасы.

Алдыңғы GPGPU әдістерімен бағдарламалаудың әлсіз жақтары мынада: бұл әдістер алдыңғы бірікпеген архитектураларда шыңы шейдер орындау бірліктерін пайдаланбайды, деректер текстураларда сақталады және экраннан тыс буферге шығарылады, ал көп өту алгоритмдері пиксельдік шейдер бірліктерін пайдаланады. GPGPU шектеулеріне мыналар жатады: аппараттық құралдар мүмкіндіктерін жеткіліксіз тиімді пайдалану, жад өткізу қабілеттілігін шектеу, шашырау операциясының болмауы (тек жинау), графикалық API міндетті түрде пайдалану.

CUDA-ның алдыңғы GPGPU әдістеріне қарағанда негізгі артықшылықтары бұл архитектураның GPU-де графикалық емес есептеулерді тиімді пайдалануға арналған және алгоритмдерді графикалық концепцияға ыңғайлы пішінге көшіруді қажет етпестен C бағдарламалау тілін пайдалану фактісіне байланысты. құбыр. CUDA графикалық API интерфейстерін пайдаланбайтын, кездейсоқ жадқа қол жеткізуді (шашырау немесе жинау) ұсынатын жаңа GPU есептеу жолын ұсынады. Мұндай архитектура GPGPU кемшіліктерінен таза және барлық орындау бірліктерін пайдаланады, сонымен қатар бүтін математика және разрядты ауыстыру операциялары арқылы мүмкіндіктерді кеңейтеді.

Сонымен қатар, CUDA графикалық API интерфейстерінде қолжетімсіз кейбір аппараттық құралдар мүмкіндіктерін ашады, мысалы, ортақ жад. Бұл ағындар блоктары қол жеткізе алатын жадтың шағын көлемі (әр мультипроцессорға 16 килобайт). Ол ең жиі қолданылатын деректерді кэштеуге мүмкіндік береді және осы тапсырма үшін текстураны алуды пайдаланудан гөрі жылдамырақ өнімділікті қамтамасыз ете алады. Бұл, өз кезегінде, көптеген қолданбаларда параллельді алгоритмдердің өткізу қабілетінің сезімталдығын төмендетеді. Мысалы, ол сызықтық алгебра, жылдам Фурье түрлендіру және кескінді өңдеу сүзгілері үшін пайдалы.

CUDA және жадқа қол жеткізуде ыңғайлы. Графикалық API коды деректерді алдын ала анықталған аймақтарда 32 бір дәлдіктегі өзгермелі нүкте мәні (RGBA мәндері сегіз көрсету мақсатына бір мезгілде) ретінде шығарады және CUDA шашырауды жазуды қолдайды - кез келген мекенжайдағы жазбалардың шектеусіз саны. Мұндай артықшылықтар графикалық API негізіндегі GPGPU әдістерін қолдану арқылы тиімді жүзеге асырылмайтын кейбір алгоритмдерді GPU-да орындауға мүмкіндік береді.

Сондай-ақ графикалық API интерфейстері міндетті түрде деректерді текстураларда сақтайды, бұл үлкен массивтерді текстураларға алдын ала орауды талап етеді, бұл алгоритмді қиындатады және арнайы адрестеуді қолдануды мәжбүр етеді. Ал CUDA кез келген мекенжайдағы деректерді оқуға мүмкіндік береді. CUDA тағы бір артықшылығы - процессор мен графикалық процессор арасындағы оңтайландырылған байланыс. Ал төменгі деңгейге қол жеткізгісі келетін әзірлеушілер үшін (мысалы, басқа бағдарламалау тілін жазғанда) CUDA төменгі деңгейлі ассемблер тілінде бағдарламалау мүмкіндігін ұсынады.

CUDA даму тарихы

CUDA әзірлеу G80 чипімен бірге 2006 жылдың қарашасында жарияланды және шығарылымы қоғамдық бета нұсқасы CUDA SDK 2007 жылдың ақпан айында өтті. 2007 жылдың маусымында 1.0 нұсқасы жоғары өнімді есептеуіш нарық үшін G80 чипіне негізделген Tesla шешімдерін шығару үшін шығарылды. Содан кейін, жылдың соңында CUDA 1.1 бета нұсқасы шығарылды, ол нұсқа санының шамалы өсуіне қарамастан, көптеген жаңа нәрселерді енгізді.

CUDA 1.1-де пайда болғаннан біз CUDA функционалдық тұрақты Nvidia бейне драйверлеріне қосылғанын атап өтуге болады. Бұл кез келген CUDA бағдарламасына қойылатын талаптарда Geforce 8 және одан жоғары сериялы бейне картасын, сондай-ақ драйвердің 169.xx ең аз нұсқасын көрсету жеткілікті екенін білдірді. Бұл әзірлеушілер үшін өте маңызды, егер бұл шарттар орындалса, CUDA бағдарламалары кез келген пайдаланушы үшін жұмыс істейді. Сондай-ақ, деректерді көшіру (тек G84, G86, G92 және одан жоғары микросхемалар үшін), деректерді бейне жадыға асинхронды тасымалдау, атомдық жадыға қол жеткізу операциялары, Windows жүйесінің 64 биттік нұсқаларын қолдау және көп файлды орнату мүмкіндігі сияқты асинхронды орындау қосылды. -чип SLI режимінде CUDA жұмысы.

Қазіргі уақытта ағымдағы нұсқа Geforce GTX 200 желісімен бірге шығарылған GT200 - CUDA 2.0 негізіндегі шешімдерге арналған.Бета нұсқасы 2008 жылдың көктемінде шығарылды. Екінші нұсқада: қос дәлдіктегі есептеулерді қолдау (тек GT200 үшін аппараттық қолдау), Windows Vista (32 және 64 биттік нұсқалары) және Mac OS X соңғы қолдауға ие, жөндеу және профильдеу құралдары қосылды, 3D текстураларына қолдау көрсетіледі, оңтайландырылған. деректерді тасымалдау.

Екі еселенген дәлдікпен есептеулерге келетін болсақ, олардың ағымдағы аппараттық генерациядағы жылдамдығы бір дәлдікке қарағанда бірнеше есе төмен. Себептер бізде талқыланады. GT200-де бұл қолдауды жүзеге асыру FP32 блоктарының нәтижелерді төрт есе баяу қарқынмен алу үшін пайдаланылмайтындығында, FP64 есептеулерін қолдау үшін Nvidia арнайы есептеу блоктарын жасауды шешті. Ал GT200-де олардың саны FP32 блоктарынан он есе аз (әр мультипроцессор үшін бір қос дәлдік блок).

Шындығында өнімділік одан да төмен болуы мүмкін, өйткені архитектура жад пен регистрлерден 32-биттік оқу үшін оңтайландырылған, сонымен қатар графикалық қосымшаларда қос дәлдік қажет емес, ал GT200-де ол жай ғана болуы мүмкін. Иә, және қазіргі төрт ядролы процессорлар нақты өнімділікті көрсетеді. Бірақ бір дәлдікке қарағанда тіпті 10 есе баяу болғандықтан, бұл тірек аралас дәлдіктегі тізбектер үшін пайдалы. Кең таралған әдістердің бірі - бастапқыда бір дәлдікте шамамен алынған нәтижелерді алу, содан кейін оларды қосарланған дәлдікте нақтылау. Енді мұны процессорға аралық деректерді жіберместен тікелей бейне картада жасауға болады.

CUDA 2.0 тағы бір пайдалы мүмкіндігінің GPU-ға ешқандай қатысы жоқ, таңқаларлық. Қазір ғана CUDA кодын орталық процессорда жылдам орындау үшін жоғары тиімді көп ағынды SSE кодына құрастыруға болады. Яғни, қазір бұл мүмкіндік отладка үшін ғана емес, сонымен қатар қолайлы нақты пайдалану Nvidia графикалық картасы жоқ жүйелерде. Ақыр соңында, CUDA-ны қалыпты кодта пайдалану Nvidia видеокарталары арнайы бейне шешімдерінің арасында ең танымал болғанымен, барлық жүйелерде қол жетімді емес екендігімен шектеледі. Ал 2.0 нұсқасына дейін мұндай жағдайларда екі түрлі код жазылуы керек еді: CUDA үшін және орталық процессор үшін бөлек. Ал енді кез келген CUDA бағдарламасын процессорда бейне чиптерге қарағанда төмен жылдамдықпен болса да жоғары тиімділікпен іске қосуға болады.

Nvidia CUDA қолдайтын шешімдер

CUDA қолдайтын барлық графикалық карталар аудио және бейне өңдеуден бастап медициналық және ғылыми зерттеулерге дейін ең күрделі тапсырмаларды жеделдетуге көмектеседі. Жалғыз нақты шектеу - көптеген CUDA бағдарламалары үшін ең аз дегенде 256 мегабайт бейне жады қажет және бұл CUDA қолданбалары үшін ең маңызды сипаттамалардың бірі.

CUDA қолдайтын өнімдердің жаңартылған тізімін мына жерден табуға болады. Осы жазбаны жазу кезінде CUDA есептеулері Geforce 200, Geforce 9 және Geforce 8 серияларының барлық өнімдерін, соның ішінде Geforce 8400M бастап мобильді өнімдерді, сондай-ақ Geforce 8100, 8200 және 8300 чипсеттерін қолдады. Қазіргі заманғы Quadro және барлық Tesla: S1070, C1060, C870, D870 және S870.

Жаңа Geforce GTX 260 және 280 видеокарталарымен бірге сәйкес жоғары өнімді есептеу шешімдері жарияланғанын атап өтеміз: Tesla C1060 және S1070 (жоғарыдағы фотода көрсетілген), оларды осы күзде сатып алуға болады. Оларда бірдей графикалық процессор қолданылады - GT200, C1060-та ол біреу, S1070-те - төрт. Бірақ, ойын шешімдерінен айырмашылығы, олар бір чипке төрт гигабайт жадты пайдаланады. Минустардың ішінде, мүмкін, бір чипке 102 гигабайт/с қамтамасыз ететін ойын карталарына қарағанда жад жиілігі мен жад өткізу қабілеттілігінің төмен болуы.

Nvidia CUDA құрамы

CUDA құрамына екі API кіреді: жоғары деңгейлі (CUDA Runtime API) және төмен деңгейлі (CUDA Driver API), екеуін бір бағдарламада бір уақытта пайдалану мүмкін болмаса да, біреуін немесе екіншісін пайдалану керек. Жоғары деңгей төменгі деңгейдің «жоғарғы жағында» жұмыс істейді, барлық орындалу уақытындағы қоңыраулар төменгі деңгейлі Driver API өңдейтін қарапайым нұсқауларға аударылады. Бірақ тіпті «жоғары деңгейлі» API Nvidia бейне чиптерінің дизайны мен жұмысы туралы білімді болжайды; онда абстракцияның тым жоғары деңгейі жоқ.

Басқа деңгей бар, одан да жоғары - екі кітапхана:

КУБЛАС- сызықтық алгебра есептерін есептеуге және GPU ресурстарына тікелей қол жеткізуге арналған BLAS (Basic Linear Algebra Subprograms) CUDA нұсқасы;

МАНЖЕТ- Сигналды өңдеуде кеңінен қолданылатын Жылдам Фурье түрлендіруін есептеуге арналған жылдам Фурье түрлендіру кітапханасының CUDA нұсқасы. Келесі түрлендіру түрлеріне қолдау көрсетіледі: күрделі-күрделі (C2C), нақты-күрделі (R2C) және күрделі-нақты (C2R).

Осы кітапханаларды толығырақ қарастырайық. CUBLAS — CUDA тіліне аударылған стандартты сызықтық алгебра алгоритмдері, қазіргі уақытта CUBLAS негізгі функцияларының белгілі бір жиынтығына ғана қолдау көрсетіледі. Кітапхананы пайдалану өте оңай: бейнежадта матрицалық және векторлық объектілерді құру, оларды деректермен толтыру, қажетті CUBLAS функцияларын шақыру және бейне жадыдан нәтижелерді жүйелік жадыға қайта жүктеу қажет. CUBLAS құрамында GPU жадында объектілерді құру және жою, сондай-ақ осы жадқа деректерді оқу және жазу үшін арнайы функциялар бар. Қолдау көрсетілетін BLAS функциялары: нақты сандар үшін 1, 2 және 3 деңгейлері, кешен үшін 1 деңгей CGEMM. 1-деңгей-вектор-векторлық амалдар, 2-деңгей-вектор-матрицалық амалдар, 3-деңгей-матрица-матрицалық амалдар.

CUFFT - Fast Furier Transform-тың CUDA нұсқасы - сигналды талдауда, сүзуде және т.б. кеңінен қолданылады және өте маңызды. CUFFT GPU үшін теңшелетін FFT әзірлеуді қажет етпей Nvidia GPU құрылғыларында тиімді FFT есептеуі үшін қарапайым интерфейсті қамтамасыз етеді. CUDA FFT нұсқасы күрделі және нақты деректердің 1D, 2D және 3D түрлендірулерін қолдайды, параллельді бірнеше 1D түрлендірулер үшін пакетті орындау, 2D және 3D түрлендірулердің өлшемдері ішінде болуы мүмкін, 1D үшін 8 миллион элементке дейінгі өлшемге қолдау көрсетіледі.

CUDA бойынша бағдарламаларды құру негіздері

Төмендегі мәтінді түсіну үшін Nvidia бейне чиптерінің негізгі архитектуралық ерекшеліктерін түсінуіңіз керек. GPU текстуралық блоктардың бірнеше кластерлерінен тұрады (Texture Processing Cluster). Әрбір кластер текстураны алудың үлкейтілген блогынан және әрқайсысы сегізден тұратын екі немесе үш ағынды мультипроцессордан тұрады. есептеуіш құрылғыларжәне екі суперфункционалды блок. Барлық нұсқаулар SIMD принципіне сәйкес орындалады, бір нұсқау бұрмадағы барлық ағындарға қолданылғанда (тоқыма өнеркәсібінен алынған термин, CUDA-да бұл 32 ағындар тобы - мультипроцессорлармен өңделген деректердің ең аз мөлшері). Бұл орындау әдісі SIMT деп аталды (бір нұсқаулық бірнеше ағындар - бір нұсқаулық және көптеген ағындар).

Көп процессорлардың әрқайсысының белгілі бір ресурстары бар. Сонымен, бір мультипроцессорға сыйымдылығы 16 килобайт болатын арнайы ортақ жады бар. Бірақ бұл кэш емес, өйткені бағдарламашы оны ұялы процессорлардың SPU ішіндегі Жергілікті дүкенге ұқсас кез келген қажеттілік үшін пайдалана алады. Бұл ортақ жад бір блоктың ағындары арасында ақпарат алмасуға мүмкіндік береді. Бір блоктың барлық ағындары әрқашан бір мультипроцессормен орындалатыны маңызды. Әртүрлі блоктардағы ағындар деректермен алмаса алмайды және бұл шектеуді есте сақтау керек. Бір жад банкіне бірнеше ағындар қатынасу жағдайларын қоспағанда, ортақ жад жиі пайдалы. Мультипроцессорлар бейне жадына да қол жеткізе алады, бірақ жоғары кідіріспен және өткізу қабілеті нашар. Қолжетімділікті тездету және бейне жадыға қол жеткізу жиілігін азайту үшін мультипроцессорларда тұрақтылар мен текстуралық деректерге арналған 8 килобайт кэш бар.

Мультипроцессор 8192-16384 (тиісінше G8x/G9x және GT2xx үшін) онда орындалатын барлық блоктардың барлық ағындарына ортақ регистрлерді пайдаланады. G8x/G9x үшін бір мультипроцессорға блоктардың максималды саны сегіз, ал деформация саны 24 (бір мультипроцессорға 768 ағын). Барлығы Geforce 8 және 9 серияларының жоғарғы бейне карталары бір уақытта 12288 ағынға дейін өңдей алады. GT200 негізіндегі GeForce GTX 280 бір мультипроцессорға 1024 ағынды ұсынады, оның 30720 ағынға дейін өңдейтін үш мультипроцессордың 10 кластері бар. Бұл шектеулерді білу қол жетімді ресурстар үшін алгоритмдерді оңтайландыруға мүмкіндік береді.

Қолданыстағы қолданбаны CUDA-ға көшірудің бірінші қадамы оны профильдеу және жұмысты баяулататын кедергілер болып табылатын код аймақтарын анықтау болып табылады. Егер мұндай бөлімдердің арасында жылдам параллельді орындау үшін қолайлылары болса, бұл функциялар GPU-да орындау үшін C және CUDA кеңейтімдеріне беріледі. Бағдарлама CPU және GPU үшін кодты жасайтын Nvidia ұсынған компилятор арқылы құрастырылған. Бағдарлама орындалғанда, орталық процессор кодтың өз бөліктерін орындайды, ал GPU ең ауыр параллельді есептеулермен CUDA кодын орындайды. GPU үшін жасалған бұл бөлік ядро ​​(ядро) деп аталады. Ядро деректермен орындалатын операцияларды анықтайды.

Бейне чип ядроны қабылдайды және әрбір деректер элементі үшін көшірмелерді жасайды. Бұл көшірмелер ағындар деп аталады. Ағында санауыш, регистрлер және күй бар. Кескінді өңдеу сияқты деректердің үлкен көлемі үшін миллиондаған ағындар іске қосылады. Жіптер warps деп аталатын 32 топта орындалады.Worps арнайы ағынды мультипроцессорларда жұмыс істеу үшін тағайындалады. Әрбір мультипроцессор сегіз ядродан тұрады – тактілік циклге бір MAD нұсқаулығын орындайтын ағындық процессорлар. Бір 32 жіпті бұрмалауды орындау үшін төрт мультипроцессорлық цикл қажет (біз 1,5 ГГц және одан жоғары шейдер доменінің жиілігі туралы айтып отырмыз).

Мультипроцессор дәстүрлі көп ядролы процессор емес, ол бір уақытта 32 бұрмаға дейін қолдау көрсететін көп ағынды үшін өте қолайлы.Әр сағаттық циклде аппараттық құрал бұрмалардың қайсысын орындау керектігін таңдайды және бірінен екіншісіне ауысады. циклдарды жоғалту. Егер орталық процессорға ұқсастық жасасақ, бұл бір уақытта 32 бағдарламаны орындау және олардың арасында контекстік ауыстырып-қосқышты жоғалтпай әр сағаттық циклде ауысу сияқты. Шындығында, процессордың ядролары бір бағдарламаның бір уақытта орындалуын қолдайды және жүздеген циклді кешіктірумен басқаларға ауысады.

CUDA бағдарламалау моделі

Қайтадан, CUDA параллельді есептеу моделін пайдаланады, мұнда SIMD процессорларының әрқайсысы әртүрлі деректер элементтеріне параллельді түрде бірдей нұсқауларды орындайды. GPU – есептеу құрылғысы, орталық процессорға (хостқа) арналған сопроцессор (құрылғы), оның өзіндік жады бар және көптеген ағындарды параллель өңдейді. Ядро (ядро) ағындармен орындалатын GPU үшін функция (3D графикасынан аналогия – шейдер).

Жоғарыда біз бейне чиптің орталық процессордан айырмашылығы ондаған мың ағындарды бір уақытта өңдей алатынын айттық, бұл әдетте жақсы параллельденген графикаға арналған. Әрбір ағын скаляр, деректерді 4 компонентті векторларға жинақтауды қажет етпейді, бұл көптеген тапсырмалар үшін ыңғайлы. Логикалық ағындар мен ағын блоктарының саны физикалық орындау бірліктерінің санынан асып түседі, бұл компания шешімдерінің барлық ауқымы үшін жақсы масштабтауға мүмкіндік береді.

CUDA-дағы бағдарламалау моделі ағындарды топтастыруды болжайды. Жіптер жіп блоктарына біріктіріледі - ортақ жад пен синхрондау нүктелерін пайдаланып бір-бірімен әрекеттесетін ағындардың бір өлшемді немесе екі өлшемді торлары. Бағдарлама (ядро) жіп блоктарының торында орындалады, төмендегі суретті қараңыз. Бір тор бір уақытта орындалады. Әрбір блок пішіні бойынша бір, екі немесе үш өлшемді болуы мүмкін және ағымдағы аппараттық құралда 512 ағыннан тұруы мүмкін.

Жіп блоктары өлшемі 32 жіпті құрайтын деформациялар деп аталатын шағын топтарда жұмыс істейді. Бұл мультипроцессорларда өңдеуге болатын деректердің ең аз мөлшері. Және бұл әрқашан қолайлы емес болғандықтан, CUDA сізге 64-тен 512-ге дейінгі ағынды қамтитын блоктармен жұмыс істеуге мүмкіндік береді.

Блоктарды торларға топтастыру шектеулерден арылуға және ядроны бір қоңырауда ағындардың көбірек санына қолдануға мүмкіндік береді. Ол сонымен қатар масштабтауға көмектеседі. Егер графикалық процессордың ресурстары жеткіліксіз болса, ол блоктарды дәйекті түрде орындайды. Әйтпесе, блоктарды параллель орындауға болады, бұл мобильді және біріктірілгеннен бастап әртүрлі деңгейдегі бейне чиптеріндегі жұмысты оңтайлы бөлу үшін маңызды.

CUDA жады моделі

CUDA жады моделі байтты адрестеу мүмкіндігімен, жинақтау мен шашырауды қолдауымен ерекшеленеді. Әрбір ағындық процессор үшін 1024 данаға дейін регистрлердің жеткілікті үлкен саны бар. Оларға қол жеткізу өте жылдам, оларда 32-биттік бүтін сандарды немесе өзгермелі нүктелерді сақтауға болады.

Әрбір ағын келесі жад түрлеріне қол жеткізе алады:

жаһандық жады- бейне чиптегі барлық мультипроцессорлар үшін қолжетімді жадтың ең үлкен көлемі, өлшемі ағымдағы шешімдер үшін 256 мегабайттан 1,5 гигабайтқа дейін (және Tesla үшін 4 Гбайтқа дейін). Ол жоғары өткізу қабілетіне ие, Nvidia үздік шешімдері үшін 100 гигабайттан астам / с, бірақ бірнеше жүз циклді өте үлкен кідірістерге ие. Кэштеу мүмкін емес, жалпы жүктеу және сақтау нұсқауларын және тұрақты жад көрсеткіштерін қолдайды.

жергілікті жадытек бір ағындық процессорға қол жеткізе алатын жадтың шағын көлемі. Бұл салыстырмалы түрде баяу - жаһандық сияқты.

Ортақ жадмультипроцессордағы барлық ағындық процессорларға ортақ рұқсаты бар 16 килобайттық (қазіргі архитектураның бейне чиптерінде) жад блогы. Бұл жад өте жылдам, регистрлер сияқты. Ол ағынның өзара әрекеттесуін қамтамасыз етеді, әзірлеуші ​​тікелей басқарады және кідірісі төмен. Ортақ жадтың артықшылықтары: бағдарламашы басқаратын бірінші деңгейлі кэш түрінде пайдалану, орындау бірліктері (ALU) деректеріне қол жеткізудегі кідірістерді азайту, жадыға ғаламдық қол жеткізу санын азайту.

Тұрақты есте сақтау- 64 килобайт жад аймағы (ағымдағы графикалық процессорлар үшін бірдей), барлық мультипроцессорлар тек оқуға арналған. Ол бір мультипроцессорға 8 килобайтта кэштелген. Өте баяу - кэште қажетті деректер болмаған кезде бірнеше жүздеген циклдің кешігуі.

текстуралық жады- барлық мультипроцессорлардың оқуы үшін қолжетімді жад блогы. Деректерді іріктеу бейне чиптің текстуралық бірліктерін пайдалану арқылы жүзеге асырылады, сондықтан деректерді сызықтық интерполяциялау мүмкіндігі қосымша шығынсыз қамтамасыз етіледі. Бір мультипроцессорға кэштелген 8 килобайт. Жаһандық сияқты баяу - кэште деректер болмаған кездегі кешігудің жүздеген циклі.

Әрине, ғаламдық, жергілікті, құрылымдық және тұрақты жад физикалық түрде бірдей жад болып табылады, ол бейне картаның жергілікті бейне жады ретінде белгілі. Олардың айырмашылықтары әртүрлі кэштеу алгоритмдерінде және қол жеткізу үлгілерінде. Орталық процессор тек жаңартуға және сұрауға болады сыртқы жады: ғаламдық, тұрақты және текстуралық.

Жоғарыда жазылғандардан CUDA процессорға арналған бағдарламаларда қабылданғанмен бірдей емес, әзірлеуге ерекше көзқарасты білдіретіні анық. Жадтың әртүрлі түрлері туралы есте сақтау керек, жергілікті және ғаламдық жад кэштелмейді және оған қол жеткізудегі кідіріс тіркелген жадқа қарағанда әлдеқайда жоғары, өйткені ол физикалық түрде жеке микросұлбаларда орналасқан.

Типтік, бірақ міндетті емес мәселені шешу үлгісі:

• тапсырма қосалқы тапсырмаларға бөлінеді;
• кіріс деректері ортақ жадыға сәйкес блоктарға бөлінеді;
• әрбір блок жіп блогымен өңделеді;
• қосалқы блок ғаламдық жадтан ортақ жадқа жүктеледі;
• ортақ жадтағы деректер бойынша сәйкес есептеулер орындалады;
• нәтижелер ортақ жадтан қайтадан жаһандық жадқа көшіріледі.

Бағдарламалау ортасы

CUDA жұмыс уақыты кітапханаларын қамтиды:

• процессор мен графикалық процессорда қолдау көрсетілетін RTL шақыруларының кірістірілген векторлық түрлерін және ішкі жиындарын қамтамасыз ететін жалпы бөлік;
• CPU құрамдас бөлігі, бір немесе бірнеше графикалық процессорларды басқару үшін;
• GPU-ға тән функционалдылықты қамтамасыз ететін GPU құрамдас бөлігі.

CUDA қосымшасының негізгі процесі жалпы процессорда (хост) жұмыс істейді, ол видеокартадағы ядролық процестердің бірнеше көшірмелерін іске қосады. Орталық процессордың коды келесі әрекеттерді орындайды: графикалық процессорды инициализациялайды, бейне карта мен жүйедегі жадты бөледі, тұрақты мәндерді видеокарта жадына көшіреді, видеокартадағы ядролық процестердің бірнеше көшірмелерін іске қосады, нәтижені компьютерден көшіреді. бейне жады, жадты босатады және шығады.

Түсіну үшін мысал ретінде CUDA-да ұсынылған векторларды қосуға арналған CPU коды:

Бейне чипімен орындалатын функциялардың келесі шектеулері бар: рекурсия жоқ, функциялар ішінде статикалық айнымалылар және аргументтердің айнымалы саны жоқ. Жадты басқарудың екі түріне қолдау көрсетіледі: 32-биттік көрсеткіштер арқылы қол жеткізілетін сызықтық жад және тек текстураны алу функциялары арқылы қол жеткізілетін CUDA массивтері.

CUDA бағдарламалары графикалық API интерфейстерімен өзара әрекеттесе алады: бағдарламада жасалған деректерді көрсету, көрсету нәтижелерін оқу және оларды CUDA құралдарының көмегімен өңдеу (мысалы, өңдеуден кейінгі сүзгілерді енгізу кезінде). Ол үшін графикалық API ресурстарын CUDA жаһандық жады кеңістігінде салыстыруға болады (ресурс мекенжайын алу). Графикалық API ресурстарының келесі түрлеріне қолдау көрсетіледі: OpenGL ішіндегі буфер нысандары (PBO / VBO), Vertex буферлері және текстуралары (2D, 3D және текше карталары) Direct3D9.

CUDA қолданбасын құрастыру қадамдары:

CUDA C бастапқы код файлдары басқа құралдарды орап, оларды шақыратын NVCC бағдарламасы арқылы құрастырылады: cudacc, g++, cl, т.б. NVCC жасайды: таза C тілінде жазылған қолданбаның қалған бөлігімен бірге құрастырылған CPU коды және Бейне чипіне арналған PTX нысан коды. CUDA коды бар орындалатын файлдар міндетті түрде CUDA орындалу кітапханасының (cudart) және CUDA негізгі кітапханасының (cuda) болуын талап етеді.

CUDA бойынша бағдарламаларды оңтайландыру

Әрине, шолу мақаласы аясында CUDA бағдарламалауында күрделі оңтайландыру мәселелерін қарастыру мүмкін емес. Сондықтан біз негізгі нәрселер туралы қысқаша сөйлесеміз. CUDA мүмкіндіктерін тиімді пайдалану үшін орталық процессорға арналған бағдарламаларды жазудың әдеттегі әдістерін ұмытып, мыңдаған ағындар үшін жақсы параллельденген алгоритмдерді пайдалану керек. Сондай-ақ деректерді сақтаудың оңтайлы орнын табу (регистрлер, ортақ жад және т.б.), орталық процессор мен графикалық процессор арасындағы мәліметтерді тасымалдауды азайту және буферлеуді пайдалану маңызды.

Жалпы алғанда, CUDA бағдарламасын оңтайландыру кезінде блоктардың өлшемі мен саны арасындағы оңтайлы тепе-теңдікке қол жеткізуге тырысу керек. Блоктағы көбірек ағындар жадтың кешігуінің әсерін азайтады, сонымен қатар регистрлердің қолжетімді санын азайтады. Сонымен қатар, 512 ағынның блогы тиімсіз, Nvidia өзі оңтайлы кідіріс пен регистрлердің санына қол жеткізу үшін ымыраға келетін мән ретінде 128 немесе 256 ағынды блоктарды пайдалануды ұсынады.

CUDA бағдарламаларын оңтайландырудың негізгі сәттерінің арасында: ортақ жадты мүмкіндігінше белсенді пайдалану, өйткені ол видеокартаның ғаламдық бейне жадынан әлдеқайда жылдамырақ; жаһандық жадтан оқу және жазу мүмкіндігінше біріктірілуі керек. Ол үшін бір операцияда бірден 32/64/128 биттік деректерді оқу және жазу үшін арнайы деректер түрлерін пайдалану қажет. Оқу әрекеттерін біріктіру қиын болса, текстураны алуды пайдаланып көруге болады.

қорытындылар

Nvidia ұсынған CUDA бейне чиптеріндегі есептеулерге арналған аппараттық-бағдарламалық қамтамасыз ету архитектурасы жоғары параллелизммен кең ауқымды тапсырмаларды шешуге өте қолайлы. CUDA көптеген Nvidia бейне чиптерінде жұмыс істейді және GPU бағдарламалау үлгісін айтарлықтай жеңілдету және ортақ жад, ағындарды синхрондау мүмкіндігі, қос дәлдіктегі есептеулер және бүтін операциялар сияқты көптеген мүмкіндіктерді қосу арқылы жақсартады.

CUDA – бұл әрбір бағдарламалық жасақтама жасаушыға қолжетімді технология, оны Си тілін білетін кез келген бағдарламашы пайдалана алады. Сізге тек параллельді есептеулерге тән басқа бағдарламалау парадигмасына үйрену керек. Бірақ егер алгоритм, негізінен, жақсы параллельді болса, онда CUDA бағдарламалауға жұмсалған зерттеу мен уақыт бірнеше өлшемде қайтарылады.

Дүние жүзінде видеокарталардың кең таралуына байланысты графикалық процессордағы параллельді есептеулердің дамуы жоғары өнімділіктегі есептеу өнеркәсібіне үлкен әсер етуі мүмкін. Бұл мүмкіндіктер қазірдің өзінде оларда ғана емес, ғылыми топтарда үлкен қызығушылық тудырды. Ақыр соңында, параллельизацияға жақсы көмектесетін алгоритмдерді жеделдету мүмкіндігі (қолжетімді жабдықта, бұл маңызды емес) бірден ондаған есе кең таралған емес.

Жалпы мақсаттағы процессорлар өте баяу дамиды, оларда мұндай өнімділік жоғарылауы болмайды. Шындығында, бұл тым қатты естілгенімен, жылдам компьютерлерді қажет ететіндердің барлығы енді үстеліне қымбат емес жеке суперкомпьютерді, кейде тіпті ақша салмай-ақ ала алады. қосымша қаражат, өйткені Nvidia бейне карталары кеңінен қолданылады. GPU өндірушілеріне ұнайтын GFLOPS/$ және GFLOPS/W бойынша тиімділік артуын айтпағанның өзінде.

Көптеген есептеулердің болашағы параллельді алгоритмдерде екені анық, барлық дерлік жаңа шешімдер мен бастамалар осы бағытқа бағытталған. Дегенмен, әзірге жаңа парадигмалардың дамуы жүріп жатыр бастапқы кезең, ағындарды қолмен жасауға және жадқа кіруді жоспарлауға тура келеді, бұл әдетте бағдарламалауда жасайтын әрекеттерден қиынырақ етеді. Бірақ CUDA технологиясы дұрыс бағытта қадам жасады және ол сәтті шешім сияқты көрінеді, әсіресе Nvidia әзірлеушілерді мүмкіндігінше оның артықшылықтары мен болашағына сендіре алса.

Бірақ, әрине, графикалық процессорлар процессорларды алмастырмайды. Қазіргі пішінінде олар бұған арналмаған. Енді бейне чиптер бірте-бірте орталық процессорға қарай жылжып, барған сайын әмбебап бола бастады (бір және екі дәлдіктегі өзгермелі нүктелермен есептеулер, бүтін сандармен есептеулер), сондықтан процессорлар барған сайын «параллель» болып, ядролардың көп санын, көп ағынды алуда. технологиялар, SIMD блоктарының және гетерогенді процессор жобаларының пайда болуын айтпағанда. Сірә, GPU мен CPU болашақта жай ғана біріктіріледі. Көптеген компаниялар, соның ішінде Intel және AMD сияқты жобалармен айналысатыны белгілі. Және GPU процессорды тұтынатыны маңызды емес немесе керісінше.

Мақалада біз негізінен CUDA артықшылықтары туралы айттық. Бірақ шыбын да бар. CUDA-ның аздаған кемшіліктерінің бірі - оның нашар тасымалдануы. Бұл архитектура тек осы компанияның бейне чиптерінде ғана жұмыс істейді, олардың барлығында емес, Geforce 8 және 9 серияларынан және сәйкес Quadro және Tesla-дан басталады. Иә, әлемде мұндай шешімдер өте көп, Nvidia 90 миллион CUDA үйлесімді бейне чиптерін береді. Бұл өте жақсы, бірақ бәсекелестер CUDA-дан ерекшеленетін өз шешімдерін ұсынады. Сонымен, AMD-де Stream Computing бар, Intel-де болашақта Ct болады.

Технологиялардың қайсысы жеңеді, кең таралады және қалғандарынан ұзағырақ өмір сүреді - тек уақыт көрсетеді. Бірақ CUDA-ның жақсы мүмкіндігі бар, өйткені Stream Computing-пен салыстырғанда, мысалы, ол әдеттегі Си тілінде анағұрлым дамыған және қолдануға оңай бағдарламалау ортасын қамтамасыз етеді. Мүмкін үшінші тарап кейбір шығару арқылы анықтауға көмектеседі ортақ шешім. Мысалы, DirectX-тің 11-нұсқасының келесі жаңартуында Microsoft корпорациясы барлығына немесе барлығына дерлік қолайлы орташа шешімге айналуы мүмкін есептеу шейдерлеріне уәде берді.

Алдын ала мәліметтерге қарағанда, бұл жаңа шейдер түрі CUDA үлгісінен көп нәрсені алады. Ал қазір осы ортада бағдарламалау арқылы сіз бірден артықшылықтар мен болашаққа қажетті дағдыларды ала аласыз. Жоғары өнімді есептеу тұрғысынан, DirectX-те портативтіліктің айқын кемшілігі бар, өйткені API Windows платформасымен шектелген. Дегенмен, басқа стандарт әзірленуде - Nvidia, AMD, Intel, IBM және басқа да көптеген компаниялар қолдайтын ашық көп платформалы OpenCL бастамасы.

Келесі CUDA мақаласында сіз Nvidia CUDA көмегімен бүкіл әлем бойынша әзірлеушілер орындайтын ғылыми және басқа графикалық емес есептеулердің нақты практикалық қолданбаларын зерттейтініңізді ұмытпаңыз.

Тарихқа оралайық - Intel және AMD ең қуатты процессор үшін бірлескен жарыста болған 2003 жылға оралайық. Бірнеше жыл ішінде бұл жарыстың нәтижесінде, әсіресе Intel Pentium 4 шығарылғаннан кейін сағат жылдамдығы айтарлықтай өсті.

Бірақ жарыс шегіне тез жақындады. Сағат жиілігіндегі орасан өсу толқынынан кейін (2001 және 2003 жылдар аралығында Pentium 4 тактілік жиілігі 1,5-тен 3 ГГц-ке дейін екі есе өсті), пайдаланушылар өндірушілер шығара алатын гигагерцтің оннан бір бөлігіне қанағаттануға мәжбүр болды (2003 жылдан 2005 жылға дейін). , тактілік жиіліктер тек 3-тен 3,8 ГГц-ке дейін өсті).

Прескотт сияқты жоғары сағаттық жылдамдықтар үшін оңтайландырылған архитектуралар да қиындықтарға тап бола бастады, бұл жолы тек өндірістік емес. Чип жасаушылар физика заңдарына тап болды. Кейбір сарапшылар тіпті Мур заңы өз жұмысын тоқтатады деп болжаған. Бірақ олай болмады. Заңның бастапқы мағынасы жиі бұрмаланады, бірақ ол кремний ядросының бетіндегі транзисторлар санына қатысты. Ұзақ уақыт бойы орталық процессордағы транзисторлар санының артуы өнімділіктің сәйкес өсуімен бірге жүрді - бұл мағынаның бұрмалануына әкелді. Бірақ кейін жағдай күрделене түсті. Орталық процессордың архитектурасын құрастырушылар пайданы азайту заңына жақындады: өнімділіктің қажетті жоғарылауы үшін қосу қажет транзисторлар саны барған сайын көбейіп, тұйыққа әкелді.

Процессор өндірушілері өз мәселелерінің шешімін табуға тырысып, шаштарын жұлып жатқанымен, графикалық процессор өндірушілері Мур заңының артықшылықтарынан айтарлықтай пайда алуды жалғастырды.

Неліктен олар CPU архитектурасын құрастырушылар сияқты тұйыққа тірелмеді? Себебі өте қарапайым: орталық процессорлар әртүрлі деректерді (бүтін сандар да, өзгермелі нүкте сандары да) өңдейтін, жедел жадқа қол жеткізуді жүзеге асыратын және т.б. нұсқаулар ағыны бойынша ең жақсы өнімділікті алуға арналған. Осы уақытқа дейін әзірлеушілер нұсқаулардың үлкен параллелизмін қамтамасыз етуге тырысты, яғни параллельді түрде мүмкіндігінше көп нұсқауларды орындау. Мысалы, белгілі бір жағдайларда сағатына екі команданы орындау мүмкін болған кезде, Pentium-да суперскалярлық орындау пайда болды. Pentium Pro инструкциялардың ретсіз орындалуын алды, бұл есептеу блоктарының өнімділігін оңтайландыруға мүмкіндік берді. Мәселе мынада, нұсқаулардың дәйекті ағынының параллель орындалуы айқын шектеулерге ие, сондықтан есептеу бірліктерінің санын соқыр ұлғайту пайда әкелмейді, өйткені олар әлі де жұмыссыз болады.

Керісінше, GPU жұмысы салыстырмалы түрде қарапайым. Ол бір жағынан көпбұрыштар тобын алып, екінші жағында пикселдер тобын құрудан тұрады. Көпбұрыштар мен пикселдер бір-бірінен тәуелсіз, сондықтан оларды параллель өңдеуге болады. Осылайша, GPU-да кристалдың үлкен бөлігін процессордан айырмашылығы, іс жүзінде қолданылатын есептеу блоктары үшін бөлуге болады.

GPU процессордан тек осымен ғана емес ерекшеленеді. GPU-дағы жадқа қол жеткізу өте біріктірілген - егер текстел оқылса, онда бірнеше циклден кейін көрші текстель оқылады; пиксел жазылған кезде, көршілес пиксель бірнеше циклден кейін жазылады. Жадты интеллектуалды ұйымдастыру арқылы сіз теориялық өткізу қабілеттілігіне жақын өнімділікке қол жеткізе аласыз. Бұл процессордан айырмашылығы GPU үлкен кэшті қажет етпейтінін білдіреді, өйткені оның рөлі текстуралау операцияларын жылдамдату болып табылады. Ол үшін екі сызықты және үш сызықты сүзгілерде қолданылатын бірнеше тексті бар бірнеше килобайт қажет.

GeForce FX өмір сүрсін!

Екі дүние ұзақ уақыт бойы бір-бірінен бөлек тұрды. Біз орталық процессорды (немесе тіпті бірнеше процессорларды) кеңсе тапсырмалары мен интернет қолданбалары үшін пайдаландық, ал GPU тек көрсетуді жылдамдату үшін жақсы болды. Бірақ бір мүмкіндік бәрін өзгертті: атап айтқанда, бағдарламаланатын GPU пайда болуы. Бастапқыда процессорларда қорқатын ештеңе болған жоқ. Алғашқы бағдарламаланатын графикалық процессорлар (NV20 және R200) ​​қауіп төндірмеді. Бағдарламадағы нұсқаулар саны шамамен 10-мен шектелді, олар 9 немесе 12 биттік бекітілген нүкте сандары сияқты өте экзотикалық деректер түрлерінде жұмыс істеді.

Бірақ Мур заңы тағы да өзінің жақсы жағын көрсетті. Транзисторлар санының көбеюі есептеу бірліктерінің санын көбейтіп қана қоймай, олардың икемділігін де жақсартты. NV30-ның пайда болуын бірнеше себептерге байланысты алға жасалған маңызды қадам деп санауға болады. Әрине, ойыншыларға NV30 карталары ұнамады, бірақ жаңа графикалық процессорлар графикалық процессорларды жай графикалық үдеткіштер ретінде қабылдауды өзгертуге арналған екі мүмкіндікке сүйене бастады.

• Бір дәлдіктегі өзгермелі нүктелік есептеулерді қолдау (тіпті ол IEEE754 стандартына сәйкес келмесе);
• мыңнан астам нұсқауларды қолдау.

Осылайша, бізде әрқашан қосымша есептеу қуатын іздейтін пионер зерттеушілерді тарта алатын барлық жағдайлар бар.

Математикалық есептеулер үшін графикалық үдеткіштерді пайдалану идеясы жаңа емес. Алғашқы әрекеттер өткен ғасырдың 90-шы жылдары жасалды. Әрине, олар өте қарапайым болды - көбінесе маршрутты іздеу немесе шығару сияқты тапсырмаларды жылдамдату үшін растризация және Z-буферлері сияқты кейбір аппараттық құралдарға негізделген мүмкіндіктерді пайдаланумен шектелді. Вороной диаграммалары .

2003 жылы дамыған шейдерлердің пайда болуымен жаңа жолаққа қол жеткізілді - бұл жолы матрицалық есептеулерді орындау. Бұл SIGGRAPH («GPU-дағы есептеулер») толық бөлімі АТ-тың жаңа саласына арналған жыл болды. Бұл ерте бастама GPGPU (GPU бойынша жалпы мақсаттағы есептеу) деп аталды. Ал пайда болуы.

BrookGPU рөлін түсіну үшін оның пайда болуына дейін бәрі қалай болғанын түсіну керек. 2003 жылы GPU ресурстарын алудың жалғыз жолы екі графикалық API бірін пайдалану болды - Direct3D немесе OpenGL. Демек, GPU қуатын есептеу үшін алғысы келетін әзірлеушілер аталған екі API-ге сенуге мәжбүр болды. Мәселе мынада, олар әрқашан графикалық карталарды бағдарламалау бойынша сарапшылар бола бермейді және бұл технологияға қол жеткізуді қиындатты. Егер 3D бағдарламашылар шейдерлермен, текстуралармен және фрагменттермен жұмыс істесе, онда параллельді бағдарламалау саласындағы мамандар ағындарға, өзектерге, шашырауларға және т.б. Сондықтан алдымен екі дүниенің арасында ұқсастықтар жасау керек болды.

• ағынбір типті элементтер ағыны болып табылады, GPU-да оны текстурамен көрсетуге болады. Негізінде, классикалық бағдарламалауда массив сияқты аналогы бар.
• Ядро- ағынның әрбір элементіне тәуелсіз қолданылатын функция; пиксельдік шейдердің баламасы болып табылады. Классикалық бағдарламалауда циклге ұқсастық беруге болады – ол элементтердің көп санына қолданылады.
• Ядроны ағынға қолдану нәтижелерін оқу үшін текстураны жасау керек. Орталық процессорда баламасы жоқ, өйткені бар толық қол жеткізуесте сақтауға.
• Жазылатын жадтағы орын (шашырау/шашырау операцияларында) шыңдық шейдер арқылы басқарылады, өйткені пиксельдік шейдер өңделген пикселдің координаттарын өзгерте алмайды.

Көріп отырғаныңыздай, жоғарыда аталған ұқсастықтарды ескерсек те, тапсырма қарапайым болып көрінбейді. Ал Брук көмекке келді. Бұл атау С тілінің кеңейтімдерін білдіреді («С ағындарымен», «С ағындарымен»), Стэнфордтағы әзірлеушілер оларды атаған. Негізінде Бруктың міндеті программалаушыдан 3D API барлық құрамдастарын жасыру болды, бұл графикалық процессорды параллельді есептеулер үшін сопроцессор ретінде көрсетуге мүмкіндік берді. Бұл үшін Брук компиляторы C++ коды және кеңейтімдері бар .br файлын өңдеді, содан кейін әртүрлі шығыстарды (DirectX, OpenGL ARB, OpenGL NV3x, x86) қолдауы бар кітапханамен байланыстырылған C++ кодын жасады.

Бруктың бірнеше несиелері бар, олардың біріншісі GPGPU-ны көлеңкеден шығару болып табылады, осылайша технология жалпы жұртшылыққа көрінуі мүмкін. Жобаны жариялағаннан кейін бірқатар IT-сайттар тым оптимистік болды, Brook-тің шығарылымы жақын арада қуаттырақ графикалық процессорлармен ауыстырылатын процессорлардың бар екендігіне күмән келтіреді. Бірақ, көріп отырғанымыздай, бес жыл өтсе де бұл болмады. Шынымды айтсам, бұл ешқашан болмайды деп ойлаймыз. Екінші жағынан, параллелизмге (көп ядролар, SMT көп ағынды технологиясы, SIMD блоктарының кеңеюі), сондай-ақ, керісінше, әмбебап болып келе жатқан графикалық процессорлардың (қолдау) көбірек бағдарланған процессорлардың сәтті эволюциясына қарап өзгермелі нүктелік есептеулер үшін). бір дәлдік, бүтін сандар, қос дәлдіктегі есептеулерді қолдау), GPU және CPU жақында жай ғана біріктірілетін сияқты. Сонда не болады? Математикалық сопроцессорлардағыдай графикалық процессорларды процессорлар жұтады ма? Әбден мүмкін. Intel және AMD бүгінде ұқсас жобалармен жұмыс істеуде. Бірақ әлі де көп нәрсе өзгеруі мүмкін.

Бірақ тақырыбымызға оралу. Бруктың артықшылығы GPGPU концепциясын танымал ету болды, ол GPU ресурстарына қол жеткізуді айтарлықтай жеңілдетті, бұл көбірек пайдаланушыларға жаңа бағдарламалау үлгісін меңгеруге мүмкіндік берді. Екінші жағынан, Бруктың барлық қасиеттеріне қарамастан, GPU ресурстарын есептеу үшін пайдалану үшін әлі көп жол бар еді.

Мәселелердің бірі абстракцияның әртүрлі деңгейлеріне, сонымен қатар, атап айтқанда, 3D API жасаған шамадан тыс қосымша жүктемеге байланысты, бұл айтарлықтай байқалуы мүмкін. Бірақ Брук әзірлеушілері ештеңе істей алмайтын үйлесімділік мәселесін неғұрлым күрделі деп санауға болады. GPU өндірушілері арасында қатал бәсекелестік бар, сондықтан олар көбінесе драйверлерін оңтайландырады. Егер мұндай оңтайландырулар көбінесе ойыншылар үшін жақсы болса, олар Brook үйлесімділігін бір сәтте жояды. Сондықтан, бұл API-ны бір жерде жұмыс істейтін өнеркәсіптік кодта пайдалануды елестету қиын. Ұзақ уақыт бойы Брук әуесқой зерттеушілер мен бағдарламашылар қатарында болды.

Дегенмен, Бруктың жетістігі ATI мен Nvidia-ның назарын аудару үшін жеткілікті болды және олар мұндай бастамаға қызығушылық танытты, өйткені ол нарықты кеңейте алады, компаниялар үшін жаңа маңызды секторды аша алады.

Бастапқыда Брук жобасына қатысқан зерттеушілер жаңа нарықты дамытудың жаһандық стратегиясын ұсыну үшін Санта-Кларадағы әзірлеу топтарына тез қосылды. Идея GPGPU тапсырмалары үшін қолайлы аппараттық және бағдарламалық құралдың комбинациясын жасау болды. NVidia әзірлеушілері графикалық процессорларының барлық құпияларын білетіндіктен, графикалық API-ге сенбей, драйвер арқылы графикалық процессормен байланысу мүмкін болды. Әрине, бұл өз проблемаларымен бірге келеді. Сонымен, CUDA (Compute Unified Device Architecture) әзірлеу тобы GPU-мен жұмыс істеу үшін бағдарламалық жасақтама қабаттарының жинағын жасады.

Диаграммада көріп отырғаныңыздай, CUDA екі API ұсынады.

• Жоғары деңгейлі API: CUDA Runtime API;
• төмен деңгейлі API: CUDA драйвер API.

Жоғары деңгейлі API төменгі деңгейдің үстіне енгізілгендіктен, әрбір орындау уақыты функциясы шақыруы Driver API өңдейтін қарапайым нұсқауларға бөлінеді. Екі API бірін-бірі жоққа шығаратынын ескеріңіз: бағдарламашы бір немесе басқа API пайдалана алады, бірақ екі API функциясының шақыруларын араластыру мүмкін емес. Жалпы, «жоғары деңгейлі API» термині салыстырмалы. Тіпті Runtime API интерфейсі соншалық, көпшілігі оны төмен деңгейлі деп санайды; дегенмен, ол әлі де инициализация немесе мәтінмәнді басқару үшін өте ыңғайлы функцияларды қамтамасыз етеді. Бірақ абстракцияның ерекше жоғары деңгейін күтпеңіз - сіз әлі де nVidia графикалық процессорлары және олардың қалай жұмыс істейтіні туралы жақсы білімге ие болуыңыз керек.

Driver API-мен жұмыс істеу одан да қиын; GPU өңдеуді іске қосу үшін сізге көбірек күш қажет. Екінші жағынан, төмен деңгейлі API икемді, қажет болған жағдайда бағдарламашыға көбірек басқару мүмкіндігін береді. Екі API OpenGL немесе Direct3D ресурстарымен жұмыс істей алады (бүгінгі күні тек тоғызыншы нұсқасы). Бұл мүмкіндіктің артықшылығы анық – CUDA графикалық API-ге берілуі мүмкін ресурстарды (геометрия, процедуралық текстуралар және т.б.) жасау үшін пайдаланылуы мүмкін немесе керісінше, 3D API көрсету нәтижелерін CUDA-ға жіберу үшін жасалуы мүмкін. бағдарлама, ол өз кезегінде кейінгі өңдеуді орындайды. Мұндай өзара әрекеттесулердің көптеген мысалдары бар және артықшылығы - ресурстар GPU жадында сақталуын жалғастыруда, оларды PCI Express шинасы арқылы тасымалдаудың қажеті жоқ, бұл әлі күнге дейін тығырық болып табылады.

Дегенмен, бейне жадыдағы ресурстарды ортақ пайдалану әрқашан идеалды емес және кейбір «бас ауруларына» әкелуі мүмкін екенін атап өткен жөн. Мысалы, ажыратымдылықты немесе түс тереңдігін өзгерткен кезде графикалық деректер басымдыққа ие болады. Сондықтан, егер кадр буферіндегі ресурстарды ұлғайту қажет болса, драйвер мұны CUDA қосымшаларының ресурстары есебінен оңай жасайды, бұл жай ғана қателікпен бұзылады. Әрине, өте талғампаз емес, бірақ бұл жағдай өте жиі болмауы керек. Біз кемшіліктер туралы айта бастағандықтан: CUDA қолданбалары үшін бірнеше GPU пайдаланғыңыз келсе, алдымен SLI режимін өшіруіңіз керек, әйтпесе CUDA қолданбалары тек бір GPU «көре» алады.

Соңында, бағдарламалық қамтамасыз етудің үшінші деңгейі кітапханаларға беріледі - дәлірек айтқанда, екі.

• GPU-да сызықтық алгебраны есептеуге қажетті блоктарды қамтитын CUBLAS;
• Фурье түрлендірулерін есептеуді қолдайтын CUFFT сигналды өңдеу саласында кеңінен қолданылатын алгоритм болып табылады.

CUDA-ға кіріспес бұрын, nVidia құжаттамасында шашыраңқы бірнеше терминдерді анықтайық. Компания өте нақты терминологияны таңдады, оған үйрену қиын. Ең алдымен мынаны атап өтеміз жіп CUDA-да бұл процессордың ағынымен бірдей мағынаға жақын емес және GPU мақалаларындағы ағынның баламасы да емес. Бұл жағдайда GPU ағыны болып табылады негізгі жиынтықөңделетін деректер. CPU ағындарынан айырмашылығы, CUDA ағындары өте «жеңіл», яғни екі ағын арасындағы контекстік ауыстырып-қосқыш ешбір жағдайда ресурсты қажет ететін операция емес.

CUDA құжаттамасында жиі кездесетін екінші термин бұрылыс. Мұнда ешқандай шатасу жоқ, өйткені орыс тілінде аналогы жоқ (егер сіз Start Trek немесе Warhammer ойындарының жанкүйері болмасаңыз). Шындығында, бұл термин тоқыма өнеркәсібінен алынған, онда тоқыма иірілген жіп тоқыма станогында созылған иірілген жіп арқылы тартылады. CUDA-дағы Warp – бұл 32 ағыннан тұратын топ және CUDA мультипроцессорларында SIMD әдісімен өңделетін деректердің ең аз мөлшері.

Бірақ мұндай «түйірлік» бағдарламашы үшін әрқашан қолайлы бола бермейді. Сондықтан, CUDA-да, ілмектермен тікелей жұмыс істеудің орнына, сіз онымен жұмыс істей аласыз блок, 64-тен 512-ге дейінгі ағындарды қамтиды.

Ақыры, бұл блоктар бір-біріне жиналды торлар. Бұл топтастырудың артықшылығы мынада: GPU бір уақытта өңдейтін блоктар саны аппараттық ресурстармен тығыз байланысты, біз төменде көреміз. Тректерді топтастыру тораптарға осы шектеуден толығымен тезірек деректеме жасауға және ядро ​​/ ядроға белгілі бір қоңырауға, бір қоңырауда, тіркелген ресурстар туралы ойланбастан қолданылады. Осының бәріне Cuda кітапханалары жауап береді. Сонымен қатар, бұл модель жақсы масштабталады. Егер GPU ресурстары аз болса, онда ол блоктарды дәйекті түрде орындайды. Есептеу процессорларының саны көп болса, блоктарды параллель орындауға болады. Яғни, бірдей код бастапқы деңгейдегі графикалық процессорларда да, жоғары деңгейлі, тіпті болашақ үлгілерде де жұмыс істей алады.

CUDA API интерфейсінде CPU дегенді білдіретін тағы бірнеше термин бар ( Хост / хост) және GPU ( құрылғы/құрылғы). Егер бұл кішкене кіріспе сізді қорқытпаса, CUDA-мен жақсырақ танысудың уақыты келді.

Егер сіз Томның аппараттық нұсқаулығын үнемі оқып отырсаңыз, онда nVidia-дан соңғы GPU архитектурасы сізге таныс. Егер жоқ болса, мақаланы оқуды ұсынамыз « nVidia GeForce GTX 260 және 280: графикалық карталардың жаңа буыны". CUDA-ға қатысты, nVidia архитектураны сәл басқаша ұсынады, бұрын жасырылған кейбір мәліметтерді көрсетеді.

Жоғарыдағы суреттен көріп отырғаныңыздай, nVidia шейдерінің өзегі текстуралық процессорлардың бірнеше кластерлерінен тұрады. (Текстуралық процессорлар кластері, TPC). Мысалы, 8800 GTX сегіз кластерді, 8800 GTS алты кластерді пайдаланды және т.б. Әрбір кластер негізінен текстура бірлігінен және екеуінен тұрады ағынды мультипроцессорлар. Соңғыларына нұсқауларды оқитын және декодтайтын, сондай-ақ оларды орындауға жіберетін құбырдың басы (алдыңғы ұшы) және сегіз есептеуіш құрылғыдан және екі суперфункционалды құрылғыдан тұратын құбырдың соңы (артқы ұшы) кіреді. SFU (Super Function Unit), мұнда нұсқаулар SIMD принципі бойынша орындалады, яғни бір нұсқау бұрмадағы барлық ағындарға қолданылады. nVidia мұны істеу әдісін шақырады Simt(бір нұсқау бірнеше ағындар, бір нұсқаулық, көп ағындар). Құбырдың соңы оның басталуынан екі есе жиі жұмыс істейтінін атап өткен жөн. Іс жүзінде бұл бөлік шын мәнінде қарағанда екі есе «кеңірек» көрінеді (яғни сегіз арнаның орнына 16 арналы SIMD блогы сияқты). Ағынды мультипроцессорлар келесідей жұмыс істейді: әрбір цикл, конвейердің басы орындауға дайын сызбаны таңдайды және нұсқауды орындауға кіріседі. Бұрмадағы барлық 32 жіпке қолданылатын нұсқау үшін құбырдың соңы төрт тактілік циклді алады, бірақ ол басталу жиілігінің екі есе жиілігінде жұмыс істейтіндіктен, оған тек екі сағат циклі қажет болады (басталуы тұрғысынан алғанда). құбыр). Сондықтан, құбырдың басталуы цикл үшін жұмыс істемейді, ал жабдық максималды түрде жүктелуі, ал идеалды жағдайда, әр цикл бойынша нұсқауларды ауыстырады - классикалық нұсқаулық және SFU үшін нұсқаулар - басқасы .

Әрбір мультипроцессорда түсінуге тұрарлық ресурстардың белгілі бір жиынтығы бар. Жадтың кішкентай ауданы бар «Ортақ жад», әр мультипроцессорға 16 КБ. Бұл ешбір жағдайда кэш жады емес: бағдарламашы оны өз қалауы бойынша пайдалана алады. Яғни, бізде ұялы процессорлардағы SPU жергілікті дүкеніне жақын нәрсе бар. Бұл егжей-тегжей өте қызықты, өйткені ол CUDA бағдарламалық және аппараттық технологиялардың үйлесімі екенін көрсетеді. Жадтың бұл аймағы пиксельдік шейдерлер үшін пайдаланылмайды, өйткені Nvidia «пикселдер бір-бірімен сөйлескенде бізге ұнамайды» деп ақылды түрде көрсетеді.

Жадтың бұл аймағы ағындар арасында ақпарат алмасу мүмкіндігін ашады. бір блокта. Бұл шектеуді атап өту маңызды: блоктағы барлық ағындардың бір мультипроцессормен орындалуына кепілдік беріледі. Керісінше, блоктардың әртүрлі мультипзиторларға байланысы мүлдем көрсетілмеген, ал әр түрлі блоктардан екі жіптер әр түрлі блоктардан екі жіптер орындалу кезінде бір-бірімен ақпарат алмасады. Яғни, ортақ жадты пайдалану оңай емес. Дегенмен, бірнеше ағындар бір жад банкіне кіруге тырысып, қақтығыс тудыратын жағдайларды қоспағанда, ортақ жад әлі де ақталған. Басқа жағдайларда ортақ жадқа қол жеткізу тіркеуге қол жеткізу сияқты жылдам.

Ортақ жад мультипроцессорлар қол жеткізе алатын жалғыз жад емес. Олар бейне жадын пайдалана алады, бірақ өткізу қабілеттілігі төмен және кідірісі жоғары. Сондықтан, осы жадқа қол жеткізу жиілігін азайту үшін NVIDIA константалар мен текстуралар мен текстураларды сақтайтын мультипроцессорлар (мультипроцессорға шаққандағы 8 КБ).

Мультипроцессорда мультипроцессорда белсенді барлық блоктардың барлық ағындарына ортақ 8192 регистр бар. Бір мультипроцессорға арналған белсенді блоктардың саны сегізден аспауы керек, ал белсенді бұрмалардың саны 24 (768 ағын) шектелген. Сондықтан 8800 GTX бір уақытта 12 288 ағынға дейін өңдей алады. Бұл шектеулердің барлығын атап өткен жөн, өйткені олар алгоритмді қолда бар ресурстар негізінде оңтайландыруға мүмкіндік береді.

Сондықтан CUDA бағдарламасын оңтайландыру блоктар саны мен олардың өлшемдері арасындағы оңтайлы тепе-теңдікті алудан тұрады. Әр блокқа көбірек ағындар жадтың кідірісін азайтуға пайдалы болар еді, бірақ бір ағынға қол жетімді регистрлердің саны да азаяды. Оның үстіне 512 ағынды блок тиімсіз болар еді, өйткені мультипроцессорда тек бір блок белсенді болуы мүмкін, нәтижесінде 256 ағын жоғалады. Сондықтан, Нвидия 128 немесе 256 жіптердің блоктарын қолдануды ұсынады, бұл төменгі кідірістер мен көптеген ядро ​​/ ядро ​​үшін тізілімдер саны.

Бағдарламалық тұрғыдан алғанда, CUDA BrookGPU-ға ұқсайтын Си тілінің кеңейтімдерінің жиынтығынан, сондай-ақ бірнеше нақты API қоңырауларынан тұрады. Кеңейтімдердің ішінде функциялар мен айнымалыларға қатысты тип спецификаторлары бар. Кілт сөзді есте сақтау маңызды __Глобал__Функцияның алдында болған кезде, яғни ядроға / ядроға қатысты екенін көрсетеді, яғни бұл функцияны CPU шақырады және ол GPU-да орындалады. Қосымша __Device__функцияның GPU-да орындалатынын көрсетеді (оны CUDA айтпақшы «құрылғы/құрылғы» деп атайды), бірақ оны тек GPU-дан шақыруға болады (басқаша айтқанда, басқа __құрылғы__ функциясынан немесе __global__ функциясынан). Ақыры, префикс __host________Қосымша, ол процессор деп аталатын және CPU-ны орындайтын функцияны білдіреді - басқаша айтқанда, тұрақты функция.

__Device__ және __global__ функцияларымен байланысты бірқатар шектеулер бар: олар реквизолардан тыс бола алмайды (яғни олар өздерін шақырады), ал олардың өзгермелі дәлелдер болуы мүмкін емес. Соңында __Device__ функциялары __u gpu жад кеңістігінде тұрады, олардың мекен-жайын алу мүмкін еместігін білдіреді. Айнымалыларда сонымен қатар жадтың орналасуын көрсететін бірқатар квалификаторлар бар. Префиксі бар айнымалы __щеде__ағынды мультипроцессордың ортақ жадында сақталатынын білдіреді. __Global__ функциясына қоңырау шалу сәл өзгеше. Мәселе мынада, қоңырау шалу кезінде сіз орындау конфигурациясын орнатуыңыз керек - дәлірек айтқанда, ядро ​​/ ядро ​​қолданылатын тордың / тордың өлшемін, сондай-ақ әрбір блоктың өлшемін. Мысалы, келесі қолтаңбасы бар ядро ​​алыңыз.

Ол деп аталады

Егер DG - бұл тор мөлшері, ал ДБ блоктың өлшемі болып табылады. Бұл екі айнымалы векторлық типтегі жаңа векторлық типке жатады.

CUDA API-де жадпен жұмыс істеу функциялары бар: varam: cudamalloc: жад, CUDAFREE, RAM және VRAM мен керісінше жадты көшіруге арналған жадты және CUDAMEMCPE-ді бөліп алу үшін.

Біз бұл шолуды өте қызықты тәсілмен аяқтаймыз, оның ішінде CUDA бағдарламасы құрастырылған: компиляция бірнеше қадаммен жасалады. Біріншіден, процессорға тән код шығарылады және стандартты компиляторға жіберіледі. GPU үшін тағайындалған код алдымен PTX аралық тіліне түрлендіріледі. Ол құрастыру тіліне ұқсас және сізге ықтимал тиімсіз бөлімдерді іздеуде кодты оқуға мүмкіндік береді. Соңында, соңғы кезең - аралық тілді GPU нақты нұсқауларына аудару және екілік ету жасау.

NVidia құжаттамасын қарау мені осы аптада CUDA-ны қолданып көргім келеді. Шынында да, өз бағдарламаңызды жасау арқылы API бағалаудан артық не болуы мүмкін? Бұл барлық проблемалар бетіне келуі керек, тіпті бәрі қағазға мінсіз көрінеді. Сонымен қатар, тәжірибе сіз қандай-да бір CUDA құжаттамасында көрсетілген барлық қағидаларды қаншалықты түсінетіндігіңізді жақсы көрсетеді.

Мұндай жобаға түсу өте оңай. Бүгінгі таңда көптеген тегін, бірақ сапалы құралдарды жүктеуге болады. Біздің сынақымыз үшін біз сізге қажет барлық нәрсе бар V ++ Express 2005 қолдандық. Ең қиыны GPU-ға өту үшін апталарды қажет етпейтін бағдарламаны табу болды, бірақ біздің күш-жігеріміз босқа кетпейтіндей қызықты болды. Соңында біз биіктік картасын алатын және сәйкес қалыпты картаны есептейтін код бөлігін таңдадық. Біз бұл функцияны егжей-тегжейлі білмейміз, өйткені бұл мақалада қызықты болады. Қысқаша айтқанда, бағдарлама аймақтардың қисықтықымен айналысады: бастапқы кескіннің әр пикселі үшін, біз алдыңғы немесе одан аз күрделі формуланы қолдана отырып, пайда болған кескіннің пайда болатын кескіннің түсін анықтайтын матрица. Бұл функцияның артықшылығы - оны параллельдеу өте оңай берілген сынақ CUDA мүмкіндіктерін жақсы көрсетеді.

Тағы бір артықшылығы - бізде процессорда іске асыру бар, сондықтан біз оның нәтижесін CUDA нұсқасымен салыстыра аламыз - дөңгелекті қайта ойлап таппаймыз.

Біз тағы бір рет қайталаймыз, сынақтың мақсаты CPU және GPU нұсқаларын салыстыру емес, CUDA SDK утилиталарымен танысу болды. Бұл CUDA бағдарламасын құрудағы бірінші әрекетіміз болғандықтан, біз жоғары өнімділікті күткен жоқпыз. Кодтың бұл бөлігі маңызды емес болғандықтан, процессор нұсқасы оңтайландырылмаған, сондықтан нәтижелерді тікелей салыстыру қызықты емес.

Өнімділік

Дегенмен, біз CUDA-ны ең өрескел енгізу кезінде де пайдаланудың артықшылығы бар-жоғын немесе GPU пайдалану кезінде пайда алу үшін ұзақ және жалықтыратын тәжірибе қажет пе екенін білу үшін орындау уақытын өлшедік. Сынақ құрылғысы біздің әзірлеу зертханамыздан алынды - Core 2 Duo T5450 процессоры бар ноутбук және Vista жүйесімен жұмыс істейтін GeForce 8600M GT графикалық картасы. Бұл суперкомпьютер болудан алыс, бірақ нәтижелер өте қызықты, өйткені сынақ GPU үшін «өткірленбеген». Нвидиядағы көптеген өткізу қабілеті бар жүйелерде үлкен жетістіктер және өткізу қабілеті бар, бірақ іс жүзінде, бірақ қазіргі компьютер нарығындағы көптеген 70 миллион CUDA қол жетімді GPU-ның көпшілігі күшті, сондықтан біздің сынақ негізделмеген.

2048 x 2048 пиксель кескіні үшін біз келесі нәтижелерге қол жеткіздік.

• CPU 1 ағыны: 1419 мс;
• CPU 2 ағындары: 749 мс;
• CPU 4 ағыны: 593 мс
• GPU (8600M GT) блоктары 256 жіптің блоктары: 109м;
• ГПУ (8600М ГТ) блоктары 128 жіптің блоктары: 94MS;
• GPU (8800 GTX) 128 жіптің блоктары / 256 жіптер: 31м.

Нәтижелерден бірнеше тұжырым жасауға болады. Бағдарламашылардың айқын жалқаулығы туралы әңгімеге қарамастан, біз бірнеше ағындар үшін орталық процессордың бастапқы нұсқасын өзгерткенімізден бастайық. Жоғарыда атап өткеніміздей, код бұл жағдай үшін өте қолайлы - талап етілетін нәрсе - бастапқы кескінді қанша ағындар болса, сонша аймақтарға бөлу. Біздің екі ядролы процессордағы бір ағыннан екі ағынға ауысқан кезде, жеделдету дерлік сызықты екенін ескеріңіз, бұл сынақ бағдарламасының параллель сипатын да көрсетеді. Күтпеген жерден төрт ағыны бар нұсқа да жылдамырақ болды, дегенмен бұл біздің процессорымыз үшін өте таңқаларлық - керісінше, қосымша ағындарды басқарудың үстеме шығындарына байланысты тиімділіктің төмендеуін күтуге болады. Мұндай нәтижені қалай түсіндіруге болады? Айту қиын, бірақ Windows жіптерін жоспарлаушы, мүмкін, кінәлі; Қалай болғанда да, нәтиже қайталанады. Кішірек текстуралармен (512x512) бөлуден түсетін пайда соншалықты айқын болмады (шамамен 35% қарсы 100%) және төрт ағынды нұсқаның әрекеті екі ағынды нұсқамен салыстырғанда ешқандай пайдасыз қисынды болды. GPU бұрынғысынша жылдамырақ болды, бірақ олай емес (8600M GT қос бұрандалы нұсқасынан үш есе жылдамырақ болды).

Екінші маңызды байқау, GPU-ның ең баяу орындалуының өзі процессордың ең жоғары өнімді нұсқасынан алты есе жылдамырақ болып шықты. Алгоритмнің бірінші бағдарламасы мен ашкөздік нұсқасы үшін нәтиже өте жақсы. Бізді білетінімізді ескеріңіз ең жақсы нәтижеШағын блоктарда, бірақ интуиция басқаша ұсына алады. Түсініктеме қарапайым - біздің бағдарламамыз әр ағынға 14 регистрді пайдаланады, ал 256 ағынды блоктармен бір блокқа 3584 регистр қажет және біз көрсеткендей толық процессорды жүктеу үшін 768 ағын қажет. Біздің жағдайда бұл үш блок немесе 10 572 регистрлер. Бірақ мультипроцессордың тек 8 192 регистрі бар, сондықтан ол тек екі блокты белсенді ұстай алады. Керісінше, 128 ағынды блоктармен бір блокқа 1792 регистр қажет; егер 8192-ні 1792-ге бөліп, ең жақын бүтін санға дейін дөңгелектесек, төрт блок аламыз. Іс жүзінде ағындар саны бірдей болады (бір мультипроцессорға 512, бірақ теориялық тұрғыдан толық жүктеу үшін 768 қажет), бірақ блоктар санын көбейту GPU-ға жадқа қол жеткізудегі икемділіктің артықшылығын береді - операция орындалып жатқанда. Үлкен кідірістермен, ол басқа блоктың нұсқауларын орындауға, нәтижелерді күтуге кірісуі мүмкін. Төрт блок кешігуді азайтады, әсіресе біздің бағдарлама жадқа бірнеше кіруді пайдаланатындықтан.

Талдау

Ақырында, жоғарыда айтылғандарға қарамастан, біз азғыруларға қарсы тұра алмадық және 8800 гТК-дағы бағдарламаны жеңе алмадық, бұл блок мөлшеріне қарамастан, 8600-ге қарағанда үш есе тез болды. Іс жүзінде сәйкес архитектураларда нәтиже төрт немесе одан да көп есе жоғары болады деп ойлауыңыз мүмкін: 128 ALU / шейдер процессорлары 32 және одан жоғары тактілік жиіліктерге қарсы (1,35 ГГц-ке қарсы 950 МГц), бірақ бұл болмады. Мүмкін, шектеуші фактор жадқа қол жеткізу болды. Нақтырақ болу үшін бастапқы суретке көп өлшемді CUDA массиві ретінде қол жеткізіледі - тек құрылымға қарағанда, өте күрделі термин. Бірақ бірнеше артықшылықтарды жеу.

• текстураның кэшіне пайдаға түседі;
• біз процессор нұсқасынан айырмашылығы кескін шекараларын өңдеуді қажет етпейтін орау режимін қолданамыз.

Сонымен қатар, біз «еркін» сүзгілеуден біздің орнына қалыптандырылған, бірақ біздің жағдайымыздан артықшылықтарға ие бола аламыз, бірақ біздің жағдайда бұл пайдалы. Өздеріңіз білетіндей, 8800 GTX үшін 32-мен салыстырғанда 8600-де 16 құрылымдық бірлік бар. Сондықтан, екі сәулет арасындағы қатынас тек екіден ғана. Жиілік айырмашылығына қосылыңыз және біз (32 x 0.575) / (16 x 0.475) = (16 x 0.475) = 2.4 - біз «үшке дейін» қатынасын аламыз. Бұл теория сонымен қатар блоктардың мөлшері G80-де неге өзгермейтінін түсіндіреді, өйткені алу әлі күнге дейін текстуралық қондырғыларға қосылады.

Перспективаны нәтижелермен қатар, біздің CUDA-мен алғашқы кездесуіміз өте жақсы өтті, олар таңдалған емес шарттарды ескере отырып, өте жақсы өтті. Vista ноутбуктарында дамып келе жатқан CUDA SDK 2.0 көмегімен, ол Beta-да, яғни 174,55 драйвері бар, ол да бета-да. Осыған қарамастан, біз қандай да бір жағымсыз тосынсыйлар туралы хабарлай алмаймыз - тек бірінші жөндеу кезіндегі бастапқы қателер, біздің бағдарлама әлі де өте «қате» бөлінген кеңістіктен тыс жадты шешуге тырысқан кезде.

Монитор жыпылықтай бастады, содан кейін экран қара түсті ... Vista жүргізуші жөндеу қызметін басқарғанша және бәрі жақсы болды. Бірақ сіз біздікі сияқты стандартты бағдарламалар бойынша әдеттегі сегментацияның ақаулығын көруге әлі таңқаларлық емес. Соңында, nVidia-ға қатысты шағын сын: CUDA үшін қол жетімді барлық құжаттамада Visual Studio үшін әзірлеу ортасын орнату жолын қадаммен көрсететін шағын нұсқаулық жоқ. Шындығында, мәселе аз, өйткені SDK бар толық жиынтығы CUDA қосымшаларын түсіну үшін зерттелетін мысалдар, бірақ бастаушы нұсқаулығы жақсы болар еді.

NVIDIA CUDA-ны GEIFORCE 8800-нің шығарылуымен таныстырды. Уәделер өте тартымды болып көрінді, бірақ біз нағыз сынаққа деген құлшынысымызды өткіздік. Шынында да, сол кезде бұл GPGPU толқынында қалу үшін аумақты белгілеу сияқты көрінді. Қол жетімді SDK болмаса, біз жұмыс істемейтін басқа маркетингтік дамияға тап болмаймыз. Бұл жақсы бастама бірінші рет айтылған емес, ол ертерек жарияланды және ол кезде қолдаудың болмауына байланысты жарыққа шықпады, әсіресе осындай бәсекеге қабілетті секторда. Енді, хабарландырудан бір жарым жыл өткен соң, nVidia өз сөзінде тұрды деп сенімді түрде айта аламыз.

SDK 2007 жылдың басында бета-артта қалды, содан бері бұл жобаның NVIDIA-ға маңыздылығын дәлелдеді. Бүгінгі таңда CUDA өте жақсы дамып келеді: SDK негізгі операциялық жүйелер үшін (Windows XP және Vista, Linux, сондай-ақ Mac OS X үшін 1.1) бета 2.0 нұсқасында қол жетімді және nVidia сайттың бүкіл бөлімін Әзірлеушілер.

Кәсіби деңгейде, алғашқы қадамдар туралы алғашқы қадамдар өте жағымды болды. Егер сіз GPU архитектурасымен таныс болсаңыз да, оны оңай анықтай аласыз. Бір қарағанда API анық көрінген кезде, сіз бірден сенімді нәтиже бересіз деп сене бастайсыз. Бірақ есептеу уақытын процессордан ГПУ-ге бірнеше аударымдардан ысырап етпейді ме? Бұл мыңдаған жіптерді қалай қолдануға болады? Біз эксперименттерімізді осы қорқыныштардың барлығын ескере отырып бастадық. Бірақ олар біздің алгоритмнің бірінші нұсқасы, өте тривиальды болса да, процессорға қарағанда айтарлықтай жылдамырақ болған кезде тез тарады.

Сонымен, CUDA университет қызметкерлерін оларға Джорсс сатып алуды сендіруді қалайтын зерттеушілер үшін құтқарушы емес. CUDA қазірдің өзінде кез келген C бағдарламашысы жаңа бағдарламалау парадигмасына үйрену үшін уақыт пен күш жұмсауға дайын болса, пайдалана алатын толық қолжетімді технология болып табылады. Егер сіздің алгоритмдеріңіз жақсы параллель болса, бұл күш-жігер босқа кетпейді. Нвидияға Н.Видияға толық және сапалы құжаттама ұсынғаны үшін алғысымызды білдіреміз, мұнда Novice CUDA бағдарламашылары жауап табады.

CUDA танылатын API болу үшін не қажет? Бір сөзбен айтқанда: портативтілік. Біз IT-ның болашағы параллельді есептеулерде екенін білеміз – бүгінде барлығы осындай өзгерістерге дайындалуда және бағдарламалық және аппараттық қамтамасыз етудің барлық бастамалары осы бағытқа бағытталған. Дегенмен, қазіргі уақытта, егер сіз парадигмалардың дамуына қарасаңыз, біз әлі бастапқы сатыдамыз: біз ағындарды қолмен жасаймыз және ортақ ресурстарға қол жеткізуді жоспарлауға тырысамыз; Мұның бәрі қандай-да бір жолмен бұрылып, бір қолдың саусақтарына санауға болады. Бірақ бірнеше жылдан кейін жүздеген процессорлардың саны болған кезде, бұл мүмкіндік енді болмайды. CUDA шығарылымымен Нвидия бұл мәселені шешудің алғашқы қадамын жасады - бірақ, әрине, бұл шешімтек осы компанияның графикалық процессорлары үшін жарамды, тіпті бәрі үшін емес. Тек GF8 және 9 (және олардың квадро / Quadro / Tesla туындылары), бүгінгі таңда CUDA бағдарламаларын іске асыра алады. Және жаңа 260/280 желісі, әрине.

Nvidia бүкіл әлем бойынша 70 миллион CUDA үйлесімді GPU сатты деп мақтана алады, бірақ бұл әлі де факто стандарты болу үшін жеткіліксіз. Бәсекелестер бос отырмағанын ескере отырып. AMD өзінің жеке SDK (Stream Computing) ұсынады және Intel шешімді (Ct) жариялады, бірақ ол әлі қол жетімді емес. Соғыс стандарттары келе жатыр, ал үш бәсекелес үшін нарықта үш бәсекелес үшін орын болмайды, бұл Microsoft сияқты басқа ойыншы ортақ API ұсынысы бар, ол, әрине, әзірлеушілер үшін өмірді жеңілдетеді.

Сондықтан, Нвидияда CUDA мақұлдау жолында көптеген қиындықтар бар. Технологиялық тұрғыдан біздің алдымызда, әрине, сәтті шешім болса да, әзірлеушілерді оның болашағына сендіру қалады - бұл оңай болмайды. Дегенмен, көптеген соңғы API хабарландырулары мен жаңалықтарына қарағанда, болашақ мүлде бұлыңғыр болып көрінбейді.

CUDA технологиясы

Владимир Фролов,[электрондық пошта қорғалған]

аннотация

Мақалада CUDA технологиясы туралы айтылады, ол бағдарламашыға видеокарталарды қуатты есептеу бірліктері ретінде пайдалануға мүмкіндік береді. Nvidia ұсынған құралдар графикалық өңдеу блогының (GPU) бағдарламаларын C++ тілінің ішкі жиынында жазуға мүмкіндік береді. Бұл бағдарламашыдан көлеңкелерді пайдаланудан және графикалық құбырдың қалай жұмыс істейтінін түсінеді. Мақалада CUDA және оптимизацияның әртүрлі әдістері арқылы бағдарламалау мысалдары келтірілген.

1. Кіріспе

Соңғы онжылдықтар ішінде есептеу технологияларын дамыту қарқынмен жүріп жатыр. Өңдеушілерді әзірлеушілер «кремнийдің жойғыштарына» баратындай тез оралды. Сағат жиілігінің шектеусіз өсуі бірқатар күрделі технологиялық себептерге байланысты мүмкін болмады.

Қазіргі заманғы есептеуіш жүйелердің барлық өндірушілері бір процессордың жиілігін арттыруға емес, процессорлар мен ядролардың санын көбейтуге ұмтылуда. Жетілдірілген жүйелердегі орталық өңдеу қондырғысының (CPU) саны 8.

Тағы бір себеп - қошқардың салыстырмалы түрде төмен жылдамдығы. Процессор қаншалықты жылдам жұмыс істесе де, қиындықтар, тәжірибе көрсеткендей, мүлде арифметикалық операциялар емес, жадқа сәтсіз кірулер - кэшті жіберіп алу.

Дегенмен, егер сіз GPU (Графикалық өңдеу блогы) бағытына қарасаңыз, онда олар параллелизм жолымен әлдеқайда ертерек жүрді. GF8800GTX сияқты бүгінгі видеокарталарда 128 процессорға дейін болуы мүмкін.Мұндай жүйелердің өнімділігі, егер шебер бағдарламаланған болса, айтарлықтай маңызды болуы мүмкін (1-сурет).

Күріш. 1. CPU және GPU үшін өзгермелі нүкте операцияларының саны

Алғашқы бейне тек қол жетімді болғанда, олар екі өлшемді деректерді кескіндеу арқылы процессор жүктемесін жоюға арналған жоғары мамандандырылған құрылғылар (процессормен салыстырғанда) салыстырмалы түрде қарапайым болды. Ойын индустриясының дамуымен және Doom (2-сурет) және Wolfenstein 3D (3-сурет) сияқты үш өлшемді ойындардың пайда болуымен 3D визуализациясының қажеттілігі туындады.

2,3-суреттер. Ойындар Doom және Wolfenstein 3D

Компания құрылған сәттен бастап Voodoo (1996) және 2001 жылға дейін GPU алғашқы 3Dfx карталары кіріс деректері бойынша операциялардың тіркелген жиынтығына ғана жүзеге асырылды.

Бағдарламалаушылардың визуализация алгоритмінде таңдауы жоқ және икемділікті арттыру үшін әрбір шың үшін немесе әрбір пиксел үшін бейне картаны іске қосатын шағын бағдарламалар болды. Олардың міндеттері шыңдардағы түрлендірулерді және көлеңкелеуді - нүктедегі жарықтандыруды есептеуді, мысалы, Фонг моделіне сәйкесті қамтиды.

Шейдерлер қазіргі уақытта өте жетілдірілген болса да, олар жоғары мамандандырылған 3D түрлендіру және растризация тапсырмалары үшін әзірленгенін түсіну керек. GPU әмбебап мультипроцессорлық жүйелерге қарай дамып жатқанымен, шейдер тілдері жоғары мамандандырылған.

Оларды FORTRAN тілімен салыстыруға болады, өйткені олар FORTRAN сияқты бірінші болды, бірақ тек бір типті мәселелерді шешуге арналған. FORTRAN сандық есептеулерге қатысы жоқ есептерді шешуге жарамсыз сияқты, шейдерлердің 3D түрлендірулері мен растризациядан басқа есептерді шешу үшін пайдасы аз.

Бүгінгі таңда кванттық механика, жасанды интеллект, физика есептеулері, криптография, физикалық дұрыс визуализация, кескіндерді қайта құру, IT тану салаларындағы қиындықтарды шешу үшін графикалық карталарды дәстүрлі емес пайдалану үрдісі бар. Бұл тапсырмаларды графикалық API (DirectX, OpenGL) ішінде шешу ыңғайсыз, өйткені бұл API интерфейстері мүлдем басқа қолданбалар үшін жасалған.

GPU бойынша General Programming (GPGPU) дамуы логикалық тұрғыдан растризацияға қарағанда кең ауқымды тапсырмаларға бағытталған технологиялардың пайда болуына әкелді. Нәтижесінде компанияда Nvidia Compute Unified Device Architecture технологиясы (қысқартылған CUDA) құрылды, бәсекелес ATI - STREAM технологиясы.

Айта кету керек, осы мақаланы жазу кезінде STREAM STREAM технологиясы CUDA әзірлеуде әлдеқайда артта қалды, сондықтан ол қарастырылмайды. Біз C++ ішкі жиынында бағдарламаларды жазуға мүмкіндік беретін CUDA – GPGPU технологиясына тоқталамыз.

2. Орталық процессор мен графикалық процессор арасындағы негізгі айырмашылық

Орталық процессор мен графикалық қолданбалардың аумақтары мен мүмкіндіктері арасындағы кейбір маңызды айырмашылықтарды қысқаша қарастырыңыз.

2.1. Мүмкіндіктер

Орталық процессор бастапқыда жоспардың жалпы мақсаттарын шешуге бейімделген және кездейсоқ адрестелетін жадымен жұмыс істейді. Орталық процессордағы бағдарлама сызықтық және біркелкі жадтың барлық ұяшықтарына тікелей қол жеткізе алады.

GPU үшін олай емес. Осы мақаланы оқу арқылы білетіндей, CUDA жадтың 6 түрі бар. Біз физикалық қол жетімді кез келген ұяшықтан оқи аламыз, бірақ бұл жазба барлық ұяшықтарда емес. Себебі GPU кез келген жағдайда белгілі бір мақсаттарға арналған нақты құрылғыны білдіреді. Бұл шектеу белгілі бір алгоритмдердің жылдамдығын арттыру және жабдықтың құнын төмендету мақсатында енгізілген.

2.2. Жад өнімділігі

Көптеген есептеу жүйелеріндегі көпжылдық мәселе - бұл жад процессорға қарағанда баяу. CPU өндірушілері оны кэштерді енгізу арқылы шешеді. Жадтың жиі қолданылатын аймақтары процессор жиілігінде жұмыс істейтін скретч-жадқа немесе кэш-жадқа орналастырылады. Бұл ең жиі қолданылатын деректерге қол жеткізу және есептеудің процессорлық жүктемесі кезінде уақытты үнемдеуге мүмкіндік береді.

Бағдарламашыға арналған кэш шын мәнінде мөлдір екенін ескеріңіз. Оқуға келетін болсақ, деректер жазу кезінде тікелей жадқа түспейді және кэштер арқылы өтеді. Бұл, атап айтқанда, жазылғаннан кейін бірден кейбір мәндерді жылдам оқуға мүмкіндік береді.

GPU-да (графикалық карта GF сегізінші сериясын білдіреді) кэштер де маңызды және олар да маңызды, бірақ бұл механизм CPU сияқты қуатты емес. Біріншіден, кэштелген жадтың барлық түрлері емес, екіншіден, кэштер тек оқуға арналған.

Параллельді есептеулерді пайдалана отырып, GPU баяу жадқа қол жеткізу жасырылады. Кейбір тапсырмалар деректерді күтсе, басқа жұмыстар есептеуге дайын. Бұл жалпы жүйенің жұмысын айтарлықтай жақсартуға мүмкіндік беретін CUDA негізгі принциптерінің бірі.

3. Негізгі CUDA

3.1. жіп үлгісі

Тұжырымдамаға негізделген CUDA есептеу архитектурасыәр түрлі деректер бойынша бір команда(Single Instruction Multiple Data, SIMD) тұжырымдамасы және мультипроцессор.

SIMD концепциясы бір нұсқау бір уақытта көптеген деректерді өңдей алатынын білдіреді. Мысалы, addps командалық процессоры Pentium 3 және одан кейінгі Pentium 4 үлгілері бір мезгілде қалқымалы нүктені бір дәлдікпен бүктеуге мүмкіндік береді.

Мультипроцессор – SIMD көп ядролы процессоры, ол берілген уақыттың кез келген нүктесінде тек бір команданың барлық ядроларын орындауға мүмкіндік береді. Әрбір көппроцессорлық ядроның скаляры, яғни, ол таза түрде векторлық операцияларды қолдамайды.

Жалғастырмас бұрын, біз бірнеше анықтамаларды енгіземіз. Көптеген бағдарламашылар үйренгендіктен, осы мақаладағы құрылғы мен хост мүлдем түсінілмейтінін ескеріңіз. Біз мұндай терминдерді CUDA құжаттамасына сәйкес келмеу үшін қолданамыз.

Осы мақаладағы құрылғының (құрылғының) астында біз CUDA драйверін қолдайтын графикалық картаны немесе CUDA көмегімен бағдарламаларды орындауға арналған басқа арнайы құрылғыны (мысалы, NVIDIA Tesla) түсінеміз. Бұл мақалада біз нақты іске асыру мәліметтерін болдырмай, логикалық құрылғы ретінде GPU ғана қарастырамыз.

Хост (хост) процессорды пайдаланатын және құрылғыны басқару үшін басқару функцияларын орындайтын қалыпты операциялық компьютер жадындағы бағдарлама деп аталады.

Шын мәнінде, сіздің бағдарламаңыздың процессорда жұмыс істейтін бөлігі - бұлхост және сіздің бейне картаңызқұрылғы. Логикалық тұрғыдан құрылғыны мультипроцессорлар жиынтығы (4-сурет) және CUDA драйвері ретінде қарастыруға болады.

Күріш. 4. Құрылғы

Біз құрылғымызда N ағындарда белгілі бір процедураны іске қосқымыз келеді делік (яғни оның жұмысын параллельдегіміз келеді). CUDA құжаттамасына сәйкес біз мұны процедураның өзегі деп атаймыз.

CUDA архитектурасының ерекшелігі көп ағынды қосымшалар үшін әдеттен тыс блок-торлы ұйым болып табылады (5-сурет). Бұл ретте CUDA драйвері ағындар арасында құрылғы ресурстарын дербес таратады.

Күріш. 5. Ағындарды ұйымдастыру

Суретте. 5. Ядро ядро ​​ретінде белгіленеді. Осы ядромен жұмыс істейтін барлық ағындар блоктарға (Блок) топтастырылған, ал блоктар, өз кезегінде, торда (Тор) біріктірілген.

5-суретте көрсетілгендей, ағындардың өлшемдік кодтарын анықтау үшін пайдаланылады. CUDA әзірлеушілеріне үш өлшемді, екі өлшемді немесе қарапайым (бірөлшемді) индекстермен жұмыс істеу мүмкіндігі беріледі, бағдарламашыға қайсысы ыңғайлы болса.

Жалпы алғанда индекстер үш өлшемді векторлар болып табылады. Әрбір ағын үшін белгілі: блок ішіндегі threadIdx ағынының индексі және blockIdx торындағы блок индексі. Іске қосу кезінде барлық ағындар тек осы индекстерде ерекшеленеді. Шын мәнінде, бұл индекстер арқылы бағдарламашы басқарады, оның деректерінің қай бөлігі әрбір ағында өңделетінін анықтайды.

Неліктен әзірлеушілер ұйымды таңдады деген сұраққа жауап тривиальды емес. Мұның бір себебі - бір блоктың орындалуына кепілдік беріледібірінде көппроцессорлы құрылғы, бірақ көппроцессор бірнеше түрлі блоктарды орындай алады. Басқа себептер мақала барысында толығырақ түсіндірілді.

Мультипроцессорлық бөліктерде немесе пулдарда жұмыс істейтін блок тапсырмалары (жіптер) бұрылыс деп аталады. Қазіргі уақытта CUDA бар графикалық карталардағы өлшемнің ауытқуы 32 ағынды құрайды. Пул ішіндегі тапсырмалар SIMD стилінде орындалады, яғни тек бір нұсқаудағы барлық бұрма ағындары бір уақытта орындалуы мүмкін.

Мұнда біз брондауымыз керек. Осы жазу кезінде қазіргі заманғы архитектура, мультипроцессор ішіндегі процессорлар саны 32 емес, 8. Бұдан шығатыны, барлық деформация бір уақытта орындалмайды, ол дәйекті түрде орындалатын төрт бөлікке бөлінеді (скалярлық процессор ретінде). ).

Бірақ, біріншіден, CUDA әзірлеушілері деформация өлшемін қатаң реттемейді. Олардың еңбектерінде олар 32 саны емес, кереует өлшемінің нұсқасын атап өтеді. Екіншіден, қисындылық тұрғысынан декорацияның минималды біріктірілген ағыны болып табылады, ол туралы оның ішіндегі барлық жіптер бір уақытта орындалады деп айта аламыз - және бұл жүйенің қалған бөлігі туралы ешқандай болжамдар жоқ.

3.1.1. тармақталу

Бірден сұрақ туындайды: егер бір уақытта иілу ішіндегі барлық жіптер бірдей нұсқауды орындаса, онда тармақтар туралы не деуге болады? Өйткені, егер бағдарламаның тармақталуының коды орын алса, нұсқаулар қазірдің өзінде басқаша. Ол стандартты SIMD бағдарламалау шешіміне қатысты (6-сурет).

Күріш. 6. SIMD филиалы

Бізде келесі код бар делік:

егер (шарт)B;

SISD (Single Instruction Single Data) жағдайында біз A операторын орындаймыз, шартты тексереміз, содан кейін B және D операторларын орындаймыз (егер шарт ақиқат болса).

Енді SIMD стилінде орындалатын 10 ағын бар делік. Барлық 10 ағында біз A операторын орындаймыз, содан кейін шартты тексереміз және 10 ағынның 9-ы дұрыс, бірақ бір ағында - жалған екенін анықтадық.

В операторын орындау үшін 9 ағынды, ал С операторын орындау үшін қалған бір ағынды іске қоса алмайтынымыз анық, өйткені барлық ағындарда бір уақытта бір ғана нұсқау орындалуы мүмкін. Бұл жағдайда мұны істеу керек: біріншіден, ешкімнің деректерін бұзбауы үшін бөлінетін ағынды «өлтіру» және қалған 9 ағынды жүзеге асыру. Содан кейін «өлтіру» 9 ағыны B операторын орындап, C операторымен бір ағынды өткізеді. Содан кейін ағындар қайтадан біріктіріліп, D операторы бір уақытта жұмыс істейді.

Бұл қайғылы нәтиже болып шықты: процессор ресурстары бөлінген ағындарда бос биттерді тегістеуге жұмсалып қана қоймайды, бірақ одан да жаманы, біз ақырында ЕКІ тармақты орындауға мәжбүр боламыз.

Дегенмен, бәрі бір қарағанда көрінетіндей жаман емес. Технологияның үлкен артықшылығы - бұл трюктердің CUDA драйверінің динамикалық түрде орындалуы және олар бағдарламашы үшін толығымен мөлдір. Сонымен қатар, қазіргі заманғы процессорлардың SSE нұсқауларымен жұмыс істегенде (дәлірек айтқанда, алгоритмнің 4 көшірмесін бір уақытта орындауға тырысқан жағдайда) бағдарламалаушының өзі егжей-тегжейлерге назар аударуы керек: деректерді төрт есе біріктіру, орындау туралауды ұмытпаңыз және әдетте ассемблерде сияқты төмен деңгейде жазыңыз.

Жоғарыда айтылғандардан өте маңызды қорытынды шығады. Тармақталу өнімділікті өздігінен төмендетпейді. Зиянды тек тармақталу, онда бір бассейн ішіндегі ағынының жылдамдығы жіптерді бұрмалайды. Сонымен қатар, егер бір блок ішінде, бірақ әртүрлі бассейндерде бөлінген ағындар иілсе немесе әртүрлі блоктар ішінде болса, бұл мүлдем әсер етпейді.

3.1.2. Жіптер арасындағы өзара әрекеттесу

Осы мақаланы жазу кезінде ағындар арасындағы кез келген өзара әрекеттесу (синхрондау және деректер алмасу) тек блок ішінде мүмкін болды. Яғни, тек құжатталған мүмкіндіктерді пайдалана отырып, әртүрлі блоктар ағындары арасындағы өзара әрекетті ұйымдастыру мүмкін емес.

Құжатсыз мүмкіндіктерге келетін болсақ, оларды пайдалану ұсынылмайды. Мұның себебі олар жүйенің нақты аппараттық мүмкіндіктеріне сүйенеді.

Құрылғыдағы барлық тапсырмаларды синхрондау __synchtreads функциясын шақыру арқылы орындалады. Деректер алмасу ортақ жад арқылы мүмкін болады, өйткені бұл блоктағы барлық тапсырмалар үшін ортақ.

3.2. Жад

CUDA-да жадтың алты түрі бар (7-сурет). Ол тіркеледі, жергілікті, ғаламдық, ортақ, тұрақты және текстуралық жады.

Мұндай молшылық видеокартаның ерекшеліктеріне және оның негізгі мақсатына, сондай-ақ әзірлеушілердің әртүрлі жағдайларда әмбебаптығын немесе жылдамдығын құрбан етіп, жүйені мүмкіндігінше арзан етуге ұмтылуымен байланысты.

Күріш. 7. CUDA жадысының түрлері

3.2.0. Тіркеулер

Мүмкін болса, компилятор барлық жергілікті айнымалыларды регистрлерге орналастыруға тырысады. Мұндай айнымалыларға қол жеткізу максималды жылдамдықпен жүзеге асырылады. Бір мультипроцессорлық сипаттамаға арналған ағымдағы архитектура 8192 32-разрядты регистрлер. Бір ағынның қанша қолжетімді тіркелетінін анықтау үшін бұл санды (8192) блок өлшеміне (ондағы ағындар саны) бөлу керек.

Ағында 64 типтік бөлуде тек 128 регистр алынады (белгілі бір объективті критерийлер бар, бірақ көптеген қосымшалар үшін орташа есеппен 64 жарамды). Шындығында, nvcc ешқашан 128 регистрді бөлмейді. Әдетте, ол 40-тан артық бермейді, ал басқа айнымалылар жергілікті жадта popadpyut. Бұл көппроцессорлық бірнеше блоктарды іске қосуға болатындықтан орын алады. Компилятор қатарлас блоктардың санын барынша көбейтуге тырысады. Жоғары тиімділік үшін 32-ден аз регистрлерді көбірек алып көру керек. Содан кейін теориялық түрде бір мультипроцессорда 4 блокты (егер ағын бір блокта 64 болса, 8 деформация) іске қосуға болады. Дегенмен, бұл жерде ағындар алатын ортақ жад көлемін де ескеру қажет, өйткені бір блок бүкіл ортақ жадты алып жатса, мұндай екі блокты мультипроцессорда бір уақытта орындау мүмкін емес.

3.2.1. жергілікті жады

Процедураның жергілікті деректері тым үлкен болған жағдайда немесе компилятор оларға қатынасу кезінде олар үшін қандай да бір тұрақты қадамды есептей алмаса, ол оларды жергілікті жадқа орналастыра алады. Бұл, мысалы, әртүрлі өлшемдердің түрлері үшін көрсеткіштерді түрлендіруді жеңілдетуге болады.

Физикалық тұрғыдан жергілікті жад аналогты жаһандық жад болып табылады және бірдей жылдамдықпен жұмыс істейді. Жазу кезінде нақты айнымалылар үшін жергілікті жад компиляторын пайдалануға нақты тыйым салатын механизмдер болған жоқ. Жергілікті жадты басқару қиын болғандықтан, оны мүлдем қолданбаған дұрыс (қараңыз. 4 «Оңтайландыру бойынша ұсыныстар» бөлімі).

3.2.2. жаһандық жады

CUDA құжаттамасында негізгі жетістіктердің бірі ретіндетехнология кез келген жаһандық жадты шешуге мүмкіндік береді. Яғни, кез келген жад орнынан оқуға болады, және сіз де еркін ұяшықта жаза аласыз (GPU-да әдетте бұл болмайды).

Дегенмен, бұл жағдайда икемділік үшін тарифтерді төлеуге тура келеді. Ғаламдық жад кэштелмеген. Ол өте баяу жұмыс істейді, жаһандық жадқа кіру саны кез келген жағдайда азайтылуы керек.

Ғаламдық жад негізінен бағдарлама нәтижелерін хостқа (қарапайым DRAM жадына) жібермес бұрын сақтау үшін қажет. Мұның себебі жаһандық жады - жадтың жалғыз түрі, ол жазуға болатын нәрсе болуы мүмкін.

__global__ квалификациясы жарияланған айнымалылар жаһандық жадқа орналастырылады. Ғаламдық жадты хостта cudaMalloc (void * mem, int өлшемі) шақыру арқылы динамикалық түрде бөлуге болады. Бұл функцияны құрылғыдан шақыру мүмкін емес. Бұдан шығатыны, жадты бөлу процессорда жұмыс істейтін хост бағдарламасына қосылуы керек. Хост құрылғысынан деректерді cudaMemcpy функционалдық шақыруына жіберуге болады:

cudaMemcpy(void* gpu_mem, void* cpu_mem, int өлшемі, cudaMemcpyHostToDevice);

Дәл осылай кері процедураны орындауға болады:

cudaMemcpy(void* cpu_mem, void* gpu_mem, int өлшемі, cudaMemcpyDeviceToHost);

Бұл қоңырау да хосттан жасалады.

Жұмыс кезінде жаһандық жады коалесинга (біріктіру) ережесін сақтау маңызды. Негізгі идея - бұл posledovaetlnym жіптер 4,8 немесе 16 байт жад ұяшықтарына сілтеме жасау керек. Сонымен қатар, ең бірінші сауда сәйкесінше 4,8 немесе 16 байт шекараға сәйкес келетін келесі мекенжайға хабарласуы керек. cudaMalloc қайтарған мекенжайлар кемінде 256 байт шекарасында тураланған.

3.2.3. Ортақ жад

Ортақ жад – бұл кэштелмейтін, бірақ жылдам жад. Басқарылатын кэш ретінде пайдалану ұсынылады. Бір мультипроцессорға тек 16 Кбайт ортақ жад қолжетімді. Бұл санды блоктағы тапсырмалар санына бөлсек, біз аламыз максималды сомаортақ жад бір ағында қолжетімді (егер сіз оны барлық ағындарда дербес пайдалануды жоспарласаңыз).

Ортақ жадтың ерекшелігі оның блок ішіндегі барлық тапсырмалар үшін бірдей адрестелуі болып табылады (7-сурет). Бұдан шығатыны, оны тек бір блоктың ағындары арасында деректер алмасу үшін пайдалануға болады.

Ортақ жад мультипроцессорында блокты орындау кезінде мазмұнның сақталуына кепілдік беріледі. Алайда, мультипроцессорлы блокқа ауыстырылғаннан кейін ескі блоктың мазмұны сақталатынына кепілдік берілмейді. Сондықтан ортақ жадта кез келген деректерді қалдырып, олардың қауіпсіздігіне үміттеніп, блоктар арасындағы тапсырмаларды синхрондауға тырыспаңыз.

__shared__ квалификациясымен жарияланған айнымалылар ортақ жадқа орналастырылады.

shared_float mem_shared;

Бірлік үшін ортақ жад екенін тағы бір рет атап өткен жөн. Сондықтан, егер сіз оны басқарылатын кэш ретінде пайдаланғыңыз келсе, массивтің әртүрлі элементтеріне сілтеме жасауыңыз керек, мысалы, келесідей:

float x = mem_shared;

Мұндағы threadIdx.x блок ішіндегі ағынның x индексі болып табылады.

3.2.4. Тұрақты есте сақтау

Тұрақты жады кэштелген, суретте көрсетілген. 4. Кэш жалғыз көп процессорға бір көшірмеде болады, бұл бірлік ішіндегі барлық тапсырмалар үшін жиынтықты білдіреді. Тұрақты жадтағы хост cudaMemcpyToSymbol шақыру арқылы бірдеңені жаза алады. Құрылғыдан тұрақты жад тек оқуға арналған.

Тұрақты жадты пайдалану өте оңай. Сіз оны кез келген деректер түріне орналастыра аласыз және оларды қарапайым тапсырма арқылы оқи аласыз.

#N 100 анықтаңыз

Constant__int gpu_buffer[N];

void host_function()

int cpu_buffer[N];

cudaMemcpyToSymbol (gpu_buffer, cpu_buffer, sizeof (int) * N);

// __global__ бұл құрылғының_ядросы - GPU-да іске қосуға болатын ядроны білдіреді

Global__void device_kernel()

int a = gpu_buffer;

int b = gpu_buffer + gpu_buffer;

// gpu_buffer = a; ҚАТЕ! тұрақты жады тек оқуға арналған

Тұрақты жад кэші пайдаланылғандықтан, оған қол жеткізу әдетте өте жылдам. Тұрақты жадтың жалғыз, бірақ өте үлкен кемшілігі оның көлемі небәрі 64 Кбайт (барлық аппаратта). Бұдан шығатыны, жад контекстінде жиі қолданылатын мәліметтердің аз ғана санын сақтаудың мәні бар.

3.2.5. текстуралық жады

Текстуралық жады кэштелген (Cурет 4). Әрбір мультипроцессорда бір ғана кэш бар, бұл бұл кэш құрылғыдағы барлық тапсырмалармен ортақ пайдаланылатынын білдіреді.

Текстуралық жадының атауы (және, өкінішке орай, функционалдығы) «текстура» және «текстуралау» ұғымдарынан мұраланған. Текстуралау - растризация процесінде текстураны (жай суреттерді) жерге қолдану процесі. Текстуралық жады үлгі 2D деректері үшін оңтайландырылған және келесі мүмкіндіктерге ие:

бір өлшемді немесе екі өлшемді массивтен бекітілген өлшемді мәндерді (байт, сөз, қос немесе төрт сөз) жылдам таңдау;

диапазондағы қалыпқа келтірілген адрестік түрдегі қалқымалы сандар. Содан кейін оларды қалыпты адрестеу арқылы таңдауға болады. Budetet сөзінің нәтиже мәні float4 түрі, интервалда көрсетіледі;

CudaMalloc ((жарамсыз **) & gpu_memory, N * sizeof (uint4)); //жадты графикалық процессорға бөлу

// paraemtrov текстурасының текстурасын орнату

Texture.addressMode = cudaAddressModeWrap; //режимі Орау

Texture.addressMode = cudaAddressModeWrap;

Texture.filterMode = cudaFilterModePoint; //blizhaysheeznachenie

texture.normalized = false; // нормаланған адрестеуді қолданбаңыз

CudaBindTexture (0, текстура, gpu _ жады, N ) // енді бұл жады текстура ретінде қарастырылады


CudaMemcpy (gpu _ жад, CPU _ буфер, N * sizeof (uint 4), cudaMemcpyHostToDevice ); // деректерді көшіруGPU

// __global__ құрылғы_ядросы параллельденетін ядро ​​екенін білдіреді

Global__void device_kernel()

uint4 a = tex1Dfetch (текстура, 0); // тек осы жолмен деректерді таңдауға болады!

uint4 b = tex1Dfetch (текстура, 1);

int c = a.x * b.y;

...

3.3. Қарапайым мысал

Қарапайым мысал CUDA SDK cppIntegration бағдарламасын қарастыруға шақырылады. Ол CUDA-мен қалай жұмыс істеу керектігін, сондай-ақ CUDA бағдарламаларын әзірлеуді айтарлықтай жеңілдететін MS Visual Studio-мен бірге nvcc (Nvidia-ның арнайы C++ ішкі жиыны компиляторы) пайдалануды көрсетеді.

4.1. Тапсырмаңызды дұрыс бөліңіз

Барлық тапсырмалар SIMD архитектурасына сәйкес келмейді. Егер сіздің мақсатыңыз сәйкес келмесе, бұл мүмкін, GPU қолданбаңыз. Бірақ егер сіз GPU пайдалануды шешсеңіз, бұл элементтердің алгоритмін бұзуға тырысуыңыз керек, осылайша оларды SIMD стилінде тиімді орындауға болады. Қаласаңыз - мәселеңізді шешу үшін алгоритмді өзгертіңіз, SIMD-де орналастыру жақсы болатын жаңасын жасаңыз. Тиісті пайдалану аймақтарының мысалы ретінде GPU іске асыру нәтижесінде пирамидалық массив элементтерін қосуға болады.

4.2. Жад түрін таңдау

Құрылғының барлық тапсырмалары бірдей жад орнына немесе жақын орналасқан станцияларға қатысты болса, деректерді текстуралық жадқа немесе тұрақтыға қойыңыз. Екі өлшемді деректерді функциялар мен text2Dfetch text2D көмегімен тиімді өңдеуге болады. Текстуралық жады екі өлшемді үлгі үшін арнайы оңтайландырылған.

Барлық тапсырмалар жадтың әртүрлі, кең аралық аймақтарына (егер ол 2D/3D деректері болса, өте әртүрлі мекенжайлары немесе координаттары бар) кездейсоқ қатынасатын болса, жалпы жадпен бірге жаһандық жадты пайдаланыңыз.

жаһандық жад => ортақ жад

Syncthreads ();

Ортақ жадтағы деректерді өңдеу

Syncthreads ();

жаһандық жады<= разделяемая память

4.3. Жад есептегіштерін қосыңыз

Flag --ptxas-options = -v компиляторы қаншалықты және қандай жадты (регистрлер, ортақ, жергілікті, тұрақты) қолданатыныңызды нақты айтуға мүмкіндік береді. Егер компилятор жергілікті жадты пайдаланса, сіз бұл туралы білесіз. Пайдаланылатын жадтың саны мен түрлері туралы деректерді талдау оңтайландыру бағдарламасында сізге көп көмектеседі.

4.4. Регистрлер мен ортақ жадты пайдалануды барынша азайтуға тырысыңыз

Неғұрлым ядро ​​ортақ жадты немесе регистрлерді пайдаланса, соғұрлым азырақ ағындар (дәлірек айтқанда, warp-лар) мультипроцессорда бір уақытта орындалуы мүмкін, өйткені көппроцессорлық ресурстар шектеулі. Сондықтан регистрлердің немесе ортақ жадтың толтырылуының шамалы ұлғаюы кейбір жағдайларда өнімділіктің екі есе төмендеуіне әкелуі мүмкін - дәл қазір мультипроцессорда бір уақытта екі есе аз бұрмалаулар орындалатындықтан.

4.5. Жергілікті емес, ортақ жад

Егер Nvidia компиляторы қандай да бір себептермен деректер жергілікті жадта орналасса (әдетте ресурс жоқ жерлерде өнімділіктің өте қатты төмендеуімен байқалады), жергілікті жадта қандай деректер бар екенін анықтаңыз және оларды ортақ жадқа орналастырыңыз. (ортақ жад).

Көбінесе компилятордың жергілікті жадында айнымалысы болады, егер ол жиі пайдаланылмайтын болса. Мысалы, циклде бірдеңе күте отырып, құнды жинайтын батареяның бұл түрі. Егер цикл код өлшемі бойынша үлкен болса (бірақ орындалу уақытында емес!), онда компилятор сіздің аккумуляторыңызды жергілікті жадқа қоя алады, себебі ол салыстырмалы түрде сирек қолданылады және аз тіркеледі. Бұл жағдайда өнімділік жоғалуы айтарлықтай болуы мүмкін.

Егер сіз айнымалыны өте сирек қолдансаңыз - оны жаһандық жадқа нақты қою жақсы.

Компиляторға мұндай айнымалы мәндерді жергілікті жадта автоматты түрде бөлу ыңғайлы болып көрінсе де, олай емес. Егер айнымалы жиірек қолданыла бастаса, бағдарламаның кейінгі модификацияларымен қиыншылықты табу оңай болмайды. Компилятор мұндай айнымалыны регистр жадына жылжыта алады немесе жылжыта алмайды. Егер __global__ модификаторы анық көрсетілсе, бағдарламашы оған көбірек көңіл бөледі.

4.6. Циклді айналдыру

Циклді ашу - көптеген жүйелердегі стандартты өнімділік трюк. Оның мәні әрбір итерацияда көбірек әрекеттерді орындау болып табылады, осылайша итерациялардың жалпы санын, демек, процессор орындауға мәжбүр болатын шартты секірулердің санын азайтады.

Массивтің қосындысын (мысалы, бүтін сан) табу үшін циклды қалай шығаруға болады:

int a[N]; инсум;

үшін (int i=0;i

Әрине, ілмектерді қолмен шығаруға болады (жоғарыда көрсетілгендей), бірақ бұл өнімсіз жұмыс. C++ үлгілерін кірістірілген функциялармен бірге пайдалану әлдеқайда жақсы.

үлгі

сынып ArraySumm

Құрылғы__ статикалық T орындау(const T* arr) (қайтару arr + ArraySumm (arr+1); )

үлгі

сынып ArraySumm<0,T>

Құрылғы__ статикалық T exec(const T* arr) (қайтару 0; )

үшін (int i=0;i

жиын+= ЖиымҚосындысы<4,int>::exec(a);

Nvcc компиляторының бір қызықты ерекшелігін атап өту керек. Компилятор әрқашан әдепкі бойынша __device__ сияқты кірістірілген функцияларды орындайды (мұны қайта анықтау үшін арнайы __noinline__ директивасы бар).

Сондықтан, жоғарыдағы мысал қарапайым мәлімдемелер тізбегіне айналатынына және тиімділігі жағынан қолмен жазылған кодтан ешбір жағдайда кем түспейтініне сенімді бола аласыз. Дегенмен, жалпы жағдайда (nvcc емес) сіз бұған сенімді бола алмайсыз, өйткені кірістірілген - бұл компилятордың елемеуі мүмкін деген нұсқауы ғана. Сондықтан, сіздің функцияларыңыз кірістірілген болатынына кепілдік берілмейді.

4.7. Деректерді туралау және 16 байт таңдау

Деректер құрылымдарын 16 байт шекарасында туралаңыз. Бұл жағдайда компилятор деректерді бірден 16 байтта жүктейтін олар үшін арнайы нұсқауларды пайдалана алады.

Құрылым 8 байт немесе одан аз болса, оны 8 байтқа туралауға болады. Бірақ бұл жағдайда екі 8 байт айнымалы мәнді біріктіру немесе көрсеткішті трансляциялау арқылы құрылымға біріктіру арқылы бірден екі айнымалы мәнді таңдауға болады. Трансляцияны сақтықпен пайдалану керек, өйткені компилятор деректерді регистрлерден гөрі жергілікті жадқа орналастыра алады.

4.8. Ортақ жад банкі қайшылықтары

Ортақ жад 4 байт қадамымен 16 (тек!) жад банктері түрінде ұйымдастырылған. Мультипроцессорда ілгерілемелі жіп пулын орындау кезінде ол ортақ жадқа кезекпен қол жеткізетін 16 ағынның екі жартысына (егер warp-өлшемі = 32 болса) бөлінеді.

Бұрманың әртүрлі жартысындағы тапсырмалар ортақ жадқа қайшы келмейді. Warp пулының жартысының тапсырмалары бірдей жад банктеріне қол жеткізетіндіктен, соқтығыстар орын алады және нәтижесінде өнімділік төмендейді. Бірдей аралықтағы тапсырмалар белгілі бір қадаммен ортақ жадтың әртүрлі аймақтарына қол жеткізе алады.

Оңтайлы қадамдар 4, 12, 28, ..., 2^n-4 байт (8-сурет).



Күріш. 8. Оңтайлы қадамдар.

Оңтайлы емес қадамдар 1, 8, 16, 32, ..., 2^n байт (9-сурет).



Күріш. 9. Оңтайлы емес қадамдар

4.9. Хост деректерінің қозғалысын азайту<=>құрылғы

Аралық нәтижелерді орталық процессормен өңдеу үшін мүмкіндігінше азырақ орналастыруға тырысыңыз. Бүкіл алгоритм болмаса, ең болмағанда оның негізгі бөлігін GPU-да іске асырыңыз, процессорды тек басқару тапсырмаларымен қалдырыңыз.

5. CPU/GPU портативті математикалық кітапханасы

Осы мақаланың авторы қарапайым кеңістіктік объектілермен жұмыс істеуге арналған MGML_MATH портативті кітапханасын жазды, оның коды құрылғыда да, хостта да жұмыс істейді.

MGML_MATH кітапханасын физикалық, графикалық немесе басқа кеңістіктік мәселелерді есептеуге арналған портативті (немесе гибридті) CPU/GPU жүйелерін жазу үшін негіз ретінде пайдалануға болады. Оның басты артықшылығы сол кодты процессорда да, графикалық процессорда да қолдануға болады, сонымен бірге кітапханаға қойылатын талаптардың басында жылдамдық қойылады.

6 . Әдебиет

Крис Касперский. Бағдарламаны оңтайландыру техникасы. Жадты тиімді пайдалану. - Санкт-Петербург: BHV-Петербург, 2003. - 464 б.: сырқат.

CUDA бағдарламалау нұсқаулығы 1.1 ( http://developer.download.nvidia.com/compute/cuda/1_1/NVIDIA_CUDA_Programming_Guide_1.1.pdf )

CUDA бағдарламалау нұсқаулығы 1.1. 14-15 бет

CUDA бағдарламалау нұсқаулығы 1.1. 48 бет