|
|
Message-ID: <a3cc580c56a18763fdb270d169a39c16@smtp.hushmail.com>
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
Confused about mailing lists and their use? Read about mailing lists on Wikipedia and check out these guidelines on proper formatting of your messages.