
#include "cg.h"

void vecset(int n, double v[], int iv[], int *nzv, int i, double val);
void sprnvc(int n, int nz, int nn1, double v[], int iv[],
                   double *tran, double amult);

//////////////////////////////////////////////////////////////////////////
// Kernels for makea()
//////////////////////////////////////////////////////////////////////////
__kernel void makea_0(__global int *arow,
                      __global int *g_acol,
                      __global double *g_aelt,
                      __global int *g_ilow,
                      __global int *g_ihigh,
                      int n,
                      int nn1,
                      double tran,
                      double amult)
{
  __global int (*acol)[NONZER+1];
  __global double (*aelt)[NONZER+1];

  int iouter, ivelt, nzv;
  int ivc[NONZER+1];
  double vc[NONZER+1];
  double temp_tran = tran;

  int myid, num_threads, ilow, ihigh;
  int work;

  acol = (__global int (*)[NONZER+1])g_acol;
  aelt = (__global double (*)[NONZER+1])g_aelt;

  //---------------------------------------------------------------------
  // Generate nonzero positions and save for the use in sparse.
  //---------------------------------------------------------------------
  num_threads = get_global_size(0);
  myid = get_global_id(0);

  work  = (n + num_threads - 1)/num_threads;
  ilow  = work * myid;
  ihigh = ilow + work;
  if (ihigh > n) ihigh = n;

  g_ilow[myid] = ilow;
  g_ihigh[myid] = ihigh;

  for (iouter = 0; iouter < ihigh; iouter++) {
    nzv = NONZER;
    sprnvc(n, nzv, nn1, vc, ivc, &temp_tran, amult);
    if (iouter >= ilow) {
      vecset(n, vc, ivc, &nzv, iouter+1, 0.5);
      arow[iouter] = nzv;
      for (ivelt = 0; ivelt < nzv; ivelt++) {
        acol[iouter][ivelt] = ivc[ivelt] - 1;
        aelt[iouter][ivelt] = vc[ivelt];
      }
    }
  }
}

