[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Message-ID: <a46b49a34d7079ed18e646bd74c90fc6@smtp.hushmail.com>
Date: Thu, 05 Jul 2012 16:23:49 +0200
From: magnum <john.magnum@...hmail.com>
To: john-dev@...ts.openwall.com
Subject: Re: common find_best_workgroup()
On 2012-07-05 16:16, Lukas Odzioba wrote:
> 2012/7/5 magnum <john.magnum@...hmail.com>:
>> Lukas,
>>
>> In the current common-opencl.c code, you never actually set
>> local_work_size = my_work_group within the find_best loop. It's 1 all
>> the time so the end result is totally random. I am about to change this.
>> Or am I confused?
> Current state is:
>
> void opencl_find_best_workgroup(struct fmt_main *pFmt)
> {
> (...)
> size_t my_work_group, optimal_work_group = 1;
> (..)
> /// Find minimum time
> for (my_work_group = wg_multiple;
> (int) my_work_group <= (int) max_group_size;
> my_work_group += wg_multiple) {
> (...)
> if ((endTime - startTime) < kernelExecTimeNs) {
> kernelExecTimeNs = endTime - startTime;
> optimal_work_group = my_work_group;
> }
> (...)
> local_work_size = optimal_work_group;
> (...)
> }
>
> It looks ok for me. Am I right?
The problem is that in the old code, you did something like:
ret_code = clEnqueueNDRangeKernel(queue_prof, crypt_kernel, 1, NULL,
&global_work_size, &my_work_group, 0, NULL, &myEvent);
in that loop. And that would run the test using my_work_group. Now you
do this:
pFmt->methods.crypt_all(pFmt->params.max_keys_per_crypt);
...and this will instead run the tests using local_work_size, which is
always 1.
I have already committed an experimental fix. Seems to work fine but
still testing (I also changed all formats except Claudio's so they use
the shared function)
> Possibly we should drop optimal_work_group variable and always set
> local_work_size. If you agree, I'll fix that. This functions is not
> perfect, for example we should introduce local_work_size for each
> device (we're not using more than two dev's per run but this should
> change in the near future - it is already in sayantan's code but he
> created his own "habitat" in code).
I'm starting to think we should drop the whole idea of
find_best_workgroup. I mostly see random results (picking totally
different values at different runs). For most kernels you will be able
to pick a good default. The more critical one is global_work_size, and I
believe that too can have a good default per format.
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.
