[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Message-ID: <9881fd379dfedb81a34edf1d8832eea1@smtp.hushmail.com>
Date: Thu, 16 Jul 2015 15:12:02 +0200
From: magnum <john.magnum@...hmail.com>
To: john-dev@...ts.openwall.com
Subject: Re: PHC: Lyra2 on GPU
On 2015-07-16 10:25, Agnieszka Bielec wrote:
> 2015-07-14 1:18 GMT+02:00 magnum <john.magnum@...hmail.com>:
>> On 2015-07-13 20:28, Agnieszka Bielec wrote:
>>> Do john support lws auto tuning? We only have get_default_workgroup()
>>> function which is called somewhere.
>>
>> Unless none were given, autotune_run() will auto-tune GWS first (using an
>> LWS of NULL during that process). Then it will auto-tune LWS using the
>> previously established GWS.
>
> works in Lyra2 but in yescrypt doesn't work
>
> I added printf here
>
> static void create_clobj(size_t gws, struct fmt_main *self)
> {
> printf("gws=%llu\n",gws);
>
>
> and the output is:
> Calculating best global worksize (GWS); max. 1s single kernel invocation.
> gws=256
> gws=0
> OpenCL error (CL_INVALID_BUFFER_SIZE) in file
> (opencl_yescrypt_fmt_plug.c) at line (163) - (Error creating device
> buffer)
>
> I changed get_default_workgroup() to return 0 in lyra and yescrypt
Does your crypt_all() handle the case where local_work_size == 0, and
supplying a null pointer (as opposed to &local_work_size) if so?
eg. in wpapsk-opencl crypt_all_benchmark():
size_t *lws = local_work_size ? &local_work_size : NULL;
...then kernel is called like this:
clEnqueueNDRangeKernel(queue[gpu_id], wpapsk_init, 1, NULL,
&global_work_size, lws, 0, NULL, multi_profilingEvent[1]);
It doesn't pass a pointer to a variable that is 0, but instead it passes
a NULL pointer. I guess your problem might be that this logic is
missing. It would be easier if the OpenCL calls could just use a 0 value
instead but they do not.
magnum
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.
