/**
	 * n - число уравнений (строк матрицы)
	 * b - диагональ, лежащая над главной (нумеруется: [0;n-2])
	 * c - главная диагональ матрицы A (нумеруется: [0;n-1])
	 * a - диагональ, лежащая под главной (нумеруется: [1;n-1])
	 * f - правая часть (столбец)
	 * x - решение, массив x будет содержать ответ
	 */
void solveMatrix (int n, double *a, double *c, double *b, double *f, double *x)
{
	double m;
	for (int i = 1; i < n; i++)
	{
		m = a[i]/c[i-1];
		c[i] = c[i] - m*b[i-1];
		f[i] = f[i] - m*f[i-1];
	}

	x[n-1] = f[n-1]/c[n-1];

	for (int i = n - 2; i >= 0; i--)
		x[i]=(f[i]-b[i]*x[i+1])/c[i];

}

__global__ void func(int n, float *a, float *b, float *c, float *f, float *x)
{
	
	int j = blockDim.x* blockIdx.x + threadIdx.x;
	int r = j*z;
	float m;
	__shared__ float *abuf;
	__shared__ float *bbuf;
	__shared__ float *cbuf;
	__shared__ float *fbuf;
	__shared__ float *xbuf;

	for (int i = 1; i < n*s; i++)
	{
		abuf[i] = a[blockIdx.x*z + i];// запись значений из глобальной памяти в разделяемую
		__syncthreads();
	
	}
	for (int i = 0; i < (n - 2)*s; i++)
	{
		
		bbuf[i] = b[blockIdx.x*z + i];
		__syncthreads();
		
	}
	for (int i = 0; i < (n - 1)*s; i++)
	{
		
		cbuf[i] = c[blockIdx.x*z + i];
		__syncthreads();
		
	}
	for (int i = 0; i < n*s; i++)
	{
		
		fbuf[i] = f[blockIdx.x*z + i];
		__syncthreads();
		
	}
	for (int i = 0; i < n*s; i++)
	{
		
		xbuf[i] = x[blockIdx.x*z + i];
		__syncthreads();
		
	}
		for (int i = 1+r ; i < n+r ; i++) //от 0 до z
	{
		m = abuf[i] / cbuf[i - 1];
		cbuf[i] = cbuf[i] - m*bbuf[i - 1];
		fbuf[i] = fbuf[i] - m*fbuf[i - 1];
	}

	xbuf[n - 1 ] = fbuf[n - 1 ] / cbuf[n - 1];

	for (int i = n - 2 +r; i >= r; i--){
		xbuf[i] = (fbuf[i] - bbuf[i] * xbuf[i + 1]) / cbuf[i];
	}
for (int i = n - 2 +r; i >= r; i--){
		x[i]=xbuf[blockIdx.x*z + i];
	}

}

const int l = 128;
__global__ void func(int n, float *a, float *b, float *c, float *f, float *x)
{

	__shared__ float bufa[l];// разделяемая память
	__shared__ float bufb[l];
	__shared__ float bufc[l];
	__shared__ float buff[l];
	__shared__ float bufx[l];
	__shared__ float alpha[l];
	__shared__ float betta[l];
	int j = threadIdx.x;
	int r = j*n;
	float m;
	for (int i = 1; i<n*l; i++)
	{
		bufa[i] = a[blockIdx.x*n + i];// запись значений из глобальной памяти в разделяемую
	}
	for (int i = 0; i<n*l - 2; i++)
	{
		bufb[i] = b[blockIdx.x*n + i];
	}
	for (int i = 0; i<n*l - 1; i++)
	{
		bufc[i] = c[blockIdx.x*n + i];
	}
	for (int i = 0; i < n*l; i++)
	{
		buff[i] = f[blockIdx.x*n + i];
	}
	for (int i = 0; i<n*l; i++)
	{
		bufx[i] = x[blockIdx.x*n+ i];

	}

	__syncthreads();

	for (int i = 1 ; i < n ; i++)
	{
	m = bufa[i+r] / bufc[i - 1+r];
	bufc[i+r] = bufc[i+r] - m*bufb[i - 1+r];
	buff[i+r] = buff[i+r] - m*buff[i - 1+r];
	}
	__syncthreads();
	bufx[n - 1 + r] = buff[n - 1 + r] / bufc[n - 1 + r];
	for (int i = n - 1 ; i >= 0; i--){
		bufx[i] = (buff[i+r] - bufb[i+r] * bufx[i + 1+r]) / bufc[i+r];
		
	}
}