Date: Sun, 07 Jul 2013 21:52:22 +0530 From: Sayantan Datta <std2048@...il.com> To: john-dev@...ts.openwall.com 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. http://www.khronos.org/message_boards/showthread.php/5822-non-blocking-call-to-clEnqueueNDRangeKernel Also the whole thing has nothing to do with execution speed. Sorry for the misleading warning. Regards, Sayantan 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.