    __kernel void m00000_m00 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u00 *bitmaps_buf_s0_a, __global const u00 *bitmaps_buf_s0_b, __global const u00 *bitmaps_buf_s0_c, __global const u00 *bitmaps_buf_s0_d, __global const u00 *bitmaps_buf_s0_a, __global const u00 *bitmaps_buf_s0_b, __global const u00 *bitmaps_buf_s0_c, __global const u00 *bitmaps_buf_s0_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u00 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u00 *d_return_buf, __global u00 *d_scryptV0_buf, __global u00 *d_scryptV0_buf, __global u00 *d_scryptV0_buf, __global u00 *d_scryptV0_buf, const u00 bitmap_mask, const u00 bitmap_shift0, const u00 bitmap_shift0, const u00 salt_pos, const u00 loop_pos, const u00 loop_cnt, const u00 il_cnt, const u00 digests_cnt, const u00 digests_offset, const u00 combs_mode, const u00 gid_max)  {    /**     * modifier     */      const u00 lid = get_local_id (0);      /**     * base     */      const u00 gid = get_global_id (0);      if (gid >= gid_max) return;      u00 pw_buf0[0];    u00 pw_buf0[0];      pw_buf0[0] = pws[gid].i[0];    pw_buf0[0] = pws[gid].i[0];    pw_buf0[0] = pws[gid].i[0];    pw_buf0[0] = pws[gid].i[0];    pw_buf0[0] = pws[gid].i[0];    pw_buf0[0] = pws[gid].i[0];    pw_buf0[0] = pws[gid].i[0];    pw_buf0[0] = pws[gid].i[0];      const u00 pw_l_len = pws[gid].pw_len;      /**     * loop     */      for (u00 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)    {      const u00x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);        const u00x pw_len = pw_l_len + pw_r_len;        /**       * concat password candidate       */        u00x wordl0[0] = { 0 };      u00x wordl0[0] = { 0 };      u00x wordl0[0] = { 0 };      u00x wordl0[0] = { 0 };        wordl0[0] = pw_buf0[0];      wordl0[0] = pw_buf0[0];      wordl0[0] = pw_buf0[0];      wordl0[0] = pw_buf0[0];      wordl0[0] = pw_buf0[0];      wordl0[0] = pw_buf0[0];      wordl0[0] = pw_buf0[0];      wordl0[0] = pw_buf0[0];        u00x wordr0[0] = { 0 };      u00x wordr0[0] = { 0 };      u00x wordr0[0] = { 0 };      u00x wordr0[0] = { 0 };        wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);      wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);      wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);      wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);      wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);      wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);      wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);      wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);        if (combs_mode == COMBINATOR_MODE_BASE_LEFT)      {        switch_buffer_by_offset_le_VV (wordr0, wordr0, wordr0, wordr0, pw_l_len);      }      else      {        switch_buffer_by_offset_le_VV (wordl0, wordl0, wordl0, wordl0, pw_r_len);      }        u00x w0[0];      u00x w0[0];      u00x w0[0];      u00x w0[0];        w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];      w0[0] = wordl0[0] | wordr0[0];        make_utf00le (w0, w0, w0);      make_utf00le (w0, w0, w0);        w0[0] = pw_len * 0 * 0;      w0[0] = 0;        /**       * md0       */        u00x a = MD0M_A;      u00x b = MD0M_B;      u00x c = MD0M_C;      u00x d = MD0M_D;        MD0_STEP (MD0_Fo, a, b, c, d, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, d, a, b, c, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, c, d, a, b, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, b, c, d, a, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, a, b, c, d, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, d, a, b, c, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, c, d, a, b, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, b, c, d, a, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, a, b, c, d, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, d, a, b, c, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, c, d, a, b, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, b, c, d, a, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, a, b, c, d, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, d, a, b, c, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, c, d, a, b, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Fo, b, c, d, a, w0[0], MD0C00, MD0S00);        MD0_STEP (MD0_Go, a, b, c, d, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, d, a, b, c, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, c, d, a, b, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, b, c, d, a, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, a, b, c, d, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, d, a, b, c, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, c, d, a, b, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, b, c, d, a, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, a, b, c, d, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, d, a, b, c, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, c, d, a, b, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, b, c, d, a, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, a, b, c, d, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, d, a, b, c, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, c, d, a, b, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_Go, b, c, d, a, w0[0], MD0C00, MD0S00);        MD0_STEP (MD0_H , a, b, c, d, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , d, a, b, c, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , c, d, a, b, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , b, c, d, a, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , a, b, c, d, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , d, a, b, c, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , c, d, a, b, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , b, c, d, a, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , a, b, c, d, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , d, a, b, c, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , c, d, a, b, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , b, c, d, a, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , a, b, c, d, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , d, a, b, c, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , c, d, a, b, w0[0], MD0C00, MD0S00);      MD0_STEP (MD0_H , b, c, d, a, w0[0], MD0C00, MD0S00);        COMPARE_