Follow @Openwall on Twitter for new release announcements and other news
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Sun, 07 Jul 2013 21:52:22 +0530
From: Sayantan Datta <>
Subject: Re: bug: GPU use in CPU-only formats

On Friday 05 July 2013 05:23 AM, magnum wrote:
> $ ../run/john -t -form:mscash2-opencl -dev=gpu
> Device 0: GeForce GTX 570
> Optimal Work Group Size:128
> Kernel Execution Speed (Higher is better):0.470798
> Optimal Global Work Size:81920
> Device 1: Tahiti (AMD Radeon HD 7900 Series)
> Optimal Work Group Size:256
> Kernel Execution Speed (Higher is better):1.565677
> Optimal Global Work Size:253952
> WARNING: Devices are NOT arranged in DECREASING order of performance and might cause degradation in performance.
> Benchmarking: mscash2-opencl, M$ Cache Hash 2 (DCC2) [PBKDF2-SHA1 OpenCL]... DONE
> Raw:    64291 c/s real, 64129 c/s virtual
> I think that "decreasing order" requirement should be taken care of by the format, not the user. When using --device=gpu you select all GPU devices available but on Bull they happen to come in increasing order of performance.

I found the root cause of this problem. Nvidia blocks the thread after 
each clEnqueueNDRangeKernel call which is not the case with AMD. So when 
nvidia is the first device it blocks the thread until all kernel 
execution is finished effectively serializing the kernel calls. But when 
AMD is called first it just en-queues the kernel and returns, then moves 
on to en queue the nvidia kernels and finally gets blocked.  Since we 
are testing only two devices, calling AMD followed by nvidia solves the 
problem. But for multiple nvidia devices this problem can't be solved 
unless we use multi threaded execution.

Also the whole thing has nothing to do with execution speed. Sorry for 
the misleading warning.


Content of type "text/html" skipped

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.