mirror of
https://github.com/hashcat/hashcat.git
synced 2025-07-18 20:49:19 +00:00
53186f0a7f
Updated all scrypt based plugins. Added condition in OpenCL/inc_hash_scrypt.cl to use local memory in case of HIP platform for a 10% speed boost.
34 lines
1.7 KiB
C
34 lines
1.7 KiB
C
/**
|
|
* Author......: See docs/credits.txt
|
|
* License.....: MIT
|
|
*/
|
|
|
|
#ifndef INC_HASH_SCRYPT_H
|
|
#define INC_HASH_SCRYPT_H
|
|
|
|
#define GET_SCRYPT_CNT(r,p) (2 * (r) * 16 * (p))
|
|
#define GET_SMIX_CNT(r,N) (2 * (r) * 16 * (N))
|
|
#define GET_STATE_CNT(r) (2 * (r) * 16)
|
|
|
|
#define SCRYPT_CNT GET_SCRYPT_CNT (SCRYPT_R, SCRYPT_P)
|
|
#define SCRYPT_CNT4 (SCRYPT_CNT / 4)
|
|
#define STATE_CNT GET_STATE_CNT (SCRYPT_R)
|
|
#define STATE_CNT4 (STATE_CNT / 4)
|
|
|
|
#define Coord(xd4,y,z) (((xd4) * ySIZE * zSIZE) + ((y) * zSIZE) + (z))
|
|
#define CO Coord(xd4,y,z)
|
|
|
|
#if defined IS_CUDA
|
|
inline __device__ uint4 operator & (const uint4 a, const u32 b) { return make_uint4 ((a.x & b ), (a.y & b ), (a.z & b ), (a.w & b )); }
|
|
inline __device__ uint4 operator << (const uint4 a, const u32 b) { return make_uint4 ((a.x << b ), (a.y << b ), (a.z << b ), (a.w << b )); }
|
|
inline __device__ uint4 operator >> (const uint4 a, const u32 b) { return make_uint4 ((a.x >> b ), (a.y >> b ), (a.z >> b ), (a.w >> b )); }
|
|
inline __device__ uint4 operator + (const uint4 a, const uint4 b) { return make_uint4 ((a.x + b.x), (a.y + b.y), (a.z + b.z), (a.w + b.w)); }
|
|
inline __device__ uint4 operator ^ (const uint4 a, const uint4 b) { return make_uint4 ((a.x ^ b.x), (a.y ^ b.y), (a.z ^ b.z), (a.w ^ b.w)); }
|
|
inline __device__ uint4 operator | (const uint4 a, const uint4 b) { return make_uint4 ((a.x | b.x), (a.y | b.y), (a.z | b.z), (a.w | b.w)); }
|
|
inline __device__ void operator ^= ( uint4 &a, const uint4 b) { a.x ^= b.x; a.y ^= b.y; a.z ^= b.z; a.w ^= b.w; }
|
|
|
|
inline __device__ uint4 rotate (const uint4 a, const int n) { return ((a << n) | ((a >> (32 - n)))); }
|
|
#endif
|
|
|
|
#endif
|