Openwall GNU/*/Linux - a small security-enhanced Linux distro for servers
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Mon, 13 Aug 2012 02:22:29 +0200
From: magnum <john.magnum@...hmail.com>
To: john-dev@...ts.openwall.com
Subject: Re: bf-opencl fails self-test on CPU

On 2012-08-12 20:33, Solar Designer wrote:
> On Sun, Aug 12, 2012 at 11:28:05PM +0530, Sayantan Datta wrote:
>> With Intel SDK bf-opencl passed the self test on i5 2500k with 1450c/s .
>> However it fails with AMD SDK on same machine. Probably its not a very good
>> idea but still I was wondering if we could we use the Intel SDK for AMD
>> cpu's too.
> 
> We can, but my point was that there could be a bug somewhere in our
> bf-opencl code that only shows up on some builds - if so, we'd want to
> find and fix the bug even if it's of little practical relevance
> currently.  Indeed, it is also possible that the bug is actually in
> AMD's SDK.  We just don't know.

I believe this patch (just committed to git) strongly indicates that
it's a bug in the SDK - it works around the problem with no ill side
effects.

magnum@...l:src [1.7.9-jumbo-6-fixes]$ ../run/john -t -fo:bf-opencl
-plat=1 -dev=1
OpenCL platform 1: AMD Accelerated Parallel Processing, 2 device(s).
Using device 1: AMD FX(tm)-8120 Eight-Core Processor
****Please see 'opencl_bf_std.h' for device specific optimizations****
Benchmarking: OpenBSD Blowfish (x32) [OpenCL]... DONE
Raw:    3056 c/s real, 386 c/s virtual

magnum

commit 2654bcfc1479423079b2e01644030741399055e2
Author: magnum <john.magnum@...hmail.com>
Date:   Mon Aug 13 02:11:36 2012 +0200

    bf-opencl: Workaround for AMD APP problems running this format on CPU. This
    fix does not affect GPU at all.

diff --git a/src/opencl/bf_kernel.cl b/src/opencl/bf_kernel.cl
index 26d3036..9acb53f 100644
--- a/src/opencl/bf_kernel.cl
+++ b/src/opencl/bf_kernel.cl
@@ -4,6 +4,13 @@
  * Redistribution and use in source and binary forms, with or without modification, are permitted.
  * Based on Solar Designer implementation of bf_std.c in jtr-v1.7.8 
  */
+
+#ifdef DEVICE_IS_CPU
+#define MAYBE_LOCAL
+#else
+#define MAYBE_LOCAL	__local
+#endif
+
 #define BF_ROUNDS          16
 
 #define WORK_GROUP_SIZE    8
@@ -126,11 +133,11 @@ __kernel void blowfish(constant uint *salt __attribute__((max_constant_size(16))
 		uint BF_key_exp[18];
 		uint BF_current_P[18];
 
-		__local uint S_Buffer[WORK_GROUP_SIZE*1024];
-		__local uint *Sptr = S_Buffer + _index_S_local;
-		__local uint *Sptr2 = Sptr + 256;
-		__local uint *Sptr3 = Sptr + 512;
-		__local uint *Sptr4 = Sptr + 768;
+		MAYBE_LOCAL uint S_Buffer[WORK_GROUP_SIZE*1024];
+		MAYBE_LOCAL uint *Sptr = S_Buffer + _index_S_local;
+		MAYBE_LOCAL uint *Sptr2 = Sptr + 256;
+		MAYBE_LOCAL uint *Sptr3 = Sptr + 512;
+		MAYBE_LOCAL uint *Sptr4 = Sptr + 768;
 
 		for(i=0;i<18;i++){ 
 			tmp0=BF_current_P_global [pos_P(i)];

Powered by blists - more mailing lists

Your e-mail address:

Powered by Openwall GNU/*/Linux - Powered by OpenVZ