arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-23 15:18:16 +00:00
Code Issues Releases Wiki Activity

updated m17210 to be cuda compatible

Browse Source
This commit is contained in:
Sein Coray 2019-05-14 14:38:55 +02:00
parent b29019ae75
commit c9c7261a05
No known key found for this signature in database
GPG Key ID: 44C4180EA69758EC
3 changed files with 42 additions and 51 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

31
OpenCL/m17210_a0-pure.cl
View File

@ -87,6 +87,7 @@ Related publication: https://scitepress.org/PublicationsDetail.aspx?ID=KLPzPqStp
#include "inc_vendor.h" #include "inc_vendor.h"
#include "inc_types.h" #include "inc_types.h"
#include "inc_platform.cl"
#include "inc_common.cl" #include "inc_common.cl"
#include "inc_simd.cl" #include "inc_simd.cl"
#include "inc_rp.h" #include "inc_rp.h"
@ -131,7 +132,7 @@ struct pkzip_hash
u32 data_length; u32 data_length;
u16 checksum_from_crc; u16 checksum_from_crc;
u16 checksum_from_timestamp; u16 checksum_from_timestamp;
u8 data[MAX_DATA]; u32 data[MAX_DATA];
} __attribute__((packed)); } __attribute__((packed));
@ -219,7 +220,7 @@ CONSTANT_AS u32a crc32tab[256] =
0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d 0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d
}; };
__kernel void m17210_sxx (KERN_ATTR_RULES_ESALT (pkzip_t)) KERNEL_FQ void m17210_sxx (KERN_ATTR_RULES_ESALT (pkzip_t))
{ {
/** /**
* modifier * modifier
@ -233,25 +234,23 @@ __kernel void m17210_sxx (KERN_ATTR_RULES_ESALT (pkzip_t))
* sbox, kbox * sbox, kbox
*/ */
LOCAL_AS u32 l_crc32tab[256]; LOCAL_VK u32 l_crc32tab[256];
for (u64 i = lid; i < 256; i += lsz) for (u64 i = lid; i < 256; i += lsz)
{ {
l_crc32tab[i] = crc32tab[i]; l_crc32tab[i] = crc32tab[i];
} }
barrier (CLK_LOCAL_MEM_FENCE); SYNC_THREADS();
__global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hash.data; LOCAL_VK u32 l_data[MAX_LOCAL];
LOCAL_AS u32 l_data[MAX_LOCAL];
for (u64 i = lid; i < MAX_LOCAL; i += lsz) for (u64 i = lid; i < MAX_LOCAL; i += lsz)
{ {
l_data[i] = data_ptr[i]; l_data[i] = esalt_bufs[digests_offset].hash.data[i];
} }
barrier (CLK_LOCAL_MEM_FENCE); SYNC_THREADS();
if (gid >= gid_max) return; if (gid >= gid_max) return;
@ -428,7 +427,7 @@ __kernel void m17210_sxx (KERN_ATTR_RULES_ESALT (pkzip_t))
} }
} }
__kernel void m17210_mxx (KERN_ATTR_RULES_ESALT (pkzip_t)) KERNEL_FQ void m17210_mxx (KERN_ATTR_RULES_ESALT (pkzip_t))
{ {
/** /**
* modifier * modifier
@ -442,25 +441,23 @@ __kernel void m17210_mxx (KERN_ATTR_RULES_ESALT (pkzip_t))
* sbox, kbox * sbox, kbox
*/ */
LOCAL_AS u32 l_crc32tab[256]; LOCAL_VK u32 l_crc32tab[256];
for (u64 i = lid; i < 256; i += lsz) for (u64 i = lid; i < 256; i += lsz)
{ {
l_crc32tab[i] = crc32tab[i]; l_crc32tab[i] = crc32tab[i];
} }
barrier (CLK_LOCAL_MEM_FENCE); SYNC_THREADS();
__global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hash.data; LOCAL_VK u32 l_data[MAX_LOCAL];
LOCAL_AS u32 l_data[MAX_LOCAL];
for (u64 i = lid; i < MAX_LOCAL; i += lsz) for (u64 i = lid; i < MAX_LOCAL; i += lsz)
{ {
l_data[i] = data_ptr[i]; l_data[i] = esalt_bufs[digests_offset].hash.data[i];
} }
barrier (CLK_LOCAL_MEM_FENCE); SYNC_THREADS();
if (gid >= gid_max) return; if (gid >= gid_max) return;

31
OpenCL/m17210_a1-pure.cl
View File

@ -87,6 +87,7 @@ Related publication: https://scitepress.org/PublicationsDetail.aspx?ID=KLPzPqStp
#include "inc_vendor.h" #include "inc_vendor.h"
#include "inc_types.h" #include "inc_types.h"
#include "inc_platform.cl"
#include "inc_common.cl" #include "inc_common.cl"
#include "inc_simd.cl" #include "inc_simd.cl"
@ -129,7 +130,7 @@ struct pkzip_hash
u32 data_length; u32 data_length;
u16 checksum_from_crc; u16 checksum_from_crc;
u16 checksum_from_timestamp; u16 checksum_from_timestamp;
u8 data[MAX_DATA]; u32 data[MAX_DATA];
} __attribute__((packed)); } __attribute__((packed));
@ -217,7 +218,7 @@ CONSTANT_AS u32a crc32tab[256] =
0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d 0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d
}; };
__kernel void m17210_sxx (KERN_ATTR_ESALT (pkzip_t)) KERNEL_FQ void m17210_sxx (KERN_ATTR_ESALT (pkzip_t))
{ {
/** /**
* modifier * modifier
@ -231,25 +232,23 @@ __kernel void m17210_sxx (KERN_ATTR_ESALT (pkzip_t))
* sbox, kbox * sbox, kbox
*/ */
LOCAL_AS u32 l_crc32tab[256]; LOCAL_VK u32 l_crc32tab[256];
for (u64 i = lid; i < 256; i += lsz) for (u64 i = lid; i < 256; i += lsz)
{ {
l_crc32tab[i] = crc32tab[i]; l_crc32tab[i] = crc32tab[i];
} }
barrier (CLK_LOCAL_MEM_FENCE); SYNC_THREADS();
__global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hash.data; LOCAL_VK u32 l_data[MAX_LOCAL];
LOCAL_AS u32 l_data[MAX_LOCAL];
for (u64 i = lid; i < MAX_LOCAL; i += lsz) for (u64 i = lid; i < MAX_LOCAL; i += lsz)
{ {
l_data[i] = data_ptr[i]; l_data[i] = esalt_bufs[digests_offset].hash.data[i];
} }
barrier (CLK_LOCAL_MEM_FENCE); SYNC_THREADS();
if (gid >= gid_max) return; if (gid >= gid_max) return;
@ -428,7 +427,7 @@ __kernel void m17210_sxx (KERN_ATTR_ESALT (pkzip_t))
} }
} }
__kernel void m17210_mxx (KERN_ATTR_ESALT (pkzip_t)) KERNEL_FQ void m17210_mxx (KERN_ATTR_ESALT (pkzip_t))
{ {
/** /**
* modifier * modifier
@ -442,25 +441,23 @@ __kernel void m17210_mxx (KERN_ATTR_ESALT (pkzip_t))
* sbox, kbox * sbox, kbox
*/ */
LOCAL_AS u32 l_crc32tab[256]; LOCAL_VK u32 l_crc32tab[256];
for (u64 i = lid; i < 256; i += lsz) for (u64 i = lid; i < 256; i += lsz)
{ {
l_crc32tab[i] = crc32tab[i]; l_crc32tab[i] = crc32tab[i];
} }
barrier (CLK_LOCAL_MEM_FENCE); SYNC_THREADS();
__global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hash.data; LOCAL_VK u32 l_data[MAX_LOCAL];
LOCAL_AS u32 l_data[MAX_LOCAL];
for (u64 i = lid; i < MAX_LOCAL; i += lsz) for (u64 i = lid; i < MAX_LOCAL; i += lsz)
{ {
l_data[i] = data_ptr[i]; l_data[i] = esalt_bufs[digests_offset].hash.data[i];
} }
barrier (CLK_LOCAL_MEM_FENCE); SYNC_THREADS();
if (gid >= gid_max) return; if (gid >= gid_max) return;

31
OpenCL/m17210_a3-pure.cl
View File

@ -87,6 +87,7 @@ Related publication: https://scitepress.org/PublicationsDetail.aspx?ID=KLPzPqStp
#include "inc_vendor.h" #include "inc_vendor.h"
#include "inc_types.h" #include "inc_types.h"
#include "inc_platform.cl"
#include "inc_common.cl" #include "inc_common.cl"
#include "inc_simd.cl" #include "inc_simd.cl"
@ -129,7 +130,7 @@ struct pkzip_hash
u32 data_length; u32 data_length;
u16 checksum_from_crc; u16 checksum_from_crc;
u16 checksum_from_timestamp; u16 checksum_from_timestamp;
u8 data[MAX_DATA]; u32 data[MAX_DATA];
} __attribute__((packed)); } __attribute__((packed));
@ -217,7 +218,7 @@ CONSTANT_AS u32a crc32tab[256] =
0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d 0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d
}; };
__kernel void m17210_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t)) KERNEL_FQ void m17210_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
{ {
/** /**
* modifier * modifier
@ -231,25 +232,23 @@ __kernel void m17210_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
* sbox, kbox * sbox, kbox
*/ */
LOCAL_AS u32 l_crc32tab[256]; LOCAL_VK u32 l_crc32tab[256];
for (u64 i = lid; i < 256; i += lsz) for (u64 i = lid; i < 256; i += lsz)
{ {
l_crc32tab[i] = crc32tab[i]; l_crc32tab[i] = crc32tab[i];
} }
barrier (CLK_LOCAL_MEM_FENCE); SYNC_THREADS();
__global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hash.data; LOCAL_VK u32 l_data[MAX_LOCAL];
LOCAL_AS u32 l_data[MAX_LOCAL];
for (u64 i = lid; i < MAX_LOCAL; i += lsz) for (u64 i = lid; i < MAX_LOCAL; i += lsz)
{ {
l_data[i] = data_ptr[i]; l_data[i] = esalt_bufs[digests_offset].hash.data[i];
} }
barrier (CLK_LOCAL_MEM_FENCE); SYNC_THREADS();
if (gid >= gid_max) return; if (gid >= gid_max) return;
@ -440,7 +439,7 @@ __kernel void m17210_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
} }
} }
__kernel void m17210_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t)) KERNEL_FQ void m17210_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
{ {
/** /**
* modifier * modifier
@ -454,25 +453,23 @@ __kernel void m17210_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
* sbox, kbox * sbox, kbox
*/ */
LOCAL_AS u32 l_crc32tab[256]; LOCAL_VK u32 l_crc32tab[256];
for (u64 i = lid; i < 256; i += lsz) for (u64 i = lid; i < 256; i += lsz)
{ {
l_crc32tab[i] = crc32tab[i]; l_crc32tab[i] = crc32tab[i];
} }
barrier (CLK_LOCAL_MEM_FENCE); SYNC_THREADS();
__global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hash.data; LOCAL_VK u32 l_data[MAX_LOCAL];
LOCAL_AS u32 l_data[MAX_LOCAL];
for (u64 i = lid; i < MAX_LOCAL; i += lsz) for (u64 i = lid; i < MAX_LOCAL; i += lsz)
{ {
l_data[i] = data_ptr[i]; l_data[i] = esalt_bufs[digests_offset].hash.data[i];
} }
barrier (CLK_LOCAL_MEM_FENCE); SYNC_THREADS();
if (gid >= gid_max) return; if (gid >= gid_max) return;