[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Tue, 21 Dec 2010 20:23:46 -0800
From: Dhiru Kholia <dhiru.kholia@...il.com>
To: john-users@...ts.openwall.com
Subject: Re: NTLM and OpenCL 1.0
On Tue, Dec 21, 2010 at 6:27 PM, Dhiru Kholia <dhiru.kholia@...il.com> wrote:
> I tried running your patch on a ATI Radeon 4870 system (with ATI
> Stream SDK 2.3, Catalyst 10.12, Linux 64-bit system, GCC 4.5.2). The
> OpenCL kernel fails to compile with "goto statement not allowed"
> errors (log attached). It seems that ATI's OpenCL implementation is
> somewhat more "strict" than NVIDIA's.
Attached patch (for OpenCL kernel) seems to work fine on both ATI 4870
and NVIDIA GT 240.
Benchmarking: NT MD4 [OpenCL 1.0]...
OpenCL Platform: <<<ATI Stream>>> and device: <<<ATI RV770>>>
DONE
Raw: 11707K c/s real, 11822K c/s virtual
CPU : AMD X3 720
Benchmarking: NT MD4 [OpenCL 1.0]...
OpenCL Platform: <<<NVIDIA CUDA>>> and device: <<<GeForce GT 240>>>
DONE
Raw: 23359K c/s real, 41391K c/s virtual
CPU : Core i5 750
The low performance numbers for 4870 don't make much sense (different
underlying CPUs?). Any guesses?
--
Cheers,
Dhiru
--- nt_opencl_kernel.c 2010-12-21 19:38:14.983334057 -0800
+++ nt_opencl_kernel_fixed.c 2010-12-21 19:52:53.120000725 -0800
@@ -55,15 +55,18 @@ __kernel void nt_crypt(const __global ui
uint key_chars = keys[i];//Coalescing access to global memory
uint cache_key = GET_CHAR(key_chars,ELEM_0);
//Extract 4 chars by cycle
+ int jump = 0;
while(cache_key)
{
md4_size++;
uint temp = GET_CHAR(key_chars,ELEM_1);
nt_buffer[nt_index] = ((temp ? temp : 0x80) << 16) | cache_key;
- if(!temp)
- goto key_cleaning;
-
+ if(!temp) {
+ jump = 1;
+ break;
+ }
+
md4_size++;
nt_index++;
cache_key = GET_CHAR(key_chars,ELEM_2);
@@ -76,8 +79,10 @@ __kernel void nt_crypt(const __global ui
temp = GET_CHAR(key_chars,ELEM_3);
nt_buffer[nt_index] = ((temp ? temp : 0x80) << 16) | cache_key;
- if(!temp)
- goto key_cleaning;
+ if(!temp) {
+ jump = 1;
+ break;
+ }
md4_size++;
nt_index++;
@@ -86,7 +91,8 @@ __kernel void nt_crypt(const __global ui
cache_key = GET_CHAR(key_chars,ELEM_0);
}
- nt_buffer[nt_index] = 0x80;
+ if(!jump)
+ nt_buffer[nt_index] = 0x80;
key_cleaning:
nt_index++;
Powered by blists - more mailing lists
Powered by Openwall GNU/*/Linux -
Powered by OpenVZ
