libgpuverify

Signature verification on GPUs (WiP)
Log | Files | Refs | README | LICENSE

commit 36724ccea92c0594381774f589bcfa9f73c6d658
parent 6b1ab235f8d80c1ef116f9624ad21910d11ff599
Author: Christian Grothoff <grothoff@gnunet.org>
Date:   Thu,  7 Dec 2023 00:25:36 +0900

-highlight brokenness of MG test

Diffstat:
Msource/montgomery.cl | 74++++++++++++++------------------------------------------------------------
1 file changed, 14 insertions(+), 60 deletions(-)

diff --git a/source/montgomery.cl b/source/montgomery.cl @@ -2483,12 +2483,14 @@ __kernel void montgomery(__global void *signature, __global unsigned long *s_off __global unsigned long *pks, unsigned long n) { - - int index = get_global_id(0); int pk = 0; - + + if (1) { + printf((__constant char *)"Shortcut on."); + return; + } while (1) { if (pks[pk] >= index) break; @@ -2519,67 +2521,19 @@ __kernel void montgomery(__global void *signature, __global unsigned long *s_off // the modulus can be assumed to be uneven – always if (mpz_even_p(m)) { - /* - mpz_t bb, x1, x2, q, powj; - mpz_init(bb); - mpz_init(x1); - mpz_init(x2); - mpz_init(q); - mpz_init(powj); - - mont_prepare_even_modulus(m, q, powj); - - // q is uneven, so we can use regular modexp - // MARK: we can improve the efficiency here by doing simple reductions - - mpz_mod(bb, b, q); // reductions like this - - mont_prepare(bb, e, q, r, r_1, ni, M, x); - mont_modexp(xx, x, e, M, q, ni, r, r_1); - mont_finish(x1, xx, q, ni, r, r_1); - - - // MARK: we can also reduce and really speed this up as well -> binary method? - mpz_powm(x2, b, e, powj); - - mpz_t y, q_1; - mpz_init(y); - mpz_init(q_1); - - mpz_sub(y, x2, x1); - - mpz_invert(q_1, q, powj); - - mpz_mul(y, y, q_1); - mpz_mod(y, y, powj); - - mpz_addmul(x1, q, y); - - mpz_set(res, x1); - - - */ - printf((__constant char *)"An even modulus is not allowed here."); - - } else { - - // MARK: prepare might not have to run individually on each kernel (prepare might even run on CPU) - mont_prepare(b, e, m, r, r_1, ni, M, x); - - - mont_modexp(xx, x, e, M, m, ni, r, r_1); - mont_finish(res, xx, m, ni, r, r_1); - + return; } - + if (mpz_even_p(m)) { + printf((__constant char *)"An odd modulus is not allowed here."); + return; + } + // MARK: prepare might not have to run individually on each kernel (prepare might even run on CPU) + mont_prepare(b, e, m, r, r_1, ni, M, x); + mont_modexp(xx, x, e, M, m, ni, r, r_1); + mont_finish(res, xx, m, ni, r, r_1); if (mpz_cmp(sig,res) != 0) { *valid += 1; - } - - - - }