[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Message-ID: <51BF5851.6060104@gmail.com>
Date: Tue, 18 Jun 2013 00:11:21 +0530
From: Sayantan Datta <std2048@...il.com>
To: john-dev@...ts.openwall.com
Subject: Re: Mask mode for GPU
On Monday 17 June 2013 10:52 PM, Solar Designer wrote:
>
> Oh. Can you investigate that and try to get around it perhaps with
> different source code layout?
>
> BTW, out of different GPUs, AMD GCN GPUs appeared to tolerate the code
> exceeding I-cache size pretty well - although perhaps not in all cases.
> I think the DES cracker referenced from that Reddit thread I mentioned
> did exceed I-cache size, but it ran pretty fast on GCN GPUs (and only on
> those...), I guess due to each fetch into the I-cache being made use of
> multiple times (by multiple wavefronts).
Yes, I was able to fool the the compiler by transferring the iteration
count from global memory instead of hardcoding it. Also it is possible
to get around the problem by putting the loop body inside a non-inline
function. Now the speeds are are nearly identical or maybe a little
better for multiple iterations.
> Oh, I overlooked the fact that with bitslice DES you obviously need to
> process a multiple of 32 candidate passwords per work-item.
>
> A drawback of your approach is that you'll probably end up doing modulo
> division of get_global_id(), which may be slow. Also, you may have to
> compute all of the GPU-generated characters (such as 2 or 3) in each
> work-item, whereas with my suggested approach you'd only update one
> character in most loop iterations.
Yes right, but since the problem is fixed I will use the approach you
mentioned. BTW how much slower is modulo division compared to normal
division ? We can always replace a modulo division with one integer
division and one subtraction. I guess this is what the compiler does.
I have a draft implementation for mask mode on the host side ready. It
is in my personal git repo( bleeding branch):
https://github.com/Sayantan2048/JohnTheRipper.git
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.
