(MIRU2010) NTT Graphic Processor Unit GPU graphi

Size: px

Start display at page:

Download "(MIRU2010) NTT Graphic Processor Unit GPU graphi"

あきひささわなか
6 years ago
Views:

1 (MIRU2010) NTT [email protected], [email protected] Graphic Processor Unit GPU graphic processor unit CUDA Fully automatic extraction of salient objects in near real-time Kazuma AKAMINE, Ken FUKUCHI, Akisato KIMURA, and Shigeru TAKAGI Department of Computer Science and System Engineering, Faculty of Engineering, Miyazaki University 1-1 Gakuen Kibanadai-Nishi, Miyazaki, Japan. Department of Information and Communication Systems Engineering, Okinawa National College of Technology Henoko 905, Nago, Okinawa, Japan. NTT Communication Science Laboratories, NTT Corporation Morinosato Wakamiya 3-1, Atsugi, Kanagawa, Japan. [email protected], [email protected] 1. (MRF) (MAP) MRF Greig [1] Boykov Interactive Graph Cuts [2] Kohli Dynamic Graph Cuts [3] MRF MAP [4] [5] Fu [6] [7]

2 Graphic Processor Unit GPU SIMD Single Instruction Multiple Data GPU CPU [8] [9] [10] GPU general-purpose GPU(GPGPU) [11] GPU C++ CUDA [12] OpenCL [13] GPU GPU CUDA GPU CUDA 2 GPU CUDA 3 [7] 4 GPU Graphic Processor Unit GPU GPU CUDA 2. 1 GPU(Graphics Processing Unit) GPU GPU SIMD GPU GPGPU GPU GPU NVIDIA CUDA [12] AMD ATi Stream [14] CPU GPU OpenCL [13] Microsoft API Direct Compute [15] 2. 2 CUDA GPU HLSL [15] GLSL [16] NVIDIA GPU CUDA GPU CUDA C GPU C GPU 1 ID ID CPU GPU GPU CPU GPU CPU GPU 2. 3 CUDA CUDA CPU GPU

3 2 CPU GPU 4 CUDA 1 host device int add(int a, int b) { 2 return a + b; 3 } 4 global void addall(int length, int a, int b, int result) { 5 int px = blockidx.x blockdim.x + threadidx; 6 if(px < length) { 7 result[px] = add(a[px], b[px]); 8 } 9 } 10 void main() { Dim3 blocks = {10, 0, 0}; 13 Dim3 threads = {32, 0, 0}; 14 int sm cap = 0; 15 addeach<<<blocks, threads, sm cap>>>(length, a, b, result); } 3 CUDA CPU GPU 1 1 CUDA host CPU CPU global CPU GPU device GPU GPU GPU device global void host device 3 3 GPU a b addall GPU add CPU GPU GPU addall <<<>>> 3 blocks 5 2 GPU CPU threads 1 sm cap 2. 4 CUDA CUDA GPU CUDA GPU KB syncthreads()

4 6 Saliency Graph Cuts 7 CPU GPU GPU 64KB GPU CPU GPU GPU CPU GPU CPU GPU CPU GPU 3. SGC [7] SGC 6 SGC ( 6(a)) [17] ( 6(b)) (Lucas-Kanade [18]) [19] / ( 6(c1))P r (O) P r (B) p P r (O; p) P r (O; p) 0 P r (B; p) = 1 P r (O; p) Boykov [20] 7 S T t-link n-link t-link (1)(2) n-link (5) C p p P r (C p O) P r (C p B) P r (O C p ) P r (B C p ) I p I q p q dist(p, q) p q R p ( obj ) = ln P r (O C p ) (1) R p ( bkg ) = ln P r (B C p ) (2) P r (O C p ) = P r(c p O)P r (O; p) P r (C p ) P r (B C p ) = P r(c p B)P r (B; p) P r (C p ) B {p,q} exp { (I p I q ) 2 } 1 2σ dist(p, q) (3) (4) (5) ( 6(c2)) (6)(8) (GMM) GMM d RGB 3 a k S k π k EM P r (C p O)

F (x, y) P (x, y) n m F k (i, j) (9) F (x, y) = n m F k (i, j) i=0 j=0 P (x

k, S k, π k ) = π k > = 0, M π k p k (x) (6) k=1 M π k = 1 (7) k=1 1 p k

GPU 3 GPU 8 GPU t-link n-link GPU GPU 1 1 P (x, y) F k (i, j) F (x, y) F k

10 10 minsrc maxsrc syncthreads 0 1 P r (O) GPU t-link EM CUDA ZONE [12]

5 F (x, y) P (x, y) n m F k (i, j) (9) F (x, y) = n m F k (i, j) i=0 j=0 P (x + i n 2, y + j m 2 ) (9) 8 GPU n P r (O) RGB P r (C p B) P r (O) RGB p(x; a k, S k, π k ) = π k > = 0, M π k p k (x) (6) k=1 M π k = 1 (7) k=1 1 p k (x) = (2π) d/2 S k 1/2 { exp 1 } 2 (x a k) T S 1 k (x a k) (8) ( 6(d)) 1 4. GPU 3 GPU 8 GPU t-link n-link GPU GPU 1 1 P (x, y) F k (i, j) F (x, y) F k (i, j) P (x, y) (filter) (src1) height width fheight fwidth result minsrc maxsrc syncthreads 0 1 P r (O) GPU t-link EM CUDA ZONE [12] Harp [21] EM k-means CPU GMM (6) 9 P k (x) (8) {(2π) d/2 S k 1/2 } 1 1 CPU GPU

1 texture<float, 1, cudareadmodeelementtype> filter; 2 texture<float, 1, cudareadmodeelementtype> src1; 3 device float Filter2DCore(texture<float, 1, cudareadmodeelementtype> fsource, int x, int y,

6 1 texture<float, 1, cudareadmodeelementtype> filter; 2 texture<float, 1, cudareadmodeelementtype> src1; 3 device float Filter2DCore(texture<float, 1, cudareadmodeelementtype> fsource, int x, int y, int height, int width, int filtersizex, int filtersizey) { 4 float sum = 0; 5 x = filtersizex/2; 6 y = filtersizey/2; 7 for(int fy = 0; fy < filtersizey; fy++) { 8 int by = y + fy; 9 if(by > 0 && by < height) { 10 by = width; 11 for(int fx = 0; fx < filtersizex; fx++) { 12 int bx = x + fx; 13 if(bx > 0 && bx < width) { 14 sum += tex1dfetch(filter, fy filtersizex + fx) tex1dfetch(fsource, by + bx); 15 } 16 } 17 } 18 } 19 return sum; 20 } 21 global void Filter2DKernel(int height, int width, int fheight, int fwidth, float result) { 22 int px = blockdim.x blockidx.x + threadidx.x; 23 if(px < height width) { 24 int x = px%width; 25 int y = px/width; 26 result[px] = Filter2DCore(src1, x, y, height, width, fwidth, fheight); 27 } 28 } 1 texture<float, 1, cudareadmodeelementtype> minsrc; 2 texture<float, 1, cudareadmodeelementtype> maxsrc; 3 global void SMRangeNormalizeKernel1(int length, float localmin, float localmax) { 4 int px = blockdim.x blockidx.x + threadidx.x; 5 6 shared float mini[32], maxi[32]; 7 if(px < length) { 8 mini[threadidx.x] = tex1dfetch(minsrc, px); 9 maxi[threadidx.x] = tex1dfetch(maxsrc, px); 10 } else { 11 mini[threadidx.x] = FLT MAX; 12 maxi[threadidx.x] = FLT MIN; 13 } 14 syncthreads(); 15 if(threadidx.x == 0) { 16 for(int i = 1; i < blockdim.x; i++) { 17 mini[0] = min(mini[0], mini[i]); 18 maxi[0] = max(maxi[0], maxi[i]); 19 } 20 localmin[blockidx.x] = mini[0]; 21 localmax[blockidx.x] = maxi[0]; 22 } 23 } 10 9 CUDA (8) exp { 1 2 (x a k) T S 1 k (x a k) } a k S 1 k n-link 2 1 GPU CUDA ZONE Vineet CUDA Cuts [22] CPU GPU x512 GPU GPU CPU Intel Core2Quad Q9550 4GB GPU NVIDIA Geforce 9800GT 512MB OS Windows XP Professional NVIDIA CUDA 2.1 OpenCV x t-link 4.2 n-link

7 352x x CUDA ZONE EM OpenCL!"! [ms] t- - link link 352 CPU GPU CPU GPU CPU GPU [us] CPU GPU CPU GPU GPU CPU GPU 2 GPU GPU SGC 11 4 GPU t-link t-link EM 2 CUDAZONE 6. GPU [1] D.Greig B.Porteous and A.Seheuit Exact maximum a posteriori estimation for binary images Royalstat Vol.B:51 No.2 pp [2] Y.Boykov and M-P.Jolly Interactive Graph Cuts for Optical Boundary & Region Segmentation of Objects in N-D Images Proc.ICCV Vol.I pp [3] P.Kohli and P.Torr Dynamic graph cuts for efficient inference in Markov random fields IEEE Trans.PAMI Vol.29 No.12 pp [4] AdaBoost Saliency Map Graph Cuts (MIRU2008) IS3-33 pp [5] PRMU pp [6] Y.Fu J.Cheng Z.Li and H.Lu Saliency cuts: Anautomatic approach to object segmentation Proc.ICPR [7] (MIRU2009) [8] TSUBAME ccwww/ tebiki/tesla/tesla.html [9] TMPGEnc 4.0 XPress html [10] MediaShow Espresso company/press-news-content.do?pid=2115 [11] GPU GPGPU [12] home new.html [13] [14] stream-technology/pages/stream-technology.aspx [15] ee663301(vs.85).aspx [16] [17] MCMC-based particle filter (MIRU2009) [18] B.D.Lucas and T.Kanade An Iterative Image Registration Technique with an Application to Stereo Vision Proceedings of the 7th International Joint Conference on Artificial Intelligence(IJCAI 81) pp August [19] L.Itti C.Koch and E.Niebur A model of saliencybased visual attention for rapid scene analysis IEEE Trans.PAMI Vol.20 No.11 pp November [20] Y.Boykov and G.F.Lea Graph cuts and efficient N- D image segmentation Proc.ICCV Vol.70 No.2 pp

8 [21] A.Harp Computational Statistics via GPU [22] V.Vineet and P.J.Narayanan Cuda Cuts: Fast Graph Cuts on the GPU CVPR Workshop on Visual Computer Vision on GPUs 2008.

07-二村幸孝・出口大輔.indd

07-二村幸孝・出口大輔.indd GPU Graphics Processing Units HPC High Performance Computing GPU GPGPU General-Purpose computation on GPU CPU GPU GPU *1 Intel Quad-Core Xeon E5472 3.0 GHz 2 6 MB L2 cache 1600 MHz FSB 80 GFlops 1 nvidia