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;
        }
    }
}



/* 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;
        }
    }

    __syncthreads();

    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;
            }
	}

        __syncthreads();

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

        __syncthreads();

        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;
        }
    }
}