From 747c4a8d6a58e778ed4a6a79eda1083b43a8103e Mon Sep 17 00:00:00 2001
From: Jens Steube <jens.steube@gmail.com>
Date: Mon, 21 Dec 2015 20:21:53 +0100
Subject: [PATCH] Fix blowfish based algos for AMD, NV needs testing

---
 OpenCL/m03200.cl | 21 +++++++--------------
 OpenCL/m09000.cl | 22 ++++++++--------------
 2 files changed, 15 insertions(+), 28 deletions(-)

diff --git a/OpenCL/m03200.cl b/OpenCL/m03200.cl
index 3b76decfb..f62fd4f7d 100644
--- a/OpenCL/m03200.cl
+++ b/OpenCL/m03200.cl
@@ -294,15 +294,6 @@ __constant u32 c_sbox3[256] =
   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
-};
-
 #define BF_ROUND(L,R,N)       \
 {                             \
   uchar4 c = as_uchar4 ((L)); \
@@ -449,12 +440,14 @@ __kernel void __attribute__((reqd_work_group_size (8, 1, 1))) m03200_init (__glo
 
   // initstate
 
-  u32 P[18];
-
-  for (u32 i = 0; i < 18; i++)
+  u32 P[18] =
   {
-    P[i] = c_pbox[i];
-  }
+    0x243f6a88, 0x85a308d3, 0x13198a2e, 0x03707344,
+    0xa4093822, 0x299f31d0, 0x082efa98, 0xec4e6c89,
+    0x452821e6, 0x38d01377, 0xbe5466cf, 0x34e90c6c,
+    0xc0ac29b7, 0xc97c50dd, 0x3f84d5b5, 0xb5470917,
+    0x9216d5d9, 0x8979fb1b
+  };
 
   for (u32 i = 0; i < 256; i++)
   {
diff --git a/OpenCL/m09000.cl b/OpenCL/m09000.cl
index 7728c3d0c..8cbf21e88 100644
--- a/OpenCL/m09000.cl
+++ b/OpenCL/m09000.cl
@@ -294,14 +294,6 @@ __constant u32 c_sbox3[256] =
   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)       \
@@ -598,12 +590,14 @@ __kernel void __attribute__((reqd_work_group_size (8, 1, 1))) m09000_init (__glo
    * blowfish setkey
    */
 
-  u32 P[18];
-
-  for (u32 i = 0; i < 18; i++)
-  {
-    P[i] = c_pbox[i];
-  }
+  u32 P[18] =
+{
+  0x243f6a88, 0x85a308d3, 0x13198a2e, 0x03707344,
+  0xa4093822, 0x299f31d0, 0x082efa98, 0xec4e6c89,
+  0x452821e6, 0x38d01377, 0xbe5466cf, 0x34e90c6c,
+  0xc0ac29b7, 0xc97c50dd, 0x3f84d5b5, 0xb5470917,
+  0x9216d5d9, 0x8979fb1b
+};
 
   __local u32 S0_all[8][256];
   __local u32 S1_all[8][256];