Distributed Computation of Feature-Detectors for Medical Image - - PowerPoint PPT Presentation

Peter Zinterhof Scientific Computing, Salzburg University

Distributed Computation of Feature-Detectors for Medical Image Processing on GPGPU and Cell Processors Peter Zinterhof

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

Task:

• automated recognition of 'regions of interest' within medical

imaging data here: recognition of kidneys within CT-slices

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

General approach: Evolve filter-sets, which generate feature vectors that can be classified by means of a Kohonen self-organizing map. Image filtering: Filter: 4 x 4 pixels, fed into perceptron Perceptron sums 16 inputs linearly, fires if threshold  is exceeded sliding-window (convolution) each filter delivers single number → for one image aggregate 8 filters into feature vector of length 8

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

Training phase: 1.Apply eight perceptron-based filters on training images (64 x 64 pixels) 2.Sum up fire-events to create feature vectors (8-bin histograms) 3.train Kohonen map 4.Repeat steps 1-3 for many 'individuals' in parallel and compute fitness of each filter-set (recognition rate). 5.Apply evolutionary strategy onto perceptrons to evolve such perceptrons that eventually generate high recognition rates. Recall phase: 1.Apply best filter-set onto image to be classified (kidney/ non-kidney) 2.Use pre-computed Kohonen map to classify feature vector

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

Solution:

• Use of clusters of General Purpose Graphics Processing Units

(GPGPUs) and Cell processors enables training within days instead of many months.

• Two levels of parallelism

coarse grained fine grained systems level image filtering Kohonen map

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

GPU mpich2 + gbic GPU + Cell mpich2 + pvm3 + gbic

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

Samplecode GPU: image filtering dim3 dimBlock (60,8); // 480 threads dim3 dimGrid (images , 1); // 12000 blocks > 5.7 Mio.. threads → __global__ void Perceptrons (unsigned char *a1, float *weights, unsigned char *border, int *ReturnVectors_dlocal) { __shared__ float w[Channels][4][4]; __shared__ unsigned char borders[Channels+Channels][4][4]; __shared__ unsigned char a[64][64]; if (threadIdx.y==0) { // fetch data from global memory aw = &w[0][0][0]; for (i=threadIdx.x; i < (1*Channels*16); i+=60) {

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

aw[i] = weights[i]; } ue = &borders[0][0][0]; for (i=threadIdx.x; i < (2*Channels*16); i+=60) { ue[i] = border[i]; } base = blockIdx.x*(4096); uw = &a[0][0]; for (i=threadIdx.x; i < 4096; i+=60) { uw[i] = a1[base+i]; } }

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

mask=threadIdx.y; // which of the 8 filters to compute locally Counter=0; for (j=0; j < 60; j++) { sum=0.0; if ((a[j+0][threadIdx.x+0]>=borders[mask][0][0])&&(a[j+0] [threadIdx.x+0]<borders[mask+Channels][0][0])) sum+=w[mask][0][0]; if ((a[j+1][threadIdx.x+0]>=borders[mask][0][1])&&(a[j+1] [threadIdx.x+0]<borders[mask+Channels][0][1])) sum+=w[mask][0][1]; if ((a[j+2][threadIdx.x+0]>=borders[mask][0][2])&&(a[j+2] [threadIdx.x+0]<borders[mask+Channels][0][2])) sum+=w[mask][0][2]; if ((a[j+3][threadIdx.x+0]>=borders[mask][0][3])&&(a[j+3] [threadIdx.x+0]<borders[mask+Channels][0][3]))

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

sum+=w[mask][0][3]; if (sum > THETA) { Counter++; } } // next j atomicAdd ( &ReturnVectors_dlocal [(blockIdx.x*Channels)+mask] , Counter); }

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

Codesample Cell: nearest-neighbor search

// ********************************************************************* // DEMONSTRATION OF spu_cmpgt and spu_sel for a nearest-neighbor search // mymap: array of Kohonen-map vectors (1/6 of total map) // testvector: random vector, whose nearest neighbor is computed // ********************************************************************* for (i=0; i < RES*DIM; i++) { sum=(vector float){0.0,0.0,0.0,0.0}; current_nr=spu_splats (i); diff = mymap[i][0] - testvector[0]; sum = spu_madd (diff,diff, sum); diff = mymap[i][1] - testvector[1]; sum = spu_madd (diff,diff, sum); ... diff = mymap[i][7] - testvector[7]; sum = spu_madd (diff,diff, sum); mask = spu_cmpgt (localbest, sum); // if localbest > sum -> arg 1, else arg 0 is new minimum localbest=spu_sel(localbest,sum,mask); localbest_nr = spu_sel (localbest_nr,current_nr,mask); }

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

Codesample Cell: Improved barrier function Transporting nearest neighbor information 'piggy-back'-style, yielding

• approx. 14 % speedup

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

Benchmarks:

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy

Peter Zinterhof Scientific Computing, Salzburg University

Conclusions:

• Kohonen map training is 12x faster on the Cell processor than on a

single i7-core (920, 2.67 GHz)

• in conjunction with a single GPGPU (used for filtering only), Cell

reaches 50 % of a GPGPU's performance, even in a cluster of 8 PS3 consoles one additional GPGPU suffices.

• GPGPUs reach an overall speedup of 22 x over a Xeon
• combining 8 GPGPUs with 8 Cells yields a speedup of ~248x over a

single i7 core 5 minutes on a cluster of 'unconventional hardware' is → equivalent to 1 day on the CPU

Third Workshop on UnConventional High Performance Computing 2010 (UCHPC 2010), Ischia – Naples, Italy