Fix -m 9000 performance for AMD

OpenCL/m09000.cl

@@ -294,6 +294,15 @@ __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)       \
 {                             \
@@ -605,14 +614,12 @@ __kernel void __attribute__((reqd_work_group_size (8, 1, 1))) m09000_init (__glo
    * 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];