arno
/
hashcat
mirror of https://github.com/hashcat/hashcat.git
1
0
Fork 0
Code Issues Releases Wiki Activity

Prepare OpenCL kernels for non-static compilation

Browse Source
pull/1971/head
jsteube 5 years ago
parent 5f5468be6f
commit c9d60c079f
766 changed files with 3830 additions and 2074 deletions
Show Stats Download Patch File Download Diff File

2
OpenCL/amp_a0.cl
Unescape View File

@ -3,10 +3,12 @@
* License.....: MIT
*/
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_rp.h"
#include "inc_rp.cl"
#endif
__kernel void amp (__global pw_t * restrict pws, __global pw_t * restrict pws_amp, __constant const kernel_rule_t * restrict rules_buf, __global const pw_t * restrict combs_buf, __global const bf_t * restrict bfs_buf, const u32 combs_mode, const u64 gid_max)
{

2
OpenCL/amp_a1.cl
Unescape View File

@ -3,9 +3,11 @@
* License.....: MIT
*/
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#endif
__kernel void amp (__global pw_t * restrict pws, __global pw_t * restrict pws_amp, __global const kernel_rule_t * restrict rules_buf, __global const pw_t * restrict combs_buf, __global const bf_t * restrict bfs_buf, const u32 combs_mode, const u64 gid_max)
{

2
OpenCL/amp_a3.cl
Unescape View File

@ -3,8 +3,10 @@
* License.....: MIT
*/
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#endif
__kernel void amp (__global pw_t * restrict pws, __global pw_t * restrict pws_amp, __global const kernel_rule_t * restrict rules_buf, __global const pw_t * restrict combs_buf, __constant const bf_t * restrict bfs_buf, const u32 combs_mode, const u64 gid_max)
{

2
OpenCL/m00000_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00000_m04 (KERN_ATTR_RULES ())
{

2
OpenCL/m00000_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00000_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00000_a1-optimized.cl
Unescape View File

@ -5,12 +5,14 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00000_m04 (KERN_ATTR_BASIC ())
{

2
OpenCL/m00000_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00000_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00000_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
#define MD5_STEP_REV(f,a,b,c,d,x,t,s) \
{ \

2
OpenCL/m00000_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00000_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00010_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00010_m04 (KERN_ATTR_RULES ())
{

2
OpenCL/m00010_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00010_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00010_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00010_m04 (KERN_ATTR_BASIC ())
{

2
OpenCL/m00010_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00010_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00010_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
#define MD5_STEP_REV(f,a,b,c,d,x,t,s) \
{ \

2
OpenCL/m00010_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00010_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00020_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00020_m04 (KERN_ATTR_RULES ())
{

2
OpenCL/m00020_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00020_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00020_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00020_m04 (KERN_ATTR_BASIC ())
{

2
OpenCL/m00020_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00020_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00020_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
DECLSPEC void m00020m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_BASIC ())
{

2
OpenCL/m00020_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00020_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00030_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00030_m04 (KERN_ATTR_RULES ())
{

2
OpenCL/m00030_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00030_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00030_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00030_m04 (KERN_ATTR_BASIC ())
{

2
OpenCL/m00030_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00030_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00030_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
#define MD5_STEP_REV(f,a,b,c,d,x,t,s) \
{ \

2
OpenCL/m00030_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00030_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00040_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00040_m04 (KERN_ATTR_RULES ())
{

2
OpenCL/m00040_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00040_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00040_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00040_m04 (KERN_ATTR_BASIC ())
{

2
OpenCL/m00040_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00040_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00040_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
DECLSPEC void m00040m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_BASIC ())
{

2
OpenCL/m00040_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00040_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00050_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
DECLSPEC void hmac_md5_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
{

2
OpenCL/m00050_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00050_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00050_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
DECLSPEC void hmac_md5_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
{

2
OpenCL/m00050_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00050_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00050_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
DECLSPEC void hmac_md5_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
{

2
OpenCL/m00050_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00050_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00060_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
DECLSPEC void hmac_md5_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
{

2
OpenCL/m00060_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00060_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00060_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
DECLSPEC void hmac_md5_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
{

2
OpenCL/m00060_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00060_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00060_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
DECLSPEC void hmac_md5_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
{

2
OpenCL/m00060_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
__kernel void m00060_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00100_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00100_m04 (KERN_ATTR_RULES ())
{

2
OpenCL/m00100_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00100_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00100_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00100_m04 (KERN_ATTR_BASIC ())
{

2
OpenCL/m00100_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00100_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00100_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
DECLSPEC void m00100m (u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
{

2
OpenCL/m00100_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00100_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00110_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00110_m04 (KERN_ATTR_RULES ())
{

2
OpenCL/m00110_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00110_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00110_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00110_m04 (KERN_ATTR_BASIC ())
{

2
OpenCL/m00110_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00110_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00110_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
DECLSPEC void m00110m (u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
{

2
OpenCL/m00110_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00110_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00120_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00120_m04 (KERN_ATTR_RULES ())
{

2
OpenCL/m00120_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00120_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00120_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00120_m04 (KERN_ATTR_BASIC ())
{

2
OpenCL/m00120_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00120_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00120_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
DECLSPEC void m00120m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_BASIC ())
{

2
OpenCL/m00120_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00120_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00130_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00130_m04 (KERN_ATTR_RULES ())
{

2
OpenCL/m00130_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00130_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00130_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00130_m04 (KERN_ATTR_BASIC ())
{

2
OpenCL/m00130_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00130_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00130_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
DECLSPEC void m00130m (u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
{

2
OpenCL/m00130_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00130_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00140_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00140_m04 (KERN_ATTR_RULES ())
{

2
OpenCL/m00140_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00140_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00140_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00140_m04 (KERN_ATTR_BASIC ())
{

2
OpenCL/m00140_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00140_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00140_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
DECLSPEC void m00140m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_BASIC ())
{

2
OpenCL/m00140_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00140_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00150_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
DECLSPEC void hmac_sha1_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
{

2
OpenCL/m00150_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00150_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00150_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
DECLSPEC void hmac_sha1_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
{

2
OpenCL/m00150_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00150_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00150_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
DECLSPEC void hmac_sha1_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
{

2
OpenCL/m00150_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00150_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00160_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
DECLSPEC void hmac_sha1_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
{

2
OpenCL/m00160_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00160_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00160_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
DECLSPEC void hmac_sha1_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
{

2
OpenCL/m00160_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00160_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00160_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
DECLSPEC void hmac_sha1_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
{

2
OpenCL/m00160_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00160_mxx (KERN_ATTR_VECTOR ())
{

2
OpenCL/m00200_a0-optimized.cl
Unescape View File

@ -6,12 +6,14 @@
//incompatible
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_rp_optimized.h"
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#endif
__kernel void m00200_m04 (KERN_ATTR_RULES ())
{

2
OpenCL/m00200_a1-optimized.cl
Unescape View File

@ -6,10 +6,12 @@
//incompatible
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#endif
__kernel void m00200_m04 (KERN_ATTR_BASIC ())
{

2
OpenCL/m00200_a3-optimized.cl
Unescape View File

@ -5,10 +5,12 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#endif
#define ROUND(v) \
{ \

2
OpenCL/m00300_a0-optimized.cl
Unescape View File

@ -5,6 +5,7 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp_optimized.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00300_m04 (KERN_ATTR_RULES ())
{

2
OpenCL/m00300_a0-pure.cl
Unescape View File

@ -5,6 +5,7 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
@ -12,6 +13,7 @@
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00300_mxx (KERN_ATTR_RULES ())
{

2
OpenCL/m00300_a1-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00300_m04 (KERN_ATTR_BASIC ())
{

2
OpenCL/m00300_a1-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_scalar.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00300_mxx (KERN_ATTR_BASIC ())
{

2
OpenCL/m00300_a3-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
DECLSPEC void m00300m (u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
{

2
OpenCL/m00300_a3-pure.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#endif
__kernel void m00300_mxx (KERN_ATTR_VECTOR ())
{

4
OpenCL/m00400-optimized.cl
Unescape View File

@ -5,11 +5,13 @@
#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
@ -230,5 +232,7 @@ __kernel void m00400_comp (KERN_ATTR_TMPS (phpass_tmp_t))
#define il_pos 0
#ifdef KERNEL_STATIC
#include COMPARE_M
#endif
}

4
OpenCL/m00400-pure.cl
Unescape View File

@ -5,11 +5,13 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"
#endif
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
@ -180,5 +182,7 @@ __kernel void m00400_comp (KERN_ATTR_TMPS (phpass_tmp_t))
#define il_pos 0
#ifdef KERNEL_STATIC
#include COMPARE_M
#endif
}

4
OpenCL/m00500-optimized.cl
Unescape View File

@ -5,10 +5,12 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_hash_md5.cl"
#endif
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
@ -1012,5 +1014,7 @@ __kernel void m00500_comp (KERN_ATTR_TMPS (md5crypt_tmp_t))
#define il_pos 0
#ifdef KERNEL_STATIC
#include COMPARE_M
#endif
}

4
OpenCL/m00500-pure.cl
Unescape View File

@ -5,10 +5,12 @@
//#define NEW_SIMD_CODE
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_common.cl"
#include "inc_hash_md5.cl"
#endif
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
@ -245,5 +247,7 @@ __kernel void m00500_comp (KERN_ATTR_TMPS (md5crypt_tmp_t))
#define il_pos 0
#ifdef KERNEL_STATIC
#include COMPARE_M
#endif
}

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

Write Preview
Loading…
Cancel
Save