projects
/
hashcat.git
/ blobdiff
commit
grep
author
committer
pickaxe
?
search:
re
summary
|
shortlog
|
log
|
commit
|
commitdiff
|
tree
raw
|
inline
| side by side
Prepare for a more dynamic #pragma unroll use
[hashcat.git]
/
OpenCL
/
m08900.cl
diff --git
a/OpenCL/m08900.cl
b/OpenCL/m08900.cl
index
5f90dc9
..
85ad68d
100644
(file)
--- a/
OpenCL/m08900.cl
+++ b/
OpenCL/m08900.cl
@@
-110,7
+110,9
@@
void sha256_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const
ROUND_STEP (0);
ROUND_STEP (0);
+ #ifdef _unroll
#pragma unroll
#pragma unroll
+ #endif
for (int i = 16; i < 64; i += 16)
{
ROUND_EXPAND (); ROUND_STEP (i);
for (int i = 16; i < 64; i += 16)
{
ROUND_EXPAND (); ROUND_STEP (i);
@@
-705,7
+707,9
@@
void scrypt_smix (uint4 *X, uint4 *T, const u32 phy, __global uint4 *V)
const u32 x = gid % xSIZE;
const u32 x = gid % xSIZE;
+ #ifdef _unroll
#pragma unroll
#pragma unroll
+ #endif
for (u32 i = 0; i < STATE_CNT4; i += 4)
{
T[0] = (uint4) (X[i + 0].x, X[i + 1].y, X[i + 2].z, X[i + 3].w);
for (u32 i = 0; i < STATE_CNT4; i += 4)
{
T[0] = (uint4) (X[i + 0].x, X[i + 1].y, X[i + 2].z, X[i + 3].w);
@@
-743,7
+747,9
@@
void scrypt_smix (uint4 *X, uint4 *T, const u32 phy, __global uint4 *V)
salsa_r (X);
}
salsa_r (X);
}
+ #ifdef _unroll
#pragma unroll
#pragma unroll
+ #endif
for (u32 i = 0; i < STATE_CNT4; i += 4)
{
T[0] = (uint4) (X[i + 0].x, X[i + 3].y, X[i + 2].z, X[i + 1].w);
for (u32 i = 0; i < STATE_CNT4; i += 4)
{
T[0] = (uint4) (X[i + 0].x, X[i + 3].y, X[i + 2].z, X[i + 1].w);
@@
-910,12
+916,16
@@
__kernel void m08900_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf
uint4 X[STATE_CNT4];
uint4 T[STATE_CNT4];
uint4 X[STATE_CNT4];
uint4 T[STATE_CNT4];
+ #ifdef _unroll
#pragma unroll
#pragma unroll
+ #endif
for (int z = 0; z < STATE_CNT4; z++) X[z] = swap32_4 (tmps[gid].P[z]);
scrypt_smix (X, T, scrypt_phy, d_scryptV_buf);
for (int z = 0; z < STATE_CNT4; z++) X[z] = swap32_4 (tmps[gid].P[z]);
scrypt_smix (X, T, scrypt_phy, d_scryptV_buf);
+ #ifdef _unroll
#pragma unroll
#pragma unroll
+ #endif
for (int z = 0; z < STATE_CNT4; z++) tmps[gid].P[z] = swap32_4 (X[z]);
#if SCRYPT_P >= 1
for (int z = 0; z < STATE_CNT4; z++) tmps[gid].P[z] = swap32_4 (X[z]);
#if SCRYPT_P >= 1