[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Message-ID: <5808e1bd4d9d6d270e0826f66834dab8@smtp.hushmail.com>
Date: Tue, 2 Oct 2012 10:18:03 +0200
From: magnum <john.magnum@...hmail.com>
To: john-dev@...ts.openwall.com
Subject: Re: OpenCL on OSX
On 2 Oct, 2012, at 4:01 , Claudio André <claudioandre.br@...il.com> wrote:
> magnum, I'm not sure why you used static/inline in kernels (for OSX): sometimes one, sometimes other.
I got silly build errors for all OpenCL functions (nb. not the kernels themselves, only helper functions) that did not use either static, inline or perhaps both. On other platforms, those keywords should be a no-op, functions are static and inline anyway afaik. Which one you use does not matter. Was there a difference?
> I sent you a pull request (split on sha256crypt). Please, say something (or just fix my crappy code).
>
> Sorry, but i had to change it. I tested here and on bull.
Sha256crypt works fine on OSX, after some build warnings. Complete output:
OpenCL platform 0: Apple, 2 device(s).
Using device 1: GeForce GT 650M
Building the kernel, this could take a while
Compilation log: In file included from <program source>:15:
/Users/magnum/src/john/src/opencl_cryptsha256.h:68:38: warning: unknown OpenCL extension 'cl_nv_pragma_unroll' - ignoring
#pragma OPENCL EXTENSION cl_nv_pragma_unroll : enable
^
(I seem to get the above warning on all nvidias so I have stopped using that enable pragma. Not sure if all documentation is wrong or why it doesn't work (and I believe defining cl_nv_pragma_unroll yet not understanding the pragma is a violation of the OpenCL spec). I still use #pragma unroll and I think/hope it works anyway)
<program source>:261:5: warning: array index of '15' indexes past the end of an array (that contains 1 element)
ctx->buffer->mem_32[15] = SWAP32(ctx->total * 8);
^ ~~
/Users/magnum/src/john/src/opencl_cryptsha256.h:89:5: note: array 'mem_32' declared here
uint32_t mem_32[1];
^
/Users/magnum/src/john/src/opencl_cryptsha256.h:22:18: note: expanded from macro 'uint32_t'
#define uint32_t unsigned int
^
<program source>:457:12: warning: unused variable 'lid'
size_t lid = get_local_id(0);
^
<program source>:484:12: warning: unused variable 'lid'
size_t lid = get_local_id(0);
^
Local work size (LWS) 128, global work size (GWS) 6144
Benchmarking: sha256crypt (rounds=5000) [OpenCL]... DONE
Raw: 2242 c/s real, 87771 c/s virtual
On CPU, I get even more warnings (very noisy compiler but sometimes that is a good thing) and it fails:
OpenCL platform 0: Apple, 2 device(s).
Using device 0: Intel(R) Core(TM) i7-3615QM CPU @ 2.30GHz
Compilation log: <program source>:70:14: warning: comparison of integers of different signs: 'int' and 'const unsigned int'
while (i < srclen) {
~ ^ ~~~~~~
<program source>:84:14: warning: comparison of integers of different signs: 'int' and 'const unsigned int'
while (i < srclen) {
~ ^ ~~~~~~
<program source>:98:14: warning: comparison of integers of different signs: 'int' and 'const unsigned int'
while (i < srclen) {
~ ^ ~~~~~~
<program source>:270:5: warning: array index of '15' indexes past the end of an array (that contains 1 element)
ctx->buffer->mem_32[15] = SWAP32(ctx->total * 8);
^ ~~
/Users/magnum/src/john/src/opencl_cryptsha256.h:89:5: note: array 'mem_32' declared here
uint32_t mem_32[1];
^
/Users/magnum/src/john/src/opencl_cryptsha256.h:22:18: note: expanded from macro 'uint32_t'
#define uint32_t unsigned int
^
Local work size (LWS) 1, global work size (GWS) 1024
Benchmarking: sha256crypt (rounds=5000) [OpenCL]... FAILED (get_hash[2](0))
The sha512crypt still doesn't work, but don't bother looking into it. I'm sure it's an Apple bug:
../run/john -t -fo:sha512crypt-opencl
OpenCL platform 0: Apple, 2 device(s).
Using device 1: GeForce GT 650M
Building the kernel, this could take a while
Compilation log: Error getting function data from server
Error building kernel. Returned build code: -11. DEVICE_INFO=130
OpenCL error (CL_BUILD_PROGRAM_FAILURE) in file (common-opencl.c) at line (146) - (clBuildProgram failed.)
And besides, it works on CPU:
OpenCL platform 0: Apple, 2 device(s).
Using device 0: Intel(R) Core(TM) i7-3615QM CPU @ 2.30GHz
Compilation log: <program source>:74:14: warning: comparison of integers of different signs: 'int' and 'const unsigned int'
while (i < srclen) {
~ ^ ~~~~~~
<program source>:88:14: warning: comparison of integers of different signs: 'int' and 'const unsigned int'
while (i < srclen) {
~ ^ ~~~~~~
<program source>:102:14: warning: comparison of integers of different signs: 'int' and 'const unsigned int'
while (i < srclen) {
~ ^ ~~~~~~
<program source>:274:5: warning: array index of '15' indexes past the end of an array (that contains 1 element)
ctx->buffer->mem_64[15] = SWAP64((uint64_t) (ctx->total * 8));
^ ~~
/Users/magnum/src/john/src/opencl_cryptsha512.h:94:5: note: array 'mem_64' declared here
uint64_t mem_64[1];
^
/Users/magnum/src/john/src/opencl_cryptsha512.h:23:18: note: expanded from macro 'uint64_t'
#define uint64_t unsigned long //Tip: unsigned long long int failed on compile (AMD).
^
Local work size (LWS) 1, global work size (GWS) 1024
Benchmarking: sha512crypt (rounds=5000) [OpenCL]... DONE
Raw: 1374 c/s real, 174 c/s virtual
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.
