mirror of
https://github.com/hashcat/hashcat.git
synced 2024-12-22 22:58:30 +00:00
Replace barrier() with SYNC_THREADS()
This commit is contained in:
parent
9faba41848
commit
00e1e32492
@ -51,7 +51,7 @@ KERNEL_FQ void m01100_m04 (KERN_ATTR_RULES ())
|
||||
s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -270,7 +270,7 @@ KERNEL_FQ void m01100_s04 (KERN_ATTR_RULES ())
|
||||
s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -49,7 +49,7 @@ KERNEL_FQ void m01100_m04 (KERN_ATTR_BASIC ())
|
||||
s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -330,7 +330,7 @@ KERNEL_FQ void m01100_s04 (KERN_ATTR_BASIC ())
|
||||
s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -541,7 +541,7 @@ KERNEL_FQ void m01100_m04 (KERN_ATTR_VECTOR ())
|
||||
s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -595,7 +595,7 @@ KERNEL_FQ void m01100_m08 (KERN_ATTR_VECTOR ())
|
||||
s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -649,7 +649,7 @@ KERNEL_FQ void m01100_m16 (KERN_ATTR_VECTOR ())
|
||||
s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -703,7 +703,7 @@ KERNEL_FQ void m01100_s04 (KERN_ATTR_VECTOR ())
|
||||
s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -757,7 +757,7 @@ KERNEL_FQ void m01100_s08 (KERN_ATTR_VECTOR ())
|
||||
s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -811,7 +811,7 @@ KERNEL_FQ void m01100_s16 (KERN_ATTR_VECTOR ())
|
||||
s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -519,7 +519,7 @@ KERNEL_FQ void m01500_mxx (KERN_ATTR_RULES ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -603,7 +603,7 @@ KERNEL_FQ void m01500_sxx (KERN_ATTR_RULES ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -517,7 +517,7 @@ KERNEL_FQ void m01500_mxx (KERN_ATTR_BASIC ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -680,7 +680,7 @@ KERNEL_FQ void m01500_sxx (KERN_ATTR_BASIC ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -682,7 +682,7 @@ KERNEL_FQ void m02500_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t)
|
||||
#ifdef IS_CUDA
|
||||
__syncthreads();
|
||||
#else
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
#endif
|
||||
|
||||
#else
|
||||
|
@ -549,7 +549,7 @@ KERNEL_FQ void m02501_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pmk_tmp_t, wpa_eapol_t))
|
||||
s_te4[i] = te4[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m02610_m04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -336,7 +336,7 @@ KERNEL_FQ void m02610_s04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m02610_mxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -155,7 +155,7 @@ KERNEL_FQ void m02610_sxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m02610_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -393,7 +393,7 @@ KERNEL_FQ void m02610_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m02610_mxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -151,7 +151,7 @@ KERNEL_FQ void m02610_sxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -616,7 +616,7 @@ KERNEL_FQ void m02610_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -686,7 +686,7 @@ KERNEL_FQ void m02610_m08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -756,7 +756,7 @@ KERNEL_FQ void m02610_m16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -826,7 +826,7 @@ KERNEL_FQ void m02610_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -896,7 +896,7 @@ KERNEL_FQ void m02610_s08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -966,7 +966,7 @@ KERNEL_FQ void m02610_s16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m02610_mxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -164,7 +164,7 @@ KERNEL_FQ void m02610_sxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m02710_m04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -421,7 +421,7 @@ KERNEL_FQ void m02710_s04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m02710_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -478,7 +478,7 @@ KERNEL_FQ void m02710_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -785,7 +785,7 @@ KERNEL_FQ void m02710_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -855,7 +855,7 @@ KERNEL_FQ void m02710_m08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -925,7 +925,7 @@ KERNEL_FQ void m02710_m16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -995,7 +995,7 @@ KERNEL_FQ void m02710_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1065,7 +1065,7 @@ KERNEL_FQ void m02710_s08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1135,7 +1135,7 @@ KERNEL_FQ void m02710_s16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m02810_m04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -420,7 +420,7 @@ KERNEL_FQ void m02810_s04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m02810_mxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -178,7 +178,7 @@ KERNEL_FQ void m02810_sxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m02810_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -477,7 +477,7 @@ KERNEL_FQ void m02810_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m02810_mxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -174,7 +174,7 @@ KERNEL_FQ void m02810_sxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -783,7 +783,7 @@ KERNEL_FQ void m02810_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -853,7 +853,7 @@ KERNEL_FQ void m02810_m08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -923,7 +923,7 @@ KERNEL_FQ void m02810_m16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -993,7 +993,7 @@ KERNEL_FQ void m02810_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1063,7 +1063,7 @@ KERNEL_FQ void m02810_s08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1133,7 +1133,7 @@ KERNEL_FQ void m02810_s16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m02810_mxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -187,7 +187,7 @@ KERNEL_FQ void m02810_sxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -529,7 +529,7 @@ KERNEL_FQ void m03000_mxx (KERN_ATTR_RULES ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -614,7 +614,7 @@ KERNEL_FQ void m03000_sxx (KERN_ATTR_RULES ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -527,7 +527,7 @@ KERNEL_FQ void m03000_mxx (KERN_ATTR_BASIC ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -691,7 +691,7 @@ KERNEL_FQ void m03000_sxx (KERN_ATTR_BASIC ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -56,7 +56,7 @@ KERNEL_FQ void m03100_m04 (KERN_ATTR_RULES ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
@ -288,7 +288,7 @@ KERNEL_FQ void m03100_s04 (KERN_ATTR_RULES ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
|
@ -54,7 +54,7 @@ KERNEL_FQ void m03100_m04 (KERN_ATTR_BASIC ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
@ -346,7 +346,7 @@ KERNEL_FQ void m03100_s04 (KERN_ATTR_BASIC ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
|
@ -452,7 +452,7 @@ KERNEL_FQ void m03100_m04 (KERN_ATTR_VECTOR ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
@ -535,7 +535,7 @@ KERNEL_FQ void m03100_m08 (KERN_ATTR_VECTOR ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
@ -622,7 +622,7 @@ KERNEL_FQ void m03100_s04 (KERN_ATTR_VECTOR ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
@ -705,7 +705,7 @@ KERNEL_FQ void m03100_s08 (KERN_ATTR_VECTOR ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m03710_m04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -367,7 +367,7 @@ KERNEL_FQ void m03710_s04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m03710_mxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -168,7 +168,7 @@ KERNEL_FQ void m03710_sxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m03710_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -424,7 +424,7 @@ KERNEL_FQ void m03710_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m03710_mxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -164,7 +164,7 @@ KERNEL_FQ void m03710_sxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -643,7 +643,7 @@ KERNEL_FQ void m03710_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -747,7 +747,7 @@ KERNEL_FQ void m03710_m08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -817,7 +817,7 @@ KERNEL_FQ void m03710_m16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -887,7 +887,7 @@ KERNEL_FQ void m03710_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -957,7 +957,7 @@ KERNEL_FQ void m03710_s08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -993,7 +993,7 @@ KERNEL_FQ void m03710_s16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m03710_mxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -177,7 +177,7 @@ KERNEL_FQ void m03710_sxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m03910_m04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -420,7 +420,7 @@ KERNEL_FQ void m03910_s04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m03910_mxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -178,7 +178,7 @@ KERNEL_FQ void m03910_sxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m03910_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -477,7 +477,7 @@ KERNEL_FQ void m03910_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m03910_mxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -174,7 +174,7 @@ KERNEL_FQ void m03910_sxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -783,7 +783,7 @@ KERNEL_FQ void m03910_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -853,7 +853,7 @@ KERNEL_FQ void m03910_m08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -923,7 +923,7 @@ KERNEL_FQ void m03910_m16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -993,7 +993,7 @@ KERNEL_FQ void m03910_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1063,7 +1063,7 @@ KERNEL_FQ void m03910_s08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1133,7 +1133,7 @@ KERNEL_FQ void m03910_s16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m03910_mxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -187,7 +187,7 @@ KERNEL_FQ void m03910_sxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m04010_m04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -393,7 +393,7 @@ KERNEL_FQ void m04010_s04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m04010_mxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -159,7 +159,7 @@ KERNEL_FQ void m04010_sxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04010_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -449,7 +449,7 @@ KERNEL_FQ void m04010_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04010_mxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -155,7 +155,7 @@ KERNEL_FQ void m04010_sxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -683,7 +683,7 @@ KERNEL_FQ void m04010_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -787,7 +787,7 @@ KERNEL_FQ void m04010_m08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -857,7 +857,7 @@ KERNEL_FQ void m04010_m16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -927,7 +927,7 @@ KERNEL_FQ void m04010_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -997,7 +997,7 @@ KERNEL_FQ void m04010_s08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1033,7 +1033,7 @@ KERNEL_FQ void m04010_s16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04010_mxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -172,7 +172,7 @@ KERNEL_FQ void m04010_sxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m04110_m04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -438,7 +438,7 @@ KERNEL_FQ void m04110_s04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m04110_mxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -172,7 +172,7 @@ KERNEL_FQ void m04110_sxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04110_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -496,7 +496,7 @@ KERNEL_FQ void m04110_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04110_mxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -168,7 +168,7 @@ KERNEL_FQ void m04110_sxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -739,7 +739,7 @@ KERNEL_FQ void m04110_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -843,7 +843,7 @@ KERNEL_FQ void m04110_m08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -913,7 +913,7 @@ KERNEL_FQ void m04110_m16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -983,7 +983,7 @@ KERNEL_FQ void m04110_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1053,7 +1053,7 @@ KERNEL_FQ void m04110_s08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1089,7 +1089,7 @@ KERNEL_FQ void m04110_s16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04110_mxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -183,7 +183,7 @@ KERNEL_FQ void m04110_sxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m04310_m04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -336,7 +336,7 @@ KERNEL_FQ void m04310_s04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m04310_mxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -155,7 +155,7 @@ KERNEL_FQ void m04310_sxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04310_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -393,7 +393,7 @@ KERNEL_FQ void m04310_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04310_mxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -151,7 +151,7 @@ KERNEL_FQ void m04310_sxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -616,7 +616,7 @@ KERNEL_FQ void m04310_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -686,7 +686,7 @@ KERNEL_FQ void m04310_m08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -756,7 +756,7 @@ KERNEL_FQ void m04310_m16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -826,7 +826,7 @@ KERNEL_FQ void m04310_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -896,7 +896,7 @@ KERNEL_FQ void m04310_s08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -966,7 +966,7 @@ KERNEL_FQ void m04310_s16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04310_mxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -164,7 +164,7 @@ KERNEL_FQ void m04310_sxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -53,7 +53,7 @@ KERNEL_FQ void m04400_m04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -363,7 +363,7 @@ KERNEL_FQ void m04400_s04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -53,7 +53,7 @@ KERNEL_FQ void m04400_mxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -150,7 +150,7 @@ KERNEL_FQ void m04400_sxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -51,7 +51,7 @@ KERNEL_FQ void m04400_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -419,7 +419,7 @@ KERNEL_FQ void m04400_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -51,7 +51,7 @@ KERNEL_FQ void m04400_mxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -146,7 +146,7 @@ KERNEL_FQ void m04400_sxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -588,7 +588,7 @@ KERNEL_FQ void m04400_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -658,7 +658,7 @@ KERNEL_FQ void m04400_m08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -728,7 +728,7 @@ KERNEL_FQ void m04400_m16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -798,7 +798,7 @@ KERNEL_FQ void m04400_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -868,7 +868,7 @@ KERNEL_FQ void m04400_s08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -938,7 +938,7 @@ KERNEL_FQ void m04400_s16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -51,7 +51,7 @@ KERNEL_FQ void m04400_mxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -159,7 +159,7 @@ KERNEL_FQ void m04400_sxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m04500_m04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -389,7 +389,7 @@ KERNEL_FQ void m04500_s04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m04500_mxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -149,7 +149,7 @@ KERNEL_FQ void m04500_sxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04500_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -445,7 +445,7 @@ KERNEL_FQ void m04500_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04500_mxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -145,7 +145,7 @@ KERNEL_FQ void m04500_sxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -647,7 +647,7 @@ KERNEL_FQ void m04500_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -717,7 +717,7 @@ KERNEL_FQ void m04500_m08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -787,7 +787,7 @@ KERNEL_FQ void m04500_m16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -857,7 +857,7 @@ KERNEL_FQ void m04500_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -927,7 +927,7 @@ KERNEL_FQ void m04500_s08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -997,7 +997,7 @@ KERNEL_FQ void m04500_s16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04500_mxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -158,7 +158,7 @@ KERNEL_FQ void m04500_sxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m04520_m04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -620,7 +620,7 @@ KERNEL_FQ void m04520_s04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -52,7 +52,7 @@ KERNEL_FQ void m04520_mxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -164,7 +164,7 @@ KERNEL_FQ void m04520_sxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04520_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -676,7 +676,7 @@ KERNEL_FQ void m04520_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04520_mxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -160,7 +160,7 @@ KERNEL_FQ void m04520_sxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -1100,7 +1100,7 @@ KERNEL_FQ void m04520_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1170,7 +1170,7 @@ KERNEL_FQ void m04520_m08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1240,7 +1240,7 @@ KERNEL_FQ void m04520_m16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1310,7 +1310,7 @@ KERNEL_FQ void m04520_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1380,7 +1380,7 @@ KERNEL_FQ void m04520_s08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1450,7 +1450,7 @@ KERNEL_FQ void m04520_s16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -50,7 +50,7 @@ KERNEL_FQ void m04520_mxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -175,7 +175,7 @@ KERNEL_FQ void m04520_sxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -53,7 +53,7 @@ KERNEL_FQ void m04700_m04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -346,7 +346,7 @@ KERNEL_FQ void m04700_s04 (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -53,7 +53,7 @@ KERNEL_FQ void m04700_mxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -145,7 +145,7 @@ KERNEL_FQ void m04700_sxx (KERN_ATTR_RULES ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -51,7 +51,7 @@ KERNEL_FQ void m04700_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -399,7 +399,7 @@ KERNEL_FQ void m04700_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -51,7 +51,7 @@ KERNEL_FQ void m04700_mxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -141,7 +141,7 @@ KERNEL_FQ void m04700_sxx (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -588,7 +588,7 @@ KERNEL_FQ void m04700_m04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -658,7 +658,7 @@ KERNEL_FQ void m04700_m08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -728,7 +728,7 @@ KERNEL_FQ void m04700_m16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -798,7 +798,7 @@ KERNEL_FQ void m04700_s04 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -868,7 +868,7 @@ KERNEL_FQ void m04700_s08 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -938,7 +938,7 @@ KERNEL_FQ void m04700_s16 (KERN_ATTR_BASIC ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -51,7 +51,7 @@ KERNEL_FQ void m04700_mxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -154,7 +154,7 @@ KERNEL_FQ void m04700_sxx (KERN_ATTR_VECTOR ())
|
||||
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -138,7 +138,7 @@ KERNEL_FQ void m05300_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -313,7 +313,7 @@ KERNEL_FQ void m05300_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -136,7 +136,7 @@ KERNEL_FQ void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -371,7 +371,7 @@ KERNEL_FQ void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -442,7 +442,7 @@ KERNEL_FQ void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -515,7 +515,7 @@ KERNEL_FQ void m05300_m08 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -588,7 +588,7 @@ KERNEL_FQ void m05300_m16 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -661,7 +661,7 @@ KERNEL_FQ void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -734,7 +734,7 @@ KERNEL_FQ void m05300_s08 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -807,7 +807,7 @@ KERNEL_FQ void m05300_s16 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -142,7 +142,7 @@ KERNEL_FQ void m05400_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -326,7 +326,7 @@ KERNEL_FQ void m05400_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -140,7 +140,7 @@ KERNEL_FQ void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -392,7 +392,7 @@ KERNEL_FQ void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -446,7 +446,7 @@ KERNEL_FQ void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -519,7 +519,7 @@ KERNEL_FQ void m05400_m08 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -592,7 +592,7 @@ KERNEL_FQ void m05400_m16 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -665,7 +665,7 @@ KERNEL_FQ void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -738,7 +738,7 @@ KERNEL_FQ void m05400_s08 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -811,7 +811,7 @@ KERNEL_FQ void m05400_s16 (KERN_ATTR_ESALT (ikepsk_t))
|
||||
s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -539,7 +539,7 @@ KERNEL_FQ void m05500_m04 (KERN_ATTR_RULES ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -752,7 +752,7 @@ KERNEL_FQ void m05500_s04 (KERN_ATTR_RULES ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -539,7 +539,7 @@ KERNEL_FQ void m05500_mxx (KERN_ATTR_RULES ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -663,7 +663,7 @@ KERNEL_FQ void m05500_sxx (KERN_ATTR_RULES ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -537,7 +537,7 @@ KERNEL_FQ void m05500_m04 (KERN_ATTR_BASIC ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -803,7 +803,7 @@ KERNEL_FQ void m05500_s04 (KERN_ATTR_BASIC ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -537,7 +537,7 @@ KERNEL_FQ void m05500_mxx (KERN_ATTR_BASIC ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -659,7 +659,7 @@ KERNEL_FQ void m05500_sxx (KERN_ATTR_BASIC ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -870,7 +870,7 @@ KERNEL_FQ void m05500_m04 (KERN_ATTR_VECTOR ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -944,7 +944,7 @@ KERNEL_FQ void m05500_m08 (KERN_ATTR_VECTOR ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1018,7 +1018,7 @@ KERNEL_FQ void m05500_m16 (KERN_ATTR_VECTOR ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1092,7 +1092,7 @@ KERNEL_FQ void m05500_s04 (KERN_ATTR_VECTOR ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1166,7 +1166,7 @@ KERNEL_FQ void m05500_s08 (KERN_ATTR_VECTOR ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -1240,7 +1240,7 @@ KERNEL_FQ void m05500_s16 (KERN_ATTR_VECTOR ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -537,7 +537,7 @@ KERNEL_FQ void m05500_mxx (KERN_ATTR_VECTOR ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -672,7 +672,7 @@ KERNEL_FQ void m05500_sxx (KERN_ATTR_VECTOR ())
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -141,7 +141,7 @@ KERNEL_FQ void m05600_m04 (KERN_ATTR_RULES_ESALT (netntlm_t))
|
||||
s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -380,7 +380,7 @@ KERNEL_FQ void m05600_s04 (KERN_ATTR_RULES_ESALT (netntlm_t))
|
||||
s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -139,7 +139,7 @@ KERNEL_FQ void m05600_m04 (KERN_ATTR_ESALT (netntlm_t))
|
||||
s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -436,7 +436,7 @@ KERNEL_FQ void m05600_s04 (KERN_ATTR_ESALT (netntlm_t))
|
||||
s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -563,7 +563,7 @@ KERNEL_FQ void m05600_m04 (KERN_ATTR_ESALT (netntlm_t))
|
||||
s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -636,7 +636,7 @@ KERNEL_FQ void m05600_m08 (KERN_ATTR_ESALT (netntlm_t))
|
||||
s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -709,7 +709,7 @@ KERNEL_FQ void m05600_m16 (KERN_ATTR_ESALT (netntlm_t))
|
||||
s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -782,7 +782,7 @@ KERNEL_FQ void m05600_s04 (KERN_ATTR_ESALT (netntlm_t))
|
||||
s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -855,7 +855,7 @@ KERNEL_FQ void m05600_s08 (KERN_ATTR_ESALT (netntlm_t))
|
||||
s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
@ -928,7 +928,7 @@ KERNEL_FQ void m05600_s16 (KERN_ATTR_ESALT (netntlm_t))
|
||||
s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -2308,7 +2308,7 @@ KERNEL_FQ void m05800_loop (KERN_ATTR_TMPS (androidpin_tmp_t))
|
||||
s_pc_len[i] = c_pc_len[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -2128,7 +2128,7 @@ KERNEL_FQ void m05800_loop (KERN_ATTR_TMPS (androidpin_tmp_t))
|
||||
s_pc_len[i] = c_pc_len[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
@ -60,7 +60,7 @@ KERNEL_FQ void m06100_m04 (KERN_ATTR_RULES ())
|
||||
s_Cl[7][i] = Cl[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
@ -200,7 +200,7 @@ KERNEL_FQ void m06100_s04 (KERN_ATTR_RULES ())
|
||||
s_Cl[7][i] = Cl[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
|
@ -55,7 +55,7 @@ KERNEL_FQ void m06100_mxx (KERN_ATTR_RULES ())
|
||||
s_Cl[7][i] = Cl[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
@ -139,7 +139,7 @@ KERNEL_FQ void m06100_sxx (KERN_ATTR_RULES ())
|
||||
s_Cl[7][i] = Cl[7][i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
|
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in New Issue
Block a user