Openwall GNU/*/Linux - a small security-enhanced Linux distro for servers
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Mon, 25 Nov 2013 21:29:02 +0100
From: magnum <john.magnum@...hmail.com>
To: john-dev@...ts.openwall.com
Subject: Re: md5crypt-opencl

On 2013-11-25 21:06, magnum wrote:
> On 2013-11-25 08:44, Lukas Odzioba wrote:
>> 2013/11/25 magnum <john.magnum@...hmail.com>:
>>> The attached patch works and is faster, but it can be made much better.
>>
>> I'll take a look on it later today, what c/s you are getting now?
>
> With that patch and a similar hack for ctx_update() I get 1881K c/s on
> Tahiti and 2383K c/s on the Titan. Like I said it can probably be done
> better, it's just a QnD hack for showing how bad 8-bit copies are.

Fwiw using #pragma unroll 42 was just detrimental. That's odd.

BTW you can also take ctx_out away from the inner loop and its value in 
the end is known:

@@ -366,17 +366,15 @@ __kernel void cryptmd5(__global const 
crypt_md5_password * inbuffer,
         ctx_update(&ctxs[7], pass.c, pass_len, &ctxs_buflen[7]);
         ctxs_buflen[7] += 16;

-
-       uchar *ctx_out;
         uint nid, cid = g[0];   //current ctx      id

         for (i = 0; i < 1000; i++) {
                 nid = g[(i + 1) % 42];  // next ctx id to process, 0-7
-               ctx_out = &((uchar *) ctxs[nid].buffer)[altpos[nid]];
                 md5_digest(&ctxs[cid], alt_result, &ctxs_buflen[cid]);
                 buf_update(&ctxs[nid], (uchar *) alt_result, altpos[nid]);
                 cid = nid;
         }
+
         for (i = 0; i < 4; i++)
-               outbuffer[idx].v[i] = ((uint *) ctx_out)[i];
+               outbuffer[idx].v[i] = ctxs[3].buffer[i];
  }


magnum

Powered by blists - more mailing lists

Your e-mail address:

Powered by Openwall GNU/*/Linux - Powered by OpenVZ