53

´ Uvod CUDA OpenCL PhysX Literatura

Hardware pro poˇc´ıtaˇcovou grafiku NPGR019 ´ Uvod do architektury CUDA

Jan Hor´aˇcek http://cgg.mff.cuni.cz/ MFF UK Praha

2012

Jan Hor´ aˇ cek

CUDA

1 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Historie V´ıcej´ adrov´ e zpracov´ an´ı CPU vs GPU

Obsah

1

´ Uvod

2

CUDA

3

OpenCL

4

PhysX

5

Literatura

Jan Hor´ aˇ cek

CUDA

2 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Historie V´ıcej´ adrov´ e zpracov´ an´ı CPU vs GPU

Historie

v´yvoj kˇrem´ıkov´ych ˇcip˚ u rychl´y - ˇz´ad´ano trhem kolem roku 2003 - pˇrestala se zvyˇsovat frekvence CPU (spotˇreba energie a zbytkov´e teplo u existuj´ıc´ıch technologi´ı pˇrestaly b´yt u ´nosn´e) souˇcasn´y trend - zvyˇsov´an´ı poˇ ctu jader - vhodn´e pro paralelizovateln´ eu ´lohy tradiˇcnˇe vˇetˇsina program˚ u a algoritm˚ u sekvenˇ cn´ı paraleln´ı programov´an´ı d´avno zn´am´e v high-performance computing komunitˇe - nyn´ı se pˇren´aˇs´ı i do spotˇrebitelsk´e sf´ery

Jan Hor´ aˇ cek

CUDA

3 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Historie V´ıcej´ adrov´ e zpracov´ an´ı CPU vs GPU

Historie GPU

v´yvoj grafick´eho HW: 1 2 3

specializovan´y jedno´ uˇcelov´y konfigurovateln´y programovateln´y

dnes dok´aˇze spustit t´emˇeˇr libovoln´ y algoritmus (aˇz na limity d´elky k´odu a pamˇeti) efektivn´ı pouze pro specifickou skupinu algoritm˚ u

Jan Hor´ aˇ cek

CUDA

4 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Historie V´ıcej´ adrov´ e zpracov´ an´ı CPU vs GPU

Technologie v´ıcej´adrov´eho zpracov´an´ı 1

1

Multicore soustˇred´ı se na bˇeh sekvenˇ cn´ıch program˚ u dvouo- a v´ıcej´adrov´e procesory (ˇr´adovˇe do nˇekolika m´alo des´ıtek) pln´ a instrukˇcn´ı sada x86 pˇr´ıklad: Core i7 - ˇctyˇri j´adra s out-of-order execution a technologi´ı hyperthreading

Jan Hor´ aˇ cek

CUDA

5 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Historie V´ıcej´ adrov´ e zpracov´ an´ı CPU vs GPU

Technologie v´ıcej´adrov´eho zpracov´an´ı 2

1

Many-core soustˇred´ı se na bˇeh paraleln´ıch program˚ u grafick´e akceler´atory → programovateln´e GPU velmi vysok´y hrub´ y v´ykon, v roce 2008 dokonce 1:10 vs CPU pˇr´ıklad: nVidia GeForce GTX 280 - 240 jader, kaˇzd´e j´adro heavily multithreaded, in-order, single-instruction issue processor, sd´ıl´ı kontrolu a instrukˇcn´ı cache se 7 dalˇs´ımi

Jan Hor´ aˇ cek

CUDA

6 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Historie V´ıcej´ adrov´ e zpracov´ an´ı CPU vs GPU

GPU vs CPU porovn´an´ı rychlosti

zdroj: nVidia CUDA Programming guide Jan Hor´ aˇ cek

CUDA

7 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Historie V´ıcej´ adrov´ e zpracov´ an´ı CPU vs GPU

GPU vs CPU porovn´an´ı ˇcipu

Jan Hor´ aˇ cek

CUDA

8 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Historie V´ıcej´ adrov´ e zpracov´ an´ı CPU vs GPU

GPU vs CPU porovn´an´ı ˇcipu 2

CPU - i7

GPU - GT200

Jan Hor´ aˇ cek

CUDA

9 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Historie V´ıcej´ adrov´ e zpracov´ an´ı CPU vs GPU

GPU vs CPU porovn´an´ı ˇcipu - pˇrehlednˇeji

Control

ALU

Cache DRAM

DRAM

CPU

GPU

Jan Hor´ aˇ cek

CUDA

10 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Historie V´ıcej´ adrov´ e zpracov´ an´ı CPU vs GPU

GPU vs CPU

CPU velk´a plocha pro cache preferuje niˇ zˇs´ı latenci

GPU rasterizaˇcn´ı algoritmy maj´ı tradiˇcnˇe koherentn´ı pˇr´ıstup do pamˇeti → skr´yv´a latenci pixel shadery jsou nav´ıc limitov´any v´ypoˇcetn´ım v´ykonem → velk´a plocha pro operace v plovouc´ı ˇr´ adov´ eˇ c´ arce preferuje pˇrenosovou rychlost nepotˇrebuje (tolik) cache

Jan Hor´ aˇ cek

CUDA

11 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Obsah

1

´ Uvod

2

CUDA

3

OpenCL

4

PhysX

5

Literatura

Jan Hor´ aˇ cek

CUDA

12 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Proˇc CUDA dˇr´ıve: GPGPU - v´ypoˇcet prob´ıhal v shaderech, data v textur´ach program´atoˇri potˇrebuj´ı jednoduˇsˇs´ı pˇr´ıstup k prostˇredk˚ um GPU 2007: architektura nVidia Tesla (G80 - GeForce 8800) obecnˇejˇs´ı model paraleln´ıho programov´an´ı hierarchie paraleln´ıch vl´aken bari´erov´a synchronizace atomick´e operace

CUDA = Compute Unified Device Architecture C for CUDA = jazyk, kter´ym se d´a programovat architektura CUDA efektivn´ı programov´an´ı ovˇsem st´ale potˇrebuje dobrou znalost HW Jan Hor´ aˇ cek

CUDA

13 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Architektura CUDA GPU organizov´ano do skupiny highly threaded streaming multiprocessors nˇekolik SM tvoˇr´ı blok (poˇcet SM v bloku r˚ uzn´y v kaˇzd´e generaci) kaˇzd´y SM obsahuje nˇekolik streaming processors G80: 1 SM = 8 SP GF100: 1 SM = 32 SP

SP v r´amci jednoho SM sd´ıl´ı kontroln´ı obvody a instrukˇcn´ı cache technologie prov´adˇen´ı k´ odu: SIMT (Single Instruction Multiple Threads)

Jan Hor´ aˇ cek

CUDA

14 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Architektura G80

GeForce 8800

Jan Hor´ aˇ cek

CUDA

15 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Architektura GF100

GeForce 480 Jan Hor´ aˇ cek

CUDA

16 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Compute capabilities 1

Jan Hor´ aˇ cek

CUDA

17 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Compute capabilities 2

Zbytek v nVidia CUDA Programming Guide Jan Hor´ aˇ cek

CUDA

18 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Datov´y paralelismus

mnoho aplikac´ı modeluj´ıc´ı re´aln´y svˇet pracuje s velk´ ym mnoˇzstv´ım dat datov´y paralelismus = na datech m˚ uˇze b´yt prov´adˇeno mnoho operac´ı soubˇ eˇ znˇ e napˇr´ıklad: n´asoben´ı matic

Jan Hor´ aˇ cek

CUDA

19 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Struktura programu jedna nebo v´ıce f´az´ı mohou b´yt puˇstˇeny na hostitelsk´em CPU nebo na GPU zdrojov´y CUDA k´ od m˚ uˇze obsahovat obˇ e ˇc´asti nvcc (nVidia C Compiler) tyto ˇc´asti oddˇeluje bˇehem kompilace podle generace GPU podpora od ANSI C aˇz po t´emˇeˇr u ´plnou podporu C++ k´od pro GPU naz´yv´an kernel kernel typicky spouˇst´ı ˇr´adovˇe tis´ıce vl´aken na rozd´ıl od CPU vl´aken jsou ˇr´adovˇe lehˇc´ı vytvoˇren´ı i pˇrep´ın´an´ı bˇehem nˇekolika cykl˚ u

Jan Hor´ aˇ cek

CUDA

20 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Struktura programu 2 kl´ıˇcov´a slova __global__, __host__, __device__ urˇcuj´ı, kde a odkud m˚ uˇze b´yt k´ od spuˇstˇen Kl´ıˇ cov´ e slovo __global__ float KernelFunc() __device__ float DeviceFunc() __host__ float HostFunc()

Bˇ eˇ z´ı na GPU GPU host

Spustiteln´ ez host GPU host

__device__ znamen´a, ˇze fce sm´ı b´yt spuˇstˇena pouze z kernelu nebo jin´e funkce na GPU __host__ a __device__ najednou znamen´a, ˇze se generuj´ı 2 verze, jedna pro CPU a druh´a pro GPU

Jan Hor´ aˇ cek

CUDA

21 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Bˇeh programu

zdroj: Wikipedia.org Jan Hor´ aˇ cek

CUDA

22 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Princip kernel˚ u

spust’ jedno vl´akno CPU k´od for(int y = 0; y < height; y++) for(int x = 0; x < width; x++) doSomething(x,y);

Jan Hor´ aˇ cek

CUDA

23 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Princip kernel˚ u

spust’ jedno vl´akno

spust’ x · y vl´aken

CPU k´od

GPU kernel

for(int y = 0; y < height; y++) for(int x = 0; x < width; x++) doSomething(x,y);

Jan Hor´ aˇ cek

1

zjisti x, y ze sv´ eho id

2

doSomething(x,y);

CUDA

23 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Uk´azka kernelu

// Definice kernelu __global__ void VecAdd(float *A, float *B, float *C) { int i = threadIdx.x; C[i] = A[i] + B[i]; } int main() { ... // Spuˇ stˇ en´ ı kernelu s N vl´ akny VecAdd<<<1, N>>>(A,B,C); }

Jan Hor´ aˇ cek

CUDA

24 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Hierarchie vl´aken

kernel je spuˇstˇen jako grid paraleln´ıch vl´aken vl´akna v gridu jsou ve dvou´ urovˇ nov´e hierarchii grid obsahuje 2D strukturu blok˚ u blok obsahuje 3D strukturu vl´aken

vˇsechny bloky v gridu stejn´e velikosti blok obsahuje max. 512 (1024) vl´aken velikosti gridu, blok˚ u a poˇcet vl´aken je zad´an pˇri spouˇstˇen´ı kernelu

Jan Hor´ aˇ cek

CUDA

25 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Hierarchie vl´aken 2

Pˇr´ıklad parametr˚ u kernelu // Nastav konfiguraci spuˇ stˇ en´ ı dim3 dimBlock(ThreadsX, ThreadsY) dim3 dimGrid(BlocksX,BlocksY) // Spust’ kernel MujKernel<<>> (param1, param2);

sch´ema hierarchie vl´aken Jan Hor´ aˇ cek

CUDA

26 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

ID vl´akna

CUDA definuje nˇekolik rozˇs´ıˇren´ı ANSI C o vestavˇen´e promˇenn´e pro identifikaci vl´aken Promˇ enn´ a threadIdx.x, threadIdx.y, threadIdx.z blockIdx.x, blockIdx.y blockDim.x, blockDim.y, blockDim.z gridDim.x, gridDim.y warpSize

Jan Hor´ aˇ cek

CUDA

v´ yznam index vl´ akna v bloku index bloku v gridu velikost bloku velikost gridu velikost warpu

27 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Synchronizace

mezivl´aknov´a komunikace pomoc´ı bari´ erov´ e synchronizace funkce __syncthreads() garantuje, ˇze vˇsechny vl´akna v bloku provedly kaˇzdou instrukci do tohoto m´ısta v pˇr´ıpadˇe if-then-else mus´ı vˇsechna vl´akna v bloku j´ıt stejnou cestou komunikace mezi bloky nen´ı moˇzn´a → bloky v libovoln´em poˇrad´ı → ˇsk´alovatelnost synchronizace mezi bloky se ˇreˇs´ı rozdˇelen´ım pr´ace do v´ıce kernel˚ u

Jan Hor´ aˇ cek

CUDA

28 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Pˇriˇrazov´an´ı vl´aken

nem´a vliv pro funkˇcnost, ale je dobr´e to zn´at pro efektivnost kaˇzd´y SM m˚ uˇze m´ıt na sobˇe aˇz 8 blok˚ u (ovlivnˇeno voln´ymi zdroji) kaˇzd´y blok rozdˇelen na jednotky o 32 vl´aknech zvan´e warp kaˇzd´y warp obsahuje 32 vl´aken s po sobˇe jdouc´ımi threadIdx maxim´alnˇe 24(48) warp˚ u na SM maxim´alnˇe 768(1536) vl´aken na SM velk´e mnoˇzstv´ı a rychl´e pˇrep´ın´an´ı vl´aken skr´yv´a operace s dlouhou latenc´ı (latency-hiding)

Jan Hor´ aˇ cek

CUDA

29 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Pamˇet’ pro kernely nˇekolik typ˚ u pamˇeti, drastick´e rozd´ıly v rychlosti per-thread registry per-thread lok´ aln´ı pamˇet’ per-block sd´ılen´ a pamˇet’ per-grid glob´ aln´ı pamˇet’ per-grid konstantn´ı pamˇet’ (pouze pro ˇcten´ı)

glob´aln´ı pamˇet’ typicky DRAM, pˇr´ıstup ˇr´adovˇe stovky cykl˚ u registry a lok´aln´ı pamˇet’ jsou na ˇcipu, pˇr´ıstup ˇr´adovˇe jednotky cykl˚ u aplikace (host) m˚ uˇze pˇresouvat data do/z glob´aln´ı a konstantn´ı pamˇeti sousedn´ı vl´akna mohou sd´ılet pˇr´ıstup do glob´aln´ı pamˇeti (memory coalescing) Jan Hor´ aˇ cek

CUDA

30 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Typov´e kvalifik´atory promˇenn´ych

Deklarace Pamˇ et’ Viditelnost autom. promˇenn´e mimo pol´ı registry vl´akno autom. pole lok´aln´ı vl´akno __shared__ sd´ılen´a blok __device__ glob´aln´ı grid __constant__ konstantn´ı grid max. velikost sd´ılen´e pamˇeti: 16KB (48KB) max. velikost konstantn´ı pamˇeti: 64KB max. poˇcet 32-bit registr˚ u na SM: 8K(32K) ’ lok´aln´ı pamˇet je pomal´ a jako glob´aln´ı Jan Hor´ aˇ cek

CUDA

31 / 53

ˇ Zivotnost kernel kernel kernel aplikace aplikace

´ Uvod CUDA OpenCL PhysX Literatura

Proˇ c CUDA Architektura Programy

Principy programov´an´ı pˇr´ıklad bˇehu programu 1 2 3 4 5

kaˇzd´e vl´akno naˇcte ˇc´ast dat z glob´aln´ı do sd´ılen´e pamˇeti __syncthreads() hlavn´ı v´ypoˇcet uloˇzen´ı v´ysledku pˇr´ıp. pˇriˇcten´ı dalˇs´ı ˇc´asti dat a pokraˇcov´an´ı v´ypoˇctu

d´at si pozor na vˇetven´ı if-then-else vˇetven´ı nezp˚ usobuje v´ykonov´y propad, pokud vˇsechna vl´akna na SM jdou stejnou cestou pokud se ovˇsem vˇetv´ı v r´amci zpracov´an´ı jednoho SM, mus´ı se spoˇc´ıtat obˇe vˇetve

Jan Hor´ aˇ cek

CUDA

32 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Obsah

1

´ Uvod

2

CUDA

3

OpenCL

4

PhysX

5

Literatura

Jan Hor´ aˇ cek

CUDA

33 / 53

´ Uvod CUDA OpenCL PhysX Literatura

OpenCL

alternativa k C for CUDA z´akladn´ı myˇslenka pˇrevzan´a z C for CUDA, v nˇekter´ych oblastech t´emˇeˇr 1:1 ekvivalence programovac´ı model pro spouˇstˇen´ı masivnˇe paraleln´ıch u ´loh na CPU, GPU, Cell, . . . podrobnˇejˇs´ı v´ybˇer v´ypoˇcetn´ıho zaˇr´ızen´ı ˇsirˇs´ı spektrum schopnost´ı zaˇr´ızen´ı dan´y algoritmus nemus´ı bˇeˇzet na kaˇzd´em HW

Jan Hor´ aˇ cek

CUDA

34 / 53

´ Uvod CUDA OpenCL PhysX Literatura

OpenCL koncepty

OpenCL kernel host program NDRange (index space) work item work group

Jan Hor´ aˇ cek

CUDA ekvivalent kernel host program grid thread block

CUDA

35 / 53

´ Uvod CUDA OpenCL PhysX Literatura

OpenCL vl´akna

OpenCL get_global_id(0) get_local_id(0) get_global_size(0) get_local_size(0)

CUDA ekvivalent blockIdx.x · blockDim.x + threadIdx.x threadIdx.x gridDim.x · blockDim.x blockDim.x

Jan Hor´ aˇ cek

CUDA

36 / 53

´ Uvod CUDA OpenCL PhysX Literatura

OpenCL pamˇet’

OpenCL global memory constant memory local memory private memory

Jan Hor´ aˇ cek

CUDA ekvivalent global memory constant memory shared memory local memory

CUDA

37 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Uk´azka OpenCL kernelu

// Definice kernelu __kernel void VecAdd(__global const float *A, __global const float *B, __global float *C) { int id = get_global_id(0); C[id] = A[id] + B[id]; }

Jan Hor´ aˇ cek

CUDA

38 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Obsah

1

´ Uvod

2

CUDA

3

OpenCL

4

PhysX

5

Literatura

Jan Hor´ aˇ cek

CUDA

39 / 53

´ Uvod CUDA OpenCL PhysX Literatura

(Velmi lehk´y) u´vod do PhysX

HW akcelerovan´e rozhran´ı pro simulaci fyzik´aln´ıch dˇej˚ u ve virtu´aln´ıch svˇetech v´yborn´e pro hern´ı simulace pro spolehliv´e stabiln´ı v´ypoˇcty n´ızk´a pˇresnost, ovˇsem i tak obˇcas pouˇziteln´e

nejˇcastˇeji vyuˇz´ıvan´e pro kolize ˇc´astic a jednoduch´ych tˇeles

Jan Hor´ aˇ cek

CUDA

40 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Historie ˇ ycarsk´a firma NovodeX AG zaˇcala vyv´ıjet platformu pro Sv´ v´ypoˇcet fyziky jako konkurenci pro Havok 2004 - akvizice firmou Ageia, NovodeX Physics SDK se stalo z´akladem pro PhysX 2.x SDK dodnes mnoho p˚ uvodn´ıch lid´ı z NovodeXu na v´yvoji PhysX

Ageia vyvinula HW platformu pro akceleraci fyziky PhysX PPU (Physics Processing Unit)

2008 - akvizice firmou nVidia implementace na architektuˇre CUDA postupn´e odstranˇen´ı podpory Ageia PPU podpora ve stovk´ach her siln´y marketing - c´ılen´e odstranˇen´ı podpory na ciz´ıch GPU

Jan Hor´ aˇ cek

CUDA

41 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Souˇcasnost CPU implementace bˇeˇz´ı tedy ”vˇsude” ve hr´ach ˇcasto simulace nˇekolika m´alo hlavn´ıch objekt˚ u, aby to zvl´adalo i CPU pˇri pouˇzit´ı GPU pˇrid´an´ı ˇc´asticov´eho ”nepoˇr´adku”, lepˇs´ı interakce s vodou, v´ıce dynamick´ych objekt˚ u (bez z´avislosti na gameplay)

pro GPU akceleraci nutn´ych alespoˇ n 32 CUDA jader a 256MB pamˇeti dostupn´e nejen na PC, ale i na PS3, Xbox 360, Wii fyzik´aln´ı syst´emy vˇseobecnˇe: aktu´alnˇe snahy o dobrou implementaci na mobiln´ı zaˇr´ızen´ı iOS Android Jan Hor´ aˇ cek

CUDA

42 / 53

´ Uvod CUDA OpenCL PhysX Literatura

PhysX 2.x rigid body dynamics (mechanika tuh´ych tˇeles) kolizn´ı primitiva - kapsule, koule, kv´adr, rovina, v´yˇskov´a mapa, konvexn´ı tˇeleso) r˚ uzn´e typy kloub˚ u ragdoll, materi´aly, tˇren´ı, . . .

deformovateln´a tˇelesa simulace l´atky (textilie) trh´an´ı, kolize sama se sebou, . . .

ˇc´astice a kapaliny jedno- a oboustrann´a interakce s l´atkami

dynamika vozidel objemov´a silov´a pole Jan Hor´ aˇ cek

CUDA

43 / 53

´ Uvod CUDA OpenCL PhysX Literatura

PhysX 3.x vylepˇsen´y syst´em kloub˚ u neuniformn´ı ˇsk´alov´an´ı tˇeles stabiln´ı ”depenetrace” nov´y solver na simulaci l´atek mnoho zlepˇsen´ı v´ykonu vylepˇsen´e cachov´an´ı pro PC a Xbox360 zlepˇsen´y multithreading

nov´y model simulace vozidla motor, pˇrevodovka, pneumatiky, kola, tlumiˇce, . . .

zmˇ eny a vyˇ ciˇstˇ en´ı API ... Jan Hor´ aˇ cek

CUDA

44 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Pˇr´ıklad implementace PhysX 2.x Inicializace // Initialize PhysicsSDK gPhysicsSDK = NxCreatePhysicsSDK(NX_PHYSICS_SDK_VERSION); if (!gPhysicsSDK) return; // Set the debug visualization parameters gPhysicsSDK->setParameter(NX_VISUALIZATION_SCALE, 1); gPhysicsSDK->setParameter(NX_VISUALIZE_COLLISION_SHAPES, 1); gPhysicsSDK->setParameter(NX_VISUALIZE_ACTOR_AXES, 1); // Set scale dependent parameters NxReal scale = 1.0f; // scale is meters per PhysX units gPhysicsSDK->setParameter(NX_SKIN_WIDTH, 0.05*(1/scale)); // ... other parameters initialization

Jan Hor´ aˇ cek

CUDA

45 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Pˇr´ıklad implementace PhysX 2.x Vytvoˇren´ı sc´eny // Create a scene NxSceneDesc sceneDesc; sceneDesc.gravity = NxVec3(0.0f, -9.81f, 0.0f); gScene = gPhysicsSDK->createScene(sceneDesc); if(gScene == NULL) return false; // Set default material NxMaterial* defaultMaterial = gScene->getMaterialFromIndex(0); defaultMaterial->setRestitution(0.0f); defaultMaterial->setStaticFriction(0.5f); defaultMaterial->setDynamicFriction(0.5f); // Create ground plane NxPlaneShapeDesc planeDesc; NxActorDesc actorDesc; actorDesc.shapes.pushBack(&planeDesc); gScene->createActor(actorDesc); Jan Hor´ aˇ cek

CUDA

46 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Pˇr´ıklad implementace PhysX 2.x Vytvoˇren´ı objektu // Create body NxBodyDesc bodyDesc; bodyDesc.angularDamping = 0.5f; if(initialVelocity) bodyDesc.linearVelocity = *initialVelocity; NxBoxShapeDesc boxDesc; boxDesc.dimensions = NxVec3((float)size, (float)size, (float)size); NxActorDesc actorDesc; actorDesc.shapes.pushBack(&boxDesc); actorDesc.body = &bodyDesc; actorDesc.density = 10.0f; actorDesc.globalPose.t = pos; gScene->createActor(actorDesc);

Jan Hor´ aˇ cek

CUDA

47 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Pˇr´ıklad implementace PhysX 2.x

Asynchronn´ı simulace // Start simulation (non blocking) gScene->simulate(1.0f/60.0f);

ˇ an´ı na v´ysledek Cek´ gScene->flushStream(); gScene->fetchResults(NX_RIGID_BODY_FINISHED, true);

Jan Hor´ aˇ cek

CUDA

48 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Pˇr´ıklad implementace PhysX 2.x

”Vykreslen´ı” int nbActors = gScene->getNbActors(); NxActor** actors = gScene->getActors(); while(nbActors--) { NxActor* actor = *actors++; // Render actor float glMat[16]; actor->getGlobalPose().getColumnMajor44(glMat); ... }

Jan Hor´ aˇ cek

CUDA

49 / 53

´ Uvod CUDA OpenCL PhysX Literatura

N´astroje vyˇsˇs´ı u´rovnˇe nVidia vytvoˇrila n´astroje vyˇsˇs´ı u ´rovnˇe - APEX 2 ˇc´asti tvorba - samostatn´e programy a pluginy pro modelov´an´ı runtime - pro jednoduchou implementaci do vlastn´ıho engine

APEX Clothing simulace kompletn´ıho obleˇcen´ı pouˇz´ıvan´e nejen ve hr´ach, ale i u n´avrh´aˇr˚ u atd.

APEX Destruction tvorba komplikovan´ych zniˇciteln´ych objekt˚ u napˇr. moˇznost rozpadnut´ı zdi na cihly

APEX Particles kompletn´ı ˇc´asticov´e efekty, kouˇr, kapaliny

APEX Turbulence kouˇr, prach, ˇc´astice s turbulentn´ım chov´an´ım eulerovsk´y solver kapalin na mˇr´ıˇzce

APEX Vegetation - ? Jan Hor´ aˇ cek

CUDA

50 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Alternativy

Havok (Havok Inc. + Intel) Open Dynamics Engine Newton Game Dynamics Physics Abstraction Layer (AMD . . . ?) ...

Jan Hor´ aˇ cek

CUDA

51 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Obsah

1

´ Uvod

2

CUDA

3

OpenCL

4

PhysX

5

Literatura

Jan Hor´ aˇ cek

CUDA

52 / 53

´ Uvod CUDA OpenCL PhysX Literatura

Literatura David B. Kirk, Wen-mei W. Hwu: Programming massive parallel processors, Morgan Kaufman, 2010, ISBN:978-0-12-381472-2 Jason Sanders, Edward Kandrot: CUDA by Example: An Introduction to General-Purpose GPU Programming, Addison-Wesley, 2010, ISBN:978-0131387683 nVidia Corporation: (CUDA C/OpenCL) Programming Guide nVidia Corporation: (CUDA C/OpenCL) Best Practices http://developer.nvidia.com http://www.opencl.org http://www.ati.com Jan Hor´ aˇ cek

CUDA

53 / 53