Newer
Older
Michael Henderson
committed
/* Signal Fix works by doing the following: Finds the current signal. If the signal is equal to 999.0 then
loops back over the signal matrix to find the previous unknown (unknown == 999.0). While looping, adds up all of
the values and puts them into sum. If the value is 1.0 then -1.0 is added. If -1.0 then 1.0 is added. The sum then
replaces the unknown value (999.0). */
__global__ void signalFixGlobal(float* signalMatrix, float* signalMatrixCopy, int ratioSize, int sigMatrixHeight) {
Michael Henderson
committed
int idx = threadIdx.x + blockDim.x * blockIdx.x;
Michael Henderson
committed
/* Bounds check: does the thread own a value in the signal matrix for
a certain k? */
if ( idx < ratioSize && idy < sigMatrixHeight ) {
float signal = signalMatrixCopy[idy * ratioSize + idx];
if (signal == 999.0) {
/* Loops backwards from the current position until it finds
another unknown value. */
float prevSignal = signalMatrixCopy[idy * ratioSize + i];
if (prevSignal == -1.0) sum = sum + 1.0;
if (prevSignal == 1.0) sum = sum - 1.0;
Michael Henderson
committed
}
/* Shared Memory Function */
__global__ void signalFixShared(float* signalMatrix, float* signalMatrixCopy, int ratioSize, int sigMatrixHeight) {
int idx = threadIdx.x + blockDim.x * blockIdx.x;
int idy = threadIdx.y + blockDim.y * blockIdx.y;
//int lastThreadIdx = (BLOCK_SIZE_SMALL + blockDim.x * blockDim.x);
int closestUnknown = -1;
__shared__ float subMatrix[BLOCK_SIZE_SMALL + 1];
if ( threadIdx.y == 0 && idx < ratioSize ) {
subMatrix[threadIdx.x] = signalMatrixCopy[idy * ratioSize + idx];
if (threadIdx.x == 0) {
subMatrix[BLOCK_SIZE_SMALL] = 0;
}
if ( idx < ratioSize && idy < sigMatrixHeight ) {
float signal = signalMatrixCopy[idy * ratioSize + idx];
/* Find the closest unknown location */
for (int i = threadIdx.x; i < BLOCK_SIZE_SMALL; i++) {
float nextSignal = subMatrix[i];
if (nextSignal == 999.0) {
closestUnknown = i;
break;
}
}
if (signal != 999.0) {
if (closestUnknown == -1) {
atomicAdd(&(subMatrix[BLOCK_SIZE_SMALL]), (int)(signal*-1));
} else {
atomicAdd(&(subMatrix[closestUnknown]), (int)(signal*-1));
}
if (signal == 999.0) {
signalMatrix[idy * ratioSize + idx] = (subMatrix[threadIdx.x] - 999.0);
}
/*if (idx == lastThreadIdx) {
atomicAdd(&(signalMatrix[idy * ratioSize + idx]), (int)subMatrix[BLOCK_SIZE_SMALL]);
}*/
}
}
/* Serial Function */
void serialSignalFix(int sigMatrixHeight, int vectorSize, float* hostSignalMatrixCopy, float* hostSignalMatrix) {
for (int s = 0; s < sigMatrixHeight; s++) {
for (int v = 0; v < vectorSize; v++) {
int pos = s*vectorSize + v;
float signal = hostSignalMatrixCopy[pos];
float sum = 0.0;
if (signal == 999.0) {
for (int l = pos - 1; l >= 0; l--) {
float prevSignal = hostSignalMatrixCopy[s*vectorSize + l];
if (prevSignal == 999.0) break;
if (prevSignal == -1.0) sum = sum + 1.0;
if (prevSignal == 1.0) sum = sum - 1.0;
}
}
hostSignalMatrix[pos] = sum;
}
}
}