Commit 43a50dd3 authored by Ferhat Mihoubi's avatar Ferhat Mihoubi
Browse files

- adds the kernel CUDA file for SparseMatrix

parent dfa744fc
/*
* Copyright (C) 2007-2012 Dynare Team
*
* This file is part of Dynare.
*
* Dynare is free software: you can redistribute it and/or modify
* it under the terms of the GNU General Public License as published by
* the Free Software Foundation, either version 3 of the License, or
* (at your option) any later version.
*
* Dynare is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU General Public License for more details.
*
* You should have received a copy of the GNU General Public License
* along with Dynare. If not, see <http://www.gnu.org/licenses/>.
*/
#ifndef SPARMATRIX_KERNEL
#define SPARMATRIX_KERNEL
// Kernel definition of vector division
__global__ void
VecDiv(double* A, double* B, double* C, int n)
{
int i = blockIdx.x * 1024 + threadIdx.x;
if (i < n)
C[i] = (B[i] != 0.0 ? A[i] / B[i] : A[i]);
}
__global__ void
VecAdd(double* res, double* r, double alpha, double* x, int n)
{
int i = blockIdx.x * 1024 + threadIdx.x;
if (i < n)
res[i] = r[i] + alpha * x[i];
}
__global__ void
VecInc(double* res, double alpha, double* x, int n)
{
int i = blockIdx.x * 1024 + threadIdx.x;
if (i < n)
res[i] += alpha * x[i];
}
__global__ void
update_x(double* x, double alpha, double* y, double omega, double *z)
{
int i = threadIdx.x;
x[i] += alpha * y[i] + omega * z[i];
}
__global__ void
Get_LU_dim(int *n, int* A_tild_i, int *A_tild_p, int *nnz_l, int *nnz_u)
{
nnz_u[0] = 0;
nnz_l[0] = 0;
for (int i = 0; i < n[0]; i++)
{
for (int j = A_tild_p[i]; j < A_tild_p[i+1]; j++)
{
if (A_tild_i[j] < i)
nnz_l[0]++;
else if (A_tild_i[j] == i)
{
nnz_u[0]++;
//nnz_l[0]++;
}
else
nnz_u[0]++;
}
}
}
__global__ void
Get_LU1_dim(int* n, int *nnz_l, int *nnz_u)
{
nnz_u[0] = 3+n[0];
nnz_l[0] = 1+n[0];
}
__global__ void
Get_L_and_U(int *n, double* A_tild_x, int* A_tild_i, int *A_tild_p, double* Lx, int* Li, int *Lp, double* Ux, int* Ui, int* Up)
{
int nnz_u = 0, nnz_l = 0;
Lp[0] = 0;
Up[0] = 0;
for (int i = 0; i < n[0]; i++)
{
for (int j = A_tild_p[i]; j < A_tild_p[i+1]; j++)
{
if (A_tild_i[j] < i)
{
Lx[nnz_l] = A_tild_x[j];
Li[nnz_l] = A_tild_i[j];
nnz_l++;
}
else if (A_tild_i[j] == i)
{
Ux[nnz_u] = A_tild_x[j];
Lx[nnz_l] = 1.0;
Li[nnz_l] = Ui[nnz_u] = A_tild_i[j];
nnz_u++;
//nnz_l++;
}
else
{
Ux[nnz_u] = A_tild_x[j];
Ui[nnz_u] = A_tild_i[j];
nnz_u++;
}
}
Lp[i+1] = nnz_l;
Up[i+1] = nnz_u;
}
}
#endif
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment