/**********************************************************************************************//**
 * @file	CUDA_Func_Library.cuh
 *
 * @brief	Cuda function library class.
 **************************************************************************************************/

#include "cuda.h"
#include "cuda_runtime.h"
#include "device_functions.h"
#include "device_launch_parameters.h"
//#include < math.h >

#ifndef CUDA_FUNC_LIBRARY_H
#define CUDA_FUNC_LIBRARY_H

typedef unsigned char BYTE;
typedef unsigned long long UINT64;
typedef unsigned int UINT32;

/**********************************************************************************************//**
 * @class	cudaInfo
 *
 * @brief	Class containing methods for getting info about CUDA device.
 *
 * @author	Tomas Bravenec
 **************************************************************************************************/
class cudaInfo
{
public:
	//output number of CUDA devices
	static int cudaDevCount();

	// return calculated number of CUDA cores depending on GPU architecture
	// number of cores per multiprocessor defined in nvidia documentation
	static int cudaGetCUDACores(int device);

	static char* cudaGetError(cudaError_t cudaStatus);
};

/**********************************************************************************************//**
 * @class	cudaCluster
 *
 * @brief	Support class for kmeans clustering method.
 *
 * @author	Tomas Bravenec
 **************************************************************************************************/
class cudaCluster
{
public:
	__host__ void randomizeCentroid();
	__host__ void setCentroid(BYTE blue, BYTE Green, BYTE Red);
public:
	// Center values of a cluster
	BYTE centroidRed;
	BYTE centroidBlue;
	BYTE centroidGreen;

	// Maximum number of sum of pixel colors	18 446 744 073 709 551 615
	UINT64 sumRed;
	UINT64 sumBlue;
	UINT64 sumGreen;

	// Maximum number of pixels in image					 4 294 967 295
	UINT32 count;
};

/**********************************************************************************************//**
 * @class	cudaKernelLauncher
 *
 * @brief	Class containing all methods for launching kernels and private data for some methods.
 *
 * @author	Tomas Bravenec
 **************************************************************************************************/
class cudaKernelLauncher
{
public:
	// constructor and destructor
	cudaKernelLauncher();
	~cudaKernelLauncher();

	// converts 24bpp RGB image to 8bpp grayscale image
	static cudaError_t cudaRGB2Gray(BYTE *rgb, BYTE *gray, int height, int width, int stride, int threads);

	// converts 8bpp grayscale to 24bpp grayscale using 24bpp RGB (8bpp grayscale can not be displayed)
	static cudaError_t cudaGray2RGB(BYTE *rgb, BYTE *gray, int height, int width, int stride, int threads);

	// detects edges in 8bpp grayscale image
	// type - filter used for edge detection - 0: Sobel, 1: Prewitt, 2: Kirsch
	static cudaError_t cudaEdgeDetection(BYTE *gray, BYTE *output, int height, int width, int type, int threads);

	// thresholding function for edge detection including find maximum using CUDA
	static cudaError_t cudaThresholdEdge(BYTE *image, int height, int width, double threshold, int threads);

	// thresholding function including find maximum using CUDA
	static cudaError_t cudaThreshold(BYTE *image, int height, int width, double top, double bottom, int type, int threads);

	// inverts colors of grayscale image
	static cudaError_t cudaNegative(BYTE *image, int height, int width, int threads);

	// adjusts brightness of grayscale image
	static cudaError_t cudaBrightness(BYTE *image, int height, int width, int brightness, int threads);

	// runs Kmeans clustering using CUDA
	cudaError_t cudaKmeans(BYTE *rgb, int div, int k, int maxIter, int *iterations, int height, int width, int stride, int threads);
private:
	void kmeansInit(int k);
	void kmeansInit(int k, int iterations);
	void kmeansSetMaxIterations(int iterations);
	void kmeansNewToOld();
	bool kmeansCheckConvergence();
	cudaCluster *newCluster;
	cudaCluster *oldCluster;
	int clusterCount;
	int maxIterations;
};

//
// CUDA Kernels
//
__global__ void rgbToGray_CUDA(BYTE *rgb, BYTE *gray, int height, int width, int stride);
__global__ void grayToRGB_CUDA(BYTE *rgb, BYTE *gray, int height, int width, int stride);
__global__ void edge_CUDA(BYTE *gray, BYTE *output, int height, int width, int *filterX, int *filterY);
__global__ void maxCol_CUDA(BYTE *input, BYTE *output, int height, int width);
__global__ void maxRow_CUDA(BYTE *input, BYTE *output, int height, int width);
__global__ void simpleThresholdEdge_CUDA(BYTE *image, int height, int width, BYTE max, double threshold);
__global__ void simpleThreshold_CUDA(BYTE *image, int height, int width, BYTE max, double threshold);
__global__ void binaryThreshold_CUDA(BYTE *image, int height, int width, BYTE max, double bottom, double top);
__global__ void grayThreshold_CUDA(BYTE *image, int height, int width, BYTE max, double bottom, double top);
__global__ void negative_CUDA(BYTE *image, int height, int width);
__global__ void adjustBrightness_CUDA(BYTE *image, int height, int width, int bright);
__global__ void kmeans_iteration_CUDA(BYTE *rgb, BYTE *clustered, cudaCluster *clusters, BYTE clusterCount, int height, int width, int stride);
__global__ void kmeans_recalculate_CUDA(cudaCluster *clusters, BYTE clusterCount);
__global__ void kmeans_clear_CUDA(cudaCluster *clusters, BYTE clusterCount);

#endif // CUDA_FUNC_LIBRARY_H
