Follow @Openwall on Twitter for new release announcements and other news
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Thu, 05 Jul 2012 16:23:49 +0200
From: magnum <>
Subject: Re: common find_best_workgroup()

On 2012-07-05 16:16, Lukas Odzioba wrote:
> 2012/7/5 magnum <>:
>> 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:


...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.


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.