<commit_before>#include "scatter_gridder.h"
#include "atomic_add.h"
#include "OskarBinReader.h"

#include <math_functions.h>

struct Pregridded
{
  short u;
  short v;
  char gcf_layer_w_plane;
  char gcf_layer_over;
  short gcf_layer_supp;
};

template <
    int over
  , int w_planes
  , int grid_size

  , int timesteps
  , int channels
  >
__inline__ __device__ void loadIntoSharedMem (
    double scale
  , int max_supp // needed for u v translation
  , const double3 uvw[timesteps * channels]
  , const Double4c vis[timesteps * channels]
  , Pregridded uvo_shared[timesteps * channels]
  , Double4c vis_shared[timesteps * channels]
  ) {
  for (int i = threadIdx.x; i < timesteps * channels; i += blockDim.x) {
    double3 coords = uvw[i];
    coords.x *= scale;
    coords.y *= scale;
    coords.z *= scale;
    int
        // We additionally translate these u v by -max_supp/2
        // because gridding procedure translate them back
        u = round(coords.x) + grid_size/2 - max_supp/2
      , v = round(coords.y) + grid_size/2 - max_supp/2
      , over_u = round(over * (coords.x - u))
      , over_v = round(over * (coords.y - v))
      , w_plane = round (coords.z / (w_planes/2))
      ;
    // Mirror if necessary for using in GCF
    if (w_plane < 0) w_plane = - w_plane;
    // uvo_shared[i] = {short(u), short(v), (char)w_plane, (char)(over_u * over + over_v), (short)get_supp(w_plane)};
    uvo_shared[i].u = short(u);
    uvo_shared[i].v = short(v);
    uvo_shared[i].gcf_layer_w_plane = (char)w_plane;
    uvo_shared[i].gcf_layer_over = (char)(over_u * over + over_v);
    uvo_shared[i].gcf_layer_supp = (short)get_supp(w_plane);
    vis_shared[i] = vis[i];
  }
}

template <
    int grid_size
  , int w_planes

  , int timesteps
  , int channels
  >
__inline__ __device__
// grid must be initialized to 0s.
void gridKernel_scatter_kernel_small(
    int max_supp
    // Should be passed 0-centered,
    // i.e. with 0-index in the middle
  , const complexd * gcf_layer[w_planes]
  , Double4c grid[grid_size][grid_size]
  , const Pregridded uvo_shared[timesteps * channels]
  , const Double4c vis[timesteps * channels]
  , int myU
  , int myV
  ) {
  complexd
      sumXX = {0, 0}
    , sumXY = {0, 0}
    , sumYX = {0, 0}
    , sumYY = {0, 0}
    ;
  int
      grid_point_u = 0
    , grid_point_v = 0;

  for (int i = 0; i < timesteps * channels; i++) {
    int myConvU, myConvV, myGridU, myGridV, u, v;
    u = uvo_shared[i].u;
    v = uvo_shared[i].v;
    myConvU = myU - u % max_supp;
    if (myConvU < 0) myConvU += max_supp;
    myConvV = myV - v % max_supp;
    if (myConvV < 0) myConvV += max_supp;
    // This translates points by max_supp/2
    // returning them back to normal (they were translates by -max_supp/2 before)
    myGridU = u + myConvU;
    myGridV = v + myConvV;

    int supp = uvo_shared[i].gcf_layer_supp;
    complexd supportPixel =
      gcf_layer[uvo_shared[i].gcf_layer_w_plane]
        [(uvo_shared[i].gcf_layer_over * supp + myConvU) * supp + myConvV];
    if (myGridU != grid_point_u || myGridV != grid_point_v) {
      atomicAdd(&grid[grid_point_u][grid_point_v], sumXX, sumXY, sumYX, sumYY);
        sumXX
      = sumXY
      = sumYX
      = sumYY
      = {0, 0};
      grid_point_u = myGridU;
      grid_point_v = myGridV;
    }
    #define __ADD_SUPP(pol) sum##pol = cuCfma(supportPixel, vis[i].pol, sum##pol)
    __ADD_SUPP(XX);
    __ADD_SUPP(XY);
    __ADD_SUPP(YX);
    __ADD_SUPP(YY);
  }
  atomicAdd(&grid[grid_point_u][grid_point_v], sumXX, sumXY, sumYX, sumYY);
}

template <
    int grid_size
  , int w_planes

  , int timesteps
  , int channels
  >
__inline__ __device__
// grid must be initialized to 0s.
void gridKernel_scatter_kernel(
    int max_supp<commit_msg>Fix negative wplanes handling.<commit_after>#include "scatter_gridder.h"
#include "atomic_add.h"
#include "OskarBinReader.h"

#include <math_functions.h>

struct Pregridded
{
  short u;
  short v;
  char gcf_layer_w_plane;
  char gcf_layer_over;
  short gcf_layer_supp;
};

template <
    int over
  , int w_planes
  , int grid_size

  , int timesteps
  , int channels
  >
__inline__ __device__ void loadIntoSharedMem (
    double scale
  , int max_supp // needed for u v translation
  , const double3 uvw[timesteps * channels]
  , const Double4c vis[timesteps * channels]
  , Pregridded uvo_shared[timesteps * channels]
  , Double4c vis_shared[timesteps * channels]
  ) {
  for (int i = threadIdx.x; i < timesteps * channels; i += blockDim.x) {
    double3 coords = uvw[i];
    coords.x *= scale;
    coords.y *= scale;
    coords.z *= scale;
    int
        // We additionally translate these u v by -max_supp/2
        // because gridding procedure translate them back
        u = round(coords.x) + grid_size/2 - max_supp/2
      , v = round(coords.y) + grid_size/2 - max_supp/2
      , over_u = round(over * (coords.x - u))
      , over_v = round(over * (coords.y - v))
      , w_plane = round (coords.z / (w_planes/2))
      ;
    // uvo_shared[i] = {short(u), short(v), (char)w_plane, (char)(over_u * over + over_v), (short)get_supp(w_plane)};
    uvo_shared[i].u = short(u);
    uvo_shared[i].v = short(v);
    uvo_shared[i].gcf_layer_w_plane = (char)w_plane;
    uvo_shared[i].gcf_layer_over = (char)(over_u * over + over_v);
    uvo_shared[i].gcf_layer_supp = (short)get_supp(w_plane);
    vis_shared[i] = vis[i];
  }
}

template <
    int grid_size
  , int w_planes

  , int timesteps
  , int channels
  >
__inline__ __device__
// grid must be initialized to 0s.
void gridKernel_scatter_kernel_small(
    int max_supp
    // Should be passed 0-centered,
    // i.e. with 0-index in the middle
  , const complexd * gcf_layer[w_planes]
  , Double4c grid[grid_size][grid_size]
  , const Pregridded uvo_shared[timesteps * channels]
  , const Double4c vis[timesteps * channels]
  , int myU
  , int myV
  ) {
  complexd
      sumXX = {0, 0}
    , sumXY = {0, 0}
    , sumYX = {0, 0}
    , sumYY = {0, 0}
    ;
  int
      grid_point_u = 0
    , grid_point_v = 0;

  for (int i = 0; i < timesteps * channels; i++) {
    int myConvU, myConvV, myGridU, myGridV, u, v;
    u = uvo_shared[i].u;
    v = uvo_shared[i].v;
    myConvU = myU - u % max_supp;
    if (myConvU < 0) myConvU += max_supp;
    myConvV = myV - v % max_supp;
    if (myConvV < 0) myConvV += max_supp;
    // This translates points by max_supp/2
    // returning them back to normal (they were translates by -max_supp/2 before)
    myGridU = u + myConvU;
    myGridV = v + myConvV;

    int supp = uvo_shared[i].gcf_layer_supp;

    int wplane, wplane_m;
    wplane = uvo_shared[i].gcf_layer_w_plane;
    if (wplane < 0) wplane_m = -wplane; else wplane_m = wplane;
    complexd supportPixel =
      gcf_layer[wplane_m]
        [(uvo_shared[i].gcf_layer_over * supp + myConvU) * supp + myConvV];
    // if mirrored make conjugation
    if (wplane < 0)
      supportPixel.y = - supportPixel.y;

    if (myGridU != grid_point_u || myGridV != grid_point_v) {
      atomicAdd(&grid[grid_point_u][grid_point_v], sumXX, sumXY, sumYX, sumYY);
        sumXX
      = sumXY
      = sumYX
      = sumYY
      = {0, 0};
      grid_point_u = myGridU;
      grid_point_v = myGridV;
    }
    #define __ADD_SUPP(pol) sum##pol = cuCfma(supportPixel, vis[i].pol, sum##pol)
    __ADD_SUPP(XX);
    __ADD_SUPP(XY);
    __ADD_SUPP(YX);
    __ADD_SUPP(YY);
  }
  atomicAdd(&grid[grid_point_u][grid_point_v], sumXX, sumXY, sumYX, sumYY);
}

template <
    int grid_size
  , int w_planes

  , int timesteps
  , int channels
  >
__inline__ __device__
// grid must be initialized to 0s.
void gridKernel_scatter_kernel(
    int max_supp
{'fraction_non_alphanumeric': 0.12018915013792197, 'fraction_numerical': 0.00906344410876133, 'mean_word_length': 2.888661899897855, 'pattern_counts': {'":': 0, '<': 22, '<?xml version=': 0, '>': 11, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '1082437', 'n_tokens_mistral': 3046, 'n_tokens_neox': 2978, 'n_words': 686}