Follow @Openwall on Twitter for new release announcements and other news
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Tue, 18 Jun 2013 00:11:21 +0530
From: Sayantan Datta <>
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):


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.