Follow @Openwall on Twitter for new release announcements and other news
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Sat, 11 Aug 2012 03:42:04 +0400
From: Solar Designer <>
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

	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.


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.