[PPT] - C a n F P G A s c o m p e t e w i t h G P U s PowerPoint Presentation

SLIDE 1

C a n F P G A s c

m

p e t e w i t h G P U s ?

J

h

n R

m

e i n , B r a m V e e n b

e

r

G P U T e c h n

l
g

y C

n

f e r e n c e ( G T C ’ 1 9 ) M a r c h 1 8

2

1 , 2 1 9

SLIDE 2

2

O u t l i n e

e

x p l a i n F P G A

– h

a r d w a r e

F

P G A v s . G P U

– p

r

g

r a m m i n g m

d

e l s ( O p e n C L )

– c

a s e s t u d i e s

m

a t r i x m u l t i p l i c a t i

n
r

a d i

a

s t r

n
m

i c a l i m a g i n g

– l

e s s

n

s l e a r n e d

a

n s w e r t h e q u e s t i

n

i n t h e t i t l e a n a l y z e p e r f

r

m a n c e & e n e r g y e fg i c i e n c y

SLIDE 3

3

W h a t i s a F i e l d

P

r

g

r a m m a b l e G a t e A r r a y ( F P G A ) ?

c
n

f i g u r a b l e p r

c

e s s

r
c
l

l e c t i

n
f

– m

u l t i p l i e r s / a d d e r s

– r

e g i s t e r s

– m

e m

r

i e s

– l

g

i c

– t

r a n s c e i v e r s

– c

l

c

k s

– i

n t e r c

n

n e c t

SLIDE 4

4

W h a t i s a n F P G A p r

g

r a m ?

c
n

n e c t e l e m e n t s

– p

e r f

r

m s f i x e d f u n c t i

n
d

a t a

f

l

w

e n g i n e

H

a r d w a r e D e s c r i p t i

n

L a n g u a g e ( H D L )

– V

e r i l

g

, V H D L

– d

i fg i c u l t

SLIDE 5

5

F P G A v s G P U

F

P G A a d v a n t a g e s

– h i g h e n e r g y e fg i c i e n c y

n
i

n s t r u c t i

n

d e c

d

i n g e t c .

u

s e / m

v

e a s f e w b i t s a s p

s

s i b l e – c

n

f i g u r a b l e I / O

h

i g h b a n d w i d t h

e

. g . , m a n y t i m e s 1 G b E

G

P U a d v a n t a g e s

– e a s i e r t

p

r

g

r a m – m

r

e f l e x i b l e – s h

r

t c

m

p i l a t i

n

t i m e s

SLIDE 6

6

N e w I n t e l ( A l t e r a ) F P G A t e c h n

l
g

i e s

1 )h i g h

l

e v e l p r

g

r a m m i n g l a n g u a g e ( O p e n C L ) 2 )h a r d F l

a

t i n g

P
i

n t U n i t s 3 )t i g h t i n t e g r a t i

n

w i t h C P U c

r

e s

➔ s

i m p l e t a s k s : l e s s p r

g

r a m m i n g e fg

r

t t h a n H D L

➔ a

l l

w

s c

m

p l e x H P C a p p l i c a t i

n

s

SLIDE 7

7

F P G A ≠ G P U

h

a r d w a r e – G P U : u s e h a r d w a r e – F P G A : c r e a t e h a r d w a r e

e

x e c u t i

n

m

d

e l – G P U : i n s t r u c t i

n

s – F P G A : d a t a f l

w

SLIDE 8

8

C

m

m

n

l a n g u a g e : O p e n C L

O

p e n C L – s i m i l a r t

C

U D A – C + e x p l i c i t p a r a l l e l i s m + s y n c h r

n

i z a t i

n

+ s

fu

w a r e

m

a n a g e d c a c h e – o fg l

a

d t

G

P U / F P G A

G

P U p r

g

r a m s n

t

s u i t a b l e f

r

F P G A s – F P G A : d a t a

f

l

w

e n g i n e

SLIDE 9

9

F P G A : d a t a

f

l

w

p i p e l i n e e x a m p l e

c

r e a t e h a r d w a r e f

r

c

m

p l e x m u l t i p l y

a

d d

C.real += A.real * B.real; C.real += -A.imag * B.imag; C.imag += A.real * B.imag; C.imag += A.imag * B.real;

n

e e d s f

u

r F P U s

n

e w i n p u t e n t e r s e v e r y c y c l e

B A A C

r

C

i r i

B

r i

C

r

C

i

F P U F P U F P U F P U

SLIDE 10

1

L

c

a l m e m

r

y

G

P U : u s e m e m

r

y , F P G A : c r e a t e m e m

r

y f l

a

t t m p [ 1 2 8 ] _ _ a t t r i b u t e _ _ ( ( …) ) ;

p

r

p

e r t i e s : – r e g i s t e r s

r

m e m

r

y b l

c

k s – # b a n k s – b a n k w i d t h – b a n k a d d r e s s s e l e c t i

n

b i t s a t t r i b u t e s – # r e a d p

r

t s – # w r i t e p

r

t s – s i n g l e / d

u

b l e p u m p e d – w i t h / w i t h

u

t a r b i t e r – r e p l i c a t i

n

f a c t

r
w

e l l d e s i g n e d p r

g

r a m : f e w r e s

u

r c e s , s t a l l

f

r e e

SLIDE 11

1 1

R u n n i n g m u l t i p l e k e r n e l s

G

P U : c

n

s e c u t i v e l y

F

P G A : c

n

c u r r e n t l y – c h a n n e l s ( O p e n C L e x t e n s i

n

) c h a n n e l f l

a

t 2 m y _ c h a n n e l _ _ a t t r i b u t e _ _ ( ( d e p t h ( 2 5 6 ) ) ) ; – r e q u i r e s l e s s m e m

r

y b a n d w i d t h – a l l k e r n e l s r e s i d e n t → m u s t f i t

G P U : F P G A :

c h a n n e l d e v i c e me mo r y k e r n e l _ 1 k e r n e l _ 1 k e r n e l _ 2 k e r n e l _ 2

SLIDE 12

1 2

P a r a l l e l c

n

s t r u c t s i n O p e n C L

w
r

k g r

u

p s , w

r

k i t e m s ( C U D A : t h r e a d b l

c

k s , t h r e a d s ) – G P U : p a r a l l e l – F P G A : d e f a u l t :

n

e w

r

k i t e m / c y c l e ; n

t

u s e f u l

F

P G A : – c

m

p i l e r a u t

p

a r a l l e l i z e s w h

l

e k e r n e l – # p r a g m a u n r

l

l

c

r e a t e h a r d w a r e f

r

c

n

c u r r e n t l

p

i t e r a t i

n

s – r e p l i c a t e k e r n e l s

1 2 # p r a g ma u n r

l

l a [ i ] = i ; f

r

( i n t i = ; i < 3 ; i + + )

B A A C

r

C

i r i

B

r i

C

r

C

i

F P U F P U F P U F P U

SLIDE 13

1 3

M

r

e G P U / F P G A d i fg e r e n c e s

c
d

e

u

t s i d e c r i t i c a l l

p

: – G P U : n

t

a n i s s u e – F P G A : c a n u s e m a n y r e s

u

r c e s

S

h i fu r e g i s t e r s – F P G A : s i n g l e c y c l e – G P U : e x p e n s i v e

init_once(); for (int i = 0; i < 100000; i ++) do_many_times(); for (int i = 256; i > 0; i --) a[i] = a[i – 1];

SLIDE 14

1 4

R e s

u

r c e

p

t i m i z a t i

n

s

c
m

p i l e r f e e d b a c k – a fu e r m i n u t e s , n

t

h

u

r s – H T M L

b

a s e d r e p

r

t s

i

n d i s p e n s i b l e t

l

SLIDE 15

1 5

F r e q u e n c y

p

t i m i z a t i

n
C
m

p i l e r ( p l a c e & r

u

t e ) d e t e r m i n e s F

m a x

– U

n l i k e G P U s

F

u l l F P G A : l

n

g e s t p a t h l i m i t s F

m a x

H

D L : f i n e

g

r a i n e d c

n

t r

l
O

p e n C L :

n

e c l

c

k f

r

f u l l d e s i g n

F

P G A : 4 5 M H z ; B S P : 4 M h z ; 1 D S P : 3 5 M h z

L

i t t l e c

m

p i l e r f e e d b a c k

R

e c

m

p i l e w i t h r a n d

m

s e e d s

C
m

p i l e k e r n e l s i n i s

l

a t i

n

t

f

i n d f r e q u e n c y l i m i t e r s

– T

e d i

u

s , b u t u s e f u l

SLIDE 16

1 6

A p p l i c a t i

n

s

d

e n s e m a t r i x m u l t i p l i c a t i

n
r

a d i

a

s t r

n
m

i c a l i m a g e r

S

i g n a l p r

c

e s s i n g ( F I R f i l t e r , F F T , c

r

r e l a t i

n

s , b e a m f

r

m i n g )

SLIDE 17

1 7

c
m

p l e x < f l

a

t >

h
r

i z

n

t a l & v e r t i c a l m e m

r

y a c c e s s e s ( → c

a

l e s c e )

r

e u s e

f

d a t a ( → c a c h e )

D e n s e m a t r i x m u l t i p l i c a t i

n

SLIDE 18

1 8

M a t r i x m u l t i p l i c a t i

n
n

F P G A

h

i e r a r c h i c a l a p p r

a

c h – s y s t

l

i c a r r a y

f

P r

c

e s s i n g E l e m e n t s ( P E ) – e a c h P E c

m

p u t e s 3 2 x 3 2 s u b m a t r i x

r e p e a t r e p e a t r e

r

d e r r e p e a t r e

r

d e r r e p e a t r e

r

d e r r e a d A ma t r i x r e

r

d e r

P E P E P E P E P E P E P E P E P E P E P E P E P E P E P E P E

r e a d B ma t r i x r e p e a t r e

r

d e r r e p e a t r e

r

d e r r e p e a t r e

r

d e r r e p e a t r e

r

d e r c

l

l e c t c

l

l e c t c

l

l e c t w r i t e C ma t r i x c

l

l e c t

SLIDE 19

1 9

M a t r i x m u l t i p l i c a t i

n

p e r f

r

m a n c e

A

r r i a 1 : – u s e s 8 9 %

f

t h e D S P s , 4 %

n
c

h i p m e m

r

y – c l

c

k ( 2 8 8 M H z ) a t 6 4 %

f

p e a k ( 4 5 M h z ) – n e a r l y s t a l l f r e e

T y p e D e v i c e P e r f

r

m a n c e ( T F l

p

/ s ) P

w

e r ( W ) E fg i c i e n c y ( G F l

p

/ W ) F P G A I n t e l A r r i a 1 . 7 7 4 3 7 2 . 9 G P U N V I D I A T i t a n X ( P a s c a l ) 1 . 1 2 6 3 3 8 . 4 G P U A M D V e g a F E 9 . 7 3 2 6 5 3 6 . 7

SLIDE 20

2

I m a g e

D
m

a i n G r i d d i n g f

r

R a d i

A

s t r

n
m

i c a l I m a g i n g

s

e e t a l k S 9 3 6 ( A s t r

n
m

i c a l I m a g i n g

n

G P U s )

SLIDE 21

2 1

I m a g e

D
m

a i n G r i d d i n g a l g

r

i t h m

#pragma parallel for s = 1...S : complex<float> subgrid[P ][N ×N ]; for i = 1...N ×N : float offset = compute_offset(s, i); for t = 1...T : float index = compute_index(s, i, t); for c = 1...C : float scale = scales[c]; float phase = offset - (index × scale); complex<float> phasor = {cos(phase), sin(phase)}; #pragma unroll for p = 1...P : // 4 polarizations complex<float> visibility = visibilities[t][c][p]; subgrid[p][i] += cmul(phasor, visibility); apply_aterm(subgrid); apply_taper(subgrid); apply_ifft(subgrid); store(subgrid);

SLIDE 22

2 2

F P G A g r i d d i n g d e s i g n

C

r e a t e a d a t a f l

w

n e t w

r

k :

– U

s e a l l a v a i l a b l e D S P s

– E

v e r y D S P p e r f

r

m s a u s e f u l c

m

p u t a t i

n

e v e r y c y c l e

– n

d

e v i c e m e m

r

y u s e ( e x c e p t k e r n e l a r g s )

SLIDE 23

2 3

F P G A O p e n C L g r i d d e r k e r n e l

__attribute__((max_global_work_dim(0))) __attribute__((autorun)) __attribute__((num_compute_units(NR_GRIDDERS))) __kernel void gridder() { int gridder = get_compute_id(0); float8 subgrid[NR_PIXELS]; for (unsigned short pixel = 0; pixel < NR_PIXELS; pixel ++) { | subgrid[pixel] = 0; } #pragma ivdep for (unsigned short vis_major = 0; vis_major < NR_VISIBILITIES; vis_major += UNROLL_FACTOR) { | float8 visibilities[UNROLL_FACTOR] __attribute__((register)); | | for (unsigned short vis_minor = 0; vis_minor < UNROLL_FACTOR; vis_minor++) { | | visibilities[vis_minor] = read_channel_intel(visibilities_channel[gridder]); | } | | for (unsigned short pixel = 0; pixel < NR_PIXELS; pixel++) { | | float8 pixel_value = subgrid[pixel]; | | float8 phasors = read_channel_intel(phasors_channel[gridder]); // { cos(phase), sin(phase) } | | | | #pragma unroll | | for (unsigned short vis_minor = 0; vis_minor < UNROLL_FACTOR; vis_minor++) { | | | pixel_value.even += phasors[vis_minor] * visibilities[vis_minor].even + -phasors[vis_minor] * visibilities[vis_minor].odd; | | | pixel_value.odd += phasors[vis_minor] * visibilities[vis_minor].odd + phasors[vis_minor] * visibilities[vis_minor].even; | | } | | | | subgrid[pixel] = pixel_value; | } } for (unsigned short pixel = 0; pixel < NR_PIXELS; pixel ++) { | write_channel_intel(pixel_channel[gridder], subgrid[pixel]); } }

SLIDE 24

2 4

S i n e / c

s

i n e

p

t i m i z a t i

n
c
m

p i l e r

g

e n e r a t e d : 8 D S P s

l

i m i t e d

p

r e c i s i

n

l

k

u p

t

a b l e : 1 D S P

– m

r

e D S P s a v a i l a b l e → r e p l i c a t e Φ 1 4 2 x

SLIDE 25

2 5

F P G A r e s

u

r c e u s a g e + F

m a x

A L U T s F F s R A M s D S P s M L A B s ɸ F

m a x

g r i d d i n g

i

p 4 3 % 3 1 % 6 4 % 1 4 3 9 ( 9 5 %) 7 1 % 1 4 2 5 8 d e g r i d d i n g

i

p 4 7 % 3 5 % 7 2 % 1 4 4 1 ( 9 5 %) 7 8 % 1 4 2 5 4 g r i d d i n g

l

u 2 7 % 3 2 % 6 1 % 1 4 9 8 ( 9 9 %) 5 7 % 2 2 5 6 d e g r i d d i n g

l

u 3 3 % 3 8 % 7 3 % 1 5 3 ( 9 9 %) 6 9 % 2 2 5 3

a

l m

s

t a l l

f

1 5 1 8 D S P s a v a i l a b l e u s e d

F

m a x

< 3 5 M h z

SLIDE 26

2 6

E x p e r i m e n t a l s e t u p

c
m

p a r e I n t e l A r r i a 1 F P G A t

c
m

p a r a b l e C P U a n d G P U

C

P U a n d G P U i m p l e m e n t a t i

n

s a r e b

t

h

p

t i m i z e d T y p e D e v i c e # F P U s P e a k B a n d w i d t h T D P P r

c

e s s C P U I n t e l X e

n

E 5

2

6 9 7 v 3 2 2 4 1 . 3 9 T F l

p

/ s 6 8 G B / s 1 4 5 W 2 8 n m ( T S M C ) F P G A N a l l a t e c h 3 8 5 A 1 5 1 8 1 . 3 7 T F l

p

/ s 3 4 G B / s 7 5 W 2 n m ( T S M C ) G P U N V I D I A G T X 7 5 T i6 4 1 . 3 9 T F l

p

/ s 8 8 G B / s 6 W 2 8 n m ( T S M C )

SLIDE 27

2 7

T h r

u

g h p u t a n d e n e r g y

e

fg i c i e n c y c

m

p a r i s

n
t

h r

u

g h p u t : n u m b e r

f

v i s i b i l i t i e s p r

c

e s s e d p e r s e c

n

d ( M v i s i b i l i t i e s / s )

e

n e r g y

e

fg i c i e n c y : n u m b e r

f

v i s i b i l i t i e s p r

c

e s s e d p e r J

u

l e ( M V i s i b i l i t i e s / J )

s

i m i l a r p e a k p e r f

r

m a n c e , w h a t c a u s e s t h e p e r f

r

m a n c e d i fg e r e n c e s ?

SLIDE 28

2 8

P e r f

r

m a n c e a n a l y s i s

C

P U : s i n / c

s

i n s

fu

w a r e t a k e s 8 %

f

t

t

a l r u n t i m e

G

P U : s i n / c

s
v

e r l a p p e d w i t h

t

h e r c

m

p u t a t i

n

s

F

P G A : s i n / c

s
v

e r l a p p e d , b u t c a n n

t

u s e a l l D S P s f

r

F M A s

SLIDE 29

2 9

F P G A s v s G P U s : l e s s

n

s l e a r n e d ( 1 / 3 )

d

a t a f l

w

v s . i m p e r a t i v e

– r

e t h i n k y

u

r a l g

r

i t h m

– d

i fg e r e n t p r

g

r a m c

d

e

o

n F P G A s :

– f

r e q u e n t u s e

f

O p e n C L e x t e n s i

n

s

– t

h i n k a b

u

t r e s

u

r c e u s a g e ,

c

c u p a n c y , t i m i n g

– l

e s s u s e

f
fg
c

h i p m e m

r

y

p

a r a l l e l i s m :

– b

t

h : k e r n e l r e p l i c a t i

n

a n d v e c t

r

i z a t i

n

– F

P G A : p i p e l i n i n g a n d l

p

u n r

l

l i n g

SLIDE 30

3

F P G A s v s G P U s : l e s s

n

s l e a r n e d ( 2 / 3 )

F

P G A ≠ G P U , y e t : s a m e

p

t i m i z a t i

n

s …

– e

x p l

i

t p a r a l l e l i s m

– m

a x i m i z e F P U u t i l i z a t i

n
h

i d e l a t e n c y

– o

p t i m i z e m e m

r

y p e r f

r

m a n c e

h

i d e l a t e n c y → p r e f e t c h

m

a x i m i z e b a n d w i d t h → a v

i

d b a n k c

n

f l i c t s , u n i t

s

t r i d e a c c e s s ( c

a

l e s c i n g )

r

e u s e d a t a ( c a c h i n g )

…

b u t a c h i e v e d i n v e r y d i fg e r e n t w a y s !

+

a r c h i t e c t u r e

s

p e c i f i c

p

t i m i z a t i

n

s

SLIDE 31

3 1

F P G A s v s G P U s : l e s s

n

s l e a r n e d ( 3 / 3 )

m

u c h s i m p l e r t h a n H D L

– F

P G A s a c c e s s i b l e t

w

i d e r a u d i e n c e

l
n

g l e a r n i n g c u r v e

– s

m a l l p e r f

r

m a n c e p e n a l t y

y

e t n

t

a s e a s y a s G P U s

– d

i s t r i b u t e r e s

u

r c e s i n c

m

p l e x d a t a f l

w

– o

p t i m i z e f

r

h i g h c l

c

k

t
l

s s t i l l m a t u r i n g

SLIDE 32

3 2

C u r r e n t / f u t u r e w

r

k

S

t r a t i x 1 v e r s u s V

l

t a / T u r i n g – S t r a t i x 1 i m a g e r p

r

t n

t

t r i v i a l ← d i fg e r e n t

p

t i m i z a t i

n

s f

r

n e w r

u

t i n g a r c h i t e c t u r e

1

G b E F P G A s u p p

r

t – s e n d / r e c e i v e U D P p a c k e t s i n O p e n C L – B S P f i r m w a r e d e v e l

p

m e n t

s

t r e a m i n g d a t a a p p l i c a t i

n

s – s i g n a l p r

c

e s s i n g ( f i l t e r i n g , c

r

r e l a t i

n

s , b e a m f

r

m i n g , . . . )

O

p e n C L F F T l i b r a r y

SLIDE 33

3 3

C

n

c l u s i

n

s

c
m

p l e x a p p l i c a t i

n

s

n

F P G A s n

w

p

s

s i b l e – h i g h

l

e v e l l a n g u a g e – h a r d F P U s – t i g h t i n t e g r a t i

n

w i t h C P U s

G

P U s v s F P G A s – 1 l a n g u a g e ; n

(

p e r f

r

m a n c e ) p

r

t a b i l i t y – v e r y d i fg e r e n t a r c h i t e c t u r e s – y e t m a n y

p

t i m i z a t i

n

s s i m i l a r

SLIDE 34

3 4

S

c

a n F P G A s c

m

p e t e w i t h G P U s ?

d

e p e n d s

n

a p p l i c a t i

n

– c

m

p u t e → G P U s – e n e r g y e fg i c i e n c y → G P U s – I / O → F P G A – f l e x i b i l i t y → C P U / G P U – p r

g

r a m m i n g e a s e → C P U , t h e n G P U , t h e n F P G A F P G A n

t

f a r b e h i n d ; C P U w a y b e h i n d

SLIDE 35

3 5

A c k n

w

l e d g e m e n t s

T

h i s w

r

k w a s f u n d e d b y – N e t h e r l a n d s e S c i e n c e C e n t e r ( T r i p l e

A

2 ) – E U H 2 2 F E T H P C ( D E E P

E

S T , G r a n t A g r e e m e n t n r . 7 5 4 3 4 ) – N W O N e t h e r l a n d s F

u

n d a t i

n

f

r

S c i e n t i f i c R e s e a r c h ( D A S

5

)

O

t h e r p e

p

l e – S u l e y m a n D e m i r s

y

( I n t e l ) , J

h

a n H i d d i n k ( N L e S C ) , A t z e v . d . P l

e

g ( N L e S C ) , D a n i ë l v . d . S c h u u r ( A S T R O N ) , M e r i j n V e r s t r a a t e n ( N L e S C ) , B e n v . W e r k h

v