signalfix.cu

/* 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) {

    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    int idy = threadIdx.y + blockDim.y * blockIdx.y;

    float sum = 0.0;
    
    /* 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. */
            for (int i = idx - 1; i >= 0; i--) {
                float prevSignal = signalMatrixCopy[idy * ratioSize + i];
                if (prevSignal == 999.0) break;
                if (prevSignal == -1.0) sum = sum + 1.0;
                if (prevSignal == 1.0) sum = sum - 1.0;
            }
            signalMatrix[idy * ratioSize + idx] = sum;
        }
    }
/*
    __syncthreads();

    if ( idx < ratioSize && idy < sigMatrixHeight ) {
        float signal = signalMatrixCopy[idy * ratioSize + idx];
        if (signal == 999.0) {
            signalMatrix[idy * ratioSize + idx] = sum;
        }
    }*/
}



/* 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];
    float curSignal = signalMatrixCopy[idy * ratioSize + idx];

    if ( threadIdx.y == 0 && idx < ratioSize ) {
        subMatrix[threadIdx.x] = curSignal;
    }

    if ( threadIdx.y == 0 && threadIdx.x == 0) {
        subMatrix[BLOCK_SIZE_SMALL] = 0;
    }

    __syncthreads();
    
    for (int i = threadIdx.x; i < BLOCK_SIZE_SMALL; i++) {
        float nextSignal = subMatrix[i];
        if (nextSignal == 999.0) {
            closestUnknown = i;
            break;
        }
    }

    __syncthreads();

    if (curSignal != 999.0) {
        if (closestUnknown == -1) {
            atomicAdd(&(subMatrix[BLOCK_SIZE_SMALL]), (int)(curSignal*-1));  
        } else {
            atomicAdd(&(subMatrix[closestUnknown]), (int)(curSignal*-1)); 
        }
    }

    __syncthreads();

    if (curSignal == 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;
        }
    }
}