/normxcorr/trunk

#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>

#include "dict_hw.h"

#include "dict_image.h"

#include "helpers.h"

#include "local_sum.cu.h"

#include "normxcorr_hw.h"
#include "normxcorr_hw.cu.h"

#define MAX_DEVICES 16

static int device_number = 0;
static int devices[MAX_DEVICES];

int dictDetectHardware() {
    int deviceCount;
    cudaDeviceProp deviceProp;

    cudaGetDeviceCount(&deviceCount);
    if (!deviceCount) return -1;

    for (int i = 0; i < deviceCount; i++) {
        cudaGetDeviceProperties(&deviceProp, i);
        if ((deviceProp.major > 1)||((deviceProp.major == 1)&&(deviceProp.minor > 2))) {
            devices[device_number++] = i;
        }
    }

    return device_number;
}

DICTContext dictCreateContext() {
    if (!device_number) {
        if (dictDetectHardware() <= 0) return NULL;
    }

    TProcessingState *pstate = pstateInit();
    return pstate;
}

void dictDestroyContext(DICTContext ctx) {
    pstateFree(ctx);
}

int dictSetLogger(DICTLogger error_reporter, DICTLogger message_writer) {
    return reportSet(error_reporter, message_writer);
}

int dictSetup(DICTContext ps, int ncp, int corr_size, int precision, int flags) {
    int base_blocks, side_blocks;

    fftFree(ps);

    ps->ncp = ncp;
    ps->corr_size = corr_size;
    ps->precision = precision;

    ps->subimage_size = ps->corr_size * 4 + 1;
    ps->fft_size = 6 * corr_size + 1;

    if (flags&DICT_FLAGS_FIXED_FFT_SIZE) {
        ps->fft_real_size = ps->fft_size;
    } else {
        ps->fft_real_size = next_power(ps->fft_size);
    }
    
    if (flags&DICT_FLAGS_MATLAB_MODE) {
	ps->matlab_mode = 1;
    } else {
	ps->matlab_mode = 0;
    }

    ps->ncp_alloc_size = calc_alloc(ps->ncp, CP_BLOCK);
    ps->side_alloc_size = calc_alloc(ps->fft_size, SIDE_BLOCK_SIZE);

    ps->fft_alloc_size = calc_alloc(ps->fft_real_size * ps->fft_real_size, BLOCK_SIZE_1D);

    ps->lsum_size = ps->corr_size * 2 + 1;
    ps->lsum_temp_size = ps->subimage_size + 2*ps->lsum_size - 1;

    ps->lsum_short_aligned_size = calc_alloc(ps->fft_size, BLOCK_SIZE_2D);
    ps->lsum_aligned_size = calc_alloc(ps->lsum_temp_size, BLOCK_SIZE_2D);
    ps->lsum_alloc_size = calc_alloc(ps->lsum_temp_size + ps->lsum_size, BLOCK_SIZE_2D);

    side_blocks = calc_blocks(2 * ps->corr_size  + 1, SIDE_BLOCK_SIZE);
    ps->side_blocks_power = get_power(side_blocks);

    base_blocks = calc_blocks(4 * ps->corr_size  + 1, BLOCK_SIZE_1D);
    ps->base_blocks_power = get_power(base_blocks);

    return fftInit(ps);
}

int dictSetTemplatePoints(DICTContext ps, const float *points_x, const float *points_y) {
    memcpy(ps->points,                      points_x, ps->ncp * sizeof(float));
    memcpy(ps->points + ps->ncp_alloc_size, points_y, ps->ncp * sizeof(float));

    return 0;
}

int dictSetDimensions(DICTContext ps, int width, int height) {
    ps->width = width;
    ps->height = height;
    
    return 0;
}

int dictSetPointsBuffer(DICTContext ps, float *point_x, float *point_y) {
    ps->res_x = point_x;
    ps->res_y = point_y;
    
    return 0;
}

int dictSetCurrentPoints(DICTContext ps, const float *points_x, const float *points_y) {
    memcpy(ps->points + 2 * ps->ncp_alloc_size, points_x, ps->ncp * sizeof(float));
    memcpy(ps->points + 3 * ps->ncp_alloc_size, points_y, ps->ncp * sizeof(float));

    ps->stored = 0;

    return 0;
}

int dictCompute(DICTContext ps) {
    return fftGetCurrentPoints(ps);
}

int dictGetCurrentPoints(DICTContext ps, float *res_x, float *res_y) {
    int err;

    err = fftGetCurrentPoints(ps);
    if (err) return err;

    if ((res_x)&&(res_x != ps->res_x)) {
        float *data_x;
	if (ps->stored) data_x = ps->res_x;
	else data_x = ps->points + 2 * ps->ncp_alloc_size;
	
	memcpy(res_x, data_x, ps->ncp * sizeof(float));
    }

    if ((res_y)&&(res_y != ps->res_y)) {
	float *data_y;
        if (ps->stored) data_y = ps->res_y;
        else data_y = ps->points + 3 * ps->ncp_alloc_size;

	memcpy(res_y, data_y, ps->ncp * sizeof(float));
    }
    
    return 0;
}


int dictLoadTemplateFragment(DICTContext ps, int icp, int ncp, const unsigned char *img) {
    return fftLoadBaseFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), img);
}

int dictLoadTemplateImage(DICTContext ps, const unsigned char *img, int width, int height) {
    int err;
    
    ps->width = width;
    ps->height = height;

    int size = 2 * ps->corr_size + 1;
    int size2 = size * size;

    int base_size = 4 * ps->corr_size + 1;
    int base_size2 = base_size * base_size;

    if (width * height > ps->ncp * size2) {
        ps->mode = 0;
    } else {
        ps->mode = 1;
    }

    // if not enoguh space for caching enable anyway ?
    if (width * height > ps->ncp * base_size2) {
        ps->base_mode = 0;
    } else {
        ps->base_mode = 1;
        if (!ps->mode) {
            ps->minx = 0;
            ps->maxx = width - 1;
            ps->miny = 0;
            ps->maxy = height - 1;
        }
    }

    for (int icp = 0; icp < ps->ncp; icp+=CP_BLOCK) {
        err = fftLoadBaseFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), img);
        if (err) break;
    }

    if ((ps->base_mode)&&(!ps->mode)) {
        // Correcting difference of area size between base and data images
        ps->minx += ps->corr_size;
        ps->miny += ps->corr_size;
        ps->maxx -= ps->corr_size;
        ps->maxy -= ps->corr_size;

        width = ceil(ps->maxx) - floor(ps->minx);
        height = ceil(ps->maxy) - floor(ps->miny);

//	    printf("%ux%u=%u %u\n", width, height, width*height, ps->ncp * size2);
        if (width * height < ps->ncp * size2) {
            ps->mode = 1;
        }
    }

    if (ps->mode) {
        reportMessage("Running in the image mode");
    } else {
        reportMessage("Running in the fragment mode");
    }

    return 0;
}


int dictLoadFragment(DICTContext ps, int icp, int ncp, const unsigned char *input) {
    int err;
    cudaStream_t stream = NULL;
    
        err = fftCopyFragment(ps, icp, ncp, input);
        if (err) return err;

        err = fftLoadFragment(ps, icp, ncp, input, stream);
        if (err) return err;

        err = fftPreprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
        if (err) return err;

        err = fftProcessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
        if (err) return err;

        err = fftPostprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
        if (err) return err;
	
	return 0;
}

int dictLoadImage(DICTContext ps, unsigned char *img) {
    int err;
    int ncp = ps->ncp;

/*
#ifdef DICT_HW_MEASURE_TIMINGS
    int time[16];
    struct timeval tv1, tv2;
    gettimeofday(&tv1, NULL);
    gettimeofday(&tv2, NULL);
    time[0] = (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
    printf("Pre: %li, Nope: %li, Comp: %li, Post: %li, Copy/Load: %li, -: %li\n", time[0], time[1], time[2], time[3], time[4], time[5]);

    cudaStream_t stream[2];
    for (int i = 0; i < 2; ++i) {
        cudaStreamCreate(&stream[i]);
    }
    for (int i = 0; i < 2; ++i) {
        cudaStreamDestroy(stream[i]);
    }
#endif
*/

    for (int icp = 0; icp < ncp; icp+=CP_BLOCK) {
	err = dictLoadFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), img);
	if (err) return err;
    }
    
/*
        err = fftCopyFragment(ps, icp, ncp, input);
        if (err) return err;

        err = fftLoadFragment(ps, icp, ncp, input, stream);
        if (err) return err;

        err = fftPreprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
        if (err) return err;

        err = fftProcessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
        if (err) return err;

        err = fftPostprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
        if (err) return err;
    }
*/


    /*
    	err = fftCopyFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), input, stream[0], NULL);
    	err = fftLoadFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), input, stream[0], NULL);
    	for (int i = 0; icp < ps->ncp; icp+=CP_BLOCK,i++) {
    	    err = fftPreprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream[i%2], NULL);
    	    err = fftComputeFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream[i%2], NULL);
    	    err = fftPostprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream[i%2], NULL);
    	    if (next_icp < ps->ncp) {
    		err = fftCopyFragment(ps, next_icp, min2(CP_BLOCK, ps->ncp - next_icp), input, stream[(i+1)%2], NULL);
    		err = fftLoadFragment(ps, next_icp, min2(CP_BLOCK, ps->ncp - next_icp), input, stream[(i+1)%2], NULL);
    	    }
    	}

    */
    return 0;
}


int dictProcessImage(DICTContext ps, unsigned char *img) {
    int err = dictLoadImage(ps, img);
    if (err) return err;
    return dictCompute(ps);
}

int dictLoadTemplateImageFile(DICTContext ps, const char *name) {
    int err = dictImageLoadTemplateImage(ps, name);
    if (err) return err;
    return dictLoadTemplateImage(ps, ps->image, ps->width, ps->height);
}
int dictLoadImageFile(DICTContext ps, const char *name) {
    int err = dictImageLoadImage(ps, name);
    if (err) return err;
    return dictLoadImage(ps, ps->image);
}

int dictProcessImageFile(DICTContext ps, const char *name) {
    int err = dictLoadImageFile(ps, name);
    if (err) return err;
    return dictCompute(ps);
}