// ***********************************************************************
//
// Demo program pro vyuku predmetu APPS (10/2021)
// Petr Olivka, katedra informatiky, FEI, VSB-TU Ostrava
// email:[email protected]
//
// Priklad pouziti CUDA technologie.
// Prevod barevneho obrazku na odstiny sede barvy
//
// ***********************************************************************

#include <cuda_runtime.h>
#include <stdio.h>

// Demo kernel pro prevod barevneho bodu na odstin sede.
__global__ void kernel_grayscale( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey )
{
	// souradnice vlakna, kontrola rozmeru obrazku
	int y = blockDim.y * blockIdx.y + threadIdx.y;
	if ( y >= sizey ) return;
	int x = blockDim.x * blockIdx.x + threadIdx.x;
	if ( x >= sizex ) return;

	uchar4 bgr = color_pic[ y * sizex + x ];

	// vsechny tri barevne slozky budou mit stejnou hodnotu
	bgr.x = bgr.y = bgr.z = bgr.x * 0.11 + bgr.y * 0.59 + bgr.z * 0.30;

	// ulozeni bodu do obrazku
	bw_pic[ y * sizex + x ] = bgr;

}




__global__ void kernel_otoc( uchar4 *color_pic, uchar4* newimg, int sizex, int sizey )
{
	// souradnice vlakna, kontrola rozmeru obrazku
	int y = blockDim.y * blockIdx.y + threadIdx.y;
	if ( y >= sizey ) return;
	int x = blockDim.x * blockIdx.x + threadIdx.x;
	if ( x > (sizex/2) ) return;

	int xDruhaPulka = ((sizex/2)-x)+(sizex/2);

	newimg[ y * sizex + x ] =color_pic[ y * sizex +xDruhaPulka ];
	newimg[ y * sizex + xDruhaPulka ] =color_pic[ y * sizex +x ];

	/*if(sizex%2 != 0)
	{
		if(x = sizex/2)
			newimg[ y * x] =color_pic[ y * +x ];
	}*/


}

__global__ void kernel_vyrez(uchar4 *in, uchar4* out, int2 size, int4 vyrez)
{
	int y = blockDim.y * blockIdx.y + threadIdx.y;
	int x = blockDim.x * blockIdx.x + threadIdx.x;

	if ( y >= size.y) return;
	if ( x >= size.x) return;
	out[size.x * y +x] = in[0];


	if ( y >= vyrez.z || y <= vyrez.y) return;
	if ( x >= vyrez.w || x <= vyrez.x) return;






/*	if ( y >= vyrez.z || y <= vyrez.y) return;
	int x = blockDim.x * blockIdx.x + threadIdx.x;
	if ( x >= vyrez.w || x <= vyrez.x) return;

	if ( y >= size.y) return;
	int x = blockDim.x * blockIdx.x + threadIdx.x;
	if ( x >= size.x) return;
*/

	int souradniceVyrezuX = x - vyrez.x;
	int souradniceVyrezuY = y - vyrez.y;


	int vyrezSizex = vyrez.w - vyrez.x;

	out[vyrezSizex * souradniceVyrezuY + souradniceVyrezuX] = in[size.x * y +x];

	//out[size.x * y +x] = in[size.x * y +x];

}

__global__ void kernel_roztrhni(uchar4 *orig, uchar4* newimg1, uchar4* newimg2, int sizex, int sizey)
{
	int y = blockDim.y * blockIdx.y + threadIdx.y;
	if ( y >= sizey ) return;
	int x = blockDim.x * blockIdx.x + threadIdx.x;
	if ( x >= sizex) return;


	if(x > sizex/2)
	{
		newimg1[ y * (sizex/2) + x - (sizex/2) ] = orig[ y * sizex + x ];
	}
	else
	{
		newimg2[ y * (sizex/2) + x ] = orig[ y * sizex + x ];
	}
}

void run_grayscale( uchar4 *original, uchar4* output, int2 size,  int4 vyrez)
{
	cudaError_t cerr;

	uchar4 *cudaIN;
	uchar4 *cudaOUT;

	cerr = cudaMalloc( &cudaIN, size.x * size.y * sizeof( uchar4 ) );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );



	cerr = cudaMalloc( &cudaOUT, (vyrez.w-vyrez.x) * (vyrez.z-vyrez.y) * sizeof( uchar4 ) );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );



/*	cerr = cudaMalloc( &cudaOUT, size.x *size.y * sizeof( uchar4 ) );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );	*/

	// prenos barevneho obrazku do videokarty
	cerr = cudaMemcpy( cudaIN, original, size.x * size.y * sizeof( uchar4 ), cudaMemcpyHostToDevice );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );




	int block = 16;
	dim3 blocks(size.x,size.y);
	dim3 threads( block, block );


	printf( "predpustenim %d %d\n", size.x, size.y);
	printf( "predpustenim %d %d, %d, %d\n", vyrez.x, vyrez.y, vyrez.w, vyrez.z);
	// vytvoreni bloku s vlakny, matice vlaken muze byt vetsi, nez samotny obrazek!
	kernel_vyrez<<< blocks, threads >>>( cudaIN, cudaOUT, size, vyrez);


/*	cerr = cudaMemcpy(output, cudaOUT, ( size.x )* (size.y) * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );	*/



	cerr = cudaMemcpy(output, cudaOUT,  (vyrez.w-vyrez.x) * (vyrez.z-vyrez.y) * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );


	cudaFree( cudaIN );
	cudaFree( cudaOUT );








	// alokace pameti ve videokarte
	/*uchar4 *cudaColorPic;
	uchar4 *cudaBWPic;
	cerr = cudaMalloc( &cudaColorPic, sizex * sizey * sizeof( uchar4 ) );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );

	cerr = cudaMalloc( &cudaBWPic, sizex * sizey * sizeof( uchar4 ) );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );

	// prenos barevneho obrazku do videokarty
	cerr = cudaMemcpy( cudaColorPic, color_pic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyHostToDevice );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );

	int block = 16;
	dim3 blocks( ((( sizex + block - 1 ) / block)/2)+1, ( sizey + block - 1 ) / block );
	dim3 threads( block, block );

	// vytvoreni bloku s vlakny, matice vlaken muze byt vetsi, nez samotny obrazek!
	kernel_otoc<<< blocks, threads >>>( cudaColorPic, cudaBWPic, sizex, sizey );

	if ( ( cerr = cudaGetLastError() ) != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );

	// prenos dat z videokarty
	cerr = cudaMemcpy( bw_pic, cudaBWPic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );




		*/

		printf( "OTOC OK" );



	// ROZTRZENI
/*	uchar4 *cudaImg1;
	cerr = cudaMalloc( &cudaImg1, sizex* sizey * sizeof( uchar4 ) );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );

	uchar4 *cudaImg2;
	cerr = cudaMalloc( &cudaImg2, sizex* sizey * sizeof( uchar4 ) );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );

		printf( "alok OK\n" );

	block = 16;
	dim3 blocks2( (( sizex + block - 1 ) / block), ( sizey + block - 1 ) / block );
	dim3 threads2( block, block );


	printf( "predpustenim\n" );
	// vytvoreni bloku s vlakny, matice vlaken muze byt vetsi, nez samotny obrazek!
	kernel_roztrhni<<< blocks2, threads2 >>>( cudaColorPic, cudaImg1, cudaImg2, sizex, sizey );

	printf( "pospusteni\n" );

	cerr = cudaMemcpy( img1, cudaImg1, (sizex/2) * sizey * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );

		printf( "ffffffff\n" );

	cerr = cudaMemcpy( img2, cudaImg2, (sizex/2) * sizey * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
	if ( cerr != cudaSuccess )
		printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );

	printf( "gggggggggg\n" );
	// uvolneni alokovane pameti ve videokarte
	cudaFree( cudaImg1 );
	cudaFree( cudaImg2 );
	printf( "hhhhhh\n" );
	cudaFree( cudaColorPic );
	cudaFree( cudaBWPic );*/
}