#define SETUP __private size_t i __attribute__((aligned)) = 0; \
  __private size_t j __attribute__((aligned)) = 0; \
\
  __private const size_t maxf __attribute__((aligned)) = ((1 << (8 * ffwidth)) - 1); \
  __private uchar a[ffwidth] __attribute__((aligned)); \
  __private uchar b[ffwidth] __attribute__((aligned)); \
  __private uchar c[ffwidth] __attribute__((aligned));

#ifndef SKIP
  #pragma message("SKIP is not defined, will define it to 1. Tests may crash on large fields.")
  #define SKIP (1)
#endif

// BEIGNET DEVELOPERS:
// Whenever I prepend any of these Test* functions with __kernel, the build
// fails as specified in the bug submission.
// There are macros used throught this file which are not currently present. If
// you really need to try it out I can send you those, too. Email me at
// sdimitrovski@gmail.com

/* __kernel */ void TestFiniteFieldValueFunctions(__global uchar* ffip_global) {
  SETUP

  printf("Value functions:\n");

  for (i = 0; i <= maxf; i += SKIP) {
    FiniteFieldValue(i, a);
    FiniteFieldCopy(a, b);

    __private size_t v = 0;
    FiniteFieldToValue(a, &v);

    if (v != i) {
      printf("  FiniteFieldToValue failed: %d %d\n", i, v);
      break;
    }

    v = 0;
    FiniteFieldToValue(b, &v);
    if (v != i) {

      printf("  FiniteFieldCopy failed: %d\n", i);
      break;
    }
  }

  printf("Done.\n");
}

void TestFiniteFieldUnionFunction(__global uchar* ffip_global) {
  SETUP

  printf("Union:\n");

  FiniteFieldValue(0, a);
  FiniteFieldUnion(a, &i);

  if (i != 0) {
    printf("  Union of zero came up non-zero.\nDone.\n");
    return;
  }

  for (i = 1; i <= maxf; i += SKIP) {
    FiniteFieldValue(i, a);

    __private size_t u = 0;
    FiniteFieldUnion(a, &u);

    if (u == 0) {
      printf(" Union of non-zero came up as zero.\n");
      break;
    }
  }

  printf("Done.\n");
}

void TestFiniteFieldShiftFunctions(__global uchar* ffip_global) {
  SETUP

  printf("Shift functions:\n");

  for (i = 0; i <= maxf; i++) {
    for (j = 0; j < 8; j++) {
      FiniteFieldValue(i, a);
      FiniteFieldShiftRight(a, j);

      __private size_t v = 0;
      FiniteFieldToValue(a, &v);

      if (v != ((i >> j) & maxf)) {
        printf("  FiniteFieldShiftRight (%u place) failed at %u got %u\n", j, i, v);
        break;
      }

      FiniteFieldValue(i, a);
      FiniteFieldShiftLeft(a, j);

      v = 0;
      FiniteFieldToValue(a, &v);

      if (v != ((i << j) & maxf)) {
        printf("  FiniteFieldShiftLeft (%u place) faield at %u got %u\n", j, i, v);
        break;
      }
    }
  }

  printf("Done.\n");
}

void TestFiniteFieldAdditionFunctions(__global uchar* ffip_global) {
  SETUP

  printf("Addition:\n");

  for (i = 0; i <= maxf; i += SKIP) {
    FiniteFieldValue(i, a);

    for (j = 0; j <= maxf; j += SKIP) {
      FiniteFieldValue(j, b);

      FiniteFieldAdd(a, b, c);

      __private size_t l = 0;
      FiniteFieldToValue(c, &l);

      FiniteFieldAdd(b, a, c);

      __private size_t r = 0;
      FiniteFieldToValue(c, &r);

      if (l != r) {
        printf("  FiniteFieldAdd is not commutative at %d + %d\n", i, j);
        break;
      }

      if (l != r && i > 0 && j > 0 && i != j) {
        printf(" FiniteFieldAdd gives an invalid result at %d + %d = %d\n", i, j, l);
        break;
      }
    }
  }

  {
    __private size_t v = 0;
    FiniteFieldValue(maxf - 1, a);
    FiniteFieldValue(2, b);
    FiniteFieldAdd(a, b, c);
    FiniteFieldToValue(c, &v);

    printf("  Example: %i + %i = %i\n", (maxf - 1), 2, v);
  }

  printf("Done.\n");
}

void TestFiniteFieldSubtractionFunctions(__global uchar* ffip_global) {
  SETUP

  printf("Subtraction:\n");

  for (i = 0; i <= maxf; i += SKIP) {
    FiniteFieldValue(i, a);

    for (j = 0; j <= maxf; j += SKIP) {
      FiniteFieldValue(j, b);

      FiniteFieldSubtract(a, b, c);

      __private size_t l = 0;
      FiniteFieldToValue(c, &l);

      FiniteFieldSubtract(b, a, c);

      __private size_t r = 0;
      FiniteFieldToValue(c, &r);

      if (l != r) {
        printf("  FiniteFieldSubtract is not commutative at %d + %d\n", i, j);
        break;
      }

      if (l != r && i > 0 && j > 0 && i != j) {
        printf(" FiniteFieldSubtract gives an invalid result at %d + %d = %d\n", i, j, l);
        break;
      }
    }
  }

  {
    __private size_t v = 0;
    FiniteFieldValue(2, a);
    FiniteFieldValue(maxf - 5, b);
    FiniteFieldSubtract(a, b, c);
    FiniteFieldToValue(c, &v);

    printf("  Example: %i + %i = %i\n", 2, (maxf - 5), v);
  }

  printf("Done.\n");
}

void TestFiniteFieldMultiplicationFunctions(__global uchar* ffip_global) {
  SETUP

  printf("Multiplication:\n");

  {
    __private size_t inverses = 0;

    for (i = 0; i <= maxf; i += SKIP) {
      FiniteFieldValue(i, a);

      for (j = 0; j <= maxf; j += SKIP) {
        FiniteFieldValue(j, b);

        FiniteFieldMultiply(a, b, c, ffip_global);

        __private size_t l = 0;
        FiniteFieldToValue(c, &l);

        FiniteFieldMultiply(b, a, c, ffip_global);

        __private size_t r = 0;
        FiniteFieldToValue(c, &r);

        if (l != r) {
          printf(" FiniteFieldMultiply is not commutative at %i * %i\n", i, j);
          break;
        }

        if (l == 1) {
          inverses += 1;
        }
      }
    }

    printf("  Found %i inverses.\n", inverses);
  }

  {
    FiniteFieldValue(3, a);
    FiniteFieldValue(maxf - 3, b);
    FiniteFieldMultiply(a, b, c, ffip_global);

    __private size_t v = 0;
    FiniteFieldToValue(c, &v);

    printf("  Example: %i * %i = %i\n", 3, (maxf - 3), v);
  }

  printf("Done.\n");
}

void TestFiniteFieldLeftAlignFunctions(__global uchar* ffip_global) {
  SETUP

  printf("Left align:\n");

  for (i = 0; i <= maxf; i++) {
    FiniteFieldValue(i, a);

    __private size_t bits = 0;
    FiniteFieldLeftAlign(a, &bits);

    __private size_t v = 0;
    FiniteFieldToValue(a, &v);

    __private size_t expected = i << bits;

    if (expected != v) {
      printf("  NO! Aligning %u, got %u expected %u\n", i, v, expected);
    }
  }

  printf("Done.\n");
}

void quot(__private uchar* a, __private uchar* b, __private uchar* r) {
  FiniteFieldQuotient(a, b, r);
}

void TestFiniteFieldQuotientFunctions(__global uchar* ffip_global) {
  SETUP

  printf("Quotient:\n");

  __private size_t sum = 0;

  for (i = 0; i <= maxf; i += SKIP) {
    FiniteFieldValue(i, a);

    for (j = 1; j <= maxf; j += SKIP) {
      FiniteFieldValue(j, b);

      FiniteFieldQuotient(a, b, c);

      __private size_t v = 0;
      FiniteFieldToValue(c, &v);

      sum = sum + v;
    }
  }

  printf("  Sum: %u\n", sum);

  sum = 0;

  for (i = 0; i <= maxf; i += SKIP) {
    FiniteFieldValue(i, a);

    for (j = 1; j <= maxf; j += SKIP) {
      FiniteFieldValue(j, b);

      FiniteFieldQuotientPrepend(a, b, c);

      __private size_t v = 0;
      FiniteFieldToValue(c, &v);

      sum = sum + v;
    }
  }

  printf("  PSum: %u\n", sum);

  printf("Done.\n");
}

void TestFiniteFieldMultiplicativeInverseFunctions(__global uchar* ffip_global) {
  SETUP

  printf("Multiplicative inverse:\n");

  __private size_t sum = 0;

  for (i = 1; i <= maxf; i += SKIP) {
    FiniteFieldValue(i, a);
    FiniteFieldMultiplicativeInverse(a, c, ffip_global);

    FiniteFieldMultiply(a, c, b, ffip_global);

    __private size_t v = 0;
    FiniteFieldToValue(c, &v);

    __private size_t o = 0;
    FiniteFieldToValue(b, &o);

    if (o != 1) {
      printf("  Inverse is fake at: %u (~%u) got %u\n", i, v, o);
      break;
    }

    sum = sum + v;
  }

  printf("  Sum: %u\nDone.\n", sum);
}

__kernel void Test(__global uchar* ffip_global) {
  TestFiniteFieldValueFunctions(ffip_global);
  TestFiniteFieldUnionFunction(ffip_global);
  TestFiniteFieldShiftFunctions(ffip_global);
  TestFiniteFieldAdditionFunctions(ffip_global);
  TestFiniteFieldSubtractionFunctions(ffip_global);
  TestFiniteFieldMultiplicationFunctions(ffip_global);
  TestFiniteFieldLeftAlignFunctions(ffip_global);
  TestFiniteFieldQuotientFunctions(ffip_global);
  TestFiniteFieldMultiplicativeInverseFunctions(ffip_global);
}