Attachment 115815 Details for Bug 90472 – Offending OpenCL source. Notes are included within.

Offending OpenCL source. Notes are included within.

finite_field_tests.cl (text/plain), 7.91 KB, created by Stojan Dimitrovski on 2015-05-15 15:19:02 UTC

(hide)

Description:

Filename:

MIME Type:

Creator: Stojan Dimitrovski

Created: 2015-05-15 15:19:02 UTC

Size: 7.91 KB

patch

obsolete

>#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);
>}

Actions: View

Attachments on bug 90472: 115815 | 115816 | 115817 | 118467 | 118468 | 118469 | 118470

Use of freedesktop.org services, including Bugzilla, is subject to our Code of Conduct. How we collect and use information is described in our Privacy Policy.