projects
/
hashcat.git
/ commitdiff
commit
grep
author
committer
pickaxe
?
search:
re
summary
|
shortlog
|
log
|
commit
| commitdiff |
tree
raw
|
patch
|
inline
| side by side (parent:
a0221cd
)
Fix -m 9000 performance for AMD
author
Jens Steube
<jens.steube@gmail.com>
Mon, 9 May 2016 07:40:57 +0000
(09:40 +0200)
committer
Jens Steube
<jens.steube@gmail.com>
Mon, 9 May 2016 07:40:57 +0000
(09:40 +0200)
OpenCL/m09000.cl
patch
|
blob
|
history
diff --git
a/OpenCL/m09000.cl
b/OpenCL/m09000.cl
index
45b4bd8
..
9e12988
100644
(file)
--- a/
OpenCL/m09000.cl
+++ b/
OpenCL/m09000.cl
@@
-294,6
+294,15
@@
__constant u32 c_sbox3[256] =
0xb74e6132, 0xce77e25b, 0x578fdfe3, 0x3ac372e6
};
0xb74e6132, 0xce77e25b, 0x578fdfe3, 0x3ac372e6
};
+__constant u32 c_pbox[18] =
+{
+ 0x243f6a88, 0x85a308d3, 0x13198a2e, 0x03707344,
+ 0xa4093822, 0x299f31d0, 0x082efa98, 0xec4e6c89,
+ 0x452821e6, 0x38d01377, 0xbe5466cf, 0x34e90c6c,
+ 0xc0ac29b7, 0xc97c50dd, 0x3f84d5b5, 0xb5470917,
+ 0x9216d5d9, 0x8979fb1b
+};
+
#ifdef IS_AMD
#define BF_ROUND(L,R,N) \
{ \
#ifdef IS_AMD
#define BF_ROUND(L,R,N) \
{ \
@@
-605,14
+614,12
@@
__kernel void __attribute__((reqd_work_group_size (8, 1, 1))) m09000_init (__glo
* blowfish setkey
*/
* blowfish setkey
*/
- u32 P[18] =
+ u32 P[18];
+
+ for (u32 i = 0; i < 18; i++)
{
{
- 0x243f6a88, 0x85a308d3, 0x13198a2e, 0x03707344,
- 0xa4093822, 0x299f31d0, 0x082efa98, 0xec4e6c89,
- 0x452821e6, 0x38d01377, 0xbe5466cf, 0x34e90c6c,
- 0xc0ac29b7, 0xc97c50dd, 0x3f84d5b5, 0xb5470917,
- 0x9216d5d9, 0x8979fb1b
- };
+ P[i] = c_pbox[i];
+ }
__local u32 S0_all[8][256];
__local u32 S1_all[8][256];
__local u32 S0_all[8][256];
__local u32 S1_all[8][256];