Skip to content
Snippets Groups Projects
signalfix.cu 3.6 KiB
Newer Older
  • Learn to ignore specific revisions
  • adisuma2's avatar
    adisuma2 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's avatar
    Michael Henderson committed
        int idy = threadIdx.y + blockDim.y * blockIdx.y;
    
        float sum = 0.0;
    
    adisuma2's avatar
    adisuma2 committed
        
        /* Bounds check: does the thread own a value in the signal matrix for
        a certain k? */
    
        if ( idx < ratioSize && idy < sigMatrixHeight ) {
    
    Michael Henderson's avatar
    Michael Henderson committed
            float signal = signalMatrixCopy[idy * ratioSize + idx];
    
    adisuma2's avatar
    adisuma2 committed
            if (signal == 999.0) {
                /* Loops backwards from the current position until it finds
                another unknown value. */
    
    Michael Henderson's avatar
    Michael Henderson committed
                for (int i = idx - 1; i >= 0; i--) {
    
    Michael Henderson's avatar
    Michael Henderson committed
                    float prevSignal = signalMatrixCopy[idy * ratioSize + i];
    
    adisuma2's avatar
    adisuma2 committed
                    if (prevSignal == 999.0) break;
    
    Michael Henderson's avatar
    Michael Henderson committed
                    if (prevSignal == -1.0) sum = sum + 1.0;
                    if (prevSignal == 1.0) sum = sum - 1.0;
    
    adisuma2's avatar
    adisuma2 committed
                signalMatrix[idy * ratioSize + idx] = sum;
    
    adisuma2's avatar
    adisuma2 committed
    /* Shared Memory Function */
    __global__ void signalFixShared(float* signalMatrix, float* signalMatrixCopy, int ratioSize, int sigMatrixHeight) {
    
    adisuma2's avatar
    adisuma2 committed
        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;
            }
    
    adisuma2's avatar
    adisuma2 committed
        }
    
        __syncthreads();
    
    
        if ( idx < ratioSize && idy < sigMatrixHeight ) {
    
            float signal = signalMatrixCopy[idy * ratioSize + idx];
    
    adisuma2's avatar
    adisuma2 committed
        
    
            /* 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;
                }
    	}
    
    adisuma2's avatar
    adisuma2 committed
    
    
            __syncthreads();
    
            if (signal != 999.0) {
                if (closestUnknown == -1) {
                    atomicAdd(&(subMatrix[BLOCK_SIZE_SMALL]), (int)(signal*-1));  
                } else {
                    atomicAdd(&(subMatrix[closestUnknown]), (int)(signal*-1)); 
                }
    
    adisuma2's avatar
    adisuma2 committed
            }
    
    
            __syncthreads();
    
    adisuma2's avatar
    adisuma2 committed
    
    
            if (signal == 999.0) {
                signalMatrix[idy * ratioSize + idx] = (subMatrix[threadIdx.x] - 999.0);
            }
    
    adisuma2's avatar
    adisuma2 committed
    
    
            /*if (idx == lastThreadIdx) {
                atomicAdd(&(signalMatrix[idy * ratioSize + idx]), (int)subMatrix[BLOCK_SIZE_SMALL]);
            }*/
        }
    
    adisuma2's avatar
    adisuma2 committed
    }
    
    
    
    /* 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;
            }
        }
    }