Openwall GNU/*/Linux - a small security-enhanced Linux distro for servers
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Sun, 23 Sep 2012 12:30:29 +0200
From: magnum <john.magnum@...hmail.com>
To: john-dev@...ts.openwall.com
Subject: Re: bitslice DES on GPU

On 22 Sep, 2012, at 13:42 , Sayantan Datta <std2048@...il.com> wrote:

> Actually the bitselect is causing problems. So I'm removing it along with the AMD sboxes for now. Without bitselect, AMD sboxes are slower than nv sboxes. I've posted the patch to unstable-jumbo branch of JtR repo.

Testing this on OSX w/ nvidia:

$ ../run/john -t -fo:des-opencl
OpenCL platform 0: Apple, 2 device(s).
Using device 1: GeForce GT 650M
Compilation log: <program source>:292:65: warning: unknown attribute 'max_constant_size' ignored
 __kernel void DES_bs_25(constant uint *index768 __attribute__((max_constant_size(3072))), __global int *index96 ,__global DES_bs_transfer *DES_bs_all,__global DES_bs_vector *B_global)
                                                                ^

Benchmarking: DES BS [128/128 BS SSE2-16]... DONE
Many salts:	2016K c/s real, 2036K c/s virtual
Only one salt:	2097K c/s real, 2114K c/s virtual

Do you think the compiler warning means we're in trouble or is it OK that the attribute is ignored on some systems?

Before this patch it was a lot slower:

$ ../run/john -t -fo:des-opencl 
OpenCL platform 0: Apple, 2 device(s).
Using device 1: GeForce GT 650M
Benchmarking: DES BS [128/128 BS SSE2-16]... DONE
Many salts:	173605 c/s real, 173893 c/s virtual
Only one salt:	173318 c/s real, 173605 c/s virtual

The [128/128 BS SSE2-16] is very confusing, it should say OpenCL instead of SSE2. Please fix that.

magnum

Powered by blists - more mailing lists

Your e-mail address:

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