Genetically Improved Software

Abstract

Genetic programming (GP) can dramatically increase computer programs’ performance. It can automatically port or refactor legacy code written by domain experts and specialist software engineers. After reviewing SBSE research on evolving software we describe an open source parallel StereoCamera image processing application in which GI optimisation gave a seven fold speedup on nVidia Tesla GPU hardware not even imagined when the original state-of-the-art CUDA GPGPU C++ code was written.

Appendix: StereoKernel Tuned for K20c Tesla

In addition to the complete Stereo Camera system StereoCamera_v1_1c.zip contains the following CUDA kernel. The modifications to the openVidia CUDA Stereo Camera code distributed by SourceForge are also described in Sect. 8.16 (pages 206 to 207).

/*******

stereoKernel

Now for the main stereo kernel: There are four parameters:

disparityPixel points to memory containing the disparity value (d)

for each pixel.

width & height are the image width & height, and out_pitch

specifies the pitch of the output data in words (i.e. the number

of floats between the start of one row and the start of the next.).

disparityMinSSD removed by GP

*********/

__attribute__((global)) void stereoKernel(

// pointer to the output memory for the disparity map

float * __restrict__ disparityPixel,

// the pitch (in pixels) of the output memory for the disparity

map const size_t out_pitch,

const int width,

const int height,

unsigned int * __restrict__ timer, //For GP timing only

int * __restrict__ sm_id //For GP timing only

)

{

FIXED_init_timings(timer,sm_id); //For GP timing only

extern __attribute__((shared)) float disparityPixel_S[];

int* const disparityMinSSD = (int*)&disparityPixel_S[ROWSper

THREAD*BLOCK_W];

// column squared difference functions

int* const col_ssd = &disparityMinSSD[ROWSperTHREAD*BLOCK_W];

float d; // disparity value

float d0,d1;

float dmin;

int diff; // difference temporary value

int ssd; // total SSD for a kernel

float x_tex; // texture coordinates for image lookup

float y_tex;

int row; // the current row in the rolling window

int i; // for index variable

const int dthreadIdx = threadIdx.x % BLOCK_W;

const int dblockIdx = threadIdx.x / BLOCK_W;

//bugfix force subsequent calculations to be signed

const int X = (__mul24(blockIdx.x,(BLOCK_W-2*RADIUS_H)) +

dthreadIdx);

const int ssdIdx = threadIdx.x;

int* const reduce_ssd = &col_ssd[(BLOCK_W )*dperblock-BLOCK_W];

const int Y = (__mul24(blockIdx.y,ROWSperTHREAD));

//int extra_read_val = 0; no longer used

//if(dthreadIdx < (2*RADIUS_H)) extra_read_val = BLOCK_W + ssdIdx;

// initialize the memory used for the disparity and the disparity

difference

//Uses first group of threads to initialise shared memory

if(threadIdx.x<BLOCK_W-2*RADIUS_H)

if(dblockIdx==0)

if(X<width )

{

for(i = 0;i<ROWSperTHREAD && Y+i < height;i++)

{

// initialize to -1 indicating no match

disparityPixel_S[i*BLOCK_W +threadIdx.x] = -1.0f;

//ssd += col_ssd[i+threadIdx.x];

disparityMinSSD[i*BLOCK_W +threadIdx.x] = MIN_SSD;

}

}

__syncthreads();

x_tex = X - RADIUS_H;

for(d0 = STEREO_MIND;d0 <= STEREO_MAXD;d0 += STEREO_DISP_STEP*

dperblock)

{

d = d0 + STEREO_DISP_STEP*dblockIdx;

col_ssd[ssdIdx] = 0;

// do the first row

y_tex = Y - RADIUS_V;

for(i = 0;i <= 2*RADIUS_V;i++)

{

diff = readLeft(x_tex,y_tex) - readRight(x_tex-d,y_tex);

col_ssd[ssdIdx] += SQ(diff);

y_tex += 1.0f;

}

__syncthreads();

// now accumulate the total

if(dthreadIdx<BLOCK_W-2*RADIUS_H)

if(X < width && Y < height)

{

ssd = 0;

for(i = 0;i<=(2*RADIUS_H);i++)

{

ssd += col_ssd[i+ssdIdx];

}

}

if(dblockIdx!=0) reduce_ssd[threadIdx.x] = ssd;

__syncthreads();

//Use first group of threads to set ssd to smallest SSD for

d1<d0+dperblock

if(threadIdx.x<BLOCK_W-2*RADIUS_H)

if(X < width && Y < height)

{

dmin = d;

d1 = d + STEREO_DISP_STEP;

for(i = threadIdx.x+BLOCK_W;i < blockDim.x;i += BLOCK_W) {

if(d1 <= STEREO_MAXD && reduce_ssd[i] < ssd) {

ssd = reduce_ssd[i];

dmin = d1;

}

d1 += STEREO_DISP_STEP;

}

//if ssd is smaller update both shared data arrays

if( ssd < disparityMinSSD[0*BLOCK_W +threadIdx.x])

{

disparityPixel_S[0*BLOCK_W +threadIdx.x] = dmin;

disparityMinSSD[0*BLOCK_W +threadIdx.x] = ssd;

}

}

__syncthreads();

// now do the remaining rows

y_tex = Y - RADIUS_V; // this is the row we will remove

#pragma unroll 11

for(row = 1;row < ROWSperTHREAD && (row+Y < (height+RADIUS_V));

row++)

{

// subtract the value of the first row from column sums

diff = readLeft(x_tex,y_tex) - readRight(x_tex-d,y_tex);

col_ssd[ssdIdx] -= SQ(diff);

// add in the value from the next row down

diff = readLeft(x_tex, y_tex + (float)(2*RADIUS_V)+1.0f) -

readRight(x_tex-d,y_tex + (float)(2*RADIUS_V)+1.0f);

col_ssd[ssdIdx] += SQ(diff);

y_tex += 1.0f;

__syncthreads();

if(dthreadIdx<BLOCK_W-2*RADIUS_H)

if(X<width && (Y+row) < height)

{

ssd = 0;

for(i = 0;i<=(2*RADIUS_H);i++)

{

ssd += col_ssd[i+ssdIdx];

}

}

if(dblockIdx!=0) reduce_ssd[threadIdx.x] = ssd;

__syncthreads();

//Use 1st group threads to set ssd/dmin to smallest SSD for

d1<d0+dperblock

if(threadIdx.x<BLOCK_W-2*RADIUS_H)

if(dblockIdx==0)

{

dmin = d;

d1 = d + STEREO_DISP_STEP;

for(i = threadIdx.x+BLOCK_W;i < blockDim.x;i += BLOCK_W) {

if(d1 <= STEREO_MAXD && reduce_ssd[i] < ssd) {

ssd = reduce_ssd[i];

dmin = d1;

}

d1 += STEREO_DISP_STEP;

}

//if smaller SSD found update shared memory

if(ssd < disparityMinSSD[row*BLOCK_W +threadIdx.x])

{

disparityPixel_S[row*BLOCK_W +threadIdx.x] = dmin;

disparityMinSSD[row*BLOCK_W +threadIdx.x] = ssd;

}

}//endif first group of thread

}// for row loop

}// for d0 loop

//Write answer in shared memory to global memory

if(threadIdx.x<BLOCK_W-2*RADIUS_H)

if(dblockIdx==0)

if(X < width) {

#pragma unroll 3

for(row = 0;row < ROWSperTHREAD && (row+Y < height);row++)

{

disparityPixel[__mul24((Y+row),out_pitch)+X] =

disparityPixel_S[row*BLOCK_W +threadIdx.x];

}

}

FIXED_report_timings(timer,sm_id); //For GP timing only

}