[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Message-ID: <d669ef6eee13332e05bdceffa6aca9f9@smtp.hushmail.com>
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
Confused about mailing lists and their use?
Read about mailing lists on Wikipedia
and check out these
guidelines on proper formatting of your messages.
