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 08:40:48 +0100
From: magnum <john.magnum@...hmail.com>
To: john-dev@...ts.openwall.com
Subject: Re: md5crypt-opencl

On 2013-11-25 07:46, magnum wrote:
> PUTCHAR need an int* destination and a char source. Obviously if both
> are aligned it's much faster to copy *dst++ = *src++ with both being ints.

Here's roughly what I mean, although it should be optimized so we avoid 
the conditional if possible (or so we can always do it 32-bit). The 
attached patch works and is faster, but it can be made much better. 
Possibly prepare somthing outside the inner loop.

magnum


commit 626ed7acd96054cd5141563b80cc3a9be66f15f5
Author: magnum <john.magnum@...hmail.com>
Date:   Mon Nov 25 08:31:10 2013 +0100

    Optimize for aligned destination

diff --git a/src/opencl/cryptmd5_kernel.cl b/src/opencl/cryptmd5_kernel.cl
index f90555c..73ab99f 100644
--- a/src/opencl/cryptmd5_kernel.cl
+++ b/src/opencl/cryptmd5_kernel.cl
@@ -98,12 +98,20 @@ typedef struct {
 __constant uchar cl_md5_salt_prefix[] = "$1$";
 __constant uchar cl_apr1_salt_prefix[] = "$apr1$";
 
-inline void buf_update(md5_ctx * ctx, uchar * string, uint len, uint buflen)
+inline void buf_update(md5_ctx * ctx, uchar * string, uint buflen)
 {
 	uint i;
-	for (i = 0; i < len; i++)
-		PUTCHAR(ctx->buffer, buflen + i, string[i]);
 
+	if (buflen & 3) {
+		// 8-bit copy
+		for (i = 0; i < 16; i++)
+			PUTCHAR(ctx->buffer, buflen + i, string[i]);
+	} else {
+		// 32-bit copy
+		buflen >>= 2;
+		for (i = 0; i < 4; i++)
+			ctx->buffer[buflen + i] = ((uint*)string)[i];
+	}
 }
 
 inline void ctx_update(md5_ctx * ctx, uchar * string, uint len,
@@ -356,7 +364,7 @@ __kernel void cryptmd5(__global const crypt_md5_password * inbuffer,
 		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, 16, altpos[nid]);
+		buf_update(&ctxs[nid], (uchar *) alt_result, altpos[nid]);
 		cid = nid;
 	}
 	for (i = 0; i < 4; i++)

Powered by blists - more mailing lists

Your e-mail address:

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