Inconsistant Values ug!

Accelerated Computing CUDA CUDA Programming and Performance
1

I have some code which is producing some inconsistant values - Most of the values are right but some arn’t when comparing to my excel spreadsheet to assert the correct values. Note that the code is written in such a way that no synchronization occurs but since the kernel makes multiple passes, inconsistancies caused by irregular block/thread execution, are hidden since fewer and fewer values are in flux as consistant values propagate across the matrix. What I don’t understand is why only some of the values and not all of them, even the last element of the array is correct so its unlikely that Im doing something obvious.

For those interested, this is the Smith-Waterman Gene alignment algorithm

[codebox]

global void SmithWaterman(int* sequenceA,

									 int* sequenceB,

									 float* dMatrix,

									 float* traceMatrix,

									 float* sMatrix,

									 float gapPenalty)

{

    ////offset for border values

//global element index

const int elementIdx = (gridDim.x * blockDim.x + 1) * (blockDim.y * blockIdx.y + 1 + threadIdx.y) + (blockDim.x * blockIdx.x + threadIdx.x + 1);

const int sharedElementIdx = blockDim.x * threadIdx.y + threadIdx.x;

//declare shared values

extern __shared__ float sharedMemory[];

float* above = &sharedMemory[0];

float* diagonal = &sharedMemory[blockDim.x * blockDim.y];

float* left = &sharedMemory[blockDim.x * blockDim.y * 2];

//copy values from global to shared, calculate indels/alignment

above[sharedElementIdx] = dMatrix[elementIdx - (gridDim.x * blockDim.x + 1)] - gapPenalty;

diagonal[sharedElementIdx] = dMatrix[elementIdx - (gridDim.x * blockDim.x + 2)] + sMatrix[sequenceB[threadIdx.y] * 4 + sequenceA[threadIdx.x]];

left[sharedElementIdx] = dMatrix[elementIdx - 1] - gapPenalty;

if(above[sharedElementIdx] > diagonal[sharedElementIdx] && above[sharedElementIdx] > left[sharedElementIdx])

	dMatrix[elementIdx] = above[sharedElementIdx];

else if(diagonal[sharedElementIdx] > above[sharedElementIdx] && diagonal[sharedElementIdx] > left[sharedElementIdx])

	dMatrix[elementIdx] = diagonal[sharedElementIdx];

else if(left[sharedElementIdx] > above[sharedElementIdx] && left[sharedElementIdx] > diagonal[sharedElementIdx])

	dMatrix[elementIdx] = left[sharedElementIdx];

}[/codebox]