mass_matrix.cpp

#define lid(N) ((int) get_local_id(N))
#define gid(N) ((int) get_group_id(N))
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
#if __OPENCL_C_VERSION__ < 120
#pragma OPENCL EXTENSION cl_khr_fp64: enable
#endif

__constant double const cnst[3 * 3] = { 0.6666666666666669, 0.16666666666666663, 0.16666666666666666, 0.16666666666666674, 0.16666666666666663, 0.6666666666666665, 0.16666666666666669, 0.6666666666666666, 0.16666666666666663 };
__constant double const cnst_0[3] = { 0.16666666666666666, 0.16666666666666666, 0.16666666666666666 };

__kernel void __attribute__ ((reqd_work_group_size(1, 1, 1))) loopy_kernel_and_loopy_kernel_and_tsfc_kernel_and_loopy_kernel(int const w_0_global_len, __global double const *__restrict__ w_0_global, __global int const *__restrict__ ltg_1, int const nelements, int const coords_global_len, __global double const *__restrict__ coords_global, __global int const *__restrict__ ltg_0, __global volatile double *__restrict__ A0_global, int const A0_size)
{
  double acc_i12;
  double sum_tmp_0[3];

  for (int i_init_0 = 0; i_init_0 <= -1 + A0_size; ++i_init_0)
    if (-1 + nelements >= 0)
    {
      double loopy_old_val;
      double loopy_new_val;
      do
      {
        loopy_old_val = A0_global[i_init_0];
        loopy_new_val = 0.0;
      }
      while (atom_cmpxchg((__global long *) &(A0_global[i_init_0]), *(long *) &loopy_old_val, *(long *) &loopy_new_val) != *(long *) &loopy_old_val);
    }
}

__constant double const cnst[3 * 3] = { 0.6666666666666669, 0.16666666666666663, 0.16666666666666666, 0.16666666666666674, 0.16666666666666663, 0.6666666666666665, 0.16666666666666669, 0.6666666666666666, 0.16666666666666663 };
__constant double const cnst_0[3] = { 0.16666666666666666, 0.16666666666666666, 0.16666666666666666 };

__kernel void __attribute__ ((reqd_work_group_size(2, 1, 1))) loopy_kernel_and_loopy_kernel_and_tsfc_kernel_and_loopy_kernel_0(int const w_0_global_len, __global double const *__restrict__ w_0_global, __global int const *__restrict__ ltg_1, int const nelements, int const coords_global_len, __global double const *__restrict__ coords_global, __global int const *__restrict__ ltg_0, __global volatile double *__restrict__ A0_global, int const A0_size)
{
  double acc_i12;
  double sum_tmp_0[3];

  if (-1 + A0_size >= 0 && -1 + -2 * gid(0) + -1 * lid(0) + nelements >= 0)
  {
    for (int i1_0 = 0; i1_0 <= 2; ++i1_0)
    {
      acc_i12 = 0.0;
      for (int i12 = 0; i12 <= 2; ++i12)
        acc_i12 = acc_i12 + cnst[3 * i12 + i1_0] * (cnst[3 * i12 + 2] * w_0_global[ltg_1[3 * (2 * gid(0) + lid(0)) + 2]] + cnst[3 * i12] * w_0_global[ltg_1[3 * (2 * gid(0) + lid(0))]] + cnst[3 * i12 + 1] * w_0_global[ltg_1[3 * (2 * gid(0) + lid(0)) + 1]]) * cnst_0[i12] * fabs((-1.0 * coords_global[2 * ltg_0[3 * (2 * gid(0) + lid(0))]] + coords_global[2 * ltg_0[3 * (2 * gid(0) + lid(0)) + 1]]) * (-1.0 * coords_global[2 * ltg_0[3 * (2 * gid(0) + lid(0))] + 1] + coords_global[2 * ltg_0[3 * (2 * gid(0) + lid(0)) + 2] + 1]) + -1.0 * (-1.0 * coords_global[2 * ltg_0[3 * (2 * gid(0) + lid(0))]] + coords_global[2 * ltg_0[3 * (2 * gid(0) + lid(0)) + 2]]) * (-1.0 * coords_global[2 * ltg_0[3 * (2 * gid(0) + lid(0))] + 1] + coords_global[2 * ltg_0[3 * (2 * gid(0) + lid(0)) + 1] + 1]));
      sum_tmp_0[i1_0] = acc_i12;
    }
    for (int ibf_gather_0 = 0; ibf_gather_0 <= 2; ++ibf_gather_0)
    {
      double loopy_old_val_0;
      double loopy_new_val_0;
      do
      {
        loopy_old_val_0 = A0_global[ltg_1[3 * (2 * gid(0) + lid(0)) + ibf_gather_0]];
        loopy_new_val_0 = loopy_old_val_0 + sum_tmp_0[ibf_gather_0];
      }
      while (atom_cmpxchg((__global long *) &(A0_global[ltg_1[3 * (2 * gid(0) + lid(0)) + ibf_gather_0]]), *(long *) &loopy_old_val_0, *(long *) &loopy_new_val_0) != *(long *) &loopy_old_val_0);
    }
  }
}

This has been updated after adding the ...gbarrier in between.