// Karty_2.cpp : Defines the entry point for the console application.
//

#include "stdafx.h"
#include <ctime>
// Macierze są pamiętane wierszami, a więc:
// M(row, col) = *(M.elements + row * M.width + col)

typedef struct {
	int width;
	int height;
	float *elements;
} Matrix;

// definiujemy rozmiar bloku wątków:
#define BLOCK_SIZE 16

// prototyp funkcji mnożącej (kernela)
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);

//	prototyp funkcji uzupelniajaca macierz
void pseudolosowe(Matrix *pointer);

// prototyp funkcji mnożącej (CPU)
void CPU_mnozenie(const Matrix, const Matrix ,Matrix);

// Zakładamy (dla uproszczenia rozważań), że wymiary macierzy są
// całkowitymi wielokrotnościami wartości BLOCK_SIZE

int main()
{
	// przygotowujemy środowisko i wywołujemy kernel
	dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
	dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
	MatMulKernel << <dimGrid, dimBlock >> >(d_A, d_B, d_C);

	// odbieramy obliczoną macierz C z pamięci globalnej urządzenia
	cudaMemcpy(C.elements, d_C.elements, size, cudaMemcpyDeviceToHost);
	//	cudaMemcpyDeviceToHost kierunek transferu danych (pamięć źródłowa w karcie, a obszar docelowy w RAM)

	//	zwalniamy pamięć
	cudaFree(d_A.elements);
	cudaFree(d_B.elements);
	cudaFree(d_C.elements);

	return 0;
}

// Funkcja mnożąca
void MatMul(const Matrix A, const Matrix B, Matrix C)
{
	// kopiujemy macierze A i B to globalnej pamięci urządzenia

	// najpierw A
	Matrix d_A;
	d_A.width = A.width;
	d_A.height = A.height;
	size_t size = A.width * A.height * sizeof(float);						//	okreslamy jakiego potrzebujemy rozmiaru
	cudaMalloc((void **)&d_A.elements, size);								//	alokuje pamiec na karcie graficznej
	pseudolosowe(d_A);
	cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice);		//	Kopiuje dane między kartą graficzną, a pamięcią RAM
	//	cudaMemcpyHostToDevice - kierunek transferu danych (pamięć źródłowa w RAM, a obszar docelowy w karcie graficznej

	// potem B
	Matrix d_B;
	d_B.width = B.width;
	d_B.height = B.height;
	size = B.width * B.height * sizeof(float);
	cudaMalloc((void **)&d_B.elements, size);
	pseudolosowe(d_B);
	cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice);

	// przydzielamy macierz C w globalnej pamięci urządzenia
	Matrix d_C;
	d_C.width = C.width;
	d_C.height = C.height;
	size = C.width * C.height * sizeof(float);
	cudaMalloc((void **)&d_C.elements, size);
}

// kernel odpowiedzialny za wymnożenie macierzy
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
	// każdy wątek oblicza jeden element macierzy C
	// akumulując wynik w zmiennej Cvalue
	float Cvalue = 0;
	int row = blockIdx.y * blockDim.y + threadIdx.y;
	int col = blockIdx.x * blockDim.x + threadIdx.x;
	for (int e = 0; e < A.width; ++e)
		Cvalue += A.elements[row * A.width + e]
		* B.elements[e * B.width + col];
	C.elements[row * C.width + col] = Cvalue;
}

void pseudolosowe(&Matrix)
{
	srand(time(NULL));
	for (int i = 0; i < pointer->width * pointer->height; i++)
	{
		pointer->elements[i] = (((std::rand() % 100000) + 1) / 10000);	//	tak uzyskujemy wartości po przecinku
	}
}

void CPU_mnozenie(Matrix A, Matrix B, Matrix C)
{
	for (int row = 0; row < A.height; row++)
	{
		for (int col = 0; col < B.width; col++)
		{
			float cValue = 0;
			for (int e = 0; e < A.width; e++)
			{
				cValue += A.elements[row * A.width + e] * B.elements[e * B.width + col];
			}
			C->elements[row * C->width + col] = cValue;
		}
	}
}