github/hashcat
1
Fork 0
mirror of https://github.com/hashcat/hashcat synced 2024-12-19 23:15:54 +01:00
Code Releases Activity

Add LOCAL_VK to make use of __shared__

Browse Source
This commit is contained in:
Jens Steube 2019-05-07 09:01:32 +02:00
parent bbed0cd67a
commit 8ff8c5d536
306 changed files with 1881 additions and 1841 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
OpenCL/inc_rp.cl
View File

@ -15,7 +15,7 @@
#ifdef REAL_SHM
#define COPY_PW(x) \
LOCAL_AS pw_t s_pws[64]; \
LOCAL_VK pw_t s_pws[64]; \
s_pws[get_local_id (0)] = (x);
#else
#define COPY_PW(x) \

3
OpenCL/inc_vendor.h
View File

@ -18,18 +18,21 @@
#define CONSTANT_VK
#define CONSTANT_AS
#define GLOBAL_AS
#define LOCAL_VK
#define LOCAL_AS
#define KERNEL_FQ
#elif defined IS_CUDA
#define CONSTANT_VK __constant__
#define CONSTANT_AS
#define GLOBAL_AS
#define LOCAL_VK __shared__
#define LOCAL_AS
#define KERNEL_FQ extern "C" __global__
#elif defined IS_OPENCL
#define CONSTANT_VK __constant
#define CONSTANT_AS __constant
#define GLOBAL_AS __global
#define LOCAL_VK __local
#define LOCAL_AS __local
#define KERNEL_FQ __kernel
#endif

4
OpenCL/m01100_a0-optimized.cl
View File

@ -43,7 +43,7 @@ KERNEL_FQ void m01100_m04 (KERN_ATTR_RULES ())
* salt
*/
LOCAL_AS salt_t s_salt_buf[1];
LOCAL_VK salt_t s_salt_buf[1];
if (lid == 0)
{
@ -262,7 +262,7 @@ KERNEL_FQ void m01100_s04 (KERN_ATTR_RULES ())
* salt
*/
LOCAL_AS salt_t s_salt_buf[1];
LOCAL_VK salt_t s_salt_buf[1];
if (lid == 0)
{

4
OpenCL/m01100_a1-optimized.cl
View File

@ -41,7 +41,7 @@ KERNEL_FQ void m01100_m04 (KERN_ATTR_BASIC ())
* salt
*/
LOCAL_AS salt_t s_salt_buf[1];
LOCAL_VK salt_t s_salt_buf[1];
if (lid == 0)
{
@ -322,7 +322,7 @@ KERNEL_FQ void m01100_s04 (KERN_ATTR_BASIC ())
* salt
*/
LOCAL_AS salt_t s_salt_buf[1];
LOCAL_VK salt_t s_salt_buf[1];
if (lid == 0)
{

12
OpenCL/m01100_a3-optimized.cl
View File

@ -533,7 +533,7 @@ KERNEL_FQ void m01100_m04 (KERN_ATTR_VECTOR ())
* salt
*/
LOCAL_AS salt_t s_salt_buf[1];
LOCAL_VK salt_t s_salt_buf[1];
if (lid == 0)
{
@ -587,7 +587,7 @@ KERNEL_FQ void m01100_m08 (KERN_ATTR_VECTOR ())
* salt
*/
LOCAL_AS salt_t s_salt_buf[1];
LOCAL_VK salt_t s_salt_buf[1];
if (lid == 0)
{
@ -641,7 +641,7 @@ KERNEL_FQ void m01100_m16 (KERN_ATTR_VECTOR ())
* salt
*/
LOCAL_AS salt_t s_salt_buf[1];
LOCAL_VK salt_t s_salt_buf[1];
if (lid == 0)
{
@ -695,7 +695,7 @@ KERNEL_FQ void m01100_s04 (KERN_ATTR_VECTOR ())
* salt
*/
LOCAL_AS salt_t s_salt_buf[1];
LOCAL_VK salt_t s_salt_buf[1];
if (lid == 0)
{
@ -749,7 +749,7 @@ KERNEL_FQ void m01100_s08 (KERN_ATTR_VECTOR ())
* salt
*/
LOCAL_AS salt_t s_salt_buf[1];
LOCAL_VK salt_t s_salt_buf[1];
if (lid == 0)
{
@ -803,7 +803,7 @@ KERNEL_FQ void m01100_s16 (KERN_ATTR_VECTOR ())
* salt
*/
LOCAL_AS salt_t s_salt_buf[1];
LOCAL_VK salt_t s_salt_buf[1];
if (lid == 0)
{

8
OpenCL/m01500_a0-pure.cl
View File

@ -496,8 +496,8 @@ KERNEL_FQ void m01500_mxx (KERN_ATTR_RULES ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -580,8 +580,8 @@ KERNEL_FQ void m01500_sxx (KERN_ATTR_RULES ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

8
OpenCL/m01500_a1-pure.cl
View File

@ -494,8 +494,8 @@ KERNEL_FQ void m01500_mxx (KERN_ATTR_BASIC ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -657,8 +657,8 @@ KERNEL_FQ void m01500_sxx (KERN_ATTR_BASIC ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

10
OpenCL/m02500-pure.cl
View File

@ -666,11 +666,11 @@ KERNEL_FQ void m02500_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t)
const u64 lid = get_local_id (0);
const u64 lsz = get_local_size (0);
LOCAL_AS u32 s_te0[256];
LOCAL_AS u32 s_te1[256];
LOCAL_AS u32 s_te2[256];
LOCAL_AS u32 s_te3[256];
LOCAL_AS u32 s_te4[256];
LOCAL_VK u32 s_te0[256];
LOCAL_VK u32 s_te1[256];
LOCAL_VK u32 s_te2[256];
LOCAL_VK u32 s_te3[256];
LOCAL_VK u32 s_te4[256];
for (u32 i = lid; i < 256; i += lsz)
{

10
OpenCL/m02501-pure.cl
View File

@ -536,11 +536,11 @@ KERNEL_FQ void m02501_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pmk_tmp_t, wpa_eapol_t))
const u64 lid = get_local_id (0);
const u64 lsz = get_local_size (0);
LOCAL_AS u32 s_te0[256];
LOCAL_AS u32 s_te1[256];
LOCAL_AS u32 s_te2[256];
LOCAL_AS u32 s_te3[256];
LOCAL_AS u32 s_te4[256];
LOCAL_VK u32 s_te0[256];
LOCAL_VK u32 s_te1[256];
LOCAL_VK u32 s_te2[256];
LOCAL_VK u32 s_te3[256];
LOCAL_VK u32 s_te4[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m02610_a0-optimized.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m02610_m04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -326,7 +326,7 @@ KERNEL_FQ void m02610_s04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m02610_a0-pure.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m02610_mxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -145,7 +145,7 @@ KERNEL_FQ void m02610_sxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m02610_a1-optimized.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m02610_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -383,7 +383,7 @@ KERNEL_FQ void m02610_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m02610_a1-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m02610_mxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -141,7 +141,7 @@ KERNEL_FQ void m02610_sxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

12
OpenCL/m02610_a3-optimized.cl
View File

@ -606,7 +606,7 @@ KERNEL_FQ void m02610_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -676,7 +676,7 @@ KERNEL_FQ void m02610_m08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -746,7 +746,7 @@ KERNEL_FQ void m02610_m16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -816,7 +816,7 @@ KERNEL_FQ void m02610_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -886,7 +886,7 @@ KERNEL_FQ void m02610_s08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -956,7 +956,7 @@ KERNEL_FQ void m02610_s16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m02610_a3-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m02610_mxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -154,7 +154,7 @@ KERNEL_FQ void m02610_sxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m02710_a0-optimized.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m02710_m04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -411,7 +411,7 @@ KERNEL_FQ void m02710_s04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m02710_a1-optimized.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m02710_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -468,7 +468,7 @@ KERNEL_FQ void m02710_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

12
OpenCL/m02710_a3-optimized.cl
View File

@ -775,7 +775,7 @@ KERNEL_FQ void m02710_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -845,7 +845,7 @@ KERNEL_FQ void m02710_m08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -915,7 +915,7 @@ KERNEL_FQ void m02710_m16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -985,7 +985,7 @@ KERNEL_FQ void m02710_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1055,7 +1055,7 @@ KERNEL_FQ void m02710_s08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1125,7 +1125,7 @@ KERNEL_FQ void m02710_s16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m02810_a0-optimized.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m02810_m04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -410,7 +410,7 @@ KERNEL_FQ void m02810_s04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m02810_a0-pure.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m02810_mxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -168,7 +168,7 @@ KERNEL_FQ void m02810_sxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m02810_a1-optimized.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m02810_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -467,7 +467,7 @@ KERNEL_FQ void m02810_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m02810_a1-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m02810_mxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -164,7 +164,7 @@ KERNEL_FQ void m02810_sxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

12
OpenCL/m02810_a3-optimized.cl
View File

@ -773,7 +773,7 @@ KERNEL_FQ void m02810_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -843,7 +843,7 @@ KERNEL_FQ void m02810_m08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -913,7 +913,7 @@ KERNEL_FQ void m02810_m16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -983,7 +983,7 @@ KERNEL_FQ void m02810_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1053,7 +1053,7 @@ KERNEL_FQ void m02810_s08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1123,7 +1123,7 @@ KERNEL_FQ void m02810_s16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m02810_a3-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m02810_mxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -177,7 +177,7 @@ KERNEL_FQ void m02810_sxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

8
OpenCL/m03000_a0-pure.cl
View File

@ -506,8 +506,8 @@ KERNEL_FQ void m03000_mxx (KERN_ATTR_RULES ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -591,8 +591,8 @@ KERNEL_FQ void m03000_sxx (KERN_ATTR_RULES ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

8
OpenCL/m03000_a1-pure.cl
View File

@ -504,8 +504,8 @@ KERNEL_FQ void m03000_mxx (KERN_ATTR_BASIC ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -668,8 +668,8 @@ KERNEL_FQ void m03000_sxx (KERN_ATTR_BASIC ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

8
OpenCL/m03100_a0-optimized.cl
View File

@ -33,8 +33,8 @@ KERNEL_FQ void m03100_m04 (KERN_ATTR_RULES ())
#ifdef REAL_SHM
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -265,8 +265,8 @@ KERNEL_FQ void m03100_s04 (KERN_ATTR_RULES ())
#ifdef REAL_SHM
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

8
OpenCL/m03100_a1-optimized.cl
View File

@ -31,8 +31,8 @@ KERNEL_FQ void m03100_m04 (KERN_ATTR_BASIC ())
#ifdef REAL_SHM
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -323,8 +323,8 @@ KERNEL_FQ void m03100_s04 (KERN_ATTR_BASIC ())
#ifdef REAL_SHM
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

16
OpenCL/m03100_a3-optimized.cl
View File

@ -429,8 +429,8 @@ KERNEL_FQ void m03100_m04 (KERN_ATTR_VECTOR ())
#ifdef REAL_SHM
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -512,8 +512,8 @@ KERNEL_FQ void m03100_m08 (KERN_ATTR_VECTOR ())
#ifdef REAL_SHM
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -599,8 +599,8 @@ KERNEL_FQ void m03100_s04 (KERN_ATTR_VECTOR ())
#ifdef REAL_SHM
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -682,8 +682,8 @@ KERNEL_FQ void m03100_s08 (KERN_ATTR_VECTOR ())
#ifdef REAL_SHM
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

48
OpenCL/m03200-pure.cl
View File

@ -444,15 +444,15 @@ KERNEL_FQ void __attribute__((reqd_work_group_size(FIXED_LOCAL_SIZE, 1, 1))) m03
* do the key setup
*/
LOCAL_AS u32 S0_all[FIXED_LOCAL_SIZE][256];
LOCAL_AS u32 S1_all[FIXED_LOCAL_SIZE][256];
LOCAL_AS u32 S2_all[FIXED_LOCAL_SIZE][256];
LOCAL_AS u32 S3_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S0_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S1_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S2_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S3_all[FIXED_LOCAL_SIZE][256];
LOCAL_AS u32 *S0 = S0_all[lid];
LOCAL_AS u32 *S1 = S1_all[lid];
LOCAL_AS u32 *S2 = S2_all[lid];
LOCAL_AS u32 *S3 = S3_all[lid];
LOCAL_VK u32 *S0 = S0_all[lid];
LOCAL_VK u32 *S1 = S1_all[lid];
LOCAL_VK u32 *S2 = S2_all[lid];
LOCAL_VK u32 *S3 = S3_all[lid];
// initstate
@ -613,15 +613,15 @@ KERNEL_FQ void __attribute__((reqd_work_group_size(FIXED_LOCAL_SIZE, 1, 1))) m03
P[i] = tmps[gid].P[i];
}
LOCAL_AS u32 S0_all[FIXED_LOCAL_SIZE][256];
LOCAL_AS u32 S1_all[FIXED_LOCAL_SIZE][256];
LOCAL_AS u32 S2_all[FIXED_LOCAL_SIZE][256];
LOCAL_AS u32 S3_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S0_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S1_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S2_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S3_all[FIXED_LOCAL_SIZE][256];
LOCAL_AS u32 *S0 = S0_all[lid];
LOCAL_AS u32 *S1 = S1_all[lid];
LOCAL_AS u32 *S2 = S2_all[lid];
LOCAL_AS u32 *S3 = S3_all[lid];
LOCAL_VK u32 *S0 = S0_all[lid];
LOCAL_VK u32 *S1 = S1_all[lid];
LOCAL_VK u32 *S2 = S2_all[lid];
LOCAL_VK u32 *S3 = S3_all[lid];
for (u32 i = 0; i < 256; i++)
{
@ -798,15 +798,15 @@ KERNEL_FQ void __attribute__((reqd_work_group_size(FIXED_LOCAL_SIZE, 1, 1))) m03
P[i] = tmps[gid].P[i];
}
LOCAL_AS u32 S0_all[FIXED_LOCAL_SIZE][256];
LOCAL_AS u32 S1_all[FIXED_LOCAL_SIZE][256];
LOCAL_AS u32 S2_all[FIXED_LOCAL_SIZE][256];
LOCAL_AS u32 S3_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S0_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S1_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S2_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S3_all[FIXED_LOCAL_SIZE][256];
LOCAL_AS u32 *S0 = S0_all[lid];
LOCAL_AS u32 *S1 = S1_all[lid];
LOCAL_AS u32 *S2 = S2_all[lid];
LOCAL_AS u32 *S3 = S3_all[lid];
LOCAL_VK u32 *S0 = S0_all[lid];
LOCAL_VK u32 *S1 = S1_all[lid];
LOCAL_VK u32 *S2 = S2_all[lid];
LOCAL_VK u32 *S3 = S3_all[lid];
for (u32 i = 0; i < 256; i++)
{

4
OpenCL/m03710_a0-optimized.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m03710_m04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -357,7 +357,7 @@ KERNEL_FQ void m03710_s04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m03710_a0-pure.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m03710_mxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -158,7 +158,7 @@ KERNEL_FQ void m03710_sxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m03710_a1-optimized.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m03710_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -414,7 +414,7 @@ KERNEL_FQ void m03710_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m03710_a1-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m03710_mxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -154,7 +154,7 @@ KERNEL_FQ void m03710_sxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

12
OpenCL/m03710_a3-optimized.cl
View File

@ -633,7 +633,7 @@ KERNEL_FQ void m03710_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -737,7 +737,7 @@ KERNEL_FQ void m03710_m08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -807,7 +807,7 @@ KERNEL_FQ void m03710_m16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -877,7 +877,7 @@ KERNEL_FQ void m03710_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -947,7 +947,7 @@ KERNEL_FQ void m03710_s08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -983,7 +983,7 @@ KERNEL_FQ void m03710_s16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m03710_a3-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m03710_mxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -167,7 +167,7 @@ KERNEL_FQ void m03710_sxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m03910_a0-optimized.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m03910_m04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -410,7 +410,7 @@ KERNEL_FQ void m03910_s04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m03910_a0-pure.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m03910_mxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -168,7 +168,7 @@ KERNEL_FQ void m03910_sxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m03910_a1-optimized.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m03910_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -467,7 +467,7 @@ KERNEL_FQ void m03910_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m03910_a1-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m03910_mxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -164,7 +164,7 @@ KERNEL_FQ void m03910_sxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

12
OpenCL/m03910_a3-optimized.cl
View File

@ -773,7 +773,7 @@ KERNEL_FQ void m03910_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -843,7 +843,7 @@ KERNEL_FQ void m03910_m08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -913,7 +913,7 @@ KERNEL_FQ void m03910_m16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -983,7 +983,7 @@ KERNEL_FQ void m03910_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1053,7 +1053,7 @@ KERNEL_FQ void m03910_s08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1123,7 +1123,7 @@ KERNEL_FQ void m03910_s16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m03910_a3-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m03910_mxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -177,7 +177,7 @@ KERNEL_FQ void m03910_sxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04010_a0-optimized.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m04010_m04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -383,7 +383,7 @@ KERNEL_FQ void m04010_s04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04010_a0-pure.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m04010_mxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -149,7 +149,7 @@ KERNEL_FQ void m04010_sxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04010_a1-optimized.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04010_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -439,7 +439,7 @@ KERNEL_FQ void m04010_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04010_a1-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04010_mxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -145,7 +145,7 @@ KERNEL_FQ void m04010_sxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

12
OpenCL/m04010_a3-optimized.cl
View File

@ -673,7 +673,7 @@ KERNEL_FQ void m04010_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -777,7 +777,7 @@ KERNEL_FQ void m04010_m08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -847,7 +847,7 @@ KERNEL_FQ void m04010_m16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -917,7 +917,7 @@ KERNEL_FQ void m04010_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -987,7 +987,7 @@ KERNEL_FQ void m04010_s08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1023,7 +1023,7 @@ KERNEL_FQ void m04010_s16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04010_a3-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04010_mxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -162,7 +162,7 @@ KERNEL_FQ void m04010_sxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04110_a0-optimized.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m04110_m04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -428,7 +428,7 @@ KERNEL_FQ void m04110_s04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04110_a0-pure.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m04110_mxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -162,7 +162,7 @@ KERNEL_FQ void m04110_sxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04110_a1-optimized.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04110_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -486,7 +486,7 @@ KERNEL_FQ void m04110_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04110_a1-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04110_mxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -158,7 +158,7 @@ KERNEL_FQ void m04110_sxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

12
OpenCL/m04110_a3-optimized.cl
View File

@ -729,7 +729,7 @@ KERNEL_FQ void m04110_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -833,7 +833,7 @@ KERNEL_FQ void m04110_m08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -903,7 +903,7 @@ KERNEL_FQ void m04110_m16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -973,7 +973,7 @@ KERNEL_FQ void m04110_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1043,7 +1043,7 @@ KERNEL_FQ void m04110_s08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1079,7 +1079,7 @@ KERNEL_FQ void m04110_s16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04110_a3-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04110_mxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -173,7 +173,7 @@ KERNEL_FQ void m04110_sxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04310_a0-optimized.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m04310_m04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -326,7 +326,7 @@ KERNEL_FQ void m04310_s04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04310_a0-pure.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m04310_mxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -145,7 +145,7 @@ KERNEL_FQ void m04310_sxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04310_a1-optimized.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04310_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -383,7 +383,7 @@ KERNEL_FQ void m04310_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04310_a1-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04310_mxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -141,7 +141,7 @@ KERNEL_FQ void m04310_sxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

12
OpenCL/m04310_a3-optimized.cl
View File

@ -606,7 +606,7 @@ KERNEL_FQ void m04310_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -676,7 +676,7 @@ KERNEL_FQ void m04310_m08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -746,7 +746,7 @@ KERNEL_FQ void m04310_m16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -816,7 +816,7 @@ KERNEL_FQ void m04310_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -886,7 +886,7 @@ KERNEL_FQ void m04310_s08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -956,7 +956,7 @@ KERNEL_FQ void m04310_s16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04310_a3-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04310_mxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -154,7 +154,7 @@ KERNEL_FQ void m04310_sxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04400_a0-optimized.cl
View File

@ -43,7 +43,7 @@ KERNEL_FQ void m04400_m04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -353,7 +353,7 @@ KERNEL_FQ void m04400_s04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04400_a0-pure.cl
View File

@ -43,7 +43,7 @@ KERNEL_FQ void m04400_mxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -140,7 +140,7 @@ KERNEL_FQ void m04400_sxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04400_a1-optimized.cl
View File

@ -41,7 +41,7 @@ KERNEL_FQ void m04400_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -409,7 +409,7 @@ KERNEL_FQ void m04400_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04400_a1-pure.cl
View File

@ -41,7 +41,7 @@ KERNEL_FQ void m04400_mxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -136,7 +136,7 @@ KERNEL_FQ void m04400_sxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

12
OpenCL/m04400_a3-optimized.cl
View File

@ -578,7 +578,7 @@ KERNEL_FQ void m04400_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -648,7 +648,7 @@ KERNEL_FQ void m04400_m08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -718,7 +718,7 @@ KERNEL_FQ void m04400_m16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -788,7 +788,7 @@ KERNEL_FQ void m04400_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -858,7 +858,7 @@ KERNEL_FQ void m04400_s08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -928,7 +928,7 @@ KERNEL_FQ void m04400_s16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04400_a3-pure.cl
View File

@ -41,7 +41,7 @@ KERNEL_FQ void m04400_mxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -149,7 +149,7 @@ KERNEL_FQ void m04400_sxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04500_a0-optimized.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m04500_m04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -379,7 +379,7 @@ KERNEL_FQ void m04500_s04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04500_a0-pure.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m04500_mxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -139,7 +139,7 @@ KERNEL_FQ void m04500_sxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04500_a1-optimized.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04500_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -435,7 +435,7 @@ KERNEL_FQ void m04500_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04500_a1-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04500_mxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -135,7 +135,7 @@ KERNEL_FQ void m04500_sxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

12
OpenCL/m04500_a3-optimized.cl
View File

@ -637,7 +637,7 @@ KERNEL_FQ void m04500_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -707,7 +707,7 @@ KERNEL_FQ void m04500_m08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -777,7 +777,7 @@ KERNEL_FQ void m04500_m16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -847,7 +847,7 @@ KERNEL_FQ void m04500_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -917,7 +917,7 @@ KERNEL_FQ void m04500_s08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -987,7 +987,7 @@ KERNEL_FQ void m04500_s16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04500_a3-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04500_mxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -148,7 +148,7 @@ KERNEL_FQ void m04500_sxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04520_a0-optimized.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m04520_m04 (KERN_ATTR_RULES ())
* shared
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -610,7 +610,7 @@ KERNEL_FQ void m04520_s04 (KERN_ATTR_RULES ())
* shared
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04520_a0-pure.cl
View File

@ -42,7 +42,7 @@ KERNEL_FQ void m04520_mxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -154,7 +154,7 @@ KERNEL_FQ void m04520_sxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04520_a1-optimized.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04520_m04 (KERN_ATTR_BASIC ())
* shared
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -666,7 +666,7 @@ KERNEL_FQ void m04520_s04 (KERN_ATTR_BASIC ())
* shared
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04520_a1-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04520_mxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -150,7 +150,7 @@ KERNEL_FQ void m04520_sxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

12
OpenCL/m04520_a3-optimized.cl
View File

@ -1090,7 +1090,7 @@ KERNEL_FQ void m04520_m04 (KERN_ATTR_BASIC ())
* shared
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1160,7 +1160,7 @@ KERNEL_FQ void m04520_m08 (KERN_ATTR_BASIC ())
* shared
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1230,7 +1230,7 @@ KERNEL_FQ void m04520_m16 (KERN_ATTR_BASIC ())
* shared
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1300,7 +1300,7 @@ KERNEL_FQ void m04520_s04 (KERN_ATTR_BASIC ())
* shared
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1370,7 +1370,7 @@ KERNEL_FQ void m04520_s08 (KERN_ATTR_BASIC ())
* shared
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -1440,7 +1440,7 @@ KERNEL_FQ void m04520_s16 (KERN_ATTR_BASIC ())
* shared
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04520_a3-pure.cl
View File

@ -40,7 +40,7 @@ KERNEL_FQ void m04520_mxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -165,7 +165,7 @@ KERNEL_FQ void m04520_sxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04700_a0-optimized.cl
View File

@ -43,7 +43,7 @@ KERNEL_FQ void m04700_m04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -336,7 +336,7 @@ KERNEL_FQ void m04700_s04 (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04700_a0-pure.cl
View File

@ -43,7 +43,7 @@ KERNEL_FQ void m04700_mxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -135,7 +135,7 @@ KERNEL_FQ void m04700_sxx (KERN_ATTR_RULES ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04700_a1-optimized.cl
View File

@ -41,7 +41,7 @@ KERNEL_FQ void m04700_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -389,7 +389,7 @@ KERNEL_FQ void m04700_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04700_a1-pure.cl
View File

@ -41,7 +41,7 @@ KERNEL_FQ void m04700_mxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -131,7 +131,7 @@ KERNEL_FQ void m04700_sxx (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

12
OpenCL/m04700_a3-optimized.cl
View File

@ -578,7 +578,7 @@ KERNEL_FQ void m04700_m04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -648,7 +648,7 @@ KERNEL_FQ void m04700_m08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -718,7 +718,7 @@ KERNEL_FQ void m04700_m16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -788,7 +788,7 @@ KERNEL_FQ void m04700_s04 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -858,7 +858,7 @@ KERNEL_FQ void m04700_s08 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -928,7 +928,7 @@ KERNEL_FQ void m04700_s16 (KERN_ATTR_BASIC ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m04700_a3-pure.cl
View File

@ -41,7 +41,7 @@ KERNEL_FQ void m04700_mxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -144,7 +144,7 @@ KERNEL_FQ void m04700_sxx (KERN_ATTR_VECTOR ())
* bin2asc table
*/
LOCAL_AS u32 l_bin2asc[256];
LOCAL_VK u32 l_bin2asc[256];
for (u32 i = lid; i < 256; i += lsz)
{

8
OpenCL/m05300_a0-optimized.cl
View File

@ -125,14 +125,14 @@ KERNEL_FQ void m05300_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -300,14 +300,14 @@ KERNEL_FQ void m05300_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{

8
OpenCL/m05300_a1-optimized.cl
View File

@ -123,14 +123,14 @@ KERNEL_FQ void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -358,14 +358,14 @@ KERNEL_FQ void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{

24
OpenCL/m05300_a3-optimized.cl
View File

@ -429,14 +429,14 @@ KERNEL_FQ void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -502,14 +502,14 @@ KERNEL_FQ void m05300_m08 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -575,14 +575,14 @@ KERNEL_FQ void m05300_m16 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -648,14 +648,14 @@ KERNEL_FQ void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -721,14 +721,14 @@ KERNEL_FQ void m05300_s08 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -794,14 +794,14 @@ KERNEL_FQ void m05300_s16 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{

8
OpenCL/m05400_a0-optimized.cl
View File

@ -129,14 +129,14 @@ KERNEL_FQ void m05400_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -313,14 +313,14 @@ KERNEL_FQ void m05400_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{

8
OpenCL/m05400_a1-optimized.cl
View File

@ -127,14 +127,14 @@ KERNEL_FQ void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -379,14 +379,14 @@ KERNEL_FQ void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{

24
OpenCL/m05400_a3-optimized.cl
View File

@ -433,14 +433,14 @@ KERNEL_FQ void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -506,14 +506,14 @@ KERNEL_FQ void m05400_m08 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -579,14 +579,14 @@ KERNEL_FQ void m05400_m16 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -652,14 +652,14 @@ KERNEL_FQ void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -725,14 +725,14 @@ KERNEL_FQ void m05400_s08 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{
@ -798,14 +798,14 @@ KERNEL_FQ void m05400_s16 (KERN_ATTR_ESALT (ikepsk_t))
* s_msg
*/
LOCAL_AS u32 s_nr_buf[16];
LOCAL_VK u32 s_nr_buf[16];
for (u32 i = lid; i < 16; i += lsz)
{
s_nr_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
}
LOCAL_AS u32 s_msg_buf[128];
LOCAL_VK u32 s_msg_buf[128];
for (u32 i = lid; i < 128; i += lsz)
{

8
OpenCL/m05500_a0-optimized.cl
View File

@ -516,8 +516,8 @@ KERNEL_FQ void m05500_m04 (KERN_ATTR_RULES ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -729,8 +729,8 @@ KERNEL_FQ void m05500_s04 (KERN_ATTR_RULES ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

8
OpenCL/m05500_a0-pure.cl
View File

@ -516,8 +516,8 @@ KERNEL_FQ void m05500_mxx (KERN_ATTR_RULES ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -640,8 +640,8 @@ KERNEL_FQ void m05500_sxx (KERN_ATTR_RULES ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

8
OpenCL/m05500_a1-optimized.cl
View File

@ -514,8 +514,8 @@ KERNEL_FQ void m05500_m04 (KERN_ATTR_BASIC ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -780,8 +780,8 @@ KERNEL_FQ void m05500_s04 (KERN_ATTR_BASIC ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

8
OpenCL/m05500_a1-pure.cl
View File

@ -514,8 +514,8 @@ KERNEL_FQ void m05500_mxx (KERN_ATTR_BASIC ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -636,8 +636,8 @@ KERNEL_FQ void m05500_sxx (KERN_ATTR_BASIC ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

24
OpenCL/m05500_a3-optimized.cl
View File

@ -847,8 +847,8 @@ KERNEL_FQ void m05500_m04 (KERN_ATTR_VECTOR ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -921,8 +921,8 @@ KERNEL_FQ void m05500_m08 (KERN_ATTR_VECTOR ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -995,8 +995,8 @@ KERNEL_FQ void m05500_m16 (KERN_ATTR_VECTOR ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -1069,8 +1069,8 @@ KERNEL_FQ void m05500_s04 (KERN_ATTR_VECTOR ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -1143,8 +1143,8 @@ KERNEL_FQ void m05500_s08 (KERN_ATTR_VECTOR ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -1217,8 +1217,8 @@ KERNEL_FQ void m05500_s16 (KERN_ATTR_VECTOR ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

8
OpenCL/m05500_a3-pure.cl
View File

@ -514,8 +514,8 @@ KERNEL_FQ void m05500_mxx (KERN_ATTR_VECTOR ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
@ -649,8 +649,8 @@ KERNEL_FQ void m05500_sxx (KERN_ATTR_VECTOR ())
* sbox, kbox
*/
LOCAL_AS u32 s_SPtrans[8][64];
LOCAL_AS u32 s_skb[8][64];
LOCAL_VK u32 s_SPtrans[8][64];
LOCAL_VK u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{

8
OpenCL/m05600_a0-optimized.cl
View File

@ -128,14 +128,14 @@ KERNEL_FQ void m05600_m04 (KERN_ATTR_RULES_ESALT (netntlm_t))
* salt
*/
LOCAL_AS u32 s_userdomain_buf[64];
LOCAL_VK u32 s_userdomain_buf[64];
for (u32 i = lid; i < 64; i += lsz)
{
s_userdomain_buf[i] = esalt_bufs[digests_offset].userdomain_buf[i];
}
LOCAL_AS u32 s_chall_buf[256];
LOCAL_VK u32 s_chall_buf[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -367,14 +367,14 @@ KERNEL_FQ void m05600_s04 (KERN_ATTR_RULES_ESALT (netntlm_t))
* salt
*/
LOCAL_AS u32 s_userdomain_buf[64];
LOCAL_VK u32 s_userdomain_buf[64];
for (u32 i = lid; i < 64; i += lsz)
{
s_userdomain_buf[i] = esalt_bufs[digests_offset].userdomain_buf[i];
}
LOCAL_AS u32 s_chall_buf[256];
LOCAL_VK u32 s_chall_buf[256];
for (u32 i = lid; i < 256; i += lsz)
{

8
OpenCL/m05600_a1-optimized.cl
View File

@ -126,14 +126,14 @@ KERNEL_FQ void m05600_m04 (KERN_ATTR_ESALT (netntlm_t))
* salt
*/
LOCAL_AS u32 s_userdomain_buf[64];
LOCAL_VK u32 s_userdomain_buf[64];
for (u32 i = lid; i < 64; i += lsz)
{
s_userdomain_buf[i] = esalt_bufs[digests_offset].userdomain_buf[i];
}
LOCAL_AS u32 s_chall_buf[256];
LOCAL_VK u32 s_chall_buf[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -423,14 +423,14 @@ KERNEL_FQ void m05600_s04 (KERN_ATTR_ESALT (netntlm_t))
* salt
*/
LOCAL_AS u32 s_userdomain_buf[64];
LOCAL_VK u32 s_userdomain_buf[64];
for (u32 i = lid; i < 64; i += lsz)
{
s_userdomain_buf[i] = esalt_bufs[digests_offset].userdomain_buf[i];
}
LOCAL_AS u32 s_chall_buf[256];
LOCAL_VK u32 s_chall_buf[256];
for (u32 i = lid; i < 256; i += lsz)
{

24
OpenCL/m05600_a3-optimized.cl
View File

@ -550,14 +550,14 @@ KERNEL_FQ void m05600_m04 (KERN_ATTR_ESALT (netntlm_t))
* salt
*/
LOCAL_AS u32 s_userdomain_buf[64];
LOCAL_VK u32 s_userdomain_buf[64];
for (u32 i = lid; i < 64; i += lsz)
{
s_userdomain_buf[i] = esalt_bufs[digests_offset].userdomain_buf[i];
}
LOCAL_AS u32 s_chall_buf[256];
LOCAL_VK u32 s_chall_buf[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -623,14 +623,14 @@ KERNEL_FQ void m05600_m08 (KERN_ATTR_ESALT (netntlm_t))
* salt
*/
LOCAL_AS u32 s_userdomain_buf[64];
LOCAL_VK u32 s_userdomain_buf[64];
for (u32 i = lid; i < 64; i += lsz)
{
s_userdomain_buf[i] = esalt_bufs[digests_offset].userdomain_buf[i];
}
LOCAL_AS u32 s_chall_buf[256];
LOCAL_VK u32 s_chall_buf[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -696,14 +696,14 @@ KERNEL_FQ void m05600_m16 (KERN_ATTR_ESALT (netntlm_t))
* salt
*/
LOCAL_AS u32 s_userdomain_buf[64];
LOCAL_VK u32 s_userdomain_buf[64];
for (u32 i = lid; i < 64; i += lsz)
{
s_userdomain_buf[i] = esalt_bufs[digests_offset].userdomain_buf[i];
}
LOCAL_AS u32 s_chall_buf[256];
LOCAL_VK u32 s_chall_buf[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -769,14 +769,14 @@ KERNEL_FQ void m05600_s04 (KERN_ATTR_ESALT (netntlm_t))
* salt
*/
LOCAL_AS u32 s_userdomain_buf[64];
LOCAL_VK u32 s_userdomain_buf[64];
for (u32 i = lid; i < 64; i += lsz)
{
s_userdomain_buf[i] = esalt_bufs[digests_offset].userdomain_buf[i];
}
LOCAL_AS u32 s_chall_buf[256];
LOCAL_VK u32 s_chall_buf[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -842,14 +842,14 @@ KERNEL_FQ void m05600_s08 (KERN_ATTR_ESALT (netntlm_t))
* salt
*/
LOCAL_AS u32 s_userdomain_buf[64];
LOCAL_VK u32 s_userdomain_buf[64];
for (u32 i = lid; i < 64; i += lsz)
{
s_userdomain_buf[i] = esalt_bufs[digests_offset].userdomain_buf[i];
}
LOCAL_AS u32 s_chall_buf[256];
LOCAL_VK u32 s_chall_buf[256];
for (u32 i = lid; i < 256; i += lsz)
{
@ -915,14 +915,14 @@ KERNEL_FQ void m05600_s16 (KERN_ATTR_ESALT (netntlm_t))
* salt
*/
LOCAL_AS u32 s_userdomain_buf[64];
LOCAL_VK u32 s_userdomain_buf[64];
for (u32 i = lid; i < 64; i += lsz)
{
s_userdomain_buf[i] = esalt_bufs[digests_offset].userdomain_buf[i];
}
LOCAL_AS u32 s_chall_buf[256];
LOCAL_VK u32 s_chall_buf[256];
for (u32 i = lid; i < 256; i += lsz)
{

4
OpenCL/m05800-optimized.cl
View File

@ -2300,8 +2300,8 @@ KERNEL_FQ void m05800_loop (KERN_ATTR_TMPS (androidpin_tmp_t))
* cache precomputed conversion table in shared memory
*/
LOCAL_AS u32 s_pc_dec[1024];
LOCAL_AS u32 s_pc_len[1024];
LOCAL_VK u32 s_pc_dec[1024];
LOCAL_VK u32 s_pc_len[1024];
for (u32 i = lid; i < 1024; i += lsz)
{

Some files were not shown because too many files have changed in this diff Show More