Date: Sun, 23 Sep 2012 17:53:42 +0530 From: Sayantan Datta <std2048@...il.com> To: john-dev@...ts.openwall.com Subject: Re: bitslice DES on GPU On Sun, Sep 23, 2012 at 4:00 PM, magnum <john.magnum@...hmail.com> wrote: > 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? > Actually it is intended for AMD GPUs but fortunately it is just ignored for any other GPUs or CPUs without causing any trouble. > > 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 > Yeah, sure. Regards, Sayantan Content of type "text/html" skipped
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.