[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Sat, 11 Aug 2012 03:42:04 +0400
From: Solar Designer <solar@...nwall.com>
To: john-dev@...ts.openwall.com
Subject: Re: Current -fixes GPU formats vs TS
Lukas -
On Sat, Aug 11, 2012 at 01:23:56AM +0200, Lukas Odzioba wrote:
> Because format fails to calculate proper values just for the first
> time after some mysterious runtime history I thought that we could do
> sth like that:
> static void init(struct fmt_main *pFmt)
> {
> host_pass = calloc(KEYS_PER_CRYPT, sizeof(pwsafe_pass));
> host_hash = calloc(KEYS_PER_CRYPT, sizeof(pwsafe_hash));
> host_salt = calloc(1, sizeof(pwsafe_salt));
> cuda_init(gpu_id);
> memset(host_pass,0,KEYS_PER_CRYPT*sizeof(pwsafe_pass));
> memset(host_hash,0,KEYS_PER_CRYPT*sizeof(pwsafe_hash));
> memset(host_salt,0,sizeof(pwsafe_salt));
> gpu_pwsafe(host_pass, host_salt, host_hash); <--- this puppy
> atexit(cleanup);
> }
> And this workaround fixes the problem. However it is at least stupid.
> I'll send final code in a moment. I am not happy about this hack, but
> as fas as it does the job we may want to use it.
Yeah, we may do that as a last resort.
In trying to trigger the problem without running xsha512-cuda first, I
added:
cudaMemset(cuda_hash, -1, PWSAFE_OUT_SIZE);
right after:
cudaMalloc(&cuda_hash, PWSAFE_OUT_SIZE);
This didn't make any difference in triggering the problem (nor in
preventing it), but surprisingly it provided a 3% speedup (approx.
106k c/s to 109k c/s on bull's GTX 570).
A memset with 0 also provides some speedup (a slightly smaller speedup?
not sure).
Any idea why? This might be a clue.
Alexander
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.
