arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-07-25 07:58:16 +00:00
Code Issues Releases Wiki Activity

using simd_shuffle on Apple Metal for argon2

Browse Source
This commit is contained in:
Gabriele Gristina 2025-07-09 08:16:00 +02:00
parent 5210ccd50d
commit adbf9d175b
No known key found for this signature in database
GPG Key ID: 9F68B59298F311F0
1 changed files with 8 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
OpenCL/inc_hash_argon2.h
View File

@ -83,17 +83,19 @@ DECLSPEC u64 hc__shfl (MAYBE_UNUSED LOCAL_AS u64 *shuffle_buf, const u64 var, co
#endif #endif
#elif defined IS_METAL #elif defined IS_METAL
#define hc__shfl_sync(shfbuf,mask,var,srcLane,argon2_thread,argon2_lsz) hc__shfl ((shfbuf),(var),(srcLane),(argon2_thread),(argon2_lsz)) #define hc__shfl_sync(shfbuf,mask,var,srcLane,argon2_thread,argon2_lsz) simd_shuffle_64 ((var),(srcLane),(argon2_lsz))
DECLSPEC u64 hc__shfl (LOCAL_AS u64 *shuffle_buf, const u64 var, const int src_lane, const u32 argon2_thread, const u32 argon2_lsz) DECLSPEC u64 simd_shuffle_64 (const u64 var, const int src_lane, const u32 argon2_lsz)
{ {
shuffle_buf[argon2_thread] = var; const u32 idx = src_lane & (argon2_lsz - 1);
SYNC_THREADS(); const u32 l32 = l32_from_64_S (var);
const u32 h32 = h32_from_64_S (var);
const u64 out = shuffle_buf[src_lane & (argon2_lsz - 1)]; u32 l32r = simd_shuffle (l32, idx);
u32 h32r = simd_shuffle (h32, idx);
SYNC_THREADS(); const u64 out = hl32_to_64_S (h32r, l32r);
return out; return out;
} }