Follow @Openwall on Twitter for new release announcements and other news
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Message-ID: <CAKGDhHUXufSW94DhdS0ANRYKQ4b+sM=wOFP6FyLkS89GBx2xHw@mail.gmail.com>
Date: Thu, 16 Jul 2015 15:20:14 +0200
From: Agnieszka Bielec <bielecagnieszka8@...il.com>
To: john-dev@...ts.openwall.com
Subject: Re: PHC: Lyra2 on GPU

2015-07-16 15:12 GMT+02:00 magnum <john.magnum@...hmail.com>:
> 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.

my code looks like this

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.