{ "metadata": { "name": "" }, "nbformat": 3, "nbformat_minor": 0, "worksheets": [ { "cells": [ { "cell_type": "markdown", "metadata": {}, "source": [ "# Programowanie masywnie r\u00f3wnoleglych procesor\u00f3w\n", "\n", "Wst\u0119p\n", "=====\n", "\n", "\n", "Trendy w rozwoju CPU GPU \n", "--------------------\n", "\n", "\n", "W XX wieku procesory graficzne s\u0142u\u017cy\u0142y do transformacji grafiki, g\u0142\u00f3wnie dwuwymiarowej. Potrzeby przemys\u0142u zwi\u0105zanego z nowoczesnymi grami wykorzystuj\u0105cymi grafik\u0119 tr\u00f3jwymiarow\u0105 doprowadzi\u0142y do intensywnego rozwoju i co r\u00f3wnie wa\u017cne - upowszechnienia si\u0119, wyspecjalizowanych procesor\u00f3w zwanych GPU (Graphics Processor Units), zdolnych renderowa\u0107 coraz bardziej skomplikowane sceny. W 2001 wraz z powstaniem technologii Pixel Shader i Vertex Shader zacz\u0119to wykorzystywa\u0107 procesor graficzny do oblicze\u0144 niezwi\u0105zanych z grafik\u0105, okre\u015blane mianem GPGPU - General-Purpose computing on Graphics Processor Units. \n", "\n", "W 2004, firma Ageia zaprezentowa\u0142a prze\u0142omowy produkt, znany pod nazw\u0105 PhysX a b\u0119d\u0105cy sprz\u0119towym akceleratorem oblicze\u0144 zwi\u0105zanych z symulacjami fizyki w grach komputerowych. Po co fizyka w grach? Rozw\u00f3j grafiki 3D szybko doprowadzi\u0142 do sytuacji w kt\u00f3rej komputer m\u00f3g\u0142 ca\u0142kiem rozs\u0105dnie wyrenderowa\u0107 \u017c\u0105dan\u0105 scen\u0119. Jednak aby to zrobi\u0142 musi wiedzie\u0107 co renderowa\u0107. Po cz\u0119\u015bci informacje o po\u0142o\u017ceniu objekt\u00f3w w grach s\u0105 zapisywane z odgrywanych przez aktor\u00f3w scen lub wprowadzane r\u0119cznie przez animatora. Jednak ma to wiele wad: jest kosztowne, zabiera duzo pami\u0119ci a powtarzalno\u015b\u0107 ruch\u00f3w negatywnie wp\u0142ywa na odbi\u00f3r akcji. Dlatego optymalnym rozwi\u0105zaniem by\u0142o generowanie scen przez pewne algorytmy. I tu przeszkod\u0105 sta\u0142 si\u0119 nasz m\u00f3zg, kt\u00f3ry jest sieci\u0105 neuronow\u0105 doskonale znaj\u0105c\u0105 podstawowe prawa fizyki! Bez trudu potrafimy odgadn\u0105\u0107 jak rusza\u0142 by si\u0119 s\u0142o\u0144, gdyby by\u0142 ze styropianu, znamy prawa odbicia, wiemy jak przep\u0142ywa ciecz itp. Nie \u0142atwo nas oszuka\u0107 - wi\u0119\u0107 programi\u015bci gier komputerowych postanowili nie \"generowa\u0107\", ale (prawie uczciwie) symulowa\u0107 \u015bwiat w grach. Jednak, aby obliczy\u0107 trajektorie na przyk\u0142ad miliona li\u015bci opadaj\u0105cych z drzewa, potrzeba ogromnych mocy obliczeniowych. Z drugiej strony gra wymaga wykonania tych oblicze\u0144 w pewnym ogranicznonym czasie. W\u0142a\u015bnie ta potrzeba, wraz z siedemdziesi\u0119cio-miliardowym biznesem gier, doprowadzi\u0142a opracowania koncepcji sprz\u0119towego akceleratora fizyki- tzw. PPU (Physics Processing Unit).\n", "\n", "Co si\u0119 sta\u0142o, \u017ce firmy Ageia nie ma ju\u017c na runku? Okaza\u0142o si\u0119, \u017ce producenci kart graficznych z procesorami GPU, zauwa\u017cyli, \u017ce jest to ich bran\u017ca. Ponadto, okaza\u0142o si\u0119, \u017ce architektura GPU jest we wielu aspektach zli\u017cona do architektury PPU. GPU wykonuje r\u00f3wnoleg\u0142e obliczenia na \"pixelach\", PPU wykonuje pewne obliczenia na wielu cz\u0105stkach. Firma Nvidia wykona\u0142a w 2006 kilka krok\u00f3w - po pierwsze udost\u0119pni\u0142a standart programowania jednostek GPU zwany CUDA i zach\u0119ci\u0142a \u015bwiat nauki to rozwijania oprogramowania. Po drugie wykupi\u0142a Agei-\u0119 i zaimplementowa\u0142a bibliotek\u0119 PhysX w CUDA, wobec czego jej akceleracja by\u0142a mo\u017cliwa nie tylko z procesorami PPU, ale na ka\u017cdej nowej karcie firmy Nvidia. Praktycznie od momentu pojawienia si\u0119 generacji GeForce serii 8, wszystkie procesory GPU z Nvidii s\u0105 kompatybilne z CUDA i mog\u0105 wykonywa\u0107 r\u00f3wnolegle napisane programy. R\u00f3\u017cnice pomi\u0119dzy procesorami polegaj\u0105 g\u0142\u00f3wnie na ilo\u015bci rdzeni, zwanych \"CUDA -cores\", kt\u00f3rych liczba waha si\u0119 od 16 do 2500! \n", "\n", "Moc obliczeniowa procesor\u00f3w GPU a tak\u017ce, co jest nies\u0142ychanie wa\u017cne - przepustowo\u015b\u0107 pami\u0119ci s\u0105 o wiele wi\u0119ksze od parametr\u00f3w dost\u0119pnych na CPU. Oczywi\u015bcie jest to kosztem ograniczenia mo\u017cliwo\u015bci procesor\u00f3w GPU. Aby zda\u0107 sobie spraw\u0119 z liczb, warto przyjrze\u0107 si\u0119 grafikom opublikowanym przez Nvidi\u0119:\n", "\n", "![Teoretyczna wydajno\u015b\u0107 obliczeniowa](files/figs/floating-point-operations-per-second.png \"Teoretyczna wydajno\u015b\u0107 obliczeniowa\") \n", "\n", "![Przepustowo\u015b\u0107 pami\u0119ci RAM ](files/figs/memory-bandwidth.png \"Przepustowo\u015b\u0107 pami\u0119ci RAM \") \n" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "## Architektura GPU \n", "\n", "\n", "Jak ju\u017c si\u0119 dowiedzieli\u015bmy, nowoczesne procesory graficzne to urz\u0105dzenia obliczeniowe, zdolne do wykonywania trylion\u00f3w operacji zmiennoprzecinkowych na sekund\u0119. \n", " \n", "Na dzie\u0144 dzisiejszy dost\u0119pne urz\u0105dzenia GPU dzia\u0142aj\u0105ce w standardzie CUDA mo\u017cna podzieli\u0107 na trzy generacje o nazwach kodowych: Tesla, Fermi i Kepler. Ka\u017cda z nich oferuje coraz bardziej zaawansowane funkcje i lepsz\u0105 wydajno\u015b\u0107.\n", "\n", "Procesor GPU jest zorganizowany wok\u00f3\u0142 koncepcji multiprocesor\u00f3w strumieniowych (MP).\n", "Takie multiprocesorowy sk\u0142adaj\u0105 si\u0119 z kilku-kilkunastu procesor\u00f3w skalarnych (SP), z kt\u00f3rych ka\u017cdy jest zdolny do wykonywania w\u0105tku w SIMT (Single Instruction, Multiple Threads). Ka\u017cdy MP posiada r\u00f3wnie\u017c ograniczon\u0105 ilo\u015b\u0107 wyspecjalizowanej pami\u0119ci on-chip: zestaw rejestr\u00f3w 32-bitowych, wsp\u00f3lny blok pami\u0119ci i pami\u0119ci podr\u0119cznej L1, cache sta\u0142ych i cache tekstur. Rejestry s\u0105 logicznie lokalne dla procesora skalarnego, ale inne typy pami\u0119ci s\u0105 wsp\u00f3lne dla wszystkich SPs w ka\u017cdym w MP. Umo\u017cliwia to m.in. wymian\u0119 danych pomi\u0119dzy w\u0105tkami operuj\u0105cymi na tym samym MP.\n", "\n", "\n", "![Schemat procesora GPU](files/figs/GF100_03.jpg \"Schemat procesora GPU\") \n", "\n", "## CUDA - jako warstwa abstakcji w dost\u0119pie do GPU\n", "\n", "Procesory graficzne mog\u0105 r\u00f3\u017cni\u0107 si\u0119 ilo\u015bci\u0105 i organizacj\u0105 rdzeni obliczeniowych (CUDA-cores). Implementacja algorytm\u00f3w w r\u00f3wnoleg\u0142y spos\u00f3b powinienna by\u0107 niezale\u017cna od konkrentego sprz\u0119tu na kt\u00f3rym si\u0119 one wykonuj\u0105. W og\u00f3lno\u015bci jest to zagadnienie - jak si\u0119 wydaje - nierozwi\u0105zywalne, ale technologia CUDA uczyni\u0142a olbrzymi krok w kierunku uzyskania niezale\u017cno\u015bci od sprz\u0119tu. \n", "\n", "Po pierwsze programuj\u0105c na GPU tworzymy du\u017co wi\u0119cej w\u0105tk\u00f3w ni\u017c jest fizycznie dost\u0119pnych rdzeni obliczeniowych. System kolejkowania wykona\u0144 zajmuje si\u0119 cz\u0119\u015bciow\u0105 serializacj\u0105 wywo\u0142a\u0144, tym intensywniejsz\u0105 im gorszy sprz\u0119t posiadamy. \n", "\n", "Po drugie dost\u0119p do sprz\u0119tu jest realizowany przez warstwe abstrakcji, zwan\u0105 w\u0142a\u015bnie CUDA, kt\u00f3ra zawiera w sobie informacje o hierarchii pami\u0119ci. Programista, nie przejmuje si\u0119 tym ile rdzeni wykonuje jego program, ale wie, \u017ce pewne grupy w\u0105tk\u00f3w maj\u0105 dost\u0119p do wsp\u00f3lnej szybkiej pami\u0119ci. \n", "\n", "Programy napisane na CUDA dramatycznie r\u00f3\u017cni\u0105 si\u0119 od wielow\u0105tkowych r\u00f3wnoleg\u0142ych odpowiednik\u00f3w napisanych na nawet kilkunastordzeniowe jednostki CPU. Z jednej strony jest to ograniczenie, skutkuj\u0105ce konieczno\u015bci\u0105 przepisania wszystkich algorytm\u00f3w. Jednak poztywnym efektem programowania na takiej architekturze jest ogromna kompatybilno\u015b\u0107 i przeno\u015bno\u015b\u0107 kodu na wsp\u00f3\u0142czesne i z du\u017cym prawdopodobie\u0144stwem przysz\u0142e urz\u0105dzenia.\n", "\n", "\n", "### Hierarchia pamieci\n", "\n", "\n", "By\u0107 mo\u017ce najbardziej istotn\u0105 cech\u0105 architektury CUDA jest hierarchia pami\u0119ci\n", "z r\u00f3\u017cnic\u0105 czas\u00f3w dost\u0119pu o 1-2 rz\u0119dy wielko\u015bci pomi\u0119dzy kolejnymi poziomami. Najwolniejsz\u0105 z punktu widzenia GPU jest pami\u0119\u0107 RAM komputera gospodarza (host). Pami\u0119\u0107 ta jest oddzielona od procesora graficznego magistral\u0105 PCIe z teoretyczn\u0105 maksymaln\u0105 przepustowo\u015bci\u0105 w jednym kierunku 16 GB/s (PCI Express 3.0,x16).\n", "\n", "Nast\u0119pn\u0105 w kolejno\u015bci jest globalna pami\u0119\u0107 urz\u0105dzenia GPU, kt\u00f3ry jest obecnie ograniczony do kilku gigabajt\u00f3w (najnowsze karty maj\u0105 juz 12GB) i o szeroko\u015bci pasma oko\u0142o 100-200 ~ Gb/s. Jest to bardzo du\u017ca warto\u015b\u0107, jednak dost\u0119p do globalnej pami\u0119ci jest operacj\u0105 wysokiej latencji wynosz\u0105c\u0105 kilkaset cykli zegara GPU.\n", "\n", "Najszybszym rodzajem pami\u0119ci dost\u0119pnym obecnie na GPU jest pami\u0119\u0107 wsp\u00f3\u0142dzielona (shared memory) znajduj\u0105ca si\u0119 bezpo\u015brednio na multiprocesorze. Jest ona obecnie ograniczona do 48 kB, ale ma przepustowo\u015b\u0107 oko\u0142o ~1,3 TB/s. Co ciekawsze, latenacja w dost\u0119pie do tej pami\u0119ci jest bardzo ma\u0142a - podobna jak dost\u0119p do rejestr\u00f3w jednostek SP.\n", "\n", "Powy\u017cszy opis \u0142atwo sugeruje strategi\u0119 pisania wydajnych program\u00f3w na CUDA, a kt\u00f3re mo\u017cna stre\u015bci\u0107 jako: przenie\u015b\u0107 jak najwi\u0119cej danych, jak to mo\u017cliwe\n", "do najszybszych rodzaju dost\u0119pnej pami\u0119ci i przechowywa\u0107 je tam tak d\u0142ugo, jak to mo\u017cliwe, przy jednoczesnej minimalizacji dost\u0119pu do wolniejszych rodzaj\u00f3w pami\u0119ci. Dodatkowo, je\u015bli dost\u0119p do wolniejszej pami\u0119ci jest konieczny, mo\u017cna pr\u00f3bowa\u0107 wykorzysta\u0107 \"wolny\" czas na wykonanie operacji arytmetycznych na pozosta\u0142ych w\u0105tkach.\n", "\n", "\n", "\n", "### W\u0105tki\n", "\n", "\n", "W programowaniu na CUDA czy OpenCL wykorzystujemy w\u0105tki. Co jest niespotykane w pisaniu kodu HPC na zwyk\u0142e procesory, w\u0105tk\u00f3w powinno si\u0119 w\u0142\u0105cza\u0107 o wiele wi\u0119cej od dost\u0119pnych fizycznych rdzeni procesora. W praktyce programowania na CPU znany jest fakt, \u017ce nadmierne rozmno\u017cenie w\u0105tk\u00f3w zazwyczaj spowalnia wykonanie programu, ze wzgl\u0119du na \"context switching\", kt\u00f3ry mo\u017ce prowadzi\u0107 do nieefektywnego wykorzystania pamieci cache jednostki centralnej a ponadto prowadzi w nieunikniony spos\u00f3b do zmniejszenia dost\u0119pnej pami\u0119ci RAM. W programowaniu w standardzie CUDA w\u0105tki s\u0105 lekkie, ich prze\u0142\u0105czanie nie zabiera wi\u0119cej ni\u017c pojedy\u0144czych cykli procesora. Co wi\u0119cej, nadmiarowe w\u0105tki potrafi\u0105 mie\u0107 pozytywny efekt na wydajno\u015b\u0107, gdy\u017c mog\u0105 przes\u0142oni\u0107 latencj\u0119 dost\u0119pu do pami\u0119ci g\u0142\u00f3wnej.\n", "\n", "Z punktu widzenia u\u017cytkownika, programy CUDA s\u0105 zorganizowane w j\u0105drach. J\u0105dro\n", "Jest to funkcja, kt\u00f3ra jest wykonywana wielokrotnie, jednocze\u015bnie na r\u00f3\u017cnych multiprocesorach. Ka\u017cda instancja j\u0105dra zwana jest w\u0105tkiem i jest przydzielana do pojedy\u0144czego procesora skalarnego (SP). W\u0105tki s\u0105 grupowane w jedno-, dwu - lub tr\u00f3jwymiarowe bloki przypisanych do multiprocesor\u00f3w jeden do jednego - czyli mamy gwarancj\u0119, \u017ce w ramach jednego bloku jeste\u015bmy na tym samym MP.\n", "\n", "### J\u0119zyk programowania i struktura kursu\n", "\n", "*Kernele* pisze si\u0119 w okrojonym dialekcie C/C++, zwanym CUDA-C. Referencyjnym i stosunkowo przyst\u0119pnie napisanym dokumentem jest CUDA C Programming Guide, kt\u00f3ry jest dost\u0119pny pod adresem: http://docs.nvidia.com/cuda/cuda-c-programming-guide/\n", "W tym kursie nie b\u0119dziemy systematycznie analizowa\u0107 wszystkich element\u00f3w CUDA-C. Post\u0105pimy inaczej. Poka\u017cemy na przyk\u0142adach typowe zastosowania programowania GPU, kt\u00f3re s\u0105 niezwykle przydatne w fizyce. Rozwi\u0105zanie w\u0142asnego problemu, proponujemy rozpocz\u0105\u0107 od znalezienia stosownego przyk\u0142adu i pr\u00f3bie samodzielnego dostosowania jego kodu. Jest to rozwi\u0105zanie typu \"crash course\", kt\u00f3re nie zast\u0105pi systemayucznego kursu. Jednak w du\u017cej wi\u0119kszo\u015bci przypadk\u00f3w, do\u015bwiadczenie nabyte podczas zabawy z zamieszczonymi przyk\u0142adami umo\u017cliwi optymalne rozwi\u0105zanie napotkanego problemu.\n", "\n", "\n", "\n", "\n", "\n", "\n", "\n" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "# Spos\u00f3b pracy\n", "\n", "## Interaktywno\u015b\u0107 - u\u017cywamy interfejsu pyCUDA\n", "\n", "\n", "W tych materia\u0142ach zostanie zaprezentowane podej\u015bcie, kt\u00f3re pozwoli maksymalnie upro\u015bci\u0107 drog\u0119 do efektywnej pracy w CUDA. W tym celu zostanie zastosowany pakiet pyCUDA, kt\u00f3ry umo\u017cliwi prac\u0119 interaktywn\u0105 z urz\u0105dzeniem CUDA bez kompromisu je\u015bli chodzi o wydajno\u015b\u0107. W pyCUDA, spos\u00f3b pracy sprowadza si\u0119 do napisania j\u0105dra obliczeniowego w j\u0119zyku C z rozszerzeniami CUDA. Wywo\u0142anie j\u0105dra, jego kompilacja oraz inspekcja danych, w tym transfer do i z urzadzenia, wykonuje si\u0119 w wygodnym j\u0119zyku python." ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "## Pierwszy program pyCUDA\n", "\n", "\n", "Przedstawimy teraz pierwszy program, kt\u00f3ry b\u0119dzie napisany w CUDA, z u\u017cyciem pyCUDA. Przyk\u0142ad ten b\u0119dzie bardzo prosty: mno\u017cenie wektora przez liczb\u0119. Na tym prostym przyk\u0142adzie poznamy spos\u00f3b pracy, kt\u00f3ry potem b\u0119dziemy mogli u\u017cy\u0107 jako szablonu do tworzenia bardziej zaawansowanych program\u00f3w. \n", "\n", "### Inicjalizacja\n", "\n", "Najprostszy spos\u00f3b inicjalizacji urz\u0105dzenia GPU tak by m\u00f3c go dalej u\u017cywa\u0107 wygl\u0105da nast\u0119puj\u0105co:\n", "\n" ] }, { "cell_type": "code", "collapsed": false, "input": [ "import numpy as np\n", "import pycuda.driver as cuda\n", "import pycuda.autoinit\n", "from pycuda.compiler import SourceModule" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 12 }, { "cell_type": "markdown", "metadata": {}, "source": [ "Opr\u00f3cz modu\u0142ow z pakietu pyCUDA, importujemy r\u00f3wnie\u017c numpy. Jest to niezwykle istotne, bowiem tablice numpy to podstawowa struktura danych kt\u00f3r\u0105 b\u0119dziemy przesy\u0142a\u0107 na urz\u0105dzenie GPU i z powrotem.\n", "\n", "Utw\u00f3rzmy wektor danych, na przyk\u0142ad wykorzystuj\u0105c `linspace`:" ] }, { "cell_type": "code", "collapsed": false, "input": [ "a = np.linspace(1,16,16).astype(np.float32)\n", "a" ], "language": "python", "metadata": {}, "outputs": [ { "metadata": {}, "output_type": "pyout", "prompt_number": 13, "text": [ "array([ 1., 2., 3., 4., 5., 6., 7., 8., 9., 10., 11.,\n", " 12., 13., 14., 15., 16.], dtype=float32)" ] } ], "prompt_number": 13 }, { "cell_type": "markdown", "metadata": {}, "source": [ "Celem naszego programu b\u0119dzie pomno\u017cenie wszystkich element\u00f3w tej tablicy przez zadan\u0105 liczb\u0119. \n", "\n", "Tablica `a` znajduje si\u0119 w pami\u0119ci systemu gospodarza. Aby wykona\u0107 operacje na niej z pomoc\u0105 GPU musimy zaalokowa\u0107 pami\u0119\u0107 na urz\u0105dzeniu GPU:" ] }, { "cell_type": "code", "collapsed": false, "input": [ "a_gpu = cuda.mem_alloc(a.nbytes)" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 14 }, { "cell_type": "markdown", "metadata": {}, "source": [ "a nast\u0119pnie skopiowa\u0107 dane do GPU:" ] }, { "cell_type": "code", "collapsed": false, "input": [ "cuda.memcpy_htod(a_gpu, a)" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 15 }, { "cell_type": "markdown", "metadata": {}, "source": [ "Mamy teraz na GPU zainicjalizowan\u0105 tablice, do ktorej mo\u017cemy si\u0119 z poziomu hosta (gospodarza) odwo\u0142ywa\u0107 przez `a_gpu`. Typow\u0105 praktyk\u0105 w programowaniu na GPU jest posiadania dw\u00f3ch kopii danych - jednej w pami\u0119ci dost\u0119pnej dla CPU a drug\u0105 w pami\u0119ci urz\u0105dzenia GPU. Warto jeszcze zaznaczy\u0107, \u017ce nasza tablica `a_gpu` znajduje si\u0119 w pami\u0119ci globlanej GPU.\n", "\n", "Teraz musimy napisa\u0107 j\u0105dro, kt\u00f3re b\u0119dziemy wywo\u0142ywa\u0107 na urz\u0105dzeniu na danych w pami\u0119ci globalnej GPU. J\u0105dro jest napisane w zmodyfikowanej wersji j\u0119zyka C, zwanej CUDA-C. \n", "\n", "J\u0105dro, kt\u00f3re b\u0119dzie mno\u017cy\u0142o wszystkie elementy tablicy przez 2 ma nast\u0119puj\u0105c\u0105 posta\u0107:\n", "\n", " __global__ void dubluj(float *a)\n", " {\n", " int idx = threadIdx.x;\n", " a[idx] *= 2;\n", " }\n", " \n", "* J\u0105dro wygl\u0105da jak zwyk\u0142a funkcja w C.\n", "* Deklaracja `__global__` jest rozszerzeniem CUDA C.\n", "* J\u0105dro b\u0119dzie uruchomione w tylu kopiach ile jest element\u00f3w wektora `a_gpu` a ka\u017cda b\u0119dzie mno\u017cy\u0142a jeden element wektora.\n", "* `threadIdx` jest struktur\u0105 okre\u015blaj\u0105c\u0105 numer w\u0105tku wewn\u0105trz bloku, kt\u00f3ry wykonuje dan\u0105 kopi\u0119. Opr\u00f3cz niego wa\u017cny jest te\u017c `blockIdx`, ale w tym przypadku odpalamy tylko jeden blok wi\u0119c nie jest nam potrzebny. Indeksy w\u0105tku, czy te\u017c wewn\u0105trz bloku czy numer bloku, to jedyny spos\u00f3b rozr\u00f3\u017cnienia w\u0105tk\u00f3w! Zmienna `idx` jest liniowym wska\u017anikiem, kt\u00f3ry b\u0119dzie si\u0119 zmienia\u0142 od `0` do `N-1`.\n", "\n", "W pyCUDA \u017ar\u00f3d\u0142o modu\u0142u podajemy jako `string` funkcji `SourceModule` w poni\u017cszy spos\u00f3b:\n" ] }, { "cell_type": "code", "collapsed": false, "input": [ "mod = SourceModule(\"\"\"\n", " __global__ void dubluj(float *a)\n", " {\n", " int idx = threadIdx.x;\n", " a[idx] *= 2;\n", " \n", " }\n", " \"\"\")" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 16 }, { "cell_type": "markdown", "metadata": {}, "source": [ "Zmienna `mod` jest teraz obiektem, kt\u00f3ry ma m.in. funkcj\u0119 `mod.get_function(\"nazwa\")` zwracaj\u0105c\u0105 funkcj\u0119 pythonow\u0105, kt\u00f3rej wywo\u0142anie uruchomi odpowiednie j\u0105dro ze \u017ar\u00f3d\u0142a uprzednio podanego do `SourceModule`. Sprawd\u017amy:" ] }, { "cell_type": "code", "collapsed": false, "input": [ "func = mod.get_function(\"dubluj\")" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 17 }, { "cell_type": "markdown", "metadata": {}, "source": [ "Wywo\u0142ajmy teraz j\u0105dro!" ] }, { "cell_type": "code", "collapsed": false, "input": [ "func(a_gpu, block=(16,1,1))" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 18 }, { "cell_type": "markdown", "metadata": {}, "source": [ "W tablicy `a_gpu` powinni\u015bmy mie\u0107 teraz " ] }, { "cell_type": "code", "collapsed": false, "input": [ "print a\n", "cuda.memcpy_dtoh(a, a_gpu)\n", "print a" ], "language": "python", "metadata": {}, "outputs": [ { "output_type": "stream", "stream": "stdout", "text": [ "[ 1. 2. 3. 4. 5. 6. 7. 8. 9. 10. 11. 12. 13. 14. 15.\n", " 16.]\n", "[ 2. 4. 6. 8. 10. 12. 14. 16. 18. 20. 22. 24. 26. 28. 30.\n", " 32.]\n" ] } ], "prompt_number": 19 }, { "cell_type": "markdown", "metadata": {}, "source": [ "Zauwa\u017cmy, \u017ce w tym programie r\u00f3wnie dobrze mo\u017cna by wykona\u0107 nast\u0119puj\u0105ce wywo\u0142anie:\n", " \n", " func(a_gpu, block=(4,4,1))\n", " \n", "ale wtedy musimy zmodyfikowa\u0107 odpowiednio obliczanie wska\u017anika `idx` w \u017ar\u00f3dle j\u0105dra, na:\n", " int idx = threadIdx.x + threadIdx.y*4;\n", "\n", "#### \u0106wiczenie: \n", "\n", "Sprawd\u017a to!" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "## Udogodnienia w pyCUDA\n", "\n", "Powy\u017cszy spos\u00f3b budowania programu w CUDA jest bardzo podobny gdyby\u015bmy u\u017cywali jedynie kompilatora oraz j\u0119zyka C/C++ zamiast pythona z pyCUDA. \n", "\n", "### InOut\n", "\n", "Warto zaznajomi\u0107 si\u0119 z kilkoma udogodnieniami, kt\u00f3re mamy wbudowane w pyCUDA. Pierwszym z nich jest zautomatyzowanie procesu alokacji i kopiowania danych. Je\u015bli chcemy wykona\u0107 po kolei: alokacje wektora na GPU, transfer danych z wektora numpy, wywo\u0142anie j\u0105dra oraz nadpisanie wyj\u015bciowego wektora numpy wynikiem z GPU, to mo\u017cemy u\u017cy\u0107 funkcji `InOut`: " ] }, { "cell_type": "code", "collapsed": false, "input": [ "func(cuda.InOut(a), block=(4, 4, 1))\n", "print a" ], "language": "python", "metadata": {}, "outputs": [ { "output_type": "stream", "stream": "stdout", "text": [ "[ 4. 8. 12. 16. 10. 12. 14. 16. 18. 20. 22. 24. 26. 28. 30.\n", " 32.]\n" ] } ], "prompt_number": 20 }, { "cell_type": "markdown", "metadata": {}, "source": [ "### `gpuarray`\n", "\n", "Modu\u0142 `gpuarray` umo\u017cliwia bardzo zwarty zapis operacji na wektorach na GPU. Kluczow\u0105 komend\u0105 jest `gpuarray.to_gpu(..`, kt\u00f3ra pozwala skopiowa\u0107 wektor numpy do urz\u0105dzenia GPU, zwracaj\u0105c uchwyt do danych na GPU. Pewne operacje mo\u017cna wykona\u0107 automatycznie na GPU po prostu wpisuj\u0105c wz\u00f3r arytmetyczny." ] }, { "cell_type": "code", "collapsed": false, "input": [ "import pycuda.gpuarray as gpuarray\n", "\n", "a = np.linspace(1,16,16).astype(np.float32)\n", "a_gpu = gpuarray.to_gpu(a.astype(numpy.float32))\n", "a_doubled = (a_gpu*2).get()\n", "\n", "print a_gpu\n", "print a_doubled" ], "language": "python", "metadata": {}, "outputs": [ { "output_type": "stream", "stream": "stdout", "text": [ "[ 1. 2. 3. 4. 5. 6. 7. 8. 9. 10. 11. 12. 13. 14. 15.\n", " 16.]\n", "[ 2. 4. 6. 8. 10. 12. 14. 16. 18. 20. 22. 24. 26. 28. 30.\n", " 32.]\n" ] } ], "prompt_number": 21 }, { "cell_type": "markdown", "metadata": {}, "source": [ "*Uwaga printf - dzia\u0142a tylko z konsoli!*" ] }, { "cell_type": "code", "collapsed": false, "input": [ "import pycuda.driver as cuda\n", "import pycuda.autoinit\n", "from pycuda.compiler import SourceModule\n", "\n", "mod = SourceModule(\"\"\"\n", " #include \n", "\n", " __global__ void printf_test()\n", " {\n", " printf(\"GONZO %d.%d.%d\\\\n\", threadIdx.x, threadIdx.y, threadIdx.z);\n", " }\n", " \"\"\")\n", "\n", "func = mod.get_function(\"printf_test\")\n", "func(block=(2,2,1))" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 22 }, { "cell_type": "code", "collapsed": false, "input": [], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 11 }, { "cell_type": "markdown", "metadata": {}, "source": [ "## Grids & Blocks \n", "W CUDA w\u0105tki mo\u017cna uruchamia\u0107 w blokach (blocks). Nale\u017cy zapami\u0119ta\u0107 regu\u0142e:\n", "\n", "**wszystkie w\u0105tki na jednym bloku zawsze b\u0119d\u0105 uruchamiane na tym samym multiprocesorze**\n", " \n", "Co za tym idzie, b\u0119d\u0105 mia\u0142y do dyspozycji t\u0105 sam\u0105 pami\u0119c typu \"shared memory\". Maksymalna liczba w\u0105tk\u00f3w w jednym bloku jest z regu\u0142y ograniczona do 512 lub 1024 w zale\u017cno\u015bci od typu urz\u0105dzenia GPU. Warto zauwazy\u0107, \u017ce jest to liczba zdecydowanie wi\u0119ksza od liczby procesor\u00f3w skalarnych w na jednym multiprocesorze. \n", "\n", "![Organizacja numeracji w\u0105tk\u00f3w](files/figs/grid_blocks.jpg) \n", "\n", "Je\u017celi chcemy odpalic np. milion w\u0105tk\u00f3w, to musimy wykorzysta\u0107 tzw. grid blok\u00f3w. Czyli odpalaj\u0105c j\u0105dro obliczeniowe specyfikujemy ile potrzeujemy w\u0105tk\u00f3w i ile jaki chcemy mie\u0107 rozmiar bloku. Chc\u0105c mie\u0107 np. $128\\times 64$ w\u0105tki, rozs\u0105dnym jest odpalenie j\u0105dra z rozmiarem gridu $128$ i rozmiarem bloku $64$. \n", "\n", "Ponadto, poniewa\u017c cz\u0119sto operujemy siatkami reprezentujemy dwu i tr\u00f3j-wymiarowe pola, CUDA ma wbudowany mechanizm, kt\u00f3ry u\u0142atwia pos\u0142ugiwaniem si\u0119 dwu i tr\u00f3j-wymiarowymi siatkami. \n", "\n", "Poni\u017cszy przyk\u0142ad pokazuje jak zbudowa\u0107 grid sk\u0142adaj\u0105cy si\u0119 z dw\u00f3ch blok\u00f3w, po pi\u0119c w\u0105tk\u00f3w.\n", "\n", "#### \u0106wiczenie. \n", "Zast\u0105p w kodzie `if(threadIdx.x==1)` przez:\n", "\n", "* `if(blockIdx.x==1)`\n", "* `if(idx==1)`\n", "\n", "i sprawd\u017a dzia\u0142anie.\n", "\n", "\n" ] }, { "cell_type": "code", "collapsed": false, "input": [ "import pycuda.driver as cuda\n", "import pycuda.autoinit\n", "from pycuda.compiler import SourceModule\n", "\n", "mod = SourceModule(\"\"\"\n", " __global__ void kernel(float *a)\n", " {\n", " int idx = threadIdx.x + blockDim.x*blockIdx.x;\n", " \n", " if(threadIdx.x==0)\n", " a[idx] = 1.0f;\n", " }\n", " \"\"\")\n", "\n", "a = np.zeros(10).astype(np.float32)\n", "func = mod.get_function(\"kernel\")\n", "print np.linspace(0,9,10)\n", "print \"----------------\"\n", "print a\n", "func(cuda.InOut(a),block=(5,1,1),grid=(2,1,1))\n", "print a" ], "language": "python", "metadata": {}, "outputs": [ { "output_type": "stream", "stream": "stdout", "text": [ "[ 0. 1. 2. 3. 4. 5. 6. 7. 8. 9.]\n", "----------------\n", "[ 0. 0. 0. 0. 0. 0. 0. 0. 0. 0.]\n", "[ 1. 0. 0. 0. 0. 1. 0. 0. 0. 0.]\n" ] } ], "prompt_number": 140 }, { "cell_type": "markdown", "metadata": {}, "source": [ "## Co dalej?\n", "\n", "Analizuj\u0105c kod powy\u017cszych przyk\u0142ad\u00f3w, jeste\u015bmy w stanie samodzielnie napisa\u0107 elementarne j\u0105dro, je wywo\u0142a\u0107 i odczyta\u0107 dane wynikowe. \n", "\n", "Proponujemy teraz przej\u015b\u0107 do przyk\u0142ad\u00f3w, kt\u00f3re umo\u017cliwi\u0105 zapoznanie si\u0119 z ogromnymi mo\u017cliwo\u015bciamy procesora GPU w zastosowaniach typowych dla nauk \u015bcis\u0142ych.\n", "\n", "\n" ] } ], "metadata": {} } ] }