Merge pull request #167 from gm4tr1x/issue84

Issue84

diff --combined include/shared.h
index 2f9cdda,2ce9727..e92abe4
--- 1/include/shared.h
--- 2/include/shared.h
+++ b/include/shared.h
@@@ -1,7 -1,5 +1,7 @@@
  /**
 - * Author......: Jens Steube <jens.steube@gmail.com>
 + * Authors.....: Jens Steube <jens.steube@gmail.com>
 + *               Gabriele Gristina <matrix@hashcat.net>
 + *
   * License.....: MIT
   */
  
@@@ -46,6 -44,20 +46,20 @@@
  
  #endif
  
+ /**
+  * libraries stuff
+  */
+ 
+ #ifdef _WIN
+ #define hc_dlopen LoadLibrary
+ #define hc_dlclose FreeLibrary
+ #define hc_dlsym GetProcAddress
+ #else
+ #define hc_dlopen dlopen
+ #define hc_dlclose dlclose
+ #define hc_dlsym dlsym
+ #endif
+ 
  /**
   * system stuff
   */
@@@ -89,12 -101,10 +103,12 @@@
  
  #define CL_VENDOR_NV          "NVIDIA Corporation"
  #define CL_VENDOR_AMD         "Advanced Micro Devices, Inc."
 +#define CL_VENDOR_APPLE       "Apple"
  #define CL_VENDOR_POCL        "The pocl project"
  
  #define VENDOR_ID_AMD         4098
  #define VENDOR_ID_NV          4318
 +#define VENDOR_ID_APPLE       16925952
  #define VENDOR_ID_GENERIC     9999
  
  #define BLOCK_SIZE            64
@@@ -169,49 -179,9 +183,49 @@@ extern hc_thread_mutex_t mux_display
  #define PW_LENGTH_MAX_7400    15
  
  /**
 - * device accel / loops macro
 + * device accel macro
   */
  
 +#ifdef OSX
 +#define KERNEL_ACCEL_5000    16
 +#define KERNEL_ACCEL_6100    1
 +#define KERNEL_ACCEL_6211    4
 +#define KERNEL_ACCEL_6231    1
 +#define KERNEL_ACCEL_6241    4
 +#define KERNEL_ACCEL_8200    1
 +#define KERNEL_ACCEL_8700    2
 +#define KERNEL_ACCEL_9500    1
 +#define KERNEL_ACCEL_9600    1
 +#define KERNEL_ACCEL_10500   4
 +#define KERNEL_ACCEL_11300   1
 +#define KERNEL_ACCEL_11600   1
 +#define KERNEL_ACCEL_11700   1
 +#define KERNEL_ACCEL_11800   1
 +#define KERNEL_ACCEL_12200   1
 +#define KERNEL_ACCEL_12400   1
 +#define KERNEL_ACCEL_12500   1
 +#define KERNEL_ACCEL_13000   1
 +#else
 +#define KERNEL_ACCEL_5000    64
 +#define KERNEL_ACCEL_6100    8
 +#define KERNEL_ACCEL_6211    16
 +#define KERNEL_ACCEL_6231    4
 +#define KERNEL_ACCEL_6241    32
 +#define KERNEL_ACCEL_8200    2
 +#define KERNEL_ACCEL_8700    8
 +#define KERNEL_ACCEL_9500    8
 +#define KERNEL_ACCEL_9600    2
 +#define KERNEL_ACCEL_10500   64
 +#define KERNEL_ACCEL_11300   2
 +#define KERNEL_ACCEL_11600   2
 +#define KERNEL_ACCEL_11700   4
 +#define KERNEL_ACCEL_11800   4
 +#define KERNEL_ACCEL_12200   2
 +#define KERNEL_ACCEL_12400   64
 +#define KERNEL_ACCEL_12500   8
 +#define KERNEL_ACCEL_13000   8
 +#endif // OSX
 +
  #define KERNEL_ACCEL_0       128
  #define KERNEL_ACCEL_10      128
  #define KERNEL_ACCEL_11      128
@@@ -294,6 -264,7 +308,6 @@@
  #define KERNEL_ACCEL_4700    64
  #define KERNEL_ACCEL_4800    128
  #define KERNEL_ACCEL_4900    64
 -#define KERNEL_ACCEL_5000    64
  #define KERNEL_ACCEL_5100    64
  #define KERNEL_ACCEL_5200    8
  #define KERNEL_ACCEL_5300    32
@@@ -303,13 -274,17 +317,13 @@@
  #define KERNEL_ACCEL_5700    64
  #define KERNEL_ACCEL_5800    8
  #define KERNEL_ACCEL_6000    64
 -#define KERNEL_ACCEL_6100    8
 -#define KERNEL_ACCEL_6211    16
  #define KERNEL_ACCEL_6212    8
  #define KERNEL_ACCEL_6213    8
  #define KERNEL_ACCEL_6221    4
  #define KERNEL_ACCEL_6222    4
  #define KERNEL_ACCEL_6223    4
 -#define KERNEL_ACCEL_6231    4
  #define KERNEL_ACCEL_6232    4
  #define KERNEL_ACCEL_6233    4
 -#define KERNEL_ACCEL_6241    32
  #define KERNEL_ACCEL_6242    16
  #define KERNEL_ACCEL_6243    16
  #define KERNEL_ACCEL_6300    8
@@@ -330,10 -305,12 +344,10 @@@
  #define KERNEL_ACCEL_7900    2
  #define KERNEL_ACCEL_8000    8
  #define KERNEL_ACCEL_8100    64
 -#define KERNEL_ACCEL_8200    2
  #define KERNEL_ACCEL_8300    64
  #define KERNEL_ACCEL_8400    64
  #define KERNEL_ACCEL_8500    64
  #define KERNEL_ACCEL_8600    8
 -#define KERNEL_ACCEL_8700    8
  #define KERNEL_ACCEL_8800    8
  #define KERNEL_ACCEL_8900    16
  #define KERNEL_ACCEL_9000    2
@@@ -341,6 -318,8 +355,6 @@@
  #define KERNEL_ACCEL_9200    2
  #define KERNEL_ACCEL_9300    2
  #define KERNEL_ACCEL_9400    8
 -#define KERNEL_ACCEL_9500    8
 -#define KERNEL_ACCEL_9600    2
  #define KERNEL_ACCEL_9700    8
  #define KERNEL_ACCEL_9710    8
  #define KERNEL_ACCEL_9720    8
@@@ -355,6 -334,7 +369,6 @@@
  #define KERNEL_ACCEL_10400   8
  #define KERNEL_ACCEL_10410   8
  #define KERNEL_ACCEL_10420   8
 -#define KERNEL_ACCEL_10500   64
  #define KERNEL_ACCEL_10600   64
  #define KERNEL_ACCEL_10700   1
  #define KERNEL_ACCEL_10800   64
@@@ -362,143 -342,25 +376,143 @@@
  #define KERNEL_ACCEL_11000   64
  #define KERNEL_ACCEL_11100   64
  #define KERNEL_ACCEL_11200   64
 -#define KERNEL_ACCEL_11300   2
  #define KERNEL_ACCEL_11400   8
  #define KERNEL_ACCEL_11500   128
 -#define KERNEL_ACCEL_11600   2
 -#define KERNEL_ACCEL_11700   4
 -#define KERNEL_ACCEL_11800   4
  #define KERNEL_ACCEL_11900   2
  #define KERNEL_ACCEL_12000   2
  #define KERNEL_ACCEL_12100   2
 -#define KERNEL_ACCEL_12200   2
  #define KERNEL_ACCEL_12300   2
 -#define KERNEL_ACCEL_12400   64
 -#define KERNEL_ACCEL_12500   8
  #define KERNEL_ACCEL_12600   32
  #define KERNEL_ACCEL_12700   64
  #define KERNEL_ACCEL_12800   64
  #define KERNEL_ACCEL_12900   8
 -#define KERNEL_ACCEL_13000   8
  
 +/**
 + * device loops macro
 + */
 +
 +#ifdef OSX
 +#define KERNEL_LOOPS_0       2
 +#define KERNEL_LOOPS_10      2
 +#define KERNEL_LOOPS_11      2
 +#define KERNEL_LOOPS_12      2
 +#define KERNEL_LOOPS_20      2
 +#define KERNEL_LOOPS_21      2
 +#define KERNEL_LOOPS_22      2
 +#define KERNEL_LOOPS_23      2
 +#define KERNEL_LOOPS_30      2
 +#define KERNEL_LOOPS_40      2
 +#define KERNEL_LOOPS_50      2
 +#define KERNEL_LOOPS_60      2
 +#define KERNEL_LOOPS_100     2
 +#define KERNEL_LOOPS_101     2
 +#define KERNEL_LOOPS_110     2
 +#define KERNEL_LOOPS_111     2
 +#define KERNEL_LOOPS_112     2
 +#define KERNEL_LOOPS_120     2
 +#define KERNEL_LOOPS_121     2
 +#define KERNEL_LOOPS_122     2
 +#define KERNEL_LOOPS_124     2
 +#define KERNEL_LOOPS_130     2
 +#define KERNEL_LOOPS_131     2
 +#define KERNEL_LOOPS_132     2
 +#define KERNEL_LOOPS_133     2
 +#define KERNEL_LOOPS_140     2
 +#define KERNEL_LOOPS_141     2
 +#define KERNEL_LOOPS_150     2
 +#define KERNEL_LOOPS_160     2
 +#define KERNEL_LOOPS_190     2
 +#define KERNEL_LOOPS_200     2
 +#define KERNEL_LOOPS_300     2
 +#define KERNEL_LOOPS_900     2
 +#define KERNEL_LOOPS_1000    2
 +#define KERNEL_LOOPS_1100    2
 +#define KERNEL_LOOPS_1400    2
 +#define KERNEL_LOOPS_1410    2
 +#define KERNEL_LOOPS_1420    2
 +#define KERNEL_LOOPS_1421    2
 +#define KERNEL_LOOPS_1430    2
 +#define KERNEL_LOOPS_1440    2
 +#define KERNEL_LOOPS_1441    2
 +#define KERNEL_LOOPS_1450    2
 +#define KERNEL_LOOPS_1460    2
 +#define KERNEL_LOOPS_1700    2
 +#define KERNEL_LOOPS_1710    2
 +#define KERNEL_LOOPS_1711    2
 +#define KERNEL_LOOPS_1720    2
 +#define KERNEL_LOOPS_1722    2
 +#define KERNEL_LOOPS_1730    2
 +#define KERNEL_LOOPS_1731    2
 +#define KERNEL_LOOPS_1740    2
 +#define KERNEL_LOOPS_1750    2
 +#define KERNEL_LOOPS_1760    2
 +#define KERNEL_LOOPS_2400    2
 +#define KERNEL_LOOPS_2410    2
 +#define KERNEL_LOOPS_2600    2
 +#define KERNEL_LOOPS_2611    2
 +#define KERNEL_LOOPS_2612    2
 +#define KERNEL_LOOPS_2711    2
 +#define KERNEL_LOOPS_2811    2
 +#define KERNEL_LOOPS_3100    2
 +#define KERNEL_LOOPS_3200    4
 +#define KERNEL_LOOPS_3710    2
 +#define KERNEL_LOOPS_3711    2
 +#define KERNEL_LOOPS_3800    2
 +#define KERNEL_LOOPS_4300    2
 +#define KERNEL_LOOPS_4400    2
 +#define KERNEL_LOOPS_4500    2
 +#define KERNEL_LOOPS_4700    2
 +#define KERNEL_LOOPS_4800    2
 +#define KERNEL_LOOPS_4900    2
 +#define KERNEL_LOOPS_5000    2
 +#define KERNEL_LOOPS_5100    2
 +#define KERNEL_LOOPS_5300    2
 +#define KERNEL_LOOPS_5400    2
 +#define KERNEL_LOOPS_5500    2
 +#define KERNEL_LOOPS_5600    2
 +#define KERNEL_LOOPS_5700    2
 +#define KERNEL_LOOPS_6000    2
 +#define KERNEL_LOOPS_6100    2
 +#define KERNEL_LOOPS_6231    2
 +#define KERNEL_LOOPS_6232    2
 +#define KERNEL_LOOPS_6233    2
 +#define KERNEL_LOOPS_6900    2
 +#define KERNEL_LOOPS_7300    2
 +#define KERNEL_LOOPS_7500    2
 +#define KERNEL_LOOPS_7600    2
 +#define KERNEL_LOOPS_7700    2
 +#define KERNEL_LOOPS_7800    2
 +#define KERNEL_LOOPS_8000    2
 +#define KERNEL_LOOPS_8100    2
 +#define KERNEL_LOOPS_8300    2
 +#define KERNEL_LOOPS_8400    2
 +#define KERNEL_LOOPS_8500    2
 +#define KERNEL_LOOPS_8600    2
 +#define KERNEL_LOOPS_8700    4
 +#define KERNEL_LOOPS_9700    2
 +#define KERNEL_LOOPS_9710    8
 +#define KERNEL_LOOPS_9720    8
 +#define KERNEL_LOOPS_9800    2
 +#define KERNEL_LOOPS_9810    2
 +#define KERNEL_LOOPS_9820    2
 +#define KERNEL_LOOPS_9900    2
 +#define KERNEL_LOOPS_10100   2
 +#define KERNEL_LOOPS_10200   2
 +#define KERNEL_LOOPS_10400   2
 +#define KERNEL_LOOPS_10410   2
 +#define KERNEL_LOOPS_10420   2
 +#define KERNEL_LOOPS_10600   2
 +#define KERNEL_LOOPS_10700   2
 +#define KERNEL_LOOPS_10800   2
 +#define KERNEL_LOOPS_11000   2
 +#define KERNEL_LOOPS_11100   2
 +#define KERNEL_LOOPS_11200   2
 +#define KERNEL_LOOPS_11400   2
 +#define KERNEL_LOOPS_11500   2
 +#define KERNEL_LOOPS_11700   8
 +#define KERNEL_LOOPS_11800   8
 +#define KERNEL_LOOPS_12600   2
 +#else
  #define KERNEL_LOOPS_0       256
  #define KERNEL_LOOPS_10      256
  #define KERNEL_LOOPS_11      256
@@@ -531,7 -393,11 +545,7 @@@
  #define KERNEL_LOOPS_190     128
  #define KERNEL_LOOPS_200     128
  #define KERNEL_LOOPS_300     64
 -#define KERNEL_LOOPS_400     256
 -#define KERNEL_LOOPS_500     256
 -#define KERNEL_LOOPS_501     256
  #define KERNEL_LOOPS_900     256
 -#define KERNEL_LOOPS_910     256
  #define KERNEL_LOOPS_1000    256
  #define KERNEL_LOOPS_1100    128
  #define KERNEL_LOOPS_1400    64
@@@ -543,6 -409,8 +557,6 @@@
  #define KERNEL_LOOPS_1441    64
  #define KERNEL_LOOPS_1450    32
  #define KERNEL_LOOPS_1460    32
 -#define KERNEL_LOOPS_1500    256
 -#define KERNEL_LOOPS_1600    256
  #define KERNEL_LOOPS_1700    32
  #define KERNEL_LOOPS_1710    32
  #define KERNEL_LOOPS_1711    32
@@@ -553,13 -421,17 +567,13 @@@
  #define KERNEL_LOOPS_1740    32
  #define KERNEL_LOOPS_1750    16
  #define KERNEL_LOOPS_1760    16
 -#define KERNEL_LOOPS_1800    16
 -#define KERNEL_LOOPS_2100    256
  #define KERNEL_LOOPS_2400    256
  #define KERNEL_LOOPS_2410    256
 -#define KERNEL_LOOPS_2500    256
  #define KERNEL_LOOPS_2600    128
  #define KERNEL_LOOPS_2611    128
  #define KERNEL_LOOPS_2612    128
  #define KERNEL_LOOPS_2711    64
  #define KERNEL_LOOPS_2811    64
 -#define KERNEL_LOOPS_3000    256
  #define KERNEL_LOOPS_3100    16
  #define KERNEL_LOOPS_3200    16
  #define KERNEL_LOOPS_3710    128
@@@ -573,101 -445,88 +587,101 @@@
  #define KERNEL_LOOPS_4900    128
  #define KERNEL_LOOPS_5000    64
  #define KERNEL_LOOPS_5100    256
 -#define KERNEL_LOOPS_5200    256
  #define KERNEL_LOOPS_5300    32
  #define KERNEL_LOOPS_5400    32
  #define KERNEL_LOOPS_5500    128
  #define KERNEL_LOOPS_5600    64
  #define KERNEL_LOOPS_5700    64
 -#define KERNEL_LOOPS_5800    256
  #define KERNEL_LOOPS_6000    64
  #define KERNEL_LOOPS_6100    64
 -#define KERNEL_LOOPS_6211    200
 -#define KERNEL_LOOPS_6212    200
 -#define KERNEL_LOOPS_6213    200
 -#define KERNEL_LOOPS_6221    200
 -#define KERNEL_LOOPS_6222    200
 -#define KERNEL_LOOPS_6223    200
  #define KERNEL_LOOPS_6231    200
  #define KERNEL_LOOPS_6232    200
  #define KERNEL_LOOPS_6233    200
 -#define KERNEL_LOOPS_6241    200
 -#define KERNEL_LOOPS_6242    200
 -#define KERNEL_LOOPS_6243    200
 -#define KERNEL_LOOPS_6300    256
 -#define KERNEL_LOOPS_6400    256
 -#define KERNEL_LOOPS_6500    256
 -#define KERNEL_LOOPS_6600    200
 -#define KERNEL_LOOPS_6700    256
 -#define KERNEL_LOOPS_6800    200
  #define KERNEL_LOOPS_6900    64
 -#define KERNEL_LOOPS_7100    256
 -#define KERNEL_LOOPS_7200    200
  #define KERNEL_LOOPS_7300    64
 -#define KERNEL_LOOPS_7400    200
  #define KERNEL_LOOPS_7500    16
  #define KERNEL_LOOPS_7600    128
  #define KERNEL_LOOPS_7700    128
  #define KERNEL_LOOPS_7800    64
 -#define KERNEL_LOOPS_7900    256
  #define KERNEL_LOOPS_8000    64
  #define KERNEL_LOOPS_8100    128
 -#define KERNEL_LOOPS_8200    200
  #define KERNEL_LOOPS_8300    64
  #define KERNEL_LOOPS_8400    64
  #define KERNEL_LOOPS_8500    16
  #define KERNEL_LOOPS_8600    16
  #define KERNEL_LOOPS_8700    16
 -#define KERNEL_LOOPS_8800    256
 -#define KERNEL_LOOPS_8900    1
 -#define KERNEL_LOOPS_9000    16
 -#define KERNEL_LOOPS_9100    256
 -#define KERNEL_LOOPS_9200    200
 -#define KERNEL_LOOPS_9300    1
 -#define KERNEL_LOOPS_9400    200
 -#define KERNEL_LOOPS_9500    200
 -#define KERNEL_LOOPS_9600    200
  #define KERNEL_LOOPS_9700    200
  #define KERNEL_LOOPS_9710    200
  #define KERNEL_LOOPS_9720    200
  #define KERNEL_LOOPS_9800    200
 -#define KERNEL_LOOPS_9810    200
  #define KERNEL_LOOPS_9820    200
 +#define KERNEL_LOOPS_9810    200
  #define KERNEL_LOOPS_9900    256
 -#define KERNEL_LOOPS_10000   200
  #define KERNEL_LOOPS_10100   512
  #define KERNEL_LOOPS_10200   64
 -#define KERNEL_LOOPS_10300   128
  #define KERNEL_LOOPS_10400   256
  #define KERNEL_LOOPS_10410   256
  #define KERNEL_LOOPS_10420   256
 -#define KERNEL_LOOPS_10500   64
  #define KERNEL_LOOPS_10600   64
  #define KERNEL_LOOPS_10700   64
  #define KERNEL_LOOPS_10800   32
 -#define KERNEL_LOOPS_10900   200
  #define KERNEL_LOOPS_11000   256
  #define KERNEL_LOOPS_11100   128
  #define KERNEL_LOOPS_11200   128
 -#define KERNEL_LOOPS_11300   256
  #define KERNEL_LOOPS_11400   128
  #define KERNEL_LOOPS_11500   256
 -#define KERNEL_LOOPS_11600   512
  #define KERNEL_LOOPS_11700   64
  #define KERNEL_LOOPS_11800   64
 +#define KERNEL_LOOPS_12600   32
 +#endif // OSX
 +
 +#define KERNEL_LOOPS_400     256
 +#define KERNEL_LOOPS_500     256
 +#define KERNEL_LOOPS_501     256
 +#define KERNEL_LOOPS_910     256
 +#define KERNEL_LOOPS_1500    256
 +#define KERNEL_LOOPS_1600    256
 +#define KERNEL_LOOPS_1800    16
 +#define KERNEL_LOOPS_2100    256
 +#define KERNEL_LOOPS_2500    256
 +#define KERNEL_LOOPS_3000    256
 +#define KERNEL_LOOPS_5200    256
 +#define KERNEL_LOOPS_5800    256
 +#define KERNEL_LOOPS_6211    200
 +#define KERNEL_LOOPS_6212    200
 +#define KERNEL_LOOPS_6213    200
 +#define KERNEL_LOOPS_6221    200
 +#define KERNEL_LOOPS_6222    200
 +#define KERNEL_LOOPS_6223    200
 +#define KERNEL_LOOPS_6241    200
 +#define KERNEL_LOOPS_6242    200
 +#define KERNEL_LOOPS_6243    200
 +#define KERNEL_LOOPS_6300    256
 +#define KERNEL_LOOPS_6400    256
 +#define KERNEL_LOOPS_6500    256
 +#define KERNEL_LOOPS_6600    200
 +#define KERNEL_LOOPS_6700    256
 +#define KERNEL_LOOPS_6800    200
 +#define KERNEL_LOOPS_7100    256
 +#define KERNEL_LOOPS_7200    200
 +#define KERNEL_LOOPS_7400    200
 +#define KERNEL_LOOPS_7900    256
 +#define KERNEL_LOOPS_8200    200
 +#define KERNEL_LOOPS_8800    256
 +#define KERNEL_LOOPS_8900    1
 +#define KERNEL_LOOPS_9000    16
 +#define KERNEL_LOOPS_9100    256
 +#define KERNEL_LOOPS_9200    200
 +#define KERNEL_LOOPS_9300    1
 +#define KERNEL_LOOPS_9400    200
 +#define KERNEL_LOOPS_9500    200
 +#define KERNEL_LOOPS_9600    200
 +#define KERNEL_LOOPS_10000   200
 +#define KERNEL_LOOPS_10300   128
 +#define KERNEL_LOOPS_10500   64
 +#define KERNEL_LOOPS_10900   200
 +#define KERNEL_LOOPS_11300   256
 +#define KERNEL_LOOPS_11600   512
  #define KERNEL_LOOPS_11900   200
  #define KERNEL_LOOPS_12000   200
  #define KERNEL_LOOPS_12100   200
@@@ -675,6 -534,7 +689,6 @@@
  #define KERNEL_LOOPS_12300   256
  #define KERNEL_LOOPS_12400   256
  #define KERNEL_LOOPS_12500   256
 -#define KERNEL_LOOPS_12600   32
  #define KERNEL_LOOPS_12700   10
  #define KERNEL_LOOPS_12800   100
  #define KERNEL_LOOPS_12900   64
@@@ -1623,14 -1483,14 +1637,14 @@@
   */
  
  #define DGST_SIZE_0                 0
 -#define DGST_SIZE_4_2               (2  * sizeof (uint))        // 8
 -#define DGST_SIZE_4_4               (4  * sizeof (uint))        // 16
 -#define DGST_SIZE_4_5               (5  * sizeof (uint))        // 20
 -#define DGST_SIZE_4_6               (6  * sizeof (uint))        // 24
 -#define DGST_SIZE_4_8               (8  * sizeof (uint))        // 32
 -#define DGST_SIZE_4_16              (16 * sizeof (uint))        // 64 !!!
 -#define DGST_SIZE_4_32              (32 * sizeof (uint))        // 128 !!!
 -#define DGST_SIZE_4_64              (64 * sizeof (uint))        // 256
 +#define DGST_SIZE_4_2               (2  * sizeof (uint))   // 8
 +#define DGST_SIZE_4_4               (4  * sizeof (uint))   // 16
 +#define DGST_SIZE_4_5               (5  * sizeof (uint))   // 20
 +#define DGST_SIZE_4_6               (6  * sizeof (uint))   // 24
 +#define DGST_SIZE_4_8               (8  * sizeof (uint))   // 32
 +#define DGST_SIZE_4_16              (16 * sizeof (uint))   // 64 !!!
 +#define DGST_SIZE_4_32              (32 * sizeof (uint))   // 128 !!!
 +#define DGST_SIZE_4_64              (64 * sizeof (uint))   // 256
  #define DGST_SIZE_8_8               (8  * sizeof (u64))    // 64 !!!
  #define DGST_SIZE_8_16              (16 * sizeof (u64))    // 128 !!!
  #define DGST_SIZE_8_25              (25 * sizeof (u64))    // 200

diff --combined include/types.h
index 4369935,cff1e02..ac7611a
--- 1/include/types.h
--- 2/include/types.h
+++ b/include/types.h
@@@ -39,22 -39,22 +39,22 @@@ typedef struc
  
  typedef struct
  {
 -  int   V;
 -  int   R;
 -  int   P;
 +  int  V;
 +  int  R;
 +  int  P;
  
 -  int   enc_md;
 +  int  enc_md;
  
 -  uint  id_buf[8];
 -  uint  u_buf[32];
 -  uint  o_buf[32];
 +  uint id_buf[8];
 +  uint u_buf[32];
 +  uint o_buf[32];
  
 -  int   id_len;
 -  int   o_len;
 -  int   u_len;
 +  int  id_len;
 +  int  o_len;
 +  int  u_len;
  
 -  uint  rc4key[2];
 -  uint  rc4data[2];
 +  uint rc4key[2];
 +  uint rc4data[2];
  
  } pdf_t;
  
@@@ -161,7 -161,7 +161,7 @@@ typedef struc
  
  typedef struct
  {
 -  u8 cipher[1040];
 +  u8   cipher[1040];
  
  } agilekey_t;
  
@@@ -239,10 -239,10 +239,10 @@@ typedef struc
    {
      uint dgst32[16];
      u64  dgst64[8];
 -  };
 +  } d;
  
 -  uint  dgst_len;
 -  uint  W_len;
 +  uint dgst_len;
 +  uint W_len;
  
  } pdf17l8_tmp_t;
  
@@@ -260,10 -260,10 +260,10 @@@ typedef struc
  
  typedef struct
  {
 -  u64 l_alt_result[8];
 +  u64  l_alt_result[8];
  
 -  u64 l_p_bytes[2];
 -  u64 l_s_bytes[2];
 +  u64  l_p_bytes[2];
 +  u64  l_s_bytes[2];
  
  } sha512crypt_tmp_t;
  
@@@ -288,7 -288,7 +288,7 @@@ typedef struc
  
  typedef struct
  {
 -  u64 dgst[8];
 +  u64  dgst[8];
  
  } bitcoin_wallet_tmp_t;
  
@@@ -362,11 -362,11 +362,11 @@@ typedef struc
  
  typedef struct
  {
 -  u64 ipad[8];
 -  u64 opad[8];
 +  u64  ipad[8];
 +  u64  opad[8];
  
 -  u64 dgst[32];
 -  u64 out[32];
 +  u64  dgst[32];
 +  u64  out[32];
  
  } tc64_tmp_t;
  
@@@ -415,11 -415,11 +415,11 @@@ typedef struc
  
  typedef struct
  {
 -  u64 ipad[8];
 -  u64 opad[8];
 +  u64  ipad[8];
 +  u64  opad[8];
  
 -  u64 dgst[8];
 -  u64 out[8];
 +  u64  dgst[8];
 +  u64  out[8];
  
  } sha512aix_tmp_t;
  
@@@ -435,7 -435,7 +435,7 @@@ typedef struc
  
  typedef struct
  {
 -  u64 digest_buf[8];
 +  u64  digest_buf[8];
  
  } drupal7_tmp_t;
  
@@@ -463,7 -463,7 +463,7 @@@ typedef struc
  
  typedef struct
  {
 -  u64 out[8];
 +  u64  out[8];
  
  } office2013_tmp_t;
  
@@@ -475,57 -475,57 +475,57 @@@ typedef struc
  
  typedef struct
  {
 -  u32 ipad[4];
 -  u32 opad[4];
 +  u32  ipad[4];
 +  u32  opad[4];
  
 -  u32 dgst[32];
 -  u32 out[32];
 +  u32  dgst[32];
 +  u32  out[32];
  
  } pbkdf2_md5_tmp_t;
  
  typedef struct
  {
 -  u32 ipad[5];
 -  u32 opad[5];
 +  u32  ipad[5];
 +  u32  opad[5];
  
 -  u32 dgst[32];
 -  u32 out[32];
 +  u32  dgst[32];
 +  u32  out[32];
  
  } pbkdf2_sha1_tmp_t;
  
  typedef struct
  {
 -  u32 ipad[8];
 -  u32 opad[8];
 +  u32  ipad[8];
 +  u32  opad[8];
  
 -  u32 dgst[32];
 -  u32 out[32];
 +  u32  dgst[32];
 +  u32  out[32];
  
  } pbkdf2_sha256_tmp_t;
  
  typedef struct
  {
 -  u64 ipad[8];
 -  u64 opad[8];
 +  u64  ipad[8];
 +  u64  opad[8];
  
 -  u64 dgst[16];
 -  u64 out[16];
 +  u64  dgst[16];
 +  u64  out[16];
  
  } pbkdf2_sha512_tmp_t;
  
  typedef struct
  {
 -  u64 out[8];
 +  u64  out[8];
  
  } ecryptfs_tmp_t;
  
  typedef struct
  {
 -  u64 ipad[8];
 -  u64 opad[8];
 +  u64  ipad[8];
 +  u64  opad[8];
  
 -  u64 dgst[16];
 -  u64 out[16];
 +  u64  dgst[16];
 +  u64  out[16];
  
  } oraclet_tmp_t;
  
@@@ -613,8 -613,8 +613,8 @@@ typedef struc
  
  typedef struct
  {
 -  uint     key;
 -  u64 val;
 +  uint key;
 +  u64  val;
  
  } hcstat_table_t;
  
@@@ -662,7 -662,7 +662,7 @@@ typedef struc
  
  typedef struct
  {
 -  u64 cnt;
 +  u64    cnt;
  
    #ifdef _POSIX
    struct stat stat;
@@@ -703,7 -703,7 +703,7 @@@ typedef struc
      u8   hc1[1][256];
      u32  hi1[1][ 64];
      u64  hl1[1][ 32];
 -  };
 +  } h;
  
    uint pw_len;
    uint alignment_placeholder_1;
@@@ -742,31 -742,31 +742,31 @@@ typedef struc
  
  typedef struct
  {
 -  u32    version_bin;
 -  char        cwd[256];
 -  u32    pid;
 +  u32  version_bin;
 +  char cwd[1024];
 +  u32  pid;
  
 -  u32    dictpos;
 -  u32    maskpos;
 +  u32  dictpos;
 +  u32  maskpos;
  
 -  u64    words_cur;
 +  u64  words_cur;
  
 -  u32    argc;
 -  char      **argv;
 +  u32  argc;
 +  char **argv;
  
  } restore_data_t;
  
  typedef struct
  {
 -  char     *file_name;
 -  long      seek;
 -  time_t    ctime;
 +  char   *file_name;
 +  long   seek;
 +  time_t ctime;
  
  } outfile_data_t;
  
  typedef struct
  {
 -  char     *buf;
 +  char *buf;
    u32  incr;
    u32  avail;
    u32  cnt;
@@@ -824,165 -824,165 +824,165 @@@ typedef struc
  #define PW_MAX1     (PW_MAX + 1)
  #define PW_DICTMAX  31
  #define PW_DICTMAX1 (PW_DICTMAX + 1)
 +#define PARAMCNT    32
  
  struct __hc_device_param
  {
    cl_device_id      device;
    cl_device_type    device_type;
  
 -  uint              device_id;
 -  uint              platform_devices_id;    // for mapping with hms devices
 +  uint    device_id;
 +  uint    platform_devices_id;   // for mapping with hms devices
 +
 +  uint    skipped;
  
 -  uint              skipped;
 +  uint    sm_major;
 +  uint    sm_minor;
 +  uint    kernel_exec_timeout;
  
 -  uint              sm_major;
 -  uint              sm_minor;
 -  uint              kernel_exec_timeout;
 +  uint    device_processors;
 +  uint    device_processor_cores;
 +  u64     device_maxmem_alloc;
 +  u64     device_global_mem;
 +  u32     device_maxclock_frequency;
  
 -  uint              device_processors;
 -  uint              device_processor_cores;
 -  u64               device_maxmem_alloc;
 -  u64               device_global_mem;
 -  u32               device_maxclock_frequency;
 +  uint    vector_width;
  
 -  uint              vector_width;
 +  uint    kernel_threads;
 +  uint    kernel_accel;
 +  uint    kernel_power;          // these both are based on their _user counterpart
 +  uint    kernel_blocks;         // but are modified by autotuner and used inside crack loops
 +  uint    kernel_power_user;
 +  uint    kernel_blocks_user;
  
 -  uint              kernel_threads;
 -  uint              kernel_accel;
 -  uint              kernel_power;          // these both are based on their _user counterpart
 -  uint              kernel_blocks;         // but are modified by autotuner and used inside crack loops
 -  uint              kernel_power_user;
 -  uint              kernel_blocks_user;
 +  uint    size_pws;
 +  uint    size_tmps;
 +  uint    size_hooks;
 +  uint    size_root_css;
 +  uint    size_markov_css;
 +  uint    size_digests;
 +  uint    size_salts;
 +  uint    size_shown;
 +  uint    size_results;
 +  uint    size_plains;
  
 -  uint              size_pws;
 -  uint              size_tmps;
 -  uint              size_hooks;
 -  uint              size_root_css;
 -  uint              size_markov_css;
 -  uint              size_digests;
 -  uint              size_salts;
 -  uint              size_shown;
 -  uint              size_results;
 -  uint              size_plains;
 +  uint  (*pw_add)       (struct __hc_device_param *, const u8 *, const uint);
  
 -  uint (*pw_add)    (struct __hc_device_param *, const u8 *, const uint);
 +  void  (*pw_transpose) (const pw_t *, pw_t *);
  
 -  void (*pw_transpose) (const pw_t *, pw_t *);
 +  FILE   *combs_fp;
 +  comb_t *combs_buf;
  
 -  FILE             *combs_fp;
 -  comb_t           *combs_buf;
 +  void   *hooks_buf;
  
 -  void             *hooks_buf;
 +  pw_cache_t *pw_caches;
  
 -  pw_cache_t       *pw_caches;
 +  pw_t   *pws_buf;
 +  uint    pws_cnt;
 +  u64     pw_cnt;
  
 -  pw_t             *pws_buf;
 -  uint              pws_cnt;
 -  u64               pw_cnt;
 +  u64     words_off;
 +  u64     words_done;
  
 -  u64               words_off;
 -  u64               words_done;
 +  uint   *result;
  
 -  uint             *result;
 +  uint    outerloop_pos;
 +  uint    outerloop_left;
  
 -  uint              outerloop_pos;
 -  uint              outerloop_left;
 +  uint    innerloop_pos;
 +  uint    innerloop_left;
  
 -  uint              innerloop_pos;
 -  uint              innerloop_left;
 +  uint    speed_pos;
 +  u64     speed_cnt[SPEED_CACHE];
 +  float   speed_ms[SPEED_CACHE];
  
 -  uint              speed_pos;
 -  u64               speed_cnt[SPEED_CACHE];
 -  float             speed_ms[SPEED_CACHE];
 -  hc_timer_t        speed_rec[SPEED_CACHE];
 +  hc_timer_t speed_rec[SPEED_CACHE];
  
 -  hc_timer_t        timer_speed;
 +  hc_timer_t timer_speed;
  
    // device specific attributes starting
  
 -  char             *device_name;
 -  char             *device_name_chksum;
 -  char             *device_version;
 -  char             *driver_version;
 -
 -  cl_uint           vendor_id;
 -
 -  cl_kernel         kernel1;
 -  cl_kernel         kernel12;
 -  cl_kernel         kernel2;
 -  cl_kernel         kernel23;
 -  cl_kernel         kernel3;
 -  cl_kernel         kernel_mp;
 -  cl_kernel         kernel_mp_l;
 -  cl_kernel         kernel_mp_r;
 -  cl_kernel         kernel_amp;
 -  cl_kernel         kernel_tb;
 -  cl_kernel         kernel_tm;
 -  cl_kernel         kernel_weak;
 -
 -  cl_context        context;
 -
 -  cl_program        program;
 -  cl_program        program_mp;
 -  cl_program        program_amp;
 -  cl_program        program_weak;
 -
 -  cl_command_queue  command_queue;
 -
 -  cl_mem            d_pws_buf;
 -  cl_mem            d_pws_amp_buf;
 -  cl_mem            d_words_buf_l;
 -  cl_mem            d_words_buf_r;
 -  cl_mem            d_rules;
 -  cl_mem            d_rules_c;
 -  cl_mem            d_combs;
 -  cl_mem            d_combs_c;
 -  cl_mem            d_bfs;
 -  cl_mem            d_bfs_c;
 -  cl_mem            d_tm_c;
 -  cl_mem            d_bitmap_s1_a;
 -  cl_mem            d_bitmap_s1_b;
 -  cl_mem            d_bitmap_s1_c;
 -  cl_mem            d_bitmap_s1_d;
 -  cl_mem            d_bitmap_s2_a;
 -  cl_mem            d_bitmap_s2_b;
 -  cl_mem            d_bitmap_s2_c;
 -  cl_mem            d_bitmap_s2_d;
 -  cl_mem            d_plain_bufs;
 -  cl_mem            d_digests_buf;
 -  cl_mem            d_digests_shown;
 -  cl_mem            d_salt_bufs;
 -  cl_mem            d_esalt_bufs;
 -  cl_mem            d_bcrypt_bufs;
 -  cl_mem            d_tmps;
 -  cl_mem            d_hooks;
 -  cl_mem            d_result;
 -  cl_mem            d_scryptV_buf;
 -  cl_mem            d_root_css_buf;
 -  cl_mem            d_markov_css_buf;
 -
 -  #define PARAMCNT 32
 -
 -  void             *kernel_params[PARAMCNT];
 -  void             *kernel_params_mp[PARAMCNT];
 -  void             *kernel_params_mp_r[PARAMCNT];
 -  void             *kernel_params_mp_l[PARAMCNT];
 -  void             *kernel_params_amp[PARAMCNT];
 -  void             *kernel_params_tb[PARAMCNT];
 -  void             *kernel_params_tm[PARAMCNT];
 -
 -  u32          kernel_params_buf32[PARAMCNT];
 -
 -  u32          kernel_params_mp_buf32[PARAMCNT];
 -  u64          kernel_params_mp_buf64[PARAMCNT];
 -
 -  u32          kernel_params_mp_r_buf32[PARAMCNT];
 -  u64          kernel_params_mp_r_buf64[PARAMCNT];
 -
 -  u32          kernel_params_mp_l_buf32[PARAMCNT];
 -  u64          kernel_params_mp_l_buf64[PARAMCNT];
 -
 -  u32          kernel_params_amp_buf32[PARAMCNT];
 +  char   *device_name;
 +  char   *device_name_chksum;
 +  char   *device_version;
 +  char   *driver_version;
 +
 +  cl_uint vendor_id;
 +
 +  cl_kernel  kernel1;
 +  cl_kernel  kernel12;
 +  cl_kernel  kernel2;
 +  cl_kernel  kernel23;
 +  cl_kernel  kernel3;
 +  cl_kernel  kernel_mp;
 +  cl_kernel  kernel_mp_l;
 +  cl_kernel  kernel_mp_r;
 +  cl_kernel  kernel_amp;
 +  cl_kernel  kernel_tb;
 +  cl_kernel  kernel_tm;
 +  cl_kernel  kernel_weak;
 +
 +  cl_context context;
 +
 +  cl_program program;
 +  cl_program program_mp;
 +  cl_program program_amp;
 +  cl_program program_weak;
 +
 +  cl_command_queue command_queue;
 +
 +  cl_mem  d_pws_buf;
 +  cl_mem  d_pws_amp_buf;
 +  cl_mem  d_words_buf_l;
 +  cl_mem  d_words_buf_r;
 +  cl_mem  d_rules;
 +  cl_mem  d_rules_c;
 +  cl_mem  d_combs;
 +  cl_mem  d_combs_c;
 +  cl_mem  d_bfs;
 +  cl_mem  d_bfs_c;
 +  cl_mem  d_tm_c;
 +  cl_mem  d_bitmap_s1_a;
 +  cl_mem  d_bitmap_s1_b;
 +  cl_mem  d_bitmap_s1_c;
 +  cl_mem  d_bitmap_s1_d;
 +  cl_mem  d_bitmap_s2_a;
 +  cl_mem  d_bitmap_s2_b;
 +  cl_mem  d_bitmap_s2_c;
 +  cl_mem  d_bitmap_s2_d;
 +  cl_mem  d_plain_bufs;
 +  cl_mem  d_digests_buf;
 +  cl_mem  d_digests_shown;
 +  cl_mem  d_salt_bufs;
 +  cl_mem  d_esalt_bufs;
 +  cl_mem  d_bcrypt_bufs;
 +  cl_mem  d_tmps;
 +  cl_mem  d_hooks;
 +  cl_mem  d_result;
 +  cl_mem  d_scryptV_buf;
 +  cl_mem  d_root_css_buf;
 +  cl_mem  d_markov_css_buf;
 +
 +  void   *kernel_params[PARAMCNT];
 +  void   *kernel_params_mp[PARAMCNT];
 +  void   *kernel_params_mp_r[PARAMCNT];
 +  void   *kernel_params_mp_l[PARAMCNT];
 +  void   *kernel_params_amp[PARAMCNT];
 +  void   *kernel_params_tb[PARAMCNT];
 +  void   *kernel_params_tm[PARAMCNT];
 +
 +  u32     kernel_params_buf32[PARAMCNT];
 +
 +  u32     kernel_params_mp_buf32[PARAMCNT];
 +  u64     kernel_params_mp_buf64[PARAMCNT];
 +
 +  u32     kernel_params_mp_r_buf32[PARAMCNT];
 +  u64     kernel_params_mp_r_buf64[PARAMCNT];
 +
 +  u32     kernel_params_mp_l_buf32[PARAMCNT];
 +  u64     kernel_params_mp_l_buf64[PARAMCNT];
 +
 +  u32     kernel_params_amp_buf32[PARAMCNT];
  };
  
  typedef struct __hc_device_param hc_device_param_t;
@@@ -1002,11 -1002,12 +1002,11 @@@ typedef struc
  
    } adapter_index;
  
 -  int od_version;
 +  int     od_version;
 +  int     fan_supported;
  
 -  int fan_supported;
 -
 -  // int busid; // used for CL_DEVICE_TOPOLOGY_AMD but broken for dual GPUs
 -  // int devid; // used for CL_DEVICE_TOPOLOGY_AMD but broken for dual GPUs
 +  // int     busid; // used for CL_DEVICE_TOPOLOGY_AMD but broken for dual GPUs
 +  // int     devid; // used for CL_DEVICE_TOPOLOGY_AMD but broken for dual GPUs
  
  } hm_attrs_t;
  #endif // HAVE_HWMON
@@@ -1017,75 -1018,79 +1017,81 @@@ typedef struc
     * threads
     */
  
 -  uint                devices_status;
 -  uint                devices_cnt;
 -  uint                devices_active;
 -  hc_device_param_t  *devices_param;
 +  uint    devices_status;
 +  uint    devices_cnt;
 +  uint    devices_active;
 +
 +  hc_device_param_t *devices_param;
  
 -  uint                kernel_blocks_all;
 +  uint    kernel_blocks_all;
  
    /**
     * attack specific
     */
  
 -  uint                wordlist_mode;
 -  uint                hashlist_mode;
 -  uint                hashlist_format;
 +  uint    wordlist_mode;
 +  uint    hashlist_mode;
 +  uint    hashlist_format;
 +
 +  uint    attack_mode;
 +  uint    attack_kern;
 +  uint    attack_exec;
  
 -  uint                attack_mode;
 -  uint                attack_kern;
 -  uint                attack_exec;
 +  uint    kernel_rules_cnt;
  
 -  uint                kernel_rules_cnt;
 -  kernel_rule_t      *kernel_rules_buf;
 +  kernel_rule_t *kernel_rules_buf;
  
 -  uint                combs_mode;
 -  uint                combs_cnt;
 +  uint    combs_mode;
 +  uint    combs_cnt;
  
 -  uint                bfs_cnt;
 +  uint    bfs_cnt;
  
 -  uint                css_cnt;
 -  cs_t               *css_buf;
 +  uint    css_cnt;
 +  cs_t   *css_buf;
  
 -  cs_t               *root_css_buf;
 -  cs_t               *markov_css_buf;
 +  cs_t   *root_css_buf;
 +  cs_t   *markov_css_buf;
  
 -  char               *rule_buf_l;
 -  char               *rule_buf_r;
 -  int                 rule_len_l;
 -  int                 rule_len_r;
 +  char   *rule_buf_l;
 +  char   *rule_buf_r;
 +  int     rule_len_l;
 +  int     rule_len_r;
  
+   /**
+    * opencl
+    */
+ 
+   void               *ocl;
+ 
    /**
     * hardware watchdog
     */
  
    #ifdef HAVE_HWMON
 -  HM_LIB              hm_dll_nv;
 -  HM_LIB              hm_dll_amd;
 -  hm_attrs_t          hm_device[DEVICES_MAX];
 +  HM_LIB  hm_dll_nv;
 +  HM_LIB  hm_dll_amd;
 +  hm_attrs_t hm_device[DEVICES_MAX];
    #endif
  
    /**
     * hashes
     */
  
 -  uint                digests_cnt;
 -  uint                digests_done;
 -  uint                digests_saved;
 +  uint    digests_cnt;
 +  uint    digests_done;
 +  uint    digests_saved;
  
 -  void               *digests_buf;
 -  uint               *digests_shown;
 -  uint               *digests_shown_tmp;
 +  void   *digests_buf;
 +  uint   *digests_shown;
 +  uint   *digests_shown_tmp;
  
 -  uint                salts_cnt;
 -  uint                salts_done;
 +  uint    salts_cnt;
 +  uint    salts_done;
  
 -  salt_t             *salts_buf;
 -  uint               *salts_shown;
 +  salt_t *salts_buf;
 +  uint   *salts_shown;
  
 -  void               *esalts_buf;
 +  void   *esalts_buf;
  
    /**
     * logging
@@@ -1100,10 -1105,10 +1106,10 @@@
     * crack-per-time
     */
  
 -  cpt_t  cpt_buf[CPT_BUF];
 -  int    cpt_pos;
 -  time_t cpt_start;
 -  u64 cpt_total;
 +  cpt_t   cpt_buf[CPT_BUF];
 +  int     cpt_pos;
 +  time_t  cpt_start;
 +  u64     cpt_total;
  
    /**
     * user
@@@ -1190,42 -1195,42 +1196,42 @@@
     * used for restore
     */
  
 -  u64 skip;
 -  u64 limit;
 +  u64     skip;
 +  u64     limit;
  
    restore_data_t *rd;
  
 -  u64 checkpoint_cur_words;  // used for the "stop at next checkpoint" feature
 +  u64     checkpoint_cur_words;     // used for the "stop at next checkpoint" feature
  
    /**
     * status, timer
     */
  
 -  time_t     runtime_start;
 -  time_t     runtime_stop;
 +  time_t  runtime_start;
 +  time_t  runtime_stop;
  
 -  time_t     proc_start;
 -  time_t     proc_stop;
 +  time_t  proc_start;
 +  time_t  proc_stop;
  
 -  u64   words_cnt;
 -  u64   words_cur;
 -  u64   words_base;
 +  u64     words_cnt;
 +  u64     words_cur;
 +  u64     words_base;
  
 -  u64  *words_progress_done;      // progress number of words done     per salt
 -  u64  *words_progress_rejected;  // progress number of words rejected per salt
 -  u64  *words_progress_restored;  // progress number of words restored per salt
 +  u64    *words_progress_done;      // progress number of words done     per salt
 +  u64    *words_progress_rejected;  // progress number of words rejected per salt
 +  u64    *words_progress_restored;  // progress number of words restored per salt
  
    hc_timer_t timer_running;         // timer on current dict
    hc_timer_t timer_paused;          // timer on current dict
  
 -  float      ms_paused;             // timer on current dict
 +  float   ms_paused;                // timer on current dict
  
    /**
      * hash_info and username
      */
  
    hashinfo_t **hash_info;
 -  uint         username;
 +  uint    username;
  
    int (*sort_by_digest) (const void *, const void *);
  

diff --combined src/Makefile
index 5da40f3,b33adc5..b92dee9
--- 1/src/Makefile
--- 2/src/Makefile
+++ b/src/Makefile
@@@ -125,7 -125,7 +125,7 @@@ VERSION_SUM              := $(shell gi
  CFLAGS                   := -O2 -pipe -W -Wall -std=c99 -Iinclude/
  
  ifeq ($(DEBUG),1)
 -CFLAGS                   += -g -ggdb -fsanitize=address -fno-omit-frame-pointer
 +CFLAGS                   += -DDEBUG -g -ggdb -fsanitize=address -fno-omit-frame-pointer
  endif
  
  ##
@@@ -142,7 -142,7 +142,7 @@@ BINARY_NATIVE            := $(PROG_NAME
  CFLAGS_NATIVE            := -D_POSIX -DOSX
  CFLAGS_NATIVE            += $(CFLAGS)
  
- LFLAGS_NATIVE            := -framework OpenCL -lpthread
+ LFLAGS_NATIVE            := -lpthread
  
  FOUND_ADL                := 0
  FOUND_NVML               := 0
@@@ -152,10 -152,9 +152,9 @@@ ifeq ($(UNAME),Linux
  CFLAGS_NATIVE            := -D_POSIX -DLINUX
  CFLAGS_NATIVE            += -s $(CFLAGS)
  
- LFLAGS_NATIVE            := -lOpenCL -lpthread
+ LFLAGS_NATIVE            := -lpthread -ldl
  
  ifneq (,$(filter 1,$(FOUND_ADL) $(FOUND_NVML)))
- LFLAGS_NATIVE            += -ldl
  CFLAGS_NATIVE            += -DHAVE_HWMON
  ifeq ($(FOUND_ADL),1)
  CFLAGS_NATIVE            += -DHAVE_ADL -I$(ADL)/include/
@@@ -206,7 -205,7 +205,7 @@@ endi
  CFLAGS_CROSS_32          := -m32
  CFLAGS_CROSS_64          := -m64
  
- LFLAGS_CROSS_LINUX       := -lpthread -lOpenCL -ldl
+ LFLAGS_CROSS_LINUX       := -lpthread -ldl
  LFLAGS_CROSS_WIN         := -lpsapi
  
  ##
@@@ -347,8 -346,8 +346,8 @@@ oclHashcat32.bin: src/oclHashcat.c $(LI
  oclHashcat64.bin: src/oclHashcat.c $(LINUX_64_OBJS)
        $(CC_LINUX_64) $(CFLAGS_CROSS_LINUX) $(CFLAGS_CROSS_64)    -o $@ $^ $(LFLAGS_CROSS_LINUX) -DCOMPTIME=$(COMPTIME) -DVERSION_TAG=\"$(VERSION_TAG)\" -DVERSION_SUM=\"$(VERSION_SUM)\" -DINSTALL_FOLDER=\"$(INSTALL_FOLDER)\" -DSHARED_FOLDER=\"$(SHARED_FOLDER)\" -DDOCUMENT_FOLDER=\"$(DOCUMENT_FOLDER)\"
  
- oclHashcat32.exe: src/oclHashcat.c $(WIN_32_OBJS) lib/libOpenCL.a
+ oclHashcat32.exe: src/oclHashcat.c $(WIN_32_OBJS)
        $(CC_WIN_32)   $(CFLAGS_CROSS_WIN)   $(CFLAGS_CROSS_32)    -o $@ $^ $(LFLAGS_CROSS_WIN)   -DCOMPTIME=$(COMPTIME) -DVERSION_TAG=\"$(VERSION_TAG)\" -DVERSION_SUM=\"$(VERSION_SUM)\" -static-libgcc
  
- oclHashcat64.exe: src/oclHashcat.c $(WIN_64_OBJS) lib/libOpenCL64.a
+ oclHashcat64.exe: src/oclHashcat.c $(WIN_64_OBJS)
        $(CC_WIN_64)   $(CFLAGS_CROSS_WIN)   $(CFLAGS_CROSS_64)    -o $@ $^ $(LFLAGS_CROSS_WIN)   -DCOMPTIME=$(COMPTIME) -DVERSION_TAG=\"$(VERSION_TAG)\" -DVERSION_SUM=\"$(VERSION_SUM)\" -static-libgcc

diff --combined src/oclHashcat.c
index 91d74e2,da3732c..5ba2c99
--- 1/src/oclHashcat.c
--- 2/src/oclHashcat.c
+++ b/src/oclHashcat.c
@@@ -1,5 -1,7 +1,7 @@@
  /**
-  * Author......: Jens Steube <jens.steube@gmail.com>
+  * Authors.....: Jens Steube <jens.steube@gmail.com>
+  *               Gabriele Gristina <matrix@hashcat.net>
+  *
   * License.....: MIT
   */
  
@@@ -12,14 -14,14 +14,14 @@@
  #include <rp_kernel_on_cpu.h>
  #include <getopt.h>
  
- const char *PROGNAME          = "oclHashcat";
- const uint  VERSION_BIN       = 210;
- const uint  RESTORE_MIN       = 210;
+ const char *PROGNAME            = "oclHashcat";
+ const uint  VERSION_BIN         = 210;
+ const uint  RESTORE_MIN         = 210;
  
- #define INCR_RULES            10000
- #define INCR_SALTS            100000
- #define INCR_MASKS            1000
- #define INCR_POT              1000
+ #define INCR_RULES              10000
+ #define INCR_SALTS              100000
+ #define INCR_MASKS              1000
+ #define INCR_POT                1000
  
  #define USAGE                   0
  #define VERSION                 0
@@@ -160,10 -162,10 +162,10 @@@ static uint default_benchmark_algorithm
    5000,
    10100,
    6000,
 -  6100,
 -  6900,
 -  11700,
 -  11800,
 +  6100, // broken in osx
 +  6900, // broken in osx
 +  11700, // broken in osx
 +  11800, // broken in osx
    400,
    8900,
    11900,
@@@ -172,55 -174,55 +174,55 @@@
    12100,
    23,
    2500,
 -  5300,
 -  5400,
 -  5500,
 -  5600,
 +  5300, // broken in osx
 +  5400, // broken in osx
 +  5500, // broken in osx
 +  5600, // broken in osx
    7300,
 -  7500,
 +  7500, // broken in osx
    8300,
 -  11100,
 +  11100, // broken in osx
    11200,
 -  11400,
 +  11400, // broken in osx
    121,
 -  2611,
 -  2711,
 -  2811,
 -  8400,
 +  2611, // broken in osx
 +  2711, // broken in osx
 +  2811, // broken in osx
 +  8400, // broken in osx
    11,
 -  2612,
 +  2612, // broken in osx
    7900,
    21,
    11000,
    124,
    10000,
 -  3711,
 -  7600,
 +  3711, // broken in osx
 +  7600, // broken in osx
    12,
    131,
    132,
    1731,
    200,
    300,
 -  3100,
 +  3100, // broken in osx
    112,
    12300,
 -  8000,
 +  8000, // broken in osx
    141,
    1441,
    1600,
 -  12600,
 +  12600, // broken in osx
    1421,
    101,
    111,
    1711,
 -  3000,
 +  3000, // broken in osx
    1000,
    1100,
    2100,
    12800,
 -  1500,
 -  12400,
 +  1500, // broken in osx
 +  12400, // broken in osx
    500,
    3200,
    7400,
@@@ -241,45 -243,45 +243,45 @@@
    501,
    5800,
    8100,
 -  8500,
 +  8500, // broken in osx
    7200,
    9900,
    7700,
    7800,
    10300,
 -  8600,
 -  8700,
 -  9100,
 +  8600, // broken in osx
 +  8700, // broken in osx
 +  9100, // broken in osx
    133,
 -  11600,
 -  12500,
 +  11600, // broken in osx
 +  12500, // broken in osx
    13000,
    6211,
    6221,
 -  6231,
 +  6231, // broken in osx
    6241,
 -  8800,
 +  8800, // broken in osx
    12900,
    12200,
 -  9700,
 -  9710,
 -  9800,
 -  9810,
 -  9400,
 -  9500,
 -  9600,
 -  10400,
 -  10410,
 +  9700, // broken in osx
 +  9710, // broken in osx
 +  9800, // broken in osx
 +  9810, // broken in osx
 +  9400, // broken in osx
 +  9500, // broken in osx
 +  9600, // broken in osx
 +  10400, // broken in osx
 +  10410, // broken in osx
    10500,
    10600,
 -  10700,
 +  10700, // broken in osx
    9000,
    5200,
 -  6800,
 -  6600,
 +  6800, // broken in osx
 +  6600, // broken in osx
    8200,
 -  11300,
 -  12700
 +  11300, // broken in osx
 +  12700 // broken in osx
  };
  
  /**
@@@ -1828,7 -1830,7 +1830,7 @@@ static void clear_prompt (
  
  static void gidd_to_pw_t (hc_device_param_t *device_param, const u64 gidd, pw_t *pw)
  {
-   hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_pws_buf, CL_TRUE, gidd * sizeof (pw_t), sizeof (pw_t), pw, 0, NULL, NULL);
+   hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, gidd * sizeof (pw_t), sizeof (pw_t), pw, 0, NULL, NULL);
  }
  
  static void check_hash (hc_device_param_t *device_param, const uint salt_pos, const uint digest_pos)
@@@ -1858,7 -1860,7 +1860,7 @@@
  
    plain_t plain;
  
-   hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_plain_bufs, CL_TRUE, idx * sizeof (plain_t), sizeof (plain_t), &plain, 0, NULL, NULL);
+   hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_plain_bufs, CL_TRUE, idx * sizeof (plain_t), sizeof (plain_t), &plain, 0, NULL, NULL);
  
    uint gidvid = plain.gidvid;
    uint il_pos = plain.il_pos;
@@@ -1881,7 -1883,7 +1883,7 @@@
  
      for (int i = 0, j = gidm; i < 16; i++, j++)
      {
 -      plain_buf[i] = pw.hi1[0][j];
 +      plain_buf[i] = pw.h.hi1[0][j];
      }
  
      plain_len = pw.pw_len;
@@@ -1930,7 -1932,7 +1932,7 @@@
  
      for (int i = 0, j = gidm; i < 16; i++, j++)
      {
 -      plain_buf[i] = pw.hi1[0][j];
 +      plain_buf[i] = pw.h.hi1[0][j];
      }
  
      plain_len = pw.pw_len;
@@@ -1991,7 -1993,7 +1993,7 @@@
  
      for (int i = 0, j = gidm; i < 16; i++, j++)
      {
 -      plain_buf[i] = pw.hi1[0][j];
 +      plain_buf[i] = pw.h.hi1[0][j];
      }
  
      plain_len = pw.pw_len;
@@@ -2025,7 -2027,7 +2027,7 @@@
  
      for (int i = 0, j = gidm; i < 16; i++, j++)
      {
 -      plain_buf[i] = pw.hi1[0][j];
 +      plain_buf[i] = pw.h.hi1[0][j];
      }
  
      plain_len = pw.pw_len;
@@@ -2175,7 -2177,7 +2177,7 @@@ static void check_cracked (hc_device_pa
  
    int found = 0;
  
-   hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
+   hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
  
    for (uint i = 0; i < KERNEL_THREADS; i++) if (device_param->result[i] == 1) found = 1;
  
@@@ -2185,7 -2187,7 +2187,7 @@@
  
      log_info_nn ("");
  
-     hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (uint), salt_buf->digests_cnt * sizeof (uint), &data.digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
+     hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (uint), salt_buf->digests_cnt * sizeof (uint), &data.digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
  
      uint cpt_cracked = 0;
  
@@@ -2240,12 -2242,12 +2242,12 @@@
  
        memset (data.digests_shown_tmp, 0, salt_buf->digests_cnt * sizeof (uint));
  
-       hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (uint), salt_buf->digests_cnt * sizeof (uint), &data.digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (uint), salt_buf->digests_cnt * sizeof (uint), &data.digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
      }
  
      memset (device_param->result, 0, device_param->size_results);
  
-     hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
+     hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
    }
  }
  
@@@ -2400,36 -2402,36 +2402,36 @@@ static void run_kernel (const uint kern
      case KERN_RUN_3:    kernel = device_param->kernel3;     break;
    }
  
-   hc_clSetKernelArg (kernel, 21, sizeof (cl_uint), device_param->kernel_params[21]);
-   hc_clSetKernelArg (kernel, 22, sizeof (cl_uint), device_param->kernel_params[22]);
-   hc_clSetKernelArg (kernel, 23, sizeof (cl_uint), device_param->kernel_params[23]);
-   hc_clSetKernelArg (kernel, 24, sizeof (cl_uint), device_param->kernel_params[24]);
-   hc_clSetKernelArg (kernel, 25, sizeof (cl_uint), device_param->kernel_params[25]);
-   hc_clSetKernelArg (kernel, 26, sizeof (cl_uint), device_param->kernel_params[26]);
-   hc_clSetKernelArg (kernel, 27, sizeof (cl_uint), device_param->kernel_params[27]);
-   hc_clSetKernelArg (kernel, 28, sizeof (cl_uint), device_param->kernel_params[28]);
-   hc_clSetKernelArg (kernel, 29, sizeof (cl_uint), device_param->kernel_params[29]);
-   hc_clSetKernelArg (kernel, 30, sizeof (cl_uint), device_param->kernel_params[30]);
-   hc_clSetKernelArg (kernel, 31, sizeof (cl_uint), device_param->kernel_params[31]);
+   hc_clSetKernelArg (data.ocl, kernel, 21, sizeof (cl_uint), device_param->kernel_params[21]);
+   hc_clSetKernelArg (data.ocl, kernel, 22, sizeof (cl_uint), device_param->kernel_params[22]);
+   hc_clSetKernelArg (data.ocl, kernel, 23, sizeof (cl_uint), device_param->kernel_params[23]);
+   hc_clSetKernelArg (data.ocl, kernel, 24, sizeof (cl_uint), device_param->kernel_params[24]);
+   hc_clSetKernelArg (data.ocl, kernel, 25, sizeof (cl_uint), device_param->kernel_params[25]);
+   hc_clSetKernelArg (data.ocl, kernel, 26, sizeof (cl_uint), device_param->kernel_params[26]);
+   hc_clSetKernelArg (data.ocl, kernel, 27, sizeof (cl_uint), device_param->kernel_params[27]);
+   hc_clSetKernelArg (data.ocl, kernel, 28, sizeof (cl_uint), device_param->kernel_params[28]);
+   hc_clSetKernelArg (data.ocl, kernel, 29, sizeof (cl_uint), device_param->kernel_params[29]);
+   hc_clSetKernelArg (data.ocl, kernel, 30, sizeof (cl_uint), device_param->kernel_params[30]);
+   hc_clSetKernelArg (data.ocl, kernel, 31, sizeof (cl_uint), device_param->kernel_params[31]);
  
    if ((data.opts_type & OPTS_TYPE_PT_BITSLICE) && (data.attack_mode == ATTACK_MODE_BF))
    {
      const size_t global_work_size[3] = { num_elements,        32, 1 };
      const size_t local_work_size[3]  = { kernel_threads / 32, 32, 1 };
  
-     hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+     hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
    }
    else
    {
      const size_t global_work_size[3] = { num_elements,   1, 1 };
      const size_t local_work_size[3]  = { kernel_threads, 1, 1 };
  
-     hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+     hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
    }
  
-   hc_clFlush (device_param->command_queue);
+   hc_clFlush (data.ocl, device_param->command_queue);
  
-   hc_clFinish (device_param->command_queue);
+   hc_clFinish (data.ocl, device_param->command_queue);
  }
  
  static void run_kernel_mp (const uint kern_run, hc_device_param_t *device_param, const uint num)
@@@ -2461,38 -2463,38 +2463,38 @@@
  
    switch (kern_run)
    {
-     case KERN_RUN_MP:   hc_clSetKernelArg (kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp[3]);
-                         hc_clSetKernelArg (kernel, 4, sizeof (cl_uint),  device_param->kernel_params_mp[4]);
-                         hc_clSetKernelArg (kernel, 5, sizeof (cl_uint),  device_param->kernel_params_mp[5]);
-                         hc_clSetKernelArg (kernel, 6, sizeof (cl_uint),  device_param->kernel_params_mp[6]);
-                         hc_clSetKernelArg (kernel, 7, sizeof (cl_uint),  device_param->kernel_params_mp[7]);
-                         hc_clSetKernelArg (kernel, 8, sizeof (cl_uint),  device_param->kernel_params_mp[8]);
+     case KERN_RUN_MP:   hc_clSetKernelArg (data.ocl, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp[3]);
+                         hc_clSetKernelArg (data.ocl, kernel, 4, sizeof (cl_uint),  device_param->kernel_params_mp[4]);
+                         hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint),  device_param->kernel_params_mp[5]);
+                         hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint),  device_param->kernel_params_mp[6]);
+                         hc_clSetKernelArg (data.ocl, kernel, 7, sizeof (cl_uint),  device_param->kernel_params_mp[7]);
+                         hc_clSetKernelArg (data.ocl, kernel, 8, sizeof (cl_uint),  device_param->kernel_params_mp[8]);
                          break;
-     case KERN_RUN_MP_R: hc_clSetKernelArg (kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_r[3]);
-                         hc_clSetKernelArg (kernel, 4, sizeof (cl_uint),  device_param->kernel_params_mp_r[4]);
-                         hc_clSetKernelArg (kernel, 5, sizeof (cl_uint),  device_param->kernel_params_mp_r[5]);
-                         hc_clSetKernelArg (kernel, 6, sizeof (cl_uint),  device_param->kernel_params_mp_r[6]);
-                         hc_clSetKernelArg (kernel, 7, sizeof (cl_uint),  device_param->kernel_params_mp_r[7]);
-                         hc_clSetKernelArg (kernel, 8, sizeof (cl_uint),  device_param->kernel_params_mp_r[8]);
+     case KERN_RUN_MP_R: hc_clSetKernelArg (data.ocl, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_r[3]);
+                         hc_clSetKernelArg (data.ocl, kernel, 4, sizeof (cl_uint),  device_param->kernel_params_mp_r[4]);
+                         hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint),  device_param->kernel_params_mp_r[5]);
+                         hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint),  device_param->kernel_params_mp_r[6]);
+                         hc_clSetKernelArg (data.ocl, kernel, 7, sizeof (cl_uint),  device_param->kernel_params_mp_r[7]);
+                         hc_clSetKernelArg (data.ocl, kernel, 8, sizeof (cl_uint),  device_param->kernel_params_mp_r[8]);
                          break;
-     case KERN_RUN_MP_L: hc_clSetKernelArg (kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_l[3]);
-                         hc_clSetKernelArg (kernel, 4, sizeof (cl_uint),  device_param->kernel_params_mp_l[4]);
-                         hc_clSetKernelArg (kernel, 5, sizeof (cl_uint),  device_param->kernel_params_mp_l[5]);
-                         hc_clSetKernelArg (kernel, 6, sizeof (cl_uint),  device_param->kernel_params_mp_l[6]);
-                         hc_clSetKernelArg (kernel, 7, sizeof (cl_uint),  device_param->kernel_params_mp_l[7]);
-                         hc_clSetKernelArg (kernel, 8, sizeof (cl_uint),  device_param->kernel_params_mp_l[8]);
-                         hc_clSetKernelArg (kernel, 9, sizeof (cl_uint),  device_param->kernel_params_mp_l[9]);
+     case KERN_RUN_MP_L: hc_clSetKernelArg (data.ocl, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_l[3]);
+                         hc_clSetKernelArg (data.ocl, kernel, 4, sizeof (cl_uint),  device_param->kernel_params_mp_l[4]);
+                         hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint),  device_param->kernel_params_mp_l[5]);
+                         hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint),  device_param->kernel_params_mp_l[6]);
+                         hc_clSetKernelArg (data.ocl, kernel, 7, sizeof (cl_uint),  device_param->kernel_params_mp_l[7]);
+                         hc_clSetKernelArg (data.ocl, kernel, 8, sizeof (cl_uint),  device_param->kernel_params_mp_l[8]);
+                         hc_clSetKernelArg (data.ocl, kernel, 9, sizeof (cl_uint),  device_param->kernel_params_mp_l[9]);
                          break;
    }
  
    const size_t global_work_size[3] = { num_elements, 1, 1 };
    const size_t local_work_size[3]  = { kernel_threads,  1, 1 };
  
-   hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+   hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
  
-   hc_clFlush (device_param->command_queue);
+   hc_clFlush (data.ocl, device_param->command_queue);
  
-   hc_clFinish (device_param->command_queue);
+   hc_clFinish (data.ocl, device_param->command_queue);
  }
  
  static void run_kernel_tb (hc_device_param_t *device_param, const uint num)
@@@ -2508,11 -2510,11 +2510,11 @@@
    const size_t global_work_size[3] = { num_elements, 1, 1 };
    const size_t local_work_size[3]  = { kernel_threads,  1, 1 };
  
-   hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+   hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
  
-   hc_clFlush (device_param->command_queue);
+   hc_clFlush (data.ocl, device_param->command_queue);
  
-   hc_clFinish (device_param->command_queue);
+   hc_clFinish (data.ocl, device_param->command_queue);
  }
  
  static void run_kernel_tm (hc_device_param_t *device_param)
@@@ -2526,11 -2528,11 +2528,11 @@@
    const size_t global_work_size[3] = { num_elements, 1, 1 };
    const size_t local_work_size[3]  = { kernel_threads,  1, 1 };
  
-   hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+   hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
  
-   hc_clFlush (device_param->command_queue);
+   hc_clFlush (data.ocl, device_param->command_queue);
  
-   hc_clFinish (device_param->command_queue);
+   hc_clFinish (data.ocl, device_param->command_queue);
  }
  
  static void run_kernel_amp (hc_device_param_t *device_param, const uint num)
@@@ -2549,17 -2551,17 +2551,17 @@@
  
    cl_kernel kernel = device_param->kernel_amp;
  
-   hc_clSetKernelArg (kernel, 5, sizeof (cl_uint), device_param->kernel_params_amp[5]);
-   hc_clSetKernelArg (kernel, 6, sizeof (cl_uint), device_param->kernel_params_amp[6]);
+   hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_amp[5]);
+   hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_amp[6]);
  
    const size_t global_work_size[3] = { num_elements, 1, 1 };
    const size_t local_work_size[3]  = { kernel_threads,  1, 1 };
  
-   hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+   hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
  
-   hc_clFlush (device_param->command_queue);
+   hc_clFlush (data.ocl, device_param->command_queue);
  
-   hc_clFinish (device_param->command_queue);
+   hc_clFinish (data.ocl, device_param->command_queue);
  }
  
  static void run_kernel_bzero (hc_device_param_t *device_param, cl_mem buf, const uint size)
@@@ -2570,7 -2572,7 +2572,7 @@@
  
      const cl_uchar zero = 0;
  
-     hc_clEnqueueFillBuffer (device_param->command_queue, buf, &zero, sizeof (cl_uchar), 0, size, 0, NULL, NULL);
+     hc_clEnqueueFillBuffer (data.ocl, device_param->command_queue, buf, &zero, sizeof (cl_uchar), 0, size, 0, NULL, NULL);
    }
    else
    {
@@@ -2589,7 -2591,7 +2591,7 @@@
  
        const int fillsz = MIN (FILLSZ, left);
  
-       hc_clEnqueueWriteBuffer (device_param->command_queue, buf, CL_TRUE, i, fillsz, tmp, 0, NULL, NULL);
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, buf, CL_TRUE, i, fillsz, tmp, 0, NULL, NULL);
      }
  
      myfree (tmp);
@@@ -2614,11 -2616,11 +2616,11 @@@ static void run_copy (hc_device_param_
  {
    if (data.attack_kern == ATTACK_KERN_STRAIGHT)
    {
-     hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
+     hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
    }
    else if (data.attack_kern == ATTACK_KERN_COMBI)
    {
-     hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
+     hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
    }
    else if (data.attack_kern == ATTACK_KERN_BF)
    {
@@@ -2856,23 -2858,23 +2858,23 @@@ static void run_cracker (hc_device_para
  
        if (data.attack_mode == ATTACK_MODE_STRAIGHT)
        {
-         hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_rules, device_param->d_rules_c, innerloop_pos * sizeof (kernel_rule_t), 0, innerloop_left * sizeof (kernel_rule_t), 0, NULL, NULL);
+         hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_rules, device_param->d_rules_c, innerloop_pos * sizeof (kernel_rule_t), 0, innerloop_left * sizeof (kernel_rule_t), 0, NULL, NULL);
        }
        else if (data.attack_mode == ATTACK_MODE_COMBI)
        {
-         hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (comb_t), device_param->combs_buf, 0, NULL, NULL);
+         hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (comb_t), device_param->combs_buf, 0, NULL, NULL);
        }
        else if (data.attack_mode == ATTACK_MODE_BF)
        {
-         hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_bfs, device_param->d_bfs_c, 0, 0, innerloop_left * sizeof (bf_t), 0, NULL, NULL);
+         hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_bfs, device_param->d_bfs_c, 0, 0, innerloop_left * sizeof (bf_t), 0, NULL, NULL);
        }
        else if (data.attack_mode == ATTACK_MODE_HYBRID1)
        {
-         hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
+         hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
        }
        else if (data.attack_mode == ATTACK_MODE_HYBRID2)
        {
-         hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
+         hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
        }
  
        if (data.attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
@@@ -2887,7 -2889,7 +2889,7 @@@
  
              run_kernel_tm (device_param);
  
-             hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_tm_c, device_param->d_bfs_c, 0, 0, size_tm, 0, NULL, NULL);
+             hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_tm_c, device_param->d_bfs_c, 0, 0, size_tm, 0, NULL, NULL);
            }
          }
  
@@@ -2939,11 -2941,11 +2941,11 @@@
          {
            run_kernel (KERN_RUN_23, device_param, pws_cnt);
  
-           hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
+           hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
  
            // do something with data
  
-           hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
+           hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
          }
  
          run_kernel (KERN_RUN_3, device_param, pws_cnt);
@@@ -3339,7 -3341,7 +3341,7 @@@ static u64 count_words (wl_data_t *wl_d
  
  static void pw_transpose_to_hi1 (const pw_t *p1, pw_t *p2)
  {
 -  memcpy (p2->hi1, p1->hi1, 64 * sizeof (uint));
 +  memcpy (p2->h.hi1, p1->h.hi1, 64 * sizeof (uint));
  }
  
  static uint pw_add_to_hc1 (hc_device_param_t *device_param, const u8 *pw_buf, const uint pw_len)
@@@ -3350,7 -3352,7 +3352,7 @@@
  
    uint cache_cnt = pw_cache->cnt;
  
 -  u8 *pw_hc1 = pw_cache->pw_buf.hc1[cache_cnt];
 +  u8 *pw_hc1 = pw_cache->pw_buf.h.hc1[cache_cnt];
  
    memcpy (pw_hc1, pw_buf, pw_len);
  
@@@ -5145,6 -5147,7 +5147,7 @@@ int main (int argc, char **argv
    #ifndef OSX
    char *cpu_affinity      = NULL;
    #endif
+   OCL_PTR *ocl            = NULL;
    char *opencl_devices    = NULL;
    char *opencl_platforms  = NULL;
    char *opencl_device_types = NULL;
@@@ -6509,6 -6512,19 +6512,19 @@@
    logfile_top_string (session);
    logfile_top_string (truecrypt_keyfiles);
  
+   /**
+    * Init OpenCL library loader
+    */
+ 
+   if (keyspace == 0)
+   {
+     ocl = (OCL_PTR *) mymalloc (sizeof (OCL_PTR));
+ 
+     ocl_init(ocl);
+ 
+     data.ocl = ocl;
+   }
+ 
    /**
     * OpenCL platform selection
     */
@@@ -11383,9 -11399,7 +11399,9 @@@
                         kernel_accel = 16;
                         break;
            case  6211:  kernel_loops = ROUNDS_TRUECRYPT_2K;
 +                       #ifndef OSX
                         kernel_accel = 64;
 +                       #endif
                         break;
            case  6212:  kernel_loops = ROUNDS_TRUECRYPT_2K;
                         kernel_accel = 32;
@@@ -11412,9 -11426,7 +11428,9 @@@
                         kernel_accel = 8;
                         break;
            case  6241:  kernel_loops = ROUNDS_TRUECRYPT_1K;
 +                       #ifndef OSX
                         kernel_accel = 128;
 +                       #endif
                         break;
            case  6242:  kernel_loops = ROUNDS_TRUECRYPT_1K;
                         kernel_accel = 64;
@@@ -11444,9 -11456,7 +11460,9 @@@
                         kernel_accel = 8;
                         break;
            case  7200:  kernel_loops = ROUNDS_GRUB;
 +                       #ifndef OSX
                         kernel_accel = 16;
 +                       #endif
                         break;
            case  7400:  kernel_loops = ROUNDS_SHA256CRYPT;
                         kernel_accel = 8;
@@@ -12491,13 -12501,16 +12507,16 @@@
  
      cl_uint platform_devices_cnt;
  
-     hc_clGetPlatformIDs (CL_PLATFORMS_MAX, platforms, &platforms_cnt);
- 
-     if (platforms_cnt == 0)
+     if (keyspace == 0)
      {
-       log_error ("ERROR: No OpenCL compatible platform found");
+       hc_clGetPlatformIDs (data.ocl, CL_PLATFORMS_MAX, platforms, &platforms_cnt);
  
-       return (-1);
+       if (platforms_cnt == 0)
+       {
+         log_error ("ERROR: No OpenCL compatible platform found");
+ 
+         return (-1);
+       }
      }
  
      /**
@@@ -12510,7 -12523,7 +12529,7 @@@
  
        char platform_vendor[INFOSZ] = { 0 };
  
-       hc_clGetPlatformInfo (platform, CL_PLATFORM_VENDOR, sizeof (platform_vendor), platform_vendor, NULL);
+       hc_clGetPlatformInfo (data.ocl, platform, CL_PLATFORM_VENDOR, sizeof (platform_vendor), platform_vendor, NULL);
  
        #ifdef HAVE_HWMON
        #if defined(HAVE_NVML) || defined(HAVE_NVAPI)
@@@ -12544,7 -12557,7 +12563,7 @@@
  
        cl_platform_id platform = platforms[platform_id];
  
-       hc_clGetDeviceIDs (platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt);
+       hc_clGetDeviceIDs (data.ocl, platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt);
  
        for (uint platform_devices_id = 0; platform_devices_id < platform_devices_cnt; platform_devices_id++)
        {
@@@ -12562,7 -12575,7 +12581,7 @@@
  
          cl_device_type device_type;
  
-         hc_clGetDeviceInfo (device_param->device, CL_DEVICE_TYPE, sizeof (device_type), &device_type, NULL);
+         hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_TYPE, sizeof (device_type), &device_type, NULL);
  
          device_type &= ~CL_DEVICE_TYPE_DEFAULT;
  
@@@ -12572,7 -12585,7 +12591,7 @@@
  
          cl_uint vendor_id = 0;
  
-         hc_clGetDeviceInfo (device_param->device, CL_DEVICE_VENDOR_ID, sizeof (vendor_id), &vendor_id, NULL);
+         hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VENDOR_ID, sizeof (vendor_id), &vendor_id, NULL);
  
          device_param->vendor_id = vendor_id;
  
@@@ -12580,7 -12593,7 +12599,7 @@@
  
          char *device_name = (char *) mymalloc (INFOSZ);
  
-         hc_clGetDeviceInfo (device_param->device, CL_DEVICE_NAME, INFOSZ, device_name, NULL);
+         hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NAME, INFOSZ, device_name, NULL);
  
          device_param->device_name = device_name;
  
@@@ -12588,7 -12601,7 +12607,7 @@@
  
          char *device_version = (char *) mymalloc (INFOSZ);
  
-         hc_clGetDeviceInfo (device_param->device, CL_DEVICE_VERSION, INFOSZ, device_version, NULL);
+         hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VERSION, INFOSZ, device_version, NULL);
  
          device_param->device_version = device_version;
  
@@@ -12608,7 -12621,7 +12627,7 @@@
  
          if (opencl_vector_width == OPENCL_VECTOR_WIDTH)
          {
-           hc_clGetDeviceInfo (device_param->device, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, sizeof (vector_width), &vector_width, NULL);
+           hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, sizeof (vector_width), &vector_width, NULL);
  
            if ((vendor_id == VENDOR_ID_NV) && (strstr (device_name, " Ti") || strstr (device_name, " TI")))
            {
@@@ -12635,7 -12648,7 +12654,7 @@@
  
          cl_uint device_processors;
  
-         hc_clGetDeviceInfo (device_param->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof (device_processors), &device_processors, NULL);
+         hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof (device_processors), &device_processors, NULL);
  
          device_param->device_processors = device_processors;
  
@@@ -12643,7 -12656,7 +12662,7 @@@
  
          cl_ulong device_maxmem_alloc;
  
-         hc_clGetDeviceInfo (device_param->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof (device_maxmem_alloc), &device_maxmem_alloc, NULL);
+         hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof (device_maxmem_alloc), &device_maxmem_alloc, NULL);
  
          device_param->device_maxmem_alloc = device_maxmem_alloc;
  
@@@ -12651,7 -12664,7 +12670,7 @@@
  
          cl_ulong device_global_mem;
  
-         hc_clGetDeviceInfo (device_param->device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof (device_global_mem), &device_global_mem, NULL);
+         hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof (device_global_mem), &device_global_mem, NULL);
  
          device_param->device_global_mem = device_global_mem;
  
@@@ -12659,7 -12672,7 +12678,7 @@@
  
          cl_uint device_maxclock_frequency;
  
-         hc_clGetDeviceInfo (device_param->device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof (device_maxclock_frequency), &device_maxclock_frequency, NULL);
+         hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof (device_maxclock_frequency), &device_maxclock_frequency, NULL);
  
          device_param->device_maxclock_frequency = device_maxclock_frequency;
  
@@@ -12674,7 -12687,7 +12693,7 @@@
  
          char *driver_version = (char *) mymalloc (INFOSZ);
  
-         hc_clGetDeviceInfo (device_param->device, CL_DRIVER_VERSION, INFOSZ, driver_version, NULL);
+         hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DRIVER_VERSION, INFOSZ, driver_version, NULL);
  
          device_param->driver_version = driver_version;
  
@@@ -12713,7 -12726,7 +12732,7 @@@
  
              #define CL_DEVICE_WAVEFRONT_WIDTH_AMD               0x4043
  
-             hc_clGetDeviceInfo (device_param->device, CL_DEVICE_WAVEFRONT_WIDTH_AMD, sizeof (device_processor_cores), &device_processor_cores, NULL);
+             hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_WAVEFRONT_WIDTH_AMD, sizeof (device_processor_cores), &device_processor_cores, NULL);
  
              device_param->device_processor_cores = device_processor_cores;
            }
@@@ -12723,7 -12736,7 +12742,7 @@@
  
              #define CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV            0x4005
  
-             hc_clGetDeviceInfo (device_param->device, CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, sizeof (kernel_exec_timeout), &kernel_exec_timeout, NULL);
+             hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, sizeof (kernel_exec_timeout), &kernel_exec_timeout, NULL);
  
              device_param->kernel_exec_timeout = kernel_exec_timeout;
  
@@@ -12731,7 -12744,7 +12750,7 @@@
  
              #define CL_DEVICE_WARP_SIZE_NV                      0x4003
  
-             hc_clGetDeviceInfo (device_param->device, CL_DEVICE_WARP_SIZE_NV, sizeof (device_processor_cores), &device_processor_cores, NULL);
+             hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_WARP_SIZE_NV, sizeof (device_processor_cores), &device_processor_cores, NULL);
  
              device_param->device_processor_cores = device_processor_cores;
  
@@@ -12741,8 -12754,8 +12760,8 @@@
              #define CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV       0x4000
              #define CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV       0x4001
  
-             hc_clGetDeviceInfo (device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, sizeof (sm_minor), &sm_minor, NULL);
-             hc_clGetDeviceInfo (device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof (sm_major), &sm_major, NULL);
+             hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, sizeof (sm_minor), &sm_minor, NULL);
+             hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof (sm_major), &sm_major, NULL);
  
              device_param->sm_minor = sm_minor;
              device_param->sm_major = sm_major;
@@@ -12864,7 -12877,7 +12883,7 @@@
        }
      }
  
-     if (devices_active == 0)
+     if (keyspace == 0 && devices_active == 0)
      {
        log_error ("ERROR: No devices found/left");
  
@@@ -13196,37 -13209,6 +13215,37 @@@
      #endif // HAVE_ADK
      #endif // HAVE_HWMON
  
 +    #ifdef OSX
 +    /*
 +     * List of OSX kernel to fix
 +     */
 +    if ((hash_mode == 6100)  || (hash_mode == 6900)  || (hash_mode == 11700) || (hash_mode == 11800) || (hash_mode == 5300)  || \
 +        (hash_mode == 5400)  || (hash_mode == 5500)  || (hash_mode == 5600)  || (hash_mode == 7500)  || (hash_mode == 11100) || \
 +        (hash_mode == 11400) || (hash_mode == 2611)  || (hash_mode == 2711)  || (hash_mode == 2811)  || (hash_mode == 8400)  || \
 +        (hash_mode == 2612)  || (hash_mode == 3711)  || (hash_mode == 7600)  || (hash_mode == 3100)  || (hash_mode == 8000)  || \
 +        (hash_mode == 12600) || (hash_mode == 3000)  || (hash_mode == 1500)  || (hash_mode == 12400) || (hash_mode == 8500)  || \
 +        (hash_mode == 8600)  || (hash_mode == 8700)  || (hash_mode == 9100)  || (hash_mode == 11600) || (hash_mode == 12500) || \
 +        (hash_mode == 6231)  || (hash_mode == 8800)  || (hash_mode == 9700)  || (hash_mode == 9710)  || (hash_mode == 9800)  || \
 +        (hash_mode == 9810)  || (hash_mode == 9400)  || (hash_mode == 9500)  || (hash_mode == 9600)  || (hash_mode == 10400) || \
 +        (hash_mode == 10410) || (hash_mode == 10700) || (hash_mode == 6800)  || (hash_mode == 6600)  || (hash_mode == 11300) || \
 +        (hash_mode == 12700))
 +    {
 +      if (force == 0)
 +      {
 +        log_info ("");
 +        log_info ("Warning: Hash mode %d is not stable in OSX.", hash_mode);
 +        log_info ("You can use --force to override this but do not post error reports if you do so");
 +        log_info ("");
 +
 +        continue;
 +      }
 +    }
 +    #endif
 +
 +    #ifdef DEBUG
 +    if (benchmark == 1) log_info ("Hashmode: %d", data.hash_mode);
 +    #endif
 +
      uint kernel_blocks_all = 0;
  
      for (uint device_id = 0; device_id < devices_cnt; device_id++)
@@@ -13255,7 -13237,7 +13274,7 @@@
         * create context for each device
         */
  
-       device_param->context = hc_clCreateContext (NULL, 1, &device_param->device, NULL, NULL);
+       device_param->context = hc_clCreateContext (data.ocl, NULL, 1, &device_param->device, NULL, NULL);
  
        /**
         * create command-queue
@@@ -13264,7 -13246,7 +13283,7 @@@
        // not supported with NV
        // device_param->command_queue = hc_clCreateCommandQueueWithProperties (device_param->context, device_param->device, NULL);
  
-       device_param->command_queue = hc_clCreateCommandQueue (device_param->context, device_param->device, 0);
+       device_param->command_queue = hc_clCreateCommandQueue (data.ocl, device_param->context, device_param->device, 0);
  
        /**
         * create input buffers on device
@@@ -13568,17 -13550,17 +13587,17 @@@
  
              load_kernel (source_file, 1, kernel_lengths, kernel_sources);
  
-             device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+             device_param->program = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
  
-             hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+             hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
  
              size_t binary_size;
  
-             clGetProgramInfo (device_param->program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+             hc_clGetProgramInfo (data.ocl, device_param->program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
  
              u8 *binary = (u8 *) mymalloc (binary_size);
  
-             clGetProgramInfo (device_param->program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
+             hc_clGetProgramInfo (data.ocl, device_param->program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
  
              writeProgramBin (cached_file, binary, binary_size);
  
@@@ -13590,9 -13572,9 +13609,9 @@@
  
              load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
  
-             device_param->program = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
+             device_param->program = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
  
-             hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+             hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
            }
          }
          else
@@@ -13601,7 -13583,7 +13620,7 @@@
  
            load_kernel (source_file, 1, kernel_lengths, kernel_sources);
  
-           device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+           device_param->program = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
  
            if (force_jit_compilation == 1500)
            {
@@@ -13612,7 -13594,7 +13631,7 @@@
              snprintf (build_opts, sizeof (build_opts) - 1, "%s -DSCRYPT_N=%d -DSCRYPT_R=%d -DSCRYPT_P=%d -DSCRYPT_TMTO=%d", build_opts, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, 1 << data.salts_buf[0].scrypt_tmto);
            }
  
-           hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+           hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
          }
  
          local_free (kernel_lengths);
@@@ -13623,13 -13605,13 +13642,13 @@@
  
          size_t ret_val_size = 0;
  
-         clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
+         hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
  
          if (ret_val_size > 2)
          {
            char *build_log = (char *) mymalloc (ret_val_size + 1);
  
-           clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
+           hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
  
            puts (build_log);
  
@@@ -13691,17 -13673,17 +13710,17 @@@
  
            load_kernel (source_file, 1, kernel_lengths, kernel_sources);
  
-           device_param->program_mp = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+           device_param->program_mp = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
  
-           hc_clBuildProgram (device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
+           hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
  
            size_t binary_size;
  
-           clGetProgramInfo (device_param->program_mp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+           hc_clGetProgramInfo (data.ocl, device_param->program_mp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
  
            u8 *binary = (u8 *) mymalloc (binary_size);
  
-           clGetProgramInfo (device_param->program_mp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
+           hc_clGetProgramInfo (data.ocl, device_param->program_mp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
  
            writeProgramBin (cached_file, binary, binary_size);
  
@@@ -13713,9 -13695,9 +13732,9 @@@
  
            load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
  
-           device_param->program_mp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
+           device_param->program_mp = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
  
-           hc_clBuildProgram (device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
+           hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
          }
  
          local_free (kernel_lengths);
@@@ -13726,13 -13708,13 +13745,13 @@@
  
          size_t ret_val_size = 0;
  
-         clGetProgramBuildInfo (device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
+         hc_clGetProgramBuildInfo (data.ocl, device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
  
          if (ret_val_size > 2)
          {
            char *build_log = (char *) mymalloc (ret_val_size + 1);
  
-           clGetProgramBuildInfo (device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
+           hc_clGetProgramBuildInfo (data.ocl, device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
  
            puts (build_log);
  
@@@ -13798,17 -13780,17 +13817,17 @@@
  
            load_kernel (source_file, 1, kernel_lengths, kernel_sources);
  
-           device_param->program_amp = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+           device_param->program_amp = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
  
-           hc_clBuildProgram (device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
+           hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
  
            size_t binary_size;
  
-           clGetProgramInfo (device_param->program_amp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+           hc_clGetProgramInfo (data.ocl, device_param->program_amp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
  
            u8 *binary = (u8 *) mymalloc (binary_size);
  
-           clGetProgramInfo (device_param->program_amp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
+           hc_clGetProgramInfo (data.ocl, device_param->program_amp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
  
            writeProgramBin (cached_file, binary, binary_size);
  
@@@ -13820,9 -13802,9 +13839,9 @@@
  
            load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
  
-           device_param->program_amp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
+           device_param->program_amp = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
  
-           hc_clBuildProgram (device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
+           hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
          }
  
          local_free (kernel_lengths);
@@@ -13833,13 -13815,13 +13852,13 @@@
  
          size_t ret_val_size = 0;
  
-         clGetProgramBuildInfo (device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
+         hc_clGetProgramBuildInfo (data.ocl, device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
  
          if (ret_val_size > 2)
          {
            char *build_log = (char *) mymalloc (ret_val_size + 1);
  
-           clGetProgramBuildInfo (device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
+           hc_clGetProgramBuildInfo (data.ocl, device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
  
            puts (build_log);
  
@@@ -13851,36 -13833,36 +13870,36 @@@
         * global buffers
         */
  
-       device_param->d_pws_buf       = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY,   size_pws,     NULL);
-       device_param->d_pws_amp_buf   = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY,   size_pws,     NULL);
-       device_param->d_tmps          = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE,  size_tmps,    NULL);
-       device_param->d_hooks         = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE,  size_hooks,   NULL);
-       device_param->d_bitmap_s1_a   = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
-       device_param->d_bitmap_s1_b   = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
-       device_param->d_bitmap_s1_c   = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
-       device_param->d_bitmap_s1_d   = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
-       device_param->d_bitmap_s2_a   = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
-       device_param->d_bitmap_s2_b   = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
-       device_param->d_bitmap_s2_c   = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
-       device_param->d_bitmap_s2_d   = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
-       device_param->d_plain_bufs    = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE,  size_plains,  NULL);
-       device_param->d_digests_buf   = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY,   size_digests, NULL);
-       device_param->d_digests_shown = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE,  size_shown,   NULL);
-       device_param->d_salt_bufs     = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY,   size_salts,   NULL);
-       device_param->d_result        = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE,  size_results, NULL);
-       device_param->d_scryptV_buf   = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE,  size_scryptV, NULL);
- 
-       hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s1_a,    CL_TRUE, 0, bitmap_size,  bitmap_s1_a,        0, NULL, NULL);
-       hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s1_b,    CL_TRUE, 0, bitmap_size,  bitmap_s1_b,        0, NULL, NULL);
-       hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s1_c,    CL_TRUE, 0, bitmap_size,  bitmap_s1_c,        0, NULL, NULL);
-       hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s1_d,    CL_TRUE, 0, bitmap_size,  bitmap_s1_d,        0, NULL, NULL);
-       hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s2_a,    CL_TRUE, 0, bitmap_size,  bitmap_s2_a,        0, NULL, NULL);
-       hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s2_b,    CL_TRUE, 0, bitmap_size,  bitmap_s2_b,        0, NULL, NULL);
-       hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s2_c,    CL_TRUE, 0, bitmap_size,  bitmap_s2_c,        0, NULL, NULL);
-       hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s2_d,    CL_TRUE, 0, bitmap_size,  bitmap_s2_d,        0, NULL, NULL);
-       hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_digests_buf,    CL_TRUE, 0, size_digests, data.digests_buf,   0, NULL, NULL);
-       hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_digests_shown,  CL_TRUE, 0, size_shown,   data.digests_shown, 0, NULL, NULL);
-       hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_salt_bufs,      CL_TRUE, 0, size_salts,   data.salts_buf,     0, NULL, NULL);
+       device_param->d_pws_buf       = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   size_pws,     NULL);
+       device_param->d_pws_amp_buf   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   size_pws,     NULL);
+       device_param->d_tmps          = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_tmps,    NULL);
+       device_param->d_hooks         = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_hooks,   NULL);
+       device_param->d_bitmap_s1_a   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
+       device_param->d_bitmap_s1_b   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
+       device_param->d_bitmap_s1_c   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
+       device_param->d_bitmap_s1_d   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
+       device_param->d_bitmap_s2_a   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
+       device_param->d_bitmap_s2_b   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
+       device_param->d_bitmap_s2_c   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
+       device_param->d_bitmap_s2_d   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   bitmap_size,  NULL);
+       device_param->d_plain_bufs    = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_plains,  NULL);
+       device_param->d_digests_buf   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   size_digests, NULL);
+       device_param->d_digests_shown = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_shown,   NULL);
+       device_param->d_salt_bufs     = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   size_salts,   NULL);
+       device_param->d_result        = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_results, NULL);
+       device_param->d_scryptV_buf   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_scryptV, NULL);
+ 
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_a,    CL_TRUE, 0, bitmap_size,  bitmap_s1_a,        0, NULL, NULL);
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_b,    CL_TRUE, 0, bitmap_size,  bitmap_s1_b,        0, NULL, NULL);
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_c,    CL_TRUE, 0, bitmap_size,  bitmap_s1_c,        0, NULL, NULL);
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_d,    CL_TRUE, 0, bitmap_size,  bitmap_s1_d,        0, NULL, NULL);
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_a,    CL_TRUE, 0, bitmap_size,  bitmap_s2_a,        0, NULL, NULL);
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_b,    CL_TRUE, 0, bitmap_size,  bitmap_s2_b,        0, NULL, NULL);
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_c,    CL_TRUE, 0, bitmap_size,  bitmap_s2_c,        0, NULL, NULL);
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_d,    CL_TRUE, 0, bitmap_size,  bitmap_s2_d,        0, NULL, NULL);
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_buf,    CL_TRUE, 0, size_digests, data.digests_buf,   0, NULL, NULL);
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown,  CL_TRUE, 0, size_shown,   data.digests_shown, 0, NULL, NULL);
+       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_salt_bufs,      CL_TRUE, 0, size_salts,   data.salts_buf,     0, NULL, NULL);
  
        run_kernel_bzero (device_param, device_param->d_pws_buf,        size_pws);
        run_kernel_bzero (device_param, device_param->d_pws_amp_buf,    size_pws);
@@@ -13895,19 -13877,19 +13914,19 @@@
  
        if (attack_kern == ATTACK_KERN_STRAIGHT)
        {
-         device_param->d_rules   = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_rules,   NULL);
-         device_param->d_rules_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_rules_c, NULL);
+         device_param->d_rules   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_rules,   NULL);
+         device_param->d_rules_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_rules_c, NULL);
  
-         hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_rules, CL_TRUE, 0, size_rules, kernel_rules_buf, 0, NULL, NULL);
+         hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_rules, CL_TRUE, 0, size_rules, kernel_rules_buf, 0, NULL, NULL);
  
          run_kernel_bzero (device_param, device_param->d_rules_c, size_rules_c);
        }
        else if (attack_kern == ATTACK_KERN_COMBI)
        {
-         device_param->d_combs           = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_combs,      NULL);
-         device_param->d_combs_c         = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_combs,      NULL);
-         device_param->d_root_css_buf    = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_root_css,   NULL);
-         device_param->d_markov_css_buf  = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
+         device_param->d_combs           = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs,      NULL);
+         device_param->d_combs_c         = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs,      NULL);
+         device_param->d_root_css_buf    = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css,   NULL);
+         device_param->d_markov_css_buf  = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
  
          run_kernel_bzero (device_param, device_param->d_combs,          size_combs);
          run_kernel_bzero (device_param, device_param->d_combs_c,        size_combs);
@@@ -13916,11 -13898,11 +13935,11 @@@
        }
        else if (attack_kern == ATTACK_KERN_BF)
        {
-         device_param->d_bfs             = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_bfs,        NULL);
-         device_param->d_bfs_c           = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_bfs,        NULL);
-         device_param->d_tm_c            = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_tm,         NULL);
-         device_param->d_root_css_buf    = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_root_css,   NULL);
-         device_param->d_markov_css_buf  = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
+         device_param->d_bfs             = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs,        NULL);
+         device_param->d_bfs_c           = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs,        NULL);
+         device_param->d_tm_c            = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_tm,         NULL);
+         device_param->d_root_css_buf    = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css,   NULL);
+         device_param->d_markov_css_buf  = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
  
          run_kernel_bzero (device_param, device_param->d_bfs,            size_bfs);
          run_kernel_bzero (device_param, device_param->d_bfs_c,          size_bfs);
@@@ -13931,9 -13913,9 +13950,9 @@@
  
        if (size_esalts)
        {
-         device_param->d_esalt_bufs = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_esalts, NULL);
+         device_param->d_esalt_bufs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_esalts, NULL);
  
-         hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_esalt_bufs, CL_TRUE, 0, size_esalts, data.esalts_buf, 0, NULL, NULL);
+         hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_esalt_bufs, CL_TRUE, 0, size_esalts, data.esalts_buf, 0, NULL, NULL);
        }
  
        /**
@@@ -14101,29 -14083,29 +14120,29 @@@
          {
            snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", kern_type, 4);
  
-           device_param->kernel1 = hc_clCreateKernel (device_param->program, kernel_name);
+           device_param->kernel1 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
  
            snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", kern_type, 8);
  
-           device_param->kernel2 = hc_clCreateKernel (device_param->program, kernel_name);
+           device_param->kernel2 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
  
            snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", kern_type, 16);
  
-           device_param->kernel3 = hc_clCreateKernel (device_param->program, kernel_name);
+           device_param->kernel3 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
          }
          else
          {
            snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_m%02d", kern_type, 4);
  
-           device_param->kernel1 = hc_clCreateKernel (device_param->program, kernel_name);
+           device_param->kernel1 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
  
            snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_m%02d", kern_type, 8);
  
-           device_param->kernel2 = hc_clCreateKernel (device_param->program, kernel_name);
+           device_param->kernel2 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
  
            snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_m%02d", kern_type, 16);
  
-           device_param->kernel3 = hc_clCreateKernel (device_param->program, kernel_name);
+           device_param->kernel3 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
          }
  
          if (data.attack_mode == ATTACK_MODE_BF)
@@@ -14132,11 -14114,11 +14151,11 @@@
            {
              snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_tb", kern_type);
  
-             device_param->kernel_tb = hc_clCreateKernel (device_param->program, kernel_name);
+             device_param->kernel_tb = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
  
              snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_tm", kern_type);
  
-             device_param->kernel_tm = hc_clCreateKernel (device_param->program, kernel_name);
+             device_param->kernel_tm = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
            }
          }
        }
@@@ -14144,71 -14126,71 +14163,71 @@@
        {
          snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_init", kern_type);
  
-         device_param->kernel1 = hc_clCreateKernel (device_param->program, kernel_name);
+         device_param->kernel1 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
  
          snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_loop", kern_type);
  
-         device_param->kernel2 = hc_clCreateKernel (device_param->program, kernel_name);
+         device_param->kernel2 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
  
          snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_comp", kern_type);
  
-         device_param->kernel3 = hc_clCreateKernel (device_param->program, kernel_name);
+         device_param->kernel3 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
  
          if (opts_type & OPTS_TYPE_HOOK12)
          {
            snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook12", kern_type);
  
-           device_param->kernel12 = hc_clCreateKernel (device_param->program, kernel_name);
+           device_param->kernel12 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
          }
  
          if (opts_type & OPTS_TYPE_HOOK23)
          {
            snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook23", kern_type);
  
-           device_param->kernel23 = hc_clCreateKernel (device_param->program, kernel_name);
+           device_param->kernel23 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
          }
        }
  
        for (uint i = 0; i <= 20; i++)
        {
-         hc_clSetKernelArg (device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]);
-         hc_clSetKernelArg (device_param->kernel2, i, sizeof (cl_mem), device_param->kernel_params[i]);
-         hc_clSetKernelArg (device_param->kernel3, i, sizeof (cl_mem), device_param->kernel_params[i]);
+         hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]);
+         hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_mem), device_param->kernel_params[i]);
+         hc_clSetKernelArg (data.ocl, device_param->kernel3, i, sizeof (cl_mem), device_param->kernel_params[i]);
  
-         if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (device_param->kernel12, i, sizeof (cl_mem), device_param->kernel_params[i]);
-         if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (device_param->kernel23, i, sizeof (cl_mem), device_param->kernel_params[i]);
+         if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (data.ocl, device_param->kernel12, i, sizeof (cl_mem), device_param->kernel_params[i]);
+         if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_mem), device_param->kernel_params[i]);
        }
  
        for (uint i = 21; i <= 31; i++)
        {
-         hc_clSetKernelArg (device_param->kernel1, i, sizeof (cl_uint), device_param->kernel_params[i]);
-         hc_clSetKernelArg (device_param->kernel2, i, sizeof (cl_uint), device_param->kernel_params[i]);
-         hc_clSetKernelArg (device_param->kernel3, i, sizeof (cl_uint), device_param->kernel_params[i]);
+         hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_uint), device_param->kernel_params[i]);
+         hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_uint), device_param->kernel_params[i]);
+         hc_clSetKernelArg (data.ocl, device_param->kernel3, i, sizeof (cl_uint), device_param->kernel_params[i]);
  
-         if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (device_param->kernel12, i, sizeof (cl_uint), device_param->kernel_params[i]);
-         if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (device_param->kernel23, i, sizeof (cl_uint), device_param->kernel_params[i]);
+         if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (data.ocl, device_param->kernel12, i, sizeof (cl_uint), device_param->kernel_params[i]);
+         if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_uint), device_param->kernel_params[i]);
        }
  
        if (attack_mode == ATTACK_MODE_BF)
        {
-         device_param->kernel_mp_l = hc_clCreateKernel (device_param->program_mp, "l_markov");
-         device_param->kernel_mp_r = hc_clCreateKernel (device_param->program_mp, "r_markov");
+         device_param->kernel_mp_l = hc_clCreateKernel (data.ocl, device_param->program_mp, "l_markov");
+         device_param->kernel_mp_r = hc_clCreateKernel (data.ocl, device_param->program_mp, "r_markov");
  
          if (opts_type & OPTS_TYPE_PT_BITSLICE)
          {
-           hc_clSetKernelArg (device_param->kernel_tb, 0, sizeof (cl_mem), device_param->kernel_params_tb[0]);
+           hc_clSetKernelArg (data.ocl, device_param->kernel_tb, 0, sizeof (cl_mem), device_param->kernel_params_tb[0]);
  
-           hc_clSetKernelArg (device_param->kernel_tm, 0, sizeof (cl_mem), device_param->kernel_params_tm[0]);
-           hc_clSetKernelArg (device_param->kernel_tm, 1, sizeof (cl_mem), device_param->kernel_params_tm[1]);
+           hc_clSetKernelArg (data.ocl, device_param->kernel_tm, 0, sizeof (cl_mem), device_param->kernel_params_tm[0]);
+           hc_clSetKernelArg (data.ocl, device_param->kernel_tm, 1, sizeof (cl_mem), device_param->kernel_params_tm[1]);
          }
        }
        else if (attack_mode == ATTACK_MODE_HYBRID1)
        {
-         device_param->kernel_mp = hc_clCreateKernel (device_param->program_mp, "C_markov");
+         device_param->kernel_mp = hc_clCreateKernel (data.ocl, device_param->program_mp, "C_markov");
        }
        else if (attack_mode == ATTACK_MODE_HYBRID2)
        {
-         device_param->kernel_mp = hc_clCreateKernel (device_param->program_mp, "C_markov");
+         device_param->kernel_mp = hc_clCreateKernel (data.ocl, device_param->program_mp, "C_markov");
        }
  
        if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
@@@ -14217,7 -14199,7 +14236,7 @@@
        }
        else
        {
-         device_param->kernel_amp = hc_clCreateKernel (device_param->program_amp, "amp");
+         device_param->kernel_amp = hc_clCreateKernel (data.ocl, device_param->program_amp, "amp");
        }
  
        if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
@@@ -14228,12 -14210,12 +14247,12 @@@
        {
          for (uint i = 0; i < 5; i++)
          {
-           hc_clSetKernelArg (device_param->kernel_amp, i, sizeof (cl_mem), device_param->kernel_params_amp[i]);
+           hc_clSetKernelArg (data.ocl, device_param->kernel_amp, i, sizeof (cl_mem), device_param->kernel_params_amp[i]);
          }
  
          for (uint i = 5; i < 7; i++)
          {
-           hc_clSetKernelArg (device_param->kernel_amp, i, sizeof (cl_uint), device_param->kernel_params_amp[i]);
+           hc_clSetKernelArg (data.ocl, device_param->kernel_amp, i, sizeof (cl_uint), device_param->kernel_params_amp[i]);
          }
        }
  
@@@ -15523,12 -15505,12 +15542,12 @@@
                device_param->kernel_params_mp_buf32[7] = 0;
              }
  
-             for (uint i = 0; i < 3; i++) hc_clSetKernelArg (device_param->kernel_mp, i, sizeof (cl_mem),   (void *) device_param->kernel_params_mp[i]);
-             for (uint i = 3; i < 4; i++) hc_clSetKernelArg (device_param->kernel_mp, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp[i]);
-             for (uint i = 4; i < 8; i++) hc_clSetKernelArg (device_param->kernel_mp, i, sizeof (cl_uint),  (void *) device_param->kernel_params_mp[i]);
+             for (uint i = 0; i < 3; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp, i, sizeof (cl_mem),   (void *) device_param->kernel_params_mp[i]);
+             for (uint i = 3; i < 4; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp[i]);
+             for (uint i = 4; i < 8; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp, i, sizeof (cl_uint),  (void *) device_param->kernel_params_mp[i]);
  
-             hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_root_css_buf,   CL_TRUE, 0, device_param->size_root_css,   root_css_buf,   0, NULL, NULL);
-             hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL);
+             hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_root_css_buf,   CL_TRUE, 0, device_param->size_root_css,   root_css_buf,   0, NULL, NULL);
+             hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL);
            }
          }
          else if (attack_mode == ATTACK_MODE_BF)
@@@ -15658,9 -15640,9 +15677,9 @@@
  
            // some more resets:
  
 -          memset (device_param->pw_caches, 0, 64 * sizeof (pw_cache_t));
 +          if (device_param->pw_caches) memset (device_param->pw_caches, 0, 64 * sizeof (pw_cache_t));
  
 -          memset (device_param->pws_buf, 0, device_param->size_pws);
 +          if (device_param->pws_buf) memset (device_param->pws_buf, 0, device_param->size_pws);
  
            device_param->pw_cnt  = 0;
            device_param->pws_cnt = 0;
@@@ -16034,16 -16016,16 +16053,16 @@@
              device_param->kernel_params_mp_r_buf32[6] = 0;
              device_param->kernel_params_mp_r_buf32[7] = 0;
  
-             for (uint i = 0; i < 3; i++) hc_clSetKernelArg (device_param->kernel_mp_l, i, sizeof (cl_mem),   (void *) device_param->kernel_params_mp_l[i]);
-             for (uint i = 3; i < 4; i++) hc_clSetKernelArg (device_param->kernel_mp_l, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_l[i]);
-             for (uint i = 4; i < 9; i++) hc_clSetKernelArg (device_param->kernel_mp_l, i, sizeof (cl_uint),  (void *) device_param->kernel_params_mp_l[i]);
+             for (uint i = 0; i < 3; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_mem),   (void *) device_param->kernel_params_mp_l[i]);
+             for (uint i = 3; i < 4; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_l[i]);
+             for (uint i = 4; i < 9; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_uint),  (void *) device_param->kernel_params_mp_l[i]);
  
-             for (uint i = 0; i < 3; i++) hc_clSetKernelArg (device_param->kernel_mp_r, i, sizeof (cl_mem),   (void *) device_param->kernel_params_mp_r[i]);
-             for (uint i = 3; i < 4; i++) hc_clSetKernelArg (device_param->kernel_mp_r, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_r[i]);
-             for (uint i = 4; i < 8; i++) hc_clSetKernelArg (device_param->kernel_mp_r, i, sizeof (cl_uint),  (void *) device_param->kernel_params_mp_r[i]);
+             for (uint i = 0; i < 3; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_mem),   (void *) device_param->kernel_params_mp_r[i]);
+             for (uint i = 3; i < 4; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_r[i]);
+             for (uint i = 4; i < 8; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_uint),  (void *) device_param->kernel_params_mp_r[i]);
  
-             hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_root_css_buf,   CL_TRUE, 0, device_param->size_root_css,   root_css_buf,   0, NULL, NULL);
-             hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL);
+             hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_root_css_buf,   CL_TRUE, 0, device_param->size_root_css,   root_css_buf,   0, NULL, NULL);
+             hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL);
            }
          }
  
@@@ -16465,53 -16447,53 +16484,53 @@@
        local_free (device_param->driver_version);
  
        if (device_param->pws_buf)            myfree                    (device_param->pws_buf);
-       if (device_param->d_pws_buf)          hc_clReleaseMemObject     (device_param->d_pws_buf);
-       if (device_param->d_pws_amp_buf)      hc_clReleaseMemObject     (device_param->d_pws_amp_buf);
-       if (device_param->d_rules)            hc_clReleaseMemObject     (device_param->d_rules);
-       if (device_param->d_rules_c)          hc_clReleaseMemObject     (device_param->d_rules_c);
-       if (device_param->d_combs)            hc_clReleaseMemObject     (device_param->d_combs);
-       if (device_param->d_combs_c)          hc_clReleaseMemObject     (device_param->d_combs_c);
-       if (device_param->d_bfs)              hc_clReleaseMemObject     (device_param->d_bfs);
-       if (device_param->d_bfs_c)            hc_clReleaseMemObject     (device_param->d_bfs_c);
-       if (device_param->d_bitmap_s1_a)      hc_clReleaseMemObject     (device_param->d_bitmap_s1_a);
-       if (device_param->d_bitmap_s1_b)      hc_clReleaseMemObject     (device_param->d_bitmap_s1_b);
-       if (device_param->d_bitmap_s1_c)      hc_clReleaseMemObject     (device_param->d_bitmap_s1_c);
-       if (device_param->d_bitmap_s1_d)      hc_clReleaseMemObject     (device_param->d_bitmap_s1_d);
-       if (device_param->d_bitmap_s2_a)      hc_clReleaseMemObject     (device_param->d_bitmap_s2_a);
-       if (device_param->d_bitmap_s2_b)      hc_clReleaseMemObject     (device_param->d_bitmap_s2_b);
-       if (device_param->d_bitmap_s2_c)      hc_clReleaseMemObject     (device_param->d_bitmap_s2_c);
-       if (device_param->d_bitmap_s2_d)      hc_clReleaseMemObject     (device_param->d_bitmap_s2_d);
-       if (device_param->d_plain_bufs)       hc_clReleaseMemObject     (device_param->d_plain_bufs);
-       if (device_param->d_digests_buf)      hc_clReleaseMemObject     (device_param->d_digests_buf);
-       if (device_param->d_digests_shown)    hc_clReleaseMemObject     (device_param->d_digests_shown);
-       if (device_param->d_salt_bufs)        hc_clReleaseMemObject     (device_param->d_salt_bufs);
-       if (device_param->d_esalt_bufs)       hc_clReleaseMemObject     (device_param->d_esalt_bufs);
-       if (device_param->d_tmps)             hc_clReleaseMemObject     (device_param->d_tmps);
-       if (device_param->d_hooks)            hc_clReleaseMemObject     (device_param->d_hooks);
-       if (device_param->d_result)           hc_clReleaseMemObject     (device_param->d_result);
-       if (device_param->d_scryptV_buf)      hc_clReleaseMemObject     (device_param->d_scryptV_buf);
-       if (device_param->d_root_css_buf)     hc_clReleaseMemObject     (device_param->d_root_css_buf);
-       if (device_param->d_markov_css_buf)   hc_clReleaseMemObject     (device_param->d_markov_css_buf);
-       if (device_param->d_tm_c)             hc_clReleaseMemObject     (device_param->d_tm_c);
- 
-       if (device_param->kernel1)            hc_clReleaseKernel        (device_param->kernel1);
-       if (device_param->kernel12)           hc_clReleaseKernel        (device_param->kernel12);
-       if (device_param->kernel2)            hc_clReleaseKernel        (device_param->kernel2);
-       if (device_param->kernel23)           hc_clReleaseKernel        (device_param->kernel23);
-       if (device_param->kernel3)            hc_clReleaseKernel        (device_param->kernel3);
-       if (device_param->kernel_mp)          hc_clReleaseKernel        (device_param->kernel_mp);
-       if (device_param->kernel_mp_l)        hc_clReleaseKernel        (device_param->kernel_mp_l);
-       if (device_param->kernel_mp_r)        hc_clReleaseKernel        (device_param->kernel_mp_r);
-       if (device_param->kernel_tb)          hc_clReleaseKernel        (device_param->kernel_tb);
-       if (device_param->kernel_tm)          hc_clReleaseKernel        (device_param->kernel_tm);
-       if (device_param->kernel_amp)         hc_clReleaseKernel        (device_param->kernel_amp);
- 
-       if (device_param->program)            hc_clReleaseProgram       (device_param->program);
-       if (device_param->program_mp)         hc_clReleaseProgram       (device_param->program_mp);
-       if (device_param->program_amp)        hc_clReleaseProgram       (device_param->program_amp);
- 
-       if (device_param->command_queue)      hc_clReleaseCommandQueue  (device_param->command_queue);
-       if (device_param->context)            hc_clReleaseContext       (device_param->context);
+       if (device_param->d_pws_buf)          hc_clReleaseMemObject     (data.ocl, device_param->d_pws_buf);
+       if (device_param->d_pws_amp_buf)      hc_clReleaseMemObject     (data.ocl, device_param->d_pws_amp_buf);
+       if (device_param->d_rules)            hc_clReleaseMemObject     (data.ocl, device_param->d_rules);
+       if (device_param->d_rules_c)          hc_clReleaseMemObject     (data.ocl, device_param->d_rules_c);
+       if (device_param->d_combs)            hc_clReleaseMemObject     (data.ocl, device_param->d_combs);
+       if (device_param->d_combs_c)          hc_clReleaseMemObject     (data.ocl, device_param->d_combs_c);
+       if (device_param->d_bfs)              hc_clReleaseMemObject     (data.ocl, device_param->d_bfs);
+       if (device_param->d_bfs_c)            hc_clReleaseMemObject     (data.ocl, device_param->d_bfs_c);
+       if (device_param->d_bitmap_s1_a)      hc_clReleaseMemObject     (data.ocl, device_param->d_bitmap_s1_a);
+       if (device_param->d_bitmap_s1_b)      hc_clReleaseMemObject     (data.ocl, device_param->d_bitmap_s1_b);
+       if (device_param->d_bitmap_s1_c)      hc_clReleaseMemObject     (data.ocl, device_param->d_bitmap_s1_c);
+       if (device_param->d_bitmap_s1_d)      hc_clReleaseMemObject     (data.ocl, device_param->d_bitmap_s1_d);
+       if (device_param->d_bitmap_s2_a)      hc_clReleaseMemObject     (data.ocl, device_param->d_bitmap_s2_a);
+       if (device_param->d_bitmap_s2_b)      hc_clReleaseMemObject     (data.ocl, device_param->d_bitmap_s2_b);
+       if (device_param->d_bitmap_s2_c)      hc_clReleaseMemObject     (data.ocl, device_param->d_bitmap_s2_c);
+       if (device_param->d_bitmap_s2_d)      hc_clReleaseMemObject     (data.ocl, device_param->d_bitmap_s2_d);
+       if (device_param->d_plain_bufs)       hc_clReleaseMemObject     (data.ocl, device_param->d_plain_bufs);
+       if (device_param->d_digests_buf)      hc_clReleaseMemObject     (data.ocl, device_param->d_digests_buf);
+       if (device_param->d_digests_shown)    hc_clReleaseMemObject     (data.ocl, device_param->d_digests_shown);
+       if (device_param->d_salt_bufs)        hc_clReleaseMemObject     (data.ocl, device_param->d_salt_bufs);
+       if (device_param->d_esalt_bufs)       hc_clReleaseMemObject     (data.ocl, device_param->d_esalt_bufs);
+       if (device_param->d_tmps)             hc_clReleaseMemObject     (data.ocl, device_param->d_tmps);
+       if (device_param->d_hooks)            hc_clReleaseMemObject     (data.ocl, device_param->d_hooks);
+       if (device_param->d_result)           hc_clReleaseMemObject     (data.ocl, device_param->d_result);
+       if (device_param->d_scryptV_buf)      hc_clReleaseMemObject     (data.ocl, device_param->d_scryptV_buf);
+       if (device_param->d_root_css_buf)     hc_clReleaseMemObject     (data.ocl, device_param->d_root_css_buf);
+       if (device_param->d_markov_css_buf)   hc_clReleaseMemObject     (data.ocl, device_param->d_markov_css_buf);
+       if (device_param->d_tm_c)             hc_clReleaseMemObject     (data.ocl, device_param->d_tm_c);
+ 
+       if (device_param->kernel1)            hc_clReleaseKernel        (data.ocl, device_param->kernel1);
+       if (device_param->kernel12)           hc_clReleaseKernel        (data.ocl, device_param->kernel12);
+       if (device_param->kernel2)            hc_clReleaseKernel        (data.ocl, device_param->kernel2);
+       if (device_param->kernel23)           hc_clReleaseKernel        (data.ocl, device_param->kernel23);
+       if (device_param->kernel3)            hc_clReleaseKernel        (data.ocl, device_param->kernel3);
+       if (device_param->kernel_mp)          hc_clReleaseKernel        (data.ocl, device_param->kernel_mp);
+       if (device_param->kernel_mp_l)        hc_clReleaseKernel        (data.ocl, device_param->kernel_mp_l);
+       if (device_param->kernel_mp_r)        hc_clReleaseKernel        (data.ocl, device_param->kernel_mp_r);
+       if (device_param->kernel_tb)          hc_clReleaseKernel        (data.ocl, device_param->kernel_tb);
+       if (device_param->kernel_tm)          hc_clReleaseKernel        (data.ocl, device_param->kernel_tm);
+       if (device_param->kernel_amp)         hc_clReleaseKernel        (data.ocl, device_param->kernel_amp);
+ 
+       if (device_param->program)            hc_clReleaseProgram       (data.ocl, device_param->program);
+       if (device_param->program_mp)         hc_clReleaseProgram       (data.ocl, device_param->program_mp);
+       if (device_param->program_amp)        hc_clReleaseProgram       (data.ocl, device_param->program_amp);
+ 
+       if (device_param->command_queue)      hc_clReleaseCommandQueue  (data.ocl, device_param->command_queue);
+       if (device_param->context)            hc_clReleaseContext       (data.ocl, device_param->context);
      }
  
      // reset default fan speed
@@@ -16792,6 -16774,8 +16811,8 @@@
    if (quiet == 0) log_info_nn ("Started: %s", ctime (&proc_start));
    if (quiet == 0) log_info_nn ("Stopped: %s", ctime (&proc_stop));
  
+   if (data.ocl) ocl_close (data.ocl);
+ 
    if (data.devices_status == STATUS_ABORTED)            return 2;
    if (data.devices_status == STATUS_QUIT)               return 2;
    if (data.devices_status == STATUS_STOP_AT_CHECKPOINT) return 2;