Prepare for VeraCrypt integration

This commit is contained in:
Jens Steube
2016-05-19 20:53:17 +02:00
parent a5b85d51b9
commit d37b6c6c30
12 changed files with 77 additions and 53 deletions

View File

@@ -656,6 +656,8 @@ __kernel void m06211_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
const u32 signature = esalt_bufs[0].signature;
u32 tmp[4];
{
@@ -666,7 +668,7 @@ __kernel void m06211_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -682,7 +684,7 @@ __kernel void m06211_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -698,7 +700,7 @@ __kernel void m06211_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);

View File

@@ -656,6 +656,8 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
const u32 signature = esalt_bufs[0].signature;
u32 tmp[4];
{
@@ -666,7 +668,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -682,7 +684,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -698,7 +700,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -737,7 +739,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey2, ukey4, tmp, tmp);
twofish256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -754,7 +756,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey2, ukey4, tmp, tmp);
aes256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -771,7 +773,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey2, ukey4, tmp, tmp);
serpent256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);

View File

@@ -656,6 +656,8 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
const u32 signature = esalt_bufs[0].signature;
u32 tmp[4];
{
@@ -666,7 +668,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -682,7 +684,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -698,7 +700,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -737,7 +739,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey2, ukey4, tmp, tmp);
twofish256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -754,7 +756,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey2, ukey4, tmp, tmp);
aes256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -771,7 +773,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey2, ukey4, tmp, tmp);
serpent256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -811,7 +813,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey2, ukey5, tmp, tmp);
serpent256_decrypt_xts (ukey1, ukey4, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -829,7 +831,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey2, ukey5, tmp, tmp);
aes256_decrypt_xts (ukey1, ukey4, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);

View File

@@ -564,6 +564,8 @@ __kernel void m06221_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
const u32 signature = esalt_bufs[0].signature;
u32 tmp[4];
{
@@ -574,7 +576,7 @@ __kernel void m06221_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -590,7 +592,7 @@ __kernel void m06221_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -606,7 +608,7 @@ __kernel void m06221_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);

View File

@@ -564,6 +564,8 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
const u32 signature = esalt_bufs[0].signature;
u32 tmp[4];
{
@@ -574,7 +576,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -590,7 +592,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -606,7 +608,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -645,7 +647,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey2, ukey4, tmp, tmp);
twofish256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -662,7 +664,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey2, ukey4, tmp, tmp);
aes256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -679,7 +681,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey2, ukey4, tmp, tmp);
serpent256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);

View File

@@ -564,6 +564,8 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
const u32 signature = esalt_bufs[0].signature;
u32 tmp[4];
{
@@ -574,7 +576,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -590,7 +592,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -606,7 +608,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -645,7 +647,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey2, ukey4, tmp, tmp);
twofish256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -662,7 +664,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey2, ukey4, tmp, tmp);
aes256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -679,7 +681,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey2, ukey4, tmp, tmp);
serpent256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -719,7 +721,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey2, ukey5, tmp, tmp);
serpent256_decrypt_xts (ukey1, ukey4, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -737,7 +739,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey2, ukey5, tmp, tmp);
aes256_decrypt_xts (ukey1, ukey4, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);

View File

@@ -2201,6 +2201,8 @@ __kernel void m06231_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
const u32 signature = esalt_bufs[0].signature;
u32 tmp[4];
{
@@ -2211,7 +2213,7 @@ __kernel void m06231_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -2227,7 +2229,7 @@ __kernel void m06231_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -2243,7 +2245,7 @@ __kernel void m06231_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);

View File

@@ -1970,6 +1970,8 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
const u32 signature = esalt_bufs[0].signature;
u32 tmp[4];
{
@@ -1980,7 +1982,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -1996,7 +1998,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -2012,7 +2014,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -2051,7 +2053,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey2, ukey4, tmp, tmp);
twofish256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -2068,7 +2070,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey2, ukey4, tmp, tmp);
aes256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -2085,7 +2087,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey2, ukey4, tmp, tmp);
serpent256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);

View File

@@ -1970,6 +1970,8 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
const u32 signature = esalt_bufs[0].signature;
u32 tmp[4];
{
@@ -1980,7 +1982,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -1996,7 +1998,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -2012,7 +2014,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -2051,7 +2053,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
aes256_decrypt_xts (ukey2, ukey4, tmp, tmp);
twofish256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -2068,7 +2070,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
serpent256_decrypt_xts (ukey2, ukey4, tmp, tmp);
aes256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -2085,7 +2087,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey2, ukey4, tmp, tmp);
serpent256_decrypt_xts (ukey1, ukey3, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -2125,7 +2127,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey2, ukey5, tmp, tmp);
serpent256_decrypt_xts (ukey1, ukey4, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);
@@ -2143,7 +2145,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
twofish256_decrypt_xts (ukey2, ukey5, tmp, tmp);
aes256_decrypt_xts (ukey1, ukey4, tmp, tmp);
if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
{
mark_hash (plains_buf, hashes_shown, 0, gid, 0);

View File

@@ -912,6 +912,7 @@ typedef struct
u32 salt_buf[16];
u32 data_buf[112];
u32 keyfile_buf[16];
u32 signature;
} tc_t;