Fakty, mity i CUDA

@M@Ti: Radeony są dobre dla graczy, ale sobie słabo radzą z OpenGL (mają bardzo słabe sterowniki do OpenGL - nawet nie chodzi o to, że działają dużo wolniej, a o to że czasami nie działa wcale (na szczęście ostatnio się to nie za często zdaża, ale w sterownikach z serii 8.x było to bardzo częste) - na wszystkich systemach operacyjnych w przeciwieństwie do NVidii)... jak pewnie dobrze wiesz praktycznie 100% programów graficznych działa na OpenGL... dodatkowo nawet stara karta graficzna typu 8600 potrafi przyspieszyć znacznie rendering w stosunku do CPU na rendererach korzystających z Cuda... do grania radeony są alternatywą, ale do grafiki jest to słaby wybór.

re: skoti

tu wiekszosc gosci to maxowcy, wiec dx to chleb powszedni. a o tej magicznej

8600 ktora potrafi renderowac to nawet sie nie bede wypowiadal. pokaz mi

renderer na gpu ktory obsluguje cala specyfikacje ".mi" a sie z tym zgodze.

co do kart graficznych to 4770 coraz trudniej dostac, za to pojawila sie tania

5670 (jakies 330 zl) z calkiem podobnymi osiagami. zapowiada sie tez 5830

wypelniajaca luke miedzy 5770 a 5870. a na nowinki od nvidia trzeba czekac do

marca (usa) / maja (reszta swiata :) ).

re: wallop

co do tej plytki 2x1366 to da sie zamontowac dwa i7... ale: i7 nie ma dodatkowego

linka qpi (tylko jeden do mostka) na bezposrednie polaczenie miedzy procesorami

(jak to ma miejsce w xeonach DP), tym samym komunikacja miedzy procesorami

odbywa sie przez mostek polnocny (jak to mialo miejsce w starszych generacjach

xeonow - bez kontrolera pamieci i interconnectu) - co rzecz jasna zabija wydajnosc

i nie jest zapewne zgodne z przykazem intela. jesli sie nie myle chip to tylersburg

wiec sie to da zrobic, choc pozostaje to w kategorii domyslow i plotek :)

PS: aha, na CES'ie wiesc gminna niosla ze nie tylko evga szykuje taka plytke

PS2: aha2, i7 980x ma miec dwa linki qpi, wiec moze bedzie dalo sie go uzyc na

platformie dwuprocesorowej - odpowiednik skulltrail 2

@dokturpotfor: w max też polecanym rendererem jest OpenGL. 8600 jest bardzo słabą kartą, ale jednak potrafi więcej niż CPU. Co do plików mental ray'a to zapewne pierwszym rendererem gpu go obsługującym będzie iray korzystający z SceneX i Cuda (oba to rozwiązania nvidii) twórców mental raya (rendermany gpgpu już są na rynku od bardzo dawna).

Co do kart graficznych z serii 56xx i 57xx to nie za bardzo warto, bo w obliczeniach gdy będzie potrzeba użyć liczb podwójnej precyzji nie będzie działać GPGPU, a do samego renderingu scen przy modelowaniu lepiej wybrać starsze tańsze rozwiązania.

dodatkowo nawet stara karta graficzna typu 8600 potrafi przyspieszyć znacznie rendering w stosunku do

Co do plików mental ray'a to zapewne pierwszym rendererem gpu go obsługującym będzie iray korzystający z SceneX i Cuda (oba to rozwiązania nvidii) twórców mental raya (rendermany gpgpu już są na rynku od bardzo dawna).

ble ble ble, "potrafi" / "bedzie (zapewne)".

"iray uses a highly optimized BSDF and EDF shading framework

(essentially, simulating the physical behavior of real-world material and

light source), as opposed to a complex collection of programmable shaders

applied through renderer specific object and scene parameters."

no pewno... po co nam "complex collection of programmable shaders", choc

moze arch materials wystarcza

"For technical audiences…:

Interpolation techniques, which trade final quality, predictability, and

simplicity of scene specification for performance, form the core of most

current global illumination renderers. Unlike them, iray rendering is based

on deterministic and consistent global illumination simulation algorithms

that converge without introducing persistent approximation artifacts."

he he, nie ma to jak sprzedac ograniczenie technologiczne i sprzetowe jako

zarabiscie porzadana zalete softu. Poza tym po demie na siggraphie gdzie

15 quadro fx5800 przez kilka minut meczylo jakis obrazek troche sie

zmartwilem ze jest to nadal w takich powijakach. Oczywiscie mental images

triumfalnie oglosilo ze na cpu liczyloby sie to 10 godzin.... tylko ze nie

dodali ze: 'liczyloby sie na cpu przez 10 godzin PRZY UZYCIU TEGO SAMEGO

ALGORYTMU'... czyli 'dla laikow technicznych' iray na gpu uzywa

nachamowego path tracera ktory jest wolny jak cholera, ale za to ma

konstrukcje cepa. na 15 quadro fx5800 scenka liczy sie 3-4 minuty. na

cpu ten sam algorytm dziala odpowiednio wolno i liczy sie kosmicznie dlugo.

i porownuje sie czasy dzialania tego samego algorytmu na cpu i gpu zamiast

zadac sobie pytanie, jak szybko uzywajac wszystkich dobrodziejstw cpu,

wyprodukowac obraz o takiej samej jakosci jak ten wyprodukowany na gpu

irayem. odpowiedz: "duzo szybciej niz uzywajac iraya na cpu" - bo mamy

photon mapping, bo mamy irradiance cache, bo mamy drzewa bsp itd...

zaproponuje przypowiesc:

to tak jakby posadzic na tandemie czterech kuternogow i kazac im sie scigac

z jednym facetem na dwoch nogach. Tylko ze zawiazmy temu zdrowemu

jedna noge, "zeby bylo sprawiedliwie". Wynik jest 4 do 1 dla czterech

kuternogow. jak ten z dwoma nogami moze uzywac ich obu to czterej kikuci

nadal wygrywaja, ale stosunek jest tylko 2:1

bo w obliczeniach gdy będzie potrzeba użyć liczb podwójnej precyzji nie będzie działać GPGPU

??? a dlaczego mialoby nie dzialac? to tak jakby podwojna precyzja pojawila

sie wczoraj... jesli sie nie myle to 'korona' najwydajniejszej karty w doublach

nadal nalezy do radeona 5970

edit: btw, na 8600 zdaje sie double nie dzialaja (ani tez na 9xxx), yep sprawdzone

tylko CUDA na gtx260/275/280/285/295 i teslach

8600 jest bardzo słabą kartą, ale jednak potrafi więcej niż CPU

niech sprobuje wywolac funkcje rekurencyjnie... gpu potrafi glownie dodawac

i mnozyc, czytac z pamieci i filtrowac tekstury... za to robi to bardzo szybko

ech... ciezko walczyc z marketingowymi machinami wielkich firm...

no pewno... po co nam "complex collection of programmable shaders", choc

moze arch materials wystarcza

O czyli jednak się nie postarali - ale jak chcesz programowalne shadery to możesz użyć jednego z rendermanów w Cuda (chociażby nVidia Gelato + rozszerzenie robiące z niego rendermana - ale łykać plików *.mi też nie będzie).

...Oczywiscie mental images

triumfalnie oglosilo ze na cpu liczyloby sie to 10 godzin.... tylko ze nie

dodali ze: 'liczyloby sie na cpu przez 10 godzin PRZY UZYCIU TEGO SAMEGO

ALGORYTMU'... czyli 'dla laikow technicznych' iray na gpu uzywa

nachamowego path tracera ktory jest wolny jak cholera, ale za to ma

konstrukcje cepa. na 15 quadro fx5800 scenka liczy sie 3-4 minuty. na

cpu ten sam algorytm dziala odpowiednio wolno i liczy sie kosmicznie dlugo.

i porownuje sie czasy dzialania tego samego algorytmu na cpu i gpu zamiast

zadac sobie pytanie, jak szybko uzywajac wszystkich dobrodziejstw cpu,

wyprodukowac obraz o takiej samej jakosci jak ten wyprodukowany na gpu

irayem. odpowiedz: "duzo szybciej niz uzywajac iraya na cpu" - bo mamy

photon mapping, bo mamy irradiance cache, bo mamy drzewa bsp itd...

Nie wiem jakich algorytmów używa dokładnie iray ale nie wierzę że to jest zwykły path tracer (i czy to jest ogólnie zwykły pathtracer a nie np. Metropolis light transport (zmodyfikowany bi-pathtracer, który jest faktycznie bardzo wolny - Maxwell, Indigo) bez żadnych optymalizacji. Co do photon mappingu (który jest implementowany bez przeszków nawet w RT w shaderach HLSL/GLSL), irradiance cache to jak najbardziej są możliwe na gpu... tylko po prostu są za szybkie (a po co używać tych metod skoro można w nie dużym czasie robić np. MLT które daje nieporównywalnie lepsze rezultaty, jednak moc CPU nie pozwalała do tej pory na szersze stosowanie), drzewa bsp, octree jak i kd-trees można zaimplementować na GPU... tylko po co się cofać i od razu implementuje się drzewa które są dziś stosowane w renderingu, a nie x lat temu (np. bvh dzięki którym blender 2.5 dostał speed up 7x jest implementowany nawet w raczkujących pathtracerach MLT na GPU jak SmallLuxGPU).

zaproponuje przypowiesc

Proponuj, ale następnym razem takie które mają co kolwiek wspólnego z tematem

??? a dlaczego mialoby nie dzialac? to tak jakby podwojna precyzja pojawila

sie wczoraj... jesli sie nie myle to 'korona' najwydajniejszej karty w doublach

nadal nalezy do radeona 5970

Dlatego miałoby nie działać bo nie ma jednostek obsługujących podwójną precyzję, bo amd je obcieła w tych seriach, które podałem http://www.geeks3d.com/public/jegx/200910/amd_opencl_supported_devices.jpg

edit: btw, na 8600 zdaje sie double nie dzialaja (ani tez na 9xxx), yep sprawdzone

tylko CUDA na gtx260/275/280/285/295 i teslach

Tak jednak plusem 8600 (jednym z niewielu) jest to że jest stara i kosztuje grosze - a jak już kupować coś nowego to lepiej nie produkt który może się możliwościami mierzyć z kartami pokroju 8600 (poza wydajnością) - więc jak już amd to 5800+ lub nVidia 200+ (a najlepiej poczekać do fermi ;p).

niech sprobuje wywolac funkcje rekurencyjnie... gpu potrafi glownie dodawac

i mnozyc, czytac z pamieci i filtrowac tekstury... za to robi to bardzo szybko

W CUDA da się zrobić rekurencję trochę na około bo z użyciem szablonów, ale się da (z tego między innymi powodu renderery używają CUDA, a nie OpenCL który działa również na AMD). Rekurencja i wskaźniki na funkcję w Cuda/OpenCL wejdą dopiero za chwilę razem z fermi (który będzie potrafił wszystko co CPU, a nawet znacznie więcej i będzie potrafił wykonać kod pisany w C++).

drzewa bsp, octree jak i kd-trees można zaimplementować na GPU... tylko po co się cofać i od razu implementuje się drzewa które są dziś stosowane w renderingu, a nie x lat temu (np. bvh dzięki którym blender 2.5 dostał speed up 7x jest implementowany nawet w raczkujących pathtracerach MLT na GPU jak SmallLuxGPU).

stary... implementacje wskaznikowych czy kursorowych struktur danych jak

drzewa, na gpu to sa dalej akademickie rozwazania, ktore maja sie do

architektur SIMD jak piesc do nosa. tak da sie je zaklepac ale ilosc

hakowania aby obejsc ograniczenia architektury skutecznie jak do tej pory

odbieraja komukolwiek ochote by tego stosowac.

Nie wiem co ma blender ale drzewo bounding volume hierarchy, (jak sie myle

to mnie popraw) ktore jako drzewo licznikowe, znane jest chyba od

paleozoiku, przytoczylem bsp jako przyklad jakich wiele.

Powiedzialem path tracer bo nie widze innej mozliwosci. czy to bedzie

2^n czy 13^n i tak jest to strasznie wolny algorytm. monte carlo czy

nie zlozonosc obliczeniowa jest gigantyczna, o stala sie nie kloce.

inny algorytm ktory widze to by byla kratka i iteracyjnie policzony transfer

swiatla (jak w demie GI cryengine 3), ale ze wzgledow pamieciowych w

jakosci lepszej niz przyblizone gi w realtimie wydaje mi sie to niemozliwe.

Proponuj, ale następnym razem takie które mają co kolwiek wspólnego z tematem

no niestety przypowiesc ma cechy metaforyczne wymagajace skupienia od

czytelnika

Dlatego miałoby nie działać bo nie ma jednostek obsługujących podwójną precyzję, bo amd je obcieła w tych seriach, które podałem http://www.geeks3d.com/public/jegx/200910/amd_opencl_supported_devices.jpg

 

Tak jednak plusem 8600 (jednym z niewielu) jest to że jest stara i kosztuje grosze - a jak już kupować coś nowego to lepiej nie produkt który może się możliwościami mierzyć z kartami pokroju 8600 (poza wydajnością) - więc jak już amd to 5800+ lub nVidia 200+ (a najlepiej poczekać do fermi ;p).

ok tu zwracam honor nie zauwazylem ze dla doubli jest support dla 4800+, 5800+

W CUDA da się zrobić rekurencję trochę na około bo z użyciem szablonów, ale się da (z tego między innymi powodu renderery używają CUDA, a nie OpenCL który działa również na AMD). Rekurencja i wskaźniki na funkcję w Cuda/OpenCL wejdą dopiero za chwilę razem z fermi (który będzie potrafił wszystko co CPU, a nawet znacznie więcej i będzie potrafił wykonać kod pisany w C++).

uzycie szablonow to zabieg dziejacy sie podczas kompilacji a nie podczas

uruchomienia programu. po co mi funkcja ktorej zagniezdzenie rekursji jest

ustalane podczas kompilacji. kolejna przypowiesc: dla parametru 5 funkcja

ABC wymaga 3 zagniezdzen rekurencyjnych, ale dla pewnosci zawsze

bedziemy wykonywac 10 tak na wszelki wypadek mimo ze punkt staly funkcji

osiagniemy juz w 3 zagniezdzeniu.

Tu dochodzimy do kolejnej sprawy, i nie zmieni tego marketing... programy

na gpu tez sa kompilowane (olaboga!)... to 'wykonywanie kodu pisanego w

c++' o ktorym wspominasz to wierutna bzdura. kod jest kompilowany w

sterowniku i na gpu trafia juz jako assembler gpu. to cale zamieszanie z

"wsparciem natywnym c++" wiaze sie jedynie z unifikacja adresowania

pamieci, nie wiem czy wiesz ale i kod C i C++ nim jakikolwiek procesor wykona

program zamienia sie na taki sam kod maszynowy, ba nawet da sie w latwy

sposob napisac funkcjonalnosc obiektowa C++ w C, ba nawet pierwsze

kompilatory C++ po prostu tlumaczyly kod na C i pozniej kompilowaly kod w

C. hmm...

podsumowujac zadne gpu nie potrafi wykonywac kodu pisanego w C++

'no spik C++ segnor', a co wiecej gpu, nawet cudo (CUDA) fermii nie ma do

ch* jasnej stosu! i nie ma szans na wydajna rekursje. A jak chcesz sobie

sam napisac stos, tylko nie wiem gdzie, w pamieci dzielonej SIMD ? jasne...

juz widze jak szybko bedzie dzialal... czy moze na jakiejs 'teksturze' -

jeszcze lepiej.

funkcje rekurencyjne to podstawa informatyki

wszystko obliczalne na maszynie Turinga ma postac funkcji rekurencyjnej

wydajna obsluga funkcji rekurencyjnych to podstawa szybkiego i

ekspresyjnego sprzetu

i nie gniewaj sie :) nie chce cie obrazic

PS: aha i z ciekawosci zerknalem na tego przyspieszonego blendera i po

pierwsze, z tego co widze blenderowi zamienili zwykle octree na drzewo

bvh, a po drugie:

"He switched to special BVH's using adaptive children count, exploiting

raycoherence using hints/LCTS, SIMD at the end. He's using self developed

heuristics to build trees and reduce the expected number of bounding box

tests by about 20%. He calls his BVH's "VBVH" and "SVBVH". VBVH is much

slower than SBVH when using non-SIMD."

wydaje mi sie ze to swiadczy glownie o tym jak slaby byl wczesniejszy

raytracer blendera :)

PS2: a tak jeszcze sie zastanowilem, upieram sie przy tej rekursji, a wczesniej

jest duzo wazniejszy problem, dynamiczny przeplyw sterowania. na gpu

nie da sie zrobic dynamicznej petli, nie da sie zrobic if'a w wydajny sposob,

nie da sie zrobic podstawowej instrukcji w stylu 'jump' (przejdz do instrukcji).

stary... implementacje wskaznikowych czy kursorowych struktur danych jak

drzewa, na gpu to sa dalej akademickie rozwazania, ktore maja sie do

architektur SIMD jak piesc do nosa. tak da sie je zaklepac ale ilosc

hakowania aby obejsc ograniczenia architektury skutecznie jak do tej pory

odbieraja komukolwiek ochote by tego stosowac.

Nie wiem czego się naczytałeś, ale korzystanie z takich struktur danych na gpu nie przysparza nikomu problemów.

Nie wiem co ma blender ale drzewo bounding volume hierarchy, (jak sie myle

to mnie popraw) ktore jako drzewo licznikowe, znane jest chyba od

paleozoiku, przytoczylem bsp jako przyklad jakich wiele.

Miał zwykłe drzewo usemkowe które w pewnych scenach jest idealne się sprawdza i wyprzedza inne drzewa - jednak w większości wypadków lepszym jest BVH (co nie zmienia faktu, że octree sprawuje się lepiej niż zwykłe binarne drzewko BSP)... no chyba że chcesz robić wszystko na GPU razem z tworzeniem drzewa co wiele prościej można zrobić jako preprocess na CPU i gotowe drzewo podać do GPU (na GPU da się to zrobić - tylko po co ;p).

Powiedzialem path tracer bo nie widze innej mozliwosci. czy to bedzie

2^n czy 13^n i tak jest to strasznie wolny algorytm. monte carlo czy

nie zlozonosc obliczeniowa jest gigantyczna, o stala sie nie kloce.

Tak złożoność w tych algorytmach jest olbrzymia i na CPU te algorytmy zdychają - dlatego na CPU lepiej użyć mniej realistycznych, ale za to szybkich metod (które na gpu liczyłyby się w mgnieniu oka).

inny algorytm ktory widze to by byla kratka i iteracyjnie policzony transfer

swiatla (jak w demie GI cryengine 3), ale ze wzgledow pamieciowych w

jakosci lepszej niz przyblizone gi w realtimie wydaje mi sie to niemozliwe.

Tak jeśli mówisz o Light Propagation Volumes to jest on szybki i na szybko jak kiedyś przejrzałem jest on bardzo podobny do Point-Based Approximate Color Bleeding stosowany przez Pixara i ILM (metody oparte na tej http://http.developer.nvidia.com/GPUGems2/gpugems2_chapter14.html). Co do tego, że jest nie możliwe, to CUDA i OpenCL mogą korzystać z RAM komputera jeśli tego programista zapragnie (straci to trochę na wydajności, bo wąskim gardłem jest przepustowość szyny, ale i tak jest to szybsze rozwiązanie niż CPU) - co do Realtime to są metody dające lepsze rezultaty tylko wtedy innych efektów nie można robić typu wolumetryczne światła (nawet najprostsze jak w cryengine 2 czyli robione przez radial blur nieba i słońca), reliefmaping już się z lepszym GI w tej chwili nie wyrobią.

no niestety przypowiesc ma cechy metaforyczne wymagajace skupienia od

czytelnika

no niestety cechą przypowieści jest też to, że ktoś kto ją opowiada powinien wiedzieć o czym mówi.

ok tu zwracam honor nie zauwazylem ze dla doubli jest support dla 4800+, 5800+

w serii 4k 4770 obsługuje fp64 - tylko w nowych obcięli.

uzycie szablonow to zabieg dziejacy sie podczas kompilacji a nie podczas

uruchomienia programu. po co mi funkcja ktorej zagniezdzenie rekursji jest

ustalane podczas kompilacji.

W ray tracingu jest to wystarczające - głębokość promieni, cieni, odbicia... ustalasz przed renderingiem, przed renderingkiem kompilujesz kernel i już. Do prawdziwej rekurencji musisz poczekać do fermi.

Tu dochodzimy do kolejnej sprawy, i nie zmieni tego marketing... programy

na gpu tez sa kompilowane (olaboga!)... to 'wykonywanie kodu pisanego w

c++' o ktorym wspominasz to wierutna bzdura. kod jest kompilowany w

sterowniku i na gpu trafia juz jako assembler gpu. to cale zamieszanie z

"wsparciem natywnym c++" wiaze sie jedynie z unifikacja adresowania

pamieci, nie wiem czy wiesz ale i kod C i C++ nim jakikolwiek procesor wykona

program zamienia sie na taki sam kod maszynowy, ba nawet da sie w latwy

sposob napisac funkcjonalnosc obiektowa C++ w C, ba nawet pierwsze

kompilatory C++ po prostu tlumaczyly kod na C i pozniej kompilowaly kod w

C. hmm...

O jakim natywnym wsparciu dla C++ piszesz? Fermi będzie to po prostu pierwszy GPU który potrafi taki kod wykonać (tu nie jest ważne czy C++/C/Java...) - wcześniejsze GPU nVidii czy AMD nie potrafiły wykonać rekurencji czy też nie miały zunifikowanego adresowania pamięci i nie potrafiły wykonać programu pisanego na CPU - Fermi to zmienia i można korzystać z "normalnego" języka, a nie okrojonego C jak w wypadku CUDA czy OpenCL 1.0

podsumowujac zadne gpu nie potrafi wykonywac kodu pisanego w C++

'no spik C++ segnor', a co wiecej gpu, nawet cudo (CUDA) fermii nie ma do

ch* jasnej stosu! i nie ma szans na wydajna rekursje. A jak chcesz sobie

sam napisac stos, tylko nie wiem gdzie, w pamieci dzielonej SIMD ? jasne...

juz widze jak szybko bedzie dzialal... czy moze na jakiejs 'teksturze' -

jeszcze lepiej.

Zapewne zaimplementują ją w pamięci globalnej, a Ty jak chcesz nawet w pamięci hosta możesz - to, że nie ma stosu to z jednej strony wada, ale z innej zaleta (tak jak w językach scala etc. stos i call jest jest implementowany na stercie dzięki czemu przy naprawdę głębokiej rekurencji nie wywali Ci przepełnienia stosu - bo umówmy się tail-call optimization w kompilatorach C++ na stosie działają raczej kiepsko).

wszystko obliczalne na maszynie Turinga ma postac funkcji rekurencyjnej

 

wydajna obsluga funkcji rekurencyjnych to podstawa szybkiego i

ekspresyjnego sprzetu

 

i nie gniewaj sie :) nie chce cie obrazic

Nie gniewam się bo tu masz rację - rekurencja jest bardzo ważna i sam z powodu jej braku w OpenCL zrezygnowałem z przepisania swojego MLT na GPU ;p - CUDA nie lubie bo one-vendor-only, a w OpenCL nie można zrobić nawet sztuczki z szablonami bo ich nie ma ;p. W Fermi i nowych wersjach Cuda i OpenCL się to zmieni i wtedy reaktywuje odprężający projekcik pisany w wolnym czasie (jednak dla firm i osób które mają dużo wolnego czasu ograniczenia tych wersji Cuda i OpenCL nie są za wielkim ograniczeniem - po prostu muszą zmienić sposób zapisu algorytmu).

wydaje mi sie ze to swiadczy glownie o tym jak slaby byl wczesniejszy

raytracer blendera :)

a czy ja mówiłem że wcześniej był super? Takiego przyspieszenia by nie uzyskali gdyby był wcześniej super.

PS2: a tak jeszcze sie zastanowilem, upieram sie przy tej rekursji, a wczesniej

jest duzo wazniejszy problem, dynamiczny przeplyw sterowania. na gpu

nie da sie zrobic dynamicznej petli, nie da sie zrobic if'a w wydajny sposob,

nie da sie zrobic podstawowej instrukcji w stylu 'jump' (przejdz do instrukcji).

Dynamiczne pętle da się zrobić (od SM3.0 jeśli dobrze pamiętam)... ale fakt, że instrukcje kontroli przepływu są strasznie wolne (wolne są i na CPU gdzie zwykły if to kilka cykli zegara, ale na GPU jest faktycznie jeszcze gorzej) - jednak w zastosowaniach takich jak raytracing i tak mimo wad GPU miażdży CPU - taka jest prawda.

Nie wiem czego się naczytałeś, ale korzystanie z takich struktur danych na gpu nie przysparza nikomu problemów.

jasne, nieregularny dostep do pamieci i nieregularny flow programu to aspekt

w ktorym gpu bryluje. a jak sobie splycisz strukture danych (a tak wlasnie

implementuje sie kddrzewa, bsp i co tam chcesz) to narzut na pamieci jest

taki duzy ze glupi zajac zajmuje 50 mega (do wydajnosci pamieci na gpu

jeszcze wrocimy). Publikacje akademickie to sa zazwyczaj 'proof of concepts',

nie wiem jak cie przekonac, moze powiem tak... wrocmy do przypadku iraya

15 quadro 5800 ma teoretyczna moc okolo 15 000 GifaFlopow na sekunde

obrazek pojawia sie po 4 minutach... core i7 920 ma jakies 50 gigaflopow

teoretycznej mocy obliczeniowej czyli 300 razy mniej. uznam ze gpu jest

wydajniejsze obliczeniowo jesli core i7 nie bedzie w stanie wypluc rownie

ladnego obrazka w ciagu 1200 minut. rozumiem jeden aspekt, ilosc zuzytego

pradu, ale mowic ze algorytmy na gpu/gpu sa wydajniejsze niz na cpu ?

rozumiem tez ze po 4 minutach juz otrzymujesz obraz, co tez moze byc

wazne, ale jak dla mnie to ma elegancje orania pola czolgiem: da sie, jest

szybko no i... ?

Tak jeśli mówisz o Light Propagation Volumes to jest on szybki i na szybko jak kiedyś przejrzałem jest on bardzo podobny do Point-Based Approximate Color Bleeding stosowany przez Pixara i ILM (metody oparte na tej http://http.developer.nvidia.com/GPUGems2/gpugems2_chapter14.html). Co do tego, że jest nie możliwe, to CUDA i OpenCL mogą korzystać z RAM komputera jeśli tego programista zapragnie (straci to trochę na wydajności, bo wąskim gardłem jest przepustowość szyny, ale i tak jest to szybsze rozwiązanie niż CPU) - co do Realtime to są metody dające lepsze rezultaty tylko wtedy innych efektów nie można robić typu wolumetryczne światła (nawet najprostsze jak w cryengine 2 czyli robione przez radial blur nieba i słońca), reliefmaping już się z lepszym GI w tej chwili nie wyrobią.

pamiec hosta i pamiec urzadzenia to dwa rozne swiaty... chcesz pojedyncze

bajty zczytywac przez DMA :) ? nawet jesli mialo by to byc mozliwe to

na szczescie implementacje openCL i CUDY tego zabraniaja... mozesz

zmapowac pamiec globalna albo stale bufory na RAM kompa, wykonac transfer

DMA, odpalic kernel, shader cokolwiek... i tyle...

w serii 4k 4770 obsługuje fp64 - tylko w nowych obcięli.

ok, whatever

W ray tracingu jest to wystarczające - głębokość promieni, cieni, odbicia... ustalasz przed renderingiem, przed renderingkiem kompilujesz kernel i już. Do prawdziwej rekurencji musisz poczekać do fermi.

stala glebokosc rekursji jeszcze przelkne (choc wyobrazam sobie zlozonosc

branchowania jesli chcesz miec choc jednego ifa), ok, ale fermi nic nie

zmieni... patrz nizej...

O jakim natywnym wsparciu dla C++ piszesz? Fermi będzie to po prostu pierwszy GPU który potrafi taki kod wykonać (tu nie jest ważne czy C++/C/Java...) - wcześniejsze GPU nVidii czy AMD nie potrafiły wykonać rekurencji czy też nie miały zunifikowanego adresowania pamięci i nie potrafiły wykonać programu pisanego na CPU - Fermi to zmienia i można korzystać z "normalnego" języka, a nie okrojonego C jak w wypadku CUDA czy OpenCL 1.0

no jesli c++ bez rekursji i bez 'jmp' na poziomie pojedynczego watka jest "nie

okrojone" to ok. Powtorze to jeszcze raz... gdzie ty chcesz ten stos i

kontrole sterowania upchnac ? primo... w 48 kilo L1 dzielonego na 32 watki ?

czy moze w 768 kilo L2 dzielonego na 512.

powtorze to kolejny raz gpu to maszyna SIMD... i chocby nvidia i ati

krzyczaly w nieboglosy inaczej, sytuacja wygloda tak:

Bobek: this is madness....

Leonidas: this is SIMD!

jak dziala cpu ? masz swoj ulubiony rejestr Instruction Pointer i co robisz ? a

teraz wykonajmy instrukcje z nastepnej komorki pamieci, a teraz z nastepnej,

a teraz skoczmy o 10 do przodu, albo pod adres 00000000... a na gpu...

znowu przypowiesc: siedzi 32 galernikow (nazywaja sie CUDA cores)

a bebniarz krzyczy: a teraz wszyscy biora rejestr r1 i mnoza przez rejestr

r2, a teraz wszyscy odejmuja 5 od rejestru r3, a teraz wszyscy zapisuja

w pamieci dzielonej w adresie ble ble ble...

oczywiscie jest branch ale zczytuje go 'bebniarz' a nie kazdy z galernikow.

oni wykonuje listy polecen, ewentualnie polecen poprzedzonych predykatami...

o tym za chwile...

Zapewne zaimplementują ją w pamięci globalnej, a Ty jak chcesz nawet w pamięci hosta możesz - to, że nie ma stosu to z jednej strony wada, ale z innej zaleta (tak jak w językach scala etc. stos i call jest jest implementowany na stercie dzięki czemu przy naprawdę głębokiej rekurencji nie wywali Ci przepełnienia stosu - bo umówmy się tail-call optimization w kompilatorach C++ na stosie działają raczej kiepsko).

sorry, znowu powiem ci... bzdura... zeby nie byc goloslownym na gtx275,

odczyt z rejestru 1 takt, odczyt z pamieci dzielonej 1 takt, odczyt z pamieci

lokalnej/globalnej.... 120-150 taktow.... stos/sterta... stos na stercie

jeden pies.... gpu nie ma podsystemu pamieci jak cpu... tam na 20 watkow jest

6-8 mega cache, tu na tysiac jest 1.5... to nie kwestia jezyka ani tego czy

stos jest na stercie (ktorej tak na serio nie ma, bo na gpu nie ma czegos takiego

jak alokacja dynamicznej pamieci... wyobrazasz sobie malloca na 1000 watkow ?).

Po prostu gpu nie ma sprzetu zeby tego robic...

Nie gniewam się bo tu masz rację - rekurencja jest bardzo ważna i sam z powodu jej braku w OpenCL zrezygnowałem z przepisania swojego MLT na GPU ;p - CUDA nie lubie bo one-vendor-only, a w OpenCL nie można zrobić nawet sztuczki z szablonami bo ich nie ma ;p. W Fermi i nowych wersjach Cuda i OpenCL się to zmieni i wtedy reaktywuje odprężający projekcik pisany w wolnym czasie (jednak dla firm i osób które mają dużo wolnego czasu ograniczenia tych wersji Cuda i OpenCL nie są za wielkim ograniczeniem - po prostu muszą zmienić sposób zapisu algorytmu).

moze nie trzeba bylo rezygnowac i cos sprobowac napisac nie gadalbys takich

farmazonow. a ty rozumiesz co ten szablon robi ? po prostu zamiast wywolywac

dynamicznie funkcje rekurencyjna po prosty przepisuje ja na chama, w uproszczeniu:

np:

abc {

print "akuku"

}

abc() {

print "ble ble" x

abc(x-1)

print "elb elb" x

}

i teraz:

abc {

print "ble ble" 5

print "ble ble" 4

print "ble ble" 3

print "ble ble" 2

print "ble ble" 1

print "akuku"

print "elb elb" 1

print "elb elb" 2

print "elb elb" 3

print "elb elb" 4

print "elb elb" 5

}

albo tworzy piec funkcji abc1, abc2, abc3, abc4, abc5 wywolujacych sie statycznie nawzajem.

masz swoje rekurencyjne gpu

Dynamiczne pętle da się zrobić (od SM3.0 jeśli dobrze pamiętam)... ale fakt, że instrukcje kontroli przepływu są strasznie wolne (wolne są i na CPU gdzie zwykły if to kilka cykli zegara, ale na GPU jest faktycznie jeszcze gorzej) - jednak w zastosowaniach takich jak raytracing i tak mimo wad GPU miażdży CPU - taka jest prawda.

petle dynamiczne nie istnieja... mozna je hakowac predykatami ale... patrz nizej...

gpu dziala w nastepujacy sposob: jakiekolwiek branching robiony jest

przez predykaty... po prostu jak masz ifa to kod:

if (a

print a

a += 10

b -= 5

} else {

print b

a *= 4

b = 0

}

zamienia sie na gpu na:

policz predykat ppp = a

wykonaj (jesli predykat ppp) print a

wykonaj (jesli predykat ppp) a += 10

wykonaj (jesli predykat ppp) b -= 5

wykonaj (jesli nie predykat ppp) print b

wykonaj (jesli nie predykat ppp) a *= 4

wykonaj (jesli nie predykat ppp) b = 0

czyli program ktory powinien wykonac 4 instrukcje na watek, na maszynie

SIMD (poza przypadkami iz kazdy watek idzie w ta sama galaz) wykona 7

instrukcji... na tym polaga branching na gpu... i fermi tego nie zmieni... no

way...

nie chce sie upierac, sam sie przylapuje na typowej programistom przypadlosci

"wszystkowiedzenia" :), ale wez napisz jakis shader w CAL'u albo PTX'ie a nie

w hdsl'u. bez kompilatora od razu zobaczysz dlaczego gpu to nie cpu.

nie twierdze ze nie znasz sie na rzeczy, ale rzucasz za duzo hasel z broszurek.

ps: przejrzyj sobie papiery do LRBni przy okazji, to bedzie calkiem fajne. Wole juz

na Larrabee czekac niz na fermi, ktore nic nowego nie wniesie.

re: nvidia

no planowalem od jakiegos czasu 'jak dziala gpu' napisac, ale nie bylo chwili.

wrocmy do przypadku iraya

15 quadro 5800 ma teoretyczna moc okolo 15 000 GifaFlopow na sekunde

obrazek pojawia sie po 4 minutach... core i7 920 ma jakies 50 gigaflopow

teoretycznej mocy obliczeniowej czyli 300 razy mniej. uznam ze gpu jest

wydajniejsze obliczeniowo jesli core i7 nie bedzie w stanie wypluc rownie

ladnego obrazka w ciagu 1200 minut. rozumiem jeden aspekt, ilosc zuzytego

pradu, ale mowic ze algorytmy na gpu/gpu sa wydajniejsze niz na cpu ?

rozumiem tez ze po 4 minutach juz otrzymujesz obraz, co tez moze byc

wazne, ale jak dla mnie to ma elegancje orania pola czolgiem: da sie, jest

szybko no i... ?

Jeśli oczekujesz różnic jakie wynikają z teoretycznych mocy obliczeniowych gpu to faktycznie możesz się zawieść, bo to uzyskasz jeśli nie będziesz nic innego robił tylko mnożył i dodawał (radeony w jednym cyklu potrafią zrobić 2x liczba rdzeni operacji muladd)... jednak co byś nie mówił szybciej rzeczy takie jak rendering policzy GPU mimo swoich ograniczeń (lub dużo taniej jeśli chciałbyś tak samo szybko liczyć na CPU i zrobić klaster).

pamiec hosta i pamiec urzadzenia to dwa rozne swiaty... chcesz pojedyncze

bajty zczytywac przez DMA :) ? nawet jesli mialo by to byc mozliwe to

na szczescie implementacje openCL i CUDY tego zabraniaja... mozesz

zmapowac pamiec globalna albo stale bufory na RAM kompa, wykonac transfer

DMA, odpalic kernel, shader cokolwiek... i tyle...

Tak to 2 różne światy, ale możesz zaalokować pamięć na GPU i przekazać do GPU jako wskaźnik na pamięć z flagami "CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR" to będziesz miał dostęp do zapisu w pamięci hosta (nie jest to wydajne, bo ciągle jest prowadzona komunikacja pomiędzy gpu a ram (co wręcz zabija wydajność), ale się da).

no jesli c++ bez rekursji i bez 'jmp' na poziomie pojedynczego watka jest "nie

okrojone" to ok. Powtorze to jeszcze raz... gdzie ty chcesz ten stos i

kontrole sterowania upchnac ? primo... w 48 kilo L1 dzielonego na 32 watki ?

czy moze w 768 kilo L2 dzielonego na 512.

 

powtorze to kolejny raz gpu to maszyna SIMD... i chocby nvidia i ati

krzyczaly w nieboglosy inaczej, sytuacja wygloda tak:

 

Bobek: this is madness....

Leonidas: this is SIMD!

C++ z rekursją w wypadku Fermi.

Co do SIMD to Fermi jest MIMD i każdy rdzeń robi co chce osobno (i kolejna przypowieść niezbyt celna - tzn celna jak się strzela do amd lub kart z serii 200-).

stos na stercie

jeden pies.... gpu nie ma podsystemu pamieci jak cpu... tam na 20 watkow jest

6-8 mega cache, tu na tysiac jest 1.5... to nie kwestia jezyka ani tego czy

stos jest na stercie (ktorej tak na serio nie ma, bo na gpu nie ma czegos takiego

jak alokacja dynamicznej pamieci... wyobrazasz sobie malloca na 1000 watkow ?).

Po prostu gpu nie ma sprzetu zeby tego robic...

Fermi ma dostęp do pamięci prawie taki jak CPU

moze nie trzeba bylo rezygnowac i cos sprobowac napisac nie gadalbys takich

farmazonow. a ty rozumiesz co ten szablon robi ? po prostu zamiast wywolywac

dynamicznie funkcje rekurencyjna po prosty przepisuje ja na chama, w uproszczeniu:

...

 

albo tworzy piec funkcji abc1, abc2, abc3, abc4, abc5 wywolujacych sie statycznie nawzajem.

 

masz swoje rekurencyjne gpu

Przecież mówię że póki co nie ma rekurencji na GPU i trzeba ją obchodzić - fermi ma wprowadzić rekurencję normalną

petle dynamiczne nie istnieja... mozna je hakowac predykatami ale... patrz nizej...

...

nie chce sie upierac, sam sie przylapuje na typowej programistom przypadlosci

"wszystkowiedzenia" :), ale wez napisz jakis shader w CAL'u albo PTX'ie a nie

w hdsl'u. bez kompilatora od razu zobaczysz dlaczego gpu to nie cpu.

Nie wiem jak w CAL, ale w Brook tak było (powodem było wykorzystywanie SM2.0+ (od kart nVidia FX) gdzie nie było dynamicznych pętli - jednak na kartach SM3.0+ pętle są dynamiczne i ich "rozmiar" zmieniać się może w trakcie działania programu (np. niezmiennik pętli przekazywany jako uniform do shadera) - więc te pętle nie mogą być rozwijane w trakcie kompilacji tak jak napisałeś).

ps: przejrzyj sobie papiery do LRBni przy okazji, to bedzie calkiem fajne. Wole juz

na Larrabee czekac niz na fermi, ktore nic nowego nie wniesie.

Od kiedy intel powiedział dwa miesiące temu, że projekt GPU Larrabee na PC został odwołany już nie czekam na tę kartę ;]

Jeśli oczekujesz różnic jakie wynikają z teoretycznych mocy obliczeniowych gpu to faktycznie możesz się zawieść, bo to uzyskasz jeśli nie będziesz nic innego robił tylko mnożył i dodawał (radeony w jednym cyklu potrafią zrobić 2x liczba rdzeni operacji muladd)... jednak co byś nie mówił szybciej rzeczy takie jak rendering policzy GPU mimo swoich ograniczeń (lub dużo taniej jeśli chciałbyś tak samo szybko liczyć na CPU i zrobić klaster).

whatever, niech nawet 10% tej mocy wykorzystaja... ide o zaklad ze w 120

minut a nie 1200 i7 wyrenderuje rownie ladny obrazek.

Tak to 2 różne światy, ale możesz zaalokować pamięć na GPU i przekazać do GPU jako wskaźnik na pamięć z flagami "CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR" to będziesz miał dostęp do zapisu w pamięci hosta (nie jest to wydajne, bo ciągle jest prowadzona komunikacja pomiędzy gpu a ram (co wręcz zabija wydajność), ale się da).

nie uzywaj terminu ktorych dzialania nie znasz... rdzen gpu ma dostep jedynie

do pamieci na gpu... koniec kropka CL_MEM_USE_HOST_PTR i na mac'u i

w streamsdk i w nvidiowej implementacji openCL przerzuca zawartosc pamieci

z hosta na device tam wykonuje kernel i znow np jesli wyniki tez byly

zmapowane na hosta przerzuca wyniki spowrotem. nie ma synchronicznej,

read write pamieci dostepnej dla cpu i gpu, ktos kto by cos takiego

zaimplementowal po szynie pci zasluzylby na ignobla.

C++ z rekursją w wypadku Fermi.

Co do SIMD to Fermi jest MIMD i każdy rdzeń robi co chce osobno (i kolejna przypowieść niezbyt celna - tzn celna jak się strzela do amd lub kart z serii 200-).

g prawda... nie wierz w plotki... nvidia reklamuje multiple kernels jako MIMD.

'kazdy rdzen robi co chce'... normalnie kaze sobie to wydrukowac na banerze

i zawiesze w biurze. podstawowa jednostka granulacji watkow jest '1 warp'

czyli 16 lub 32 (nie liczac double dispatcha) watki wykonujace ta sama

instrukcje S.I.M.D. jakbys chcial miec pelny MIMD, to fermi z maszyny 512

'potokowej' zamienia sie w 16 potokowa... chyba nie o to nam chodzilo.

Fermi ma dostęp do pamięci prawie taki jak CPU

prawie robi roznice.

512 aktywnych watkow, ~2K niekatywnych, 768 kilo L2, 768 kilo L1 na GPU

vs.

5 aktywnych watkow 50 nieaktywnych i 256kilo L1, 1 MB L2, 6-8-12 MB L3

Przecież mówię że póki co nie ma rekurencji na GPU i trzeba ją obchodzić - fermi ma wprowadzić rekurencję normalną

w przeciwienstwie do nienormalnej... normalna czyli: wspolna dla warpa

i bez stosu ?

Nie wiem jak w CAL, ale w Brook tak było (powodem było wykorzystywanie SM2.0+ (od kart nVidia FX) gdzie nie było dynamicznych pętli - jednak na kartach SM3.0+ pętle są dynamiczne i ich "rozmiar" zmieniać się może w trakcie działania programu (np. niezmiennik pętli przekazywany jako uniform do shadera) - więc te pętle nie mogą być rozwijane w trakcie kompilacji tak jak napisałeś).

hmm... wszystkie petle moga byc rozwijane dziala to w nastepujacy sposob:

po prostu miedzy kolejnymi iteracjami wkladany jest dynamiczny test

(srodkowa czesc for'a) ktory moze wyskoczyc calym warpem na sam koniec.

Ma to calkowity sens, przeciez mozesz np. wiedziec ze dynamiczny test

wyznaczy na oko 20 iteracji, dlaczego by wiec nie unrollowac np 5.

CAL IL i PTX to nie sa jezyki klienckie (jak brook, cg, hdsl itd...) tylko ISA,

wewnetrze assemblery kart graficznych. narzedziem robienia branchowania

w hardwarze (nawet w hd5000, fermii, gtx200 itd..) sa predykowane instrukcje.

dynamiczne petle (czyli rozne ilosci iteracji dla roznych watkow)

wykonywane sa w nastepujacy sposob... wszystkie watki wykonuja najdluzsza

potrzebna petle... (nim wszystkie zbranczuja na zewnatrz) a te ktore nie

musza wykonywac kolejnych iteracji po prostu wykonuja jalowe instrukcje

ktore predykatami sa olewane... nie przeskoczysz SIMD na poziomie warpa

(czyli 16, 32 watkow)

Od kiedy intel powiedział dwa miesiące temu, że projekt GPU Larrabee na PC został odwołany już nie czekam na tę kartę ;]

Larrabee to trzy produkty, intel wycofal sie z jednego. nie bedzie

konsumerskiej karty graficznej w pierwszej (wypuszczonej) generacji LRB.

ok... nie wiem co dalej... dyskusja schodzi troche na psy...

widze ze sie interesujesz tematyka, zapewne studiujesz informatyke albo

piszesz gry, ale w twojej wiedzy znajduja sie podstawowe dziury. i nie mowie

tu o wiedzy typu czy nowa karta ma to czy tamto, albo czy nowe sterowniki

nvidii wspieraja costam, bo ja np za tym nie nadazam. A raczej o braku

zrozumienia podstawowych zagadnien w stylu: model obliczen. Ja mam

(choc nie twierdze ze to cos znaczy) wysztalcenie informatyczne -

matematyczne, wiec sam tez mam tendencje do olewania stalej .... ale

nie mozna mowic o wydajnosci / zlozonosci bez dobrego okreslenia

modelu obliczen. Na twoim miejscu rowna uwage poswiecilbym nie tylko

na nowe techniki postprocesingu i raytracingu na gpu, ale na zapoznanie

sie z tym co w samym sprzecie piszczy - bo w tej dzialce akurat to ma

ogromne znaczenie (w przeciwienstwie do wielu dziedzin w ktorych

programowanie wysokopoziomowe sprawia sie bardzo dobrze).

whatever, niech nawet 10% tej mocy wykorzystaja... ide o zaklad ze w 120

minut a nie 1200 i7 wyrenderuje rownie ladny obrazek.

To bardzo możliwe - jednak 120 min to nie 4 min.

nie uzywaj terminu ktorych dzialania nie znasz... rdzen gpu ma dostep jedynie

do pamieci na gpu... koniec kropka CL_MEM_USE_HOST_PTR i na mac'u i

w streamsdk i w nvidiowej implementacji openCL przerzuca zawartosc pamieci

z hosta na device tam wykonuje kernel i znow np jesli wyniki tez byly

zmapowane na hosta przerzuca wyniki spowrotem. nie ma synchronicznej,

read write pamieci dostepnej dla cpu i gpu, ktos kto by cos takiego

zaimplementowal po szynie pci zasluzylby na ignobla.

Nie mówiłem o bezpośrednim zapisie - pamięć Ram jest cashowane do GPU tam odczytywane i zapisywane go pamięci GPU i idzie spowrotem.

g prawda... nie wierz w plotki... nvidia reklamuje multiple kernels jako MIMD.

'kazdy rdzen robi co chce'... normalnie kaze sobie to wydrukowac na banerze

i zawiesze w biurze. podstawowa jednostka granulacji watkow jest '1 warp'

czyli 16 lub 32 (nie liczac double dispatcha) watki wykonujace ta sama

instrukcje S.I.M.D. jakbys chcial miec pelny MIMD, to fermi z maszyny 512

'potokowej' zamienia sie w 16 potokowa... chyba nie o to nam chodzilo.

nVidia nie nazywa multiple kernels jako MIMD, tylko przynajmniej w papierkach pisze, że jeden cuda core może wykonać na raz 32 instrukcje (każdy procesor zupełnie niezależnie by pracował wtedy) - jak to będzie na prawdę zobaczymy jak się fermi wreszcie pojawi i nie będzie się trzeba informować ze źródeł zmanipulowanych "trochę" marketingiem.

prawie robi roznice.

 

512 aktywnych watkow, ~2K niekatywnych, 768 kilo L2, 768 kilo L1 na GPU

vs.

5 aktywnych watkow 50 nieaktywnych i 256kilo L1, 1 MB L2, 6-8-12 MB L3

Akurat tu mówiłem prawie do dram (gdzie wskaźnik po prostu wskazuje w jakiekolwiek miejsce w pamięci ram karty i może odczytać).

w przeciwienstwie do nienormalnej... normalna czyli: wspolna dla warpa

i bez stosu ?

Jeśli fermi będzie tak jak piszą mimd to nie dla wspólnego warpa, a stos mogą zaimplementować chociażby w ram karty. Trochę bełkotu marketingowego o rekurencji w fermi (nie za wiele on mówi, ale potwierdza, że fermi będzie obsługiwało rekurencję, więc jakoś z implementacją stosu sobie muszą poradzić):

GF100 is the first GPU to support recursion in hardware, enabling efficient GF100 Pathtracing Performance ray tracing and a host of other graphics 400% algorithms.

hmm... wszystkie petle moga byc rozwijane dziala to w nastepujacy sposob:

po prostu miedzy kolejnymi iteracjami wkladany jest dynamiczny test

(srodkowa czesc for'a) ktory moze wyskoczyc calym warpem na sam koniec.

Ale są robione przez normalne skoki jak tu w ptx:

	.func (.reg .u32 %rv1) main (.reg .u32 %ra1, .reg .u64 %rda2)
{
.reg .u32 %r;
.reg .u64 %rda1;
.reg .pred %p;
.loc	2	2	0
$LBB1_main:
mov.s32 	%r1, %ra1;
mov.u32 	%r2, 0;
setp.le.s32 	%p1, %r1, %r2;
mov.s32 	%r3, 0;
@%p1 bra 	$Lt_0_2306;
mov.s32 	%r4, %r1;
mov.s32 	%r5, 0;
mov.s32 	%r6, %r4;
$Lt_0_1794:
// Loop body line 2, nesting depth: 1, estimated iterations: unknown
.loc	2	6	0
add.s32 	%r3, %r5, %r3;
add.s32 	%r5, %r5, 1;
setp.ne.s32 	%p2, %r1, %r5;
@%p2 bra 	$Lt_0_1794;
bra.uni 	$Lt_0_1282;
$Lt_0_2306:
$Lt_0_1282:
.loc	2	8	0
mov.s32 	%rv1, %r3;
ret;

CAL IL i PTX to nie sa jezyki klienckie (jak brook, cg, hdsl itd...) tylko ISA,

wewnetrze assemblery kart graficznych. narzedziem robienia branchowania

w hardwarze (nawet w hd5000, fermii, gtx200 itd..) sa predykowane instrukcje.

dynamiczne petle (czyli rozne ilosci iteracji dla roznych watkow)

wykonywane sa w nastepujacy sposob... wszystkie watki wykonuja najdluzsza

potrzebna petle...

Wiem co to jest CAL i PTX. W Fermi ma być nowa wersja PTX 2.0. Nie chce mi się sprawdzać czy w Cal są skoki, ale raczej są bo wymaga od kart SM3.0

Co do dynamicznych pętli to się nie zrozumieliśmy - Ty rozumiesz je jako dynamiczna ilość iteracji dla różnych wątków (co w SIMD jest niemożliwe), a ja zrozumiałem jako po prostu dynamiczną ilość iteracji.

PS. już drugi raz napisałeś hdsl (myślałem, że to literówka dlatego nie poprawiałem) - masz na myśli chyba hlsl (inna nazwa cg - teraz jednak się rozwinął i już hlsl != cg)

Larrabee to trzy produkty, intel wycofal sie z jednego. nie bedzie

konsumerskiej karty graficznej w pierwszej (wypuszczonej) generacji LRB.

Tak z tym, że dla mnie to jedyny interesujący rynek dla kart pozwalających obliczenia GPGPU. A do kolejnych generacji jeszcze daleko (nawet do pierwszej, jak dokończą te projekty które robią) jeśli w ogóle się pojawią, więc na razie nie ma na co czekać IMO.

widze ze sie interesujesz tematyka, zapewne studiujesz informatyke albo

piszesz gry, ale w twojej wiedzy znajduja sie podstawowe dziury. i nie mowie

tu o wiedzy typu czy nowa karta ma to czy tamto, albo czy nowe sterowniki

nvidii wspieraja costam, bo ja np za tym nie nadazam. A raczej o braku

zrozumienia podstawowych zagadnien w stylu: model obliczen. Ja mam

(choc nie twierdze ze to cos znaczy) wysztalcenie informatyczne -

matematyczne, wiec sam tez mam tendencje do olewania stalej .... ale

nie mozna mowic o wydajnosci / zlozonosci bez dobrego okreslenia

modelu obliczen. Na twoim miejscu rowna uwage poswiecilbym nie tylko

na nowe techniki postprocesingu i raytracingu na gpu, ale na zapoznanie

sie z tym co w samym sprzecie piszczy - bo w tej dzialce akurat to ma

ogromne znaczenie (w przeciwienstwie do wielu dziedzin w ktorych

programowanie wysokopoziomowe sprawia sie bardzo dobrze).

Prawie zgadłeś - studiowałem informatykę i faktycznie piszę gry. Co do co w sprzęcie piszczy to zdecydowanie bardziej się skupiam właśnie na tym i dobrze znam ograniczenia sprzętu (tym bardziej, że często trzeba się ostro nagimnastykować, żeby jakiś algorytm przenieść z CPU na GPU w shaderach). Jednak widać, że Ty masz podejście typowo akademickie bo wszystko musi być ładnie i czytelnie rozwiązane, a wydajność jest na drugim miejscu, ja mam inne priorytety tak jak twórcy rendererów (oraz użytkownicy, którzy mają gdzieś, że GPU mimo teoretycznej mocy robi coś w tak długim czasie jak "4min", a CPU i7 jest tylko trochę wolniejsze bo wyrabia się z tym samym w "120min") - w bardzo wielu wypadkach CPU nawet jednordzeniowe może być o wiele wydajniejsze niż GPU (tu dopiero jest różnica pomiędzy teoretyczną wydajnością, a praktyką ;p), ale w dziedzinie mnie interesującej oraz na tym forum zgromadzonych obliczenia na GPU dają wielkie przyspieszenie.

przeczytałem co obaj panowie macie do powiedzenia i nie wiem co powiedzieć hehe. trochę jak przysłuchiwanie się rozmowom działu IT :]

dawno równie ciekawej i okraszonej rzeczowymi komentarzami "kłótni" na forum nie było. kontynuujcie, proszę ;)

Nie mówiłem o bezpośrednim zapisie - pamięć Ram jest cashowane do GPU tam odczytywane i zapisywane go pamięci GPU i idzie spowrotem.

to w jaki niby sposob mam tej pamieci uzywac do budowania stosu (od czego zaczela sie

w ogole dyskusja o pamieci global i hosta) ? o hoscie zapomnijmy bo to nie tedy droga.

pamiec global jest waskim gardlem wielu aplikacji na gpu. pamiec globalna ma przykre

wlasciwosci na gpu, slabo dziala przy nieregularnym dostepie (cashy nie ma w ogole,

lub sa kosmicznie male jak na ilosc watkow ktore walcza o pamiec), wymaga lokalnosci

operacji (nie wiem jak lepiej nazwac coalescent access) i jako pamiec 'random access'

narzuca rzad jesli nie dwa wielkosci wolniejszy dostep. glowna zasada pisania wydajnego

kodu: "keep ALU:TEX high!" to nie bierze sie z nikad. rozumiem budowanie stosu w

pamieci shared, ale do tej pory bylo jej tez smiesznie malo 16kb, a w fermii kiedy sie

wreszcie ukaze bedziemy mieli dostep do 48kilo _na wszystkie te watki dzialajace na

jednym SM_

A co do SIMD/MIMD nie przekonam cie zapewne samym opowiadaniem, zacytuje,

papiery fermi:

"

Hardware Execution

CUDA’s hierarchy of threads maps to a hierarchy of processors on the GPU; a

GPU executes one or more kernel grids; a streaming multiprocessor (SM)

executes one or more thread blocks; and CUDA cores and other execution

units in the SM execute threads. The SM executes threads in groups of 32

threads called a warp. While programmers can generally ignore warp

execution for functional correctness and think of programming one thread,

they can greatly improve performance by having threads in a warp execute

the same code path and access memory in nearby addresses.

"

"

Dual Warp Scheduler

The SM schedules threads in groups of 32 parallel threads called warps. Each

SM features two warp schedulers and two instruction dispatch units, allowing

two warps to be issued and executed concurrently. Fermi’s dual warp

scheduler selects two warps, and issues one instruction from each warp to a

group of sixteen cores, sixteen load/store units, or four SFUs. Because

warps execute independently, Fermi’s scheduler does not need to check for

dependencies from within the instruction stream. Using this elegant model of

dual-issue, Fermi achieves near peak hardware performance.

"

to nie jest MIMD, granulacja na poziomie kernela to 16/32 watki wykonywane

jako SIMD, oczywiscie techniki jak Out of order execution, czy rownolegle

wykonywanie ALU, TEX, SFU pozwalaja _calymi warpami_ wykonywac kilka

instrukcji na raz.

na poziomie jednego kernela/shadera gpu dziala jedynie jako calkiem szerokie

SIMD. na poziomie kilku kerneli rozne SM moga wykonywac rozne kernele -

to jest jedyne MIMD na fermi. jeden SM moze tez miec kilka warpow ktore

aktualnie obsluguje i np jesli jeden warp jedzie po ALU, to drugi moze

rownolegle meczyc TEX, albo SFU. Ale znow to nie jest MIMD, to zwykly

pipeline / vliw

tak, tak, hlsl :) brain fart...

btw... wyobrazasz sobie chip ktory mialby 512 fetchy i dekoderow?

jeszcze pare tekstow:

"

Divergent Branches. The PTX machine model defines a

divergent branch as a branch instruction where some threads

within the same warp branch while others fall through. In

this study, we report the branch divergence of an application

as the ratio of divergent branches to total branches.

As control flow divergence results in two separate sets of

threads whose execution must be serialized, the location of

the synchronization point has a profound impact on thread

activity and the number of dynamic instructions executed. In

[19], Fung et al. show that the earliest (ideal) location of thread

reconvergence is the immediate post dominator of the branch

instruction - that is, the nearest successor basic block that must

be executed if the branch instruction is executed regardless of

control path taken. We also investigated an alternative method

in which divergent threads are not reconverged until explicit

synchronization barriers are encountered by ignoring compiler

inserted reconverge points.

"

"

At every instruction issue time, the SIMT unit selects a warp that is

ready to execute and issues the next instruction to the active threads

of the warp. A warp executes one common instruction at a time, so full

efficiency is realized when all threads of a warp agree on their execution

path. If threads of a warp diverge via a data-dependent conditional branch,

the warp __serially__ executes each branch path taken, disabling threads that are

not on that path, and when all paths complete, the threads converge back to

the same execution path. Branch divergence occurs only within a warp;

different warps execute independently regardless of whether they are

executing common or disjointed code paths.

"

"

SIMT architecture is akin to SIMD (Single Instruction, Multiple Data) vector

organizations in that a single instruction controls multiple processing elements.

A key difference is that SIMD vector organizations expose the SIMD width to

the software, whereas SIMT instructions specify the execution and branching

behavior of a single thread. In contrast with SIMD vector machines, SIMT

enables programmers to write thread-level parallel code for independent,

scalar threads, as well as data-parallel code for coordinated threads. For the

purposes of correctness, the programmer can essentially ignore the SIMT

behavior; however, substantial performance improvements can be realized by

taking care that the code seldom requires threads in a warp to diverge.

"

fermi nic tu nie zmieni... fermi to nadal 16 procesorow SIMD o szerokosci

16/32 watki, model obliczen jest taki sam, oczywiscie dodano udogodnienia

w stylu jednoczesne wykonywanie kilku kerneli, cache przed kazda pamiecia,

wydajniejsze ale to nie zmienia modelu.

wyobraz sobie raytracing na gpu w ktorym 16/32 promienie ida w jednej

paczce przez jakas strukture danych (octree, kdtree, cokolwiek), trafiaja

w koncu na jakas gemetrie. To byla przestrzennie spojna paczka promieni

(kilka sasiednich pixeli np) zalozmy ze trafily w jeden trojkat (to sie zdarza

dla primary rays calkiem czesto), a teraz co? jeden sie odbije w jedna strone,

drugi w druga, trzeci bedzie mial sss, czwarty przeniknie... na pierwszym

bouncu moze jeszcze pojawia sie 'peczki' promieni, mozna temu nawet pomoc

i puszczac od razu promienie zmultiplikowane (po prostu duzo probek), ale

to nie zmienia faktu ze z kazdym poziomem tej pseudo-rekursji wykladniczo

rosnie zlozonosc branchowania. majac 16/32 watki na warpa szybko

zauwazysz ze liczac np gi, kazdy jest w innej odnodze sterowania - co

czyni kod seryjnym (czyli nie mozemy pominac zadnych instrukcji).

oczywiscie swiat sie kreci i powstaja hybrydowe rozwiazania, ale wg mnie

SIMD nadaje sie do rasteryzacji, nadaje sie do numeryki, ale do raytracingu

to jest tak jak powiedzialem: "oranie pola czolgiem", dlatego tak zaciekawilo

mnie larrabee (racja ze rodzi sie w bolach, ale moze sie doczekamy) -

egzemplarz 48 way MIMD z 512 bitowymi operacjami wektorowymi @ 2ghz

chetnie bym ogladnal.

mam nalecialosci akademickie (caly czas probuje sie w koncu zdoktoryzowac),

ale w tym roku stuknie mi 10 lat pracy zawodowej, na wielu architekturach

nisko i wysokopoziomowych. ostatnie 2 lata to tylko DX, wczesniej openGL.

ale jedynym powodem mojej dociekliwosci jest wlasnie proba wycisniecia

ze sprzetu tyle ile sie da... a nie da sie tego zrobic nie wiedzac jak on dziala.

re: tweety

smiej sie smiej a my tu prawimy co sie w twoim rendererze dzieje... :P

re: tweety

 

smiej sie smiej a my tu prawimy co sie w twoim rendererze dzieje... :P

czy ja się śmieję? bynajmniej! prawcie dalej, ja z chęcią dowiem się co mi podczas renderingu psizczy.

A co do SIMD/MIMD nie przekonam cie zapewne samym opowiadaniem, zacytuje,

papiery fermi:

...

to nie jest MIMD, granulacja na poziomie kernela to 16/32 watki wykonywane

jako SIMD, oczywiscie techniki jak Out of order execution, czy rownolegle

wykonywanie ALU, TEX, SFU pozwalaja _calymi warpami_ wykonywac kilka

instrukcji na raz.

A jednak mnie przekonałeś ;p. Jeszcze chętnie bym zobaczył linki do papierków. Ale jest to trochę dziwne bo stoi to w sprzeczności z wczesnymi papierkami, które mówiły o MIMD każdego cuda core - nie wiem czy to była pomyłka, czy celowy FUD ukierunkowany przeciw LRB

wyobraz sobie raytracing na gpu w ktorym 16/32 promienie ida w jednej

paczce przez jakas strukture danych (octree, kdtree, cokolwiek), trafiaja

w koncu na jakas gemetrie. To byla przestrzennie spojna paczka promieni

(kilka sasiednich pixeli np) zalozmy ze trafily w jeden trojkat (to sie zdarza

dla primary rays calkiem czesto), a teraz co? jeden sie odbije w jedna strone,

drugi w druga, trzeci bedzie mial sss, czwarty przeniknie... na pierwszym

bouncu moze jeszcze pojawia sie 'peczki' promieni, mozna temu nawet pomoc

i puszczac od razu promienie zmultiplikowane (po prostu duzo probek), ale

to nie zmienia faktu ze z kazdym poziomem tej pseudo-rekursji wykladniczo

rosnie zlozonosc branchowania. majac 16/32 watki na warpa szybko

zauwazysz ze liczac np gi, kazdy jest w innej odnodze sterowania - co

czyni kod seryjnym (czyli nie mozemy pominac zadnych instrukcji).

Tu jeszcze sobie wyobrażam bardzo dobrze, bo w zwykłym raytracerze takie promienie odbiłyby się o tą samą (prawie tą samą) normalną i dalej miałyby duże szanse na trafienie w ten sam tri (lub przynajmniej w tą samą przestrzeń w drzewie). Problem jest jak chce się mieć GI/AO przez rzucanie sampli w hemispherze do normalnej (wtedy odbite od powierzchni sample mogą być względem siebie zupełnie skierowane plecami (180 stopni)), jednak mimo to karta poradzi sobie dużo szybciej niż CPU (Można też zrobić rendering w kilku przebiegach zapisując potrzebne informacje na kilku teksturach i wysyłać paczki z samplami blisko siebie (żeby nie tracić czasu, a mieć szybko jeszcze lepszą jakość przez powiększenie liczby sample), a później złożyć wszysko (można nawet dopiero wtedy dodać cienie i oświetlenie jak w Deferred Shading)).

oczywiscie swiat sie kreci i powstaja hybrydowe rozwiazania, ale wg mnie

SIMD nadaje sie do rasteryzacji, nadaje sie do numeryki, ale do raytracingu

to jest tak jak powiedzialem: "oranie pola czolgiem", dlatego tak zaciekawilo

mnie larrabee (racja ze rodzi sie w bolach, ale moze sie doczekamy) -

egzemplarz 48 way MIMD z 512 bitowymi operacjami wektorowymi @ 2ghz

chetnie bym ogladnal.

Może i oranie pola czołgiem, ale tak szybkim, tanim oraz dostępnym czołgiem, że i tak się to opłaca.

Tu jeszcze sobie wyobrażam bardzo dobrze, bo w zwykłym raytracerze takie promienie odbiłyby się o tą samą (prawie tą samą) normalną i dalej miałyby duże szanse na trafienie w ten sam tri (lub przynajmniej w tą samą przestrzeń w drzewie). Problem jest jak chce się mieć GI/AO przez rzucanie sampli w hemispherze do normalnej (wtedy odbite od powierzchni sample mogą być względem siebie zupełnie skierowane plecami (180 stopni)), jednak mimo to karta poradzi sobie dużo szybciej niż CPU (Można też zrobić rendering w kilku przebiegach zapisując potrzebne informacje na kilku teksturach i wysyłać paczki z samplami blisko siebie (żeby nie tracić czasu, a mieć szybko jeszcze lepszą jakość przez powiększenie liczby sample), a później złożyć wszysko (można nawet dopiero wtedy dodać cienie i oświetlenie jak w Deferred Shading)).

od CES'a dostepne sa biale papiery fermii:

http://www.nvidia.com/content/PDF/fermi_white_papers/NVIDIA_Fermi_Compute_Architecture_Whitepaper.pdf

no tak, ale rzadko kiedy mamy doczynienia z idealnymi lustrami :) materialy

maja dosc dowolne parametry: odbic specular, zalaman, odbic diffuse, dodaj

do tego glossa, dodaj sss... w ogolnosci BSDF... lustra i szklanki (oraz im

podobne) ktore niezle zachowuja po odbiciu spojnosc przestrzenna to bardzo

waski wycinek. a z kazdym odbiciem rozrzut sie powieksza.

edit: btw. zauwazyles ze ulubionym tematem dem gpu raytracing sa blyszczace

lustrzane samochody ? :)

wiele przejsc algorytmu to oczywiscie jakies rozwiazanie, ale... sortowanie,

grupowanie promieni kolejnych poziomow jest, kosztowne: bo generuje sie

ich ogromna ilosc, a po drugie... nie wiadomo jak tego dokonac... a priori nie

mamy wiedzy ktore promienie ida w to samo miejsce, gdybysmy mieli to

po co w ogole je tracowac? wystarczy znalezc 'przedstawiciela' takiej grupki

(w uproszczeniu) policzyc gdzie trafi, a pozniej z duza dokladnoscia okreslic

gdzie trafia sasiedzi. pozostaja heurezy, mozna podzielic promienie na 'kubelki'

wg. punktu startowego i kierunku i takimi paczkami sobie dalej dzialac. Jasne

bedzie szybciej... ale to nie jakosciowa zmiana.

zastanowmy sie nad ta sztuczka z puszczaniem 16 probek w to samo miejsce

(albo bardzo blisko), to dobrze dziala na gpu, bo te probki ida ta sama sciezka

kodu i nie wykonuja nadmiarowych instrukcji na branchujacy kod. trafiaja

w jeden trojkat z malutkim rozrzutem i jakos tam sie rozrzucaja dalej (i dla

dobra dyskusji powiedzmy ze tam dalej to juz jest gorzej bo leca w roznych

kierunkach). fan raytracingu na gpu powie: no tak, ale pierwsze 16 primary

rayow zalatwilismy od tak *pstryk* moca gpu... a fan cpu powie: no tak

ale tylko dlatego ze gpu nie moze sobie z jednego promienia wyspawnowac

16 (co sie dzieje w monte carlo na cpu bardzo czesto) bo by zdechlo, wiec

musi na poczatku wykonac 16 razy wiecej praktycznie takich samych obliczen

zeby pierwsze promienie dotarly do pierwszego zdarzenia.

a trzeba pamietac ze branch nie zdarza sie tylko na 'trafieniu' i shaderach

opisujacych material... wystarczy drobny rozrzut i promienie lecace

z nieodleglych punktów w podobnych kierunkach zaczna bladzic roznymi

sciezkami w drzewkach.

ps: coraz bardziej mi sie wydaje ze na obecnej klasie sprzetu droga ktora

da lepsze wyniki bedzie radiance transfer... zbudowac siatke opisujaca

przeplyw swiatla (chocby grid, a moze inne cudo - losowy podzial? podzial

delaunaya ), naladowac punkty startowe od zrodel swiatla, a pozniej

zapuscic rownanie Poissona i iterowac, iterowac...

no ja też przeczytałem z wypiekami na twarzy, jak nic powinni obydwaj fronta dostać ! czekam na więcej "przytulony" do tweetiego (bo trochę się boję co będzie dalej).

"wanna fight?! you wanna fight?!" :D

najs, ale szkoda energii na dyskusje o GPU, prawda jest taka ze dopiero wielu programistow usiluje sobie wyobrazic jak zaimplementowac niektore algorytmy

w ukladzie strumieniowym, a producenci GPU maja zlewke, gdyby sie przejmowali,

musieliby zbudowac duzo drozsza i mniej wydajna architekture... a to chyba

nikomu narazie nie jest na reke bo nie ma na to rynku, wiec

przecigna nas jeszcze na tych swoich biedusich rozwiazaniach kilka lat

edit: btw. zauwazyles ze ulubionym tematem dem gpu raytracing sa blyszczace

lustrzane samochody ? :)

Nie tylko GPU raytracing - Cell raytracing też gustuje w takich scenach (jednostki wektorowe Cell (SPU) to też SIMD ;p). Chociaż demo fermi z Bugatti jest tu wyjątkiem, bo to pathtracer był pokazany i rzucał sample normalnie.

a priori nie mamy wiedzy ktore promienie ida w to samo miejsce, gdybysmy mieli to

po co w ogole je tracowac?

Nie mamy wiedzy gdzie idą, ale mamy wiedzę jak się rozchodzą sample (niezależnie czy to materiał glossy czy liczysz światło odbite od innych obiektów) - przecież nie losujesz tych wektorów z każdym zderzeniem tylko jako preproces (ilość psełdolosowych próbek które możesz upakować w paczki gdzie pomiędzy wektorami jest mały kąt - później nie musisz już losować i sortować tylko zorientować wektory, żeby znajdowały się w hemisferze normalnej, a kąty pomiędzy próbkami się nie zmienią , tylko ich kierunek). Dodatkowo trzeba pamiętać, że rzadko zdarza się głęboko liczyć GI/Glossy w zwykłych pathtracerach, bo nie daje to dużych zysków jakościowych, a spore minusy wydajnościowe.

do połowy dyskusji jeszcze coś rozumiałem.... ;)

Obserwowałem dyskusje od początku i również muszę przyznać, że takich rzeczowych "kłótni" więcej. Aż szkoda, że już wszystko sobie wyjaśniliście.

Dlatego tez przeniosłem dyskusje do General Discussion i przeniosłem również początkowe posty dla ciągłości.

Jeśli chcielibyście dokonać kontynuacji kłótnie w "formie artykułu" dajcie znać :) "Fakty, mity GPU - całkiem ciekawy temat".

też bym chętnie przeczytał taki artykuł.

jako student 2 roku informatyki też się zgubiłem gdzieś w połowie;)

no i spłoszyła się ptaszyna ... ;)

to teraz czas aby SWRSC Morphe Lis 1515 i inni masterzy sobie zrobili "łona fajt ? łona fajt ? " z ustawień vraya i postrocesingu w szopie. Już przecieram okulary.