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.