Commit 6c241b71 authored by Ondrej Mosnáček's avatar Ondrej Mosnáček

Add Argon2id functionality

parent 2c579e5e
......@@ -28,8 +28,9 @@ typedef ptrdiff_t intptr_t;
#endif
#endif /* __OPENCL_VERSION__ */
#define ARGON2_D 0
#define ARGON2_I 1
#define ARGON2_D 0
#define ARGON2_I 1
#define ARGON2_ID 2
#define ARGON2_VERSION_10 0x10
#define ARGON2_VERSION_13 0x13
......@@ -179,7 +180,7 @@ void fill_block_xor(__global const struct block_g *restrict ref_block,
}
#endif
#if ARGON2_TYPE == ARGON2_I
#if ARGON2_TYPE == ARGON2_I || ARGON2_TYPE == ARGON2_ID
void next_addresses(uint thread_input,
__local struct block_l *restrict addr,
__local struct block_l *restrict tmp,
......@@ -235,14 +236,12 @@ __kernel void argon2_kernel_segment(
/* select job's memory region: */
memory += job_id * lanes * lane_blocks;
__local struct block_l local_curr, local_prev;
__local struct block_l local_curr, local_prev, local_addr;
__local struct block_l *curr = &local_curr;
__local struct block_l *prev = &local_prev;
#if ARGON2_TYPE == ARGON2_I
__local struct block_l local_addr;
uint thread_input;
#if ARGON2_TYPE == ARGON2_I || ARGON2_TYPE == ARGON2_ID
switch (thread) {
case 0:
thread_input = pass;
......@@ -260,7 +259,7 @@ __kernel void argon2_kernel_segment(
thread_input = passes;
break;
case 5:
thread_input = ARGON2_I;
thread_input = ARGON2_TYPE;
break;
default:
thread_input = 0;
......@@ -302,24 +301,25 @@ __kernel void argon2_kernel_segment(
for (uint offset = start_offset; offset < segment_blocks; ++offset) {
uint pseudo_rand_lo, pseudo_rand_hi;
#if ARGON2_TYPE == ARGON2_I
uint addr_index = offset % ARGON2_QWORDS_IN_BLOCK;
if (addr_index == 0) {
if (thread == 6) {
++thread_input;
if (ARGON2_TYPE == ARGON2_I || (ARGON2_TYPE == ARGON2_ID && pass == 0 &&
slice < ARGON2_SYNC_POINTS / 2)) {
uint addr_index = offset % ARGON2_QWORDS_IN_BLOCK;
if (addr_index == 0) {
if (thread == 6) {
++thread_input;
}
next_addresses(thread_input, &local_addr, curr, thread);
}
next_addresses(thread_input, &local_addr, curr, thread);
uint addr_index_x = addr_index % 16;
uint addr_index_y = addr_index / 16;
addr_index = addr_index_y * 16 +
(addr_index_x + (addr_index_y / 2) * 4) % 16;
pseudo_rand_lo = local_addr.lo[addr_index];
pseudo_rand_hi = local_addr.hi[addr_index];
} else {
pseudo_rand_lo = prev->lo[0];
pseudo_rand_hi = prev->hi[0];
}
uint addr_index_x = addr_index % 16;
uint addr_index_y = addr_index / 16;
addr_index = addr_index_y * 16 +
(addr_index_x + (addr_index_y / 2) * 4) % 16;
pseudo_rand_lo = local_addr.lo[addr_index];
pseudo_rand_hi = local_addr.hi[addr_index];
#else
pseudo_rand_lo = prev->lo[0];
pseudo_rand_hi = prev->hi[0];
#endif
uint ref_lane = pseudo_rand_hi % lanes;
......@@ -384,7 +384,7 @@ __kernel void argon2_kernel_segment(
}
}
#if ARGON2_TYPE == ARGON2_I
#if ARGON2_TYPE == ARGON2_I || ARGON2_TYPE == ARGON2_ID
#define SHARED_BLOCKS 3
#else
#define SHARED_BLOCKS 2
......@@ -407,10 +407,12 @@ __kernel void argon2_kernel_oneshot(
__local struct block_l *restrict curr = &shared[0];
__local struct block_l *restrict prev = &shared[1];
#if ARGON2_TYPE == ARGON2_I
__local struct block_l *restrict addr = &shared[2];
__local struct block_l *restrict addr;
uint thread_input;
#if ARGON2_TYPE == ARGON2_I || ARGON2_TYPE == ARGON2_ID
addr = &shared[2];
switch (thread) {
case 1:
thread_input = lane;
......@@ -422,7 +424,7 @@ __kernel void argon2_kernel_oneshot(
thread_input = passes;
break;
case 5:
thread_input = ARGON2_I;
thread_input = ARGON2_TYPE;
break;
default:
thread_input = 0;
......@@ -457,24 +459,25 @@ __kernel void argon2_kernel_oneshot(
}
uint pseudo_rand_lo, pseudo_rand_hi;
#if ARGON2_TYPE == ARGON2_I
uint addr_index = offset % ARGON2_QWORDS_IN_BLOCK;
if (addr_index == 0) {
if (thread == 6) {
++thread_input;
if (ARGON2_TYPE == ARGON2_I || (ARGON2_TYPE == ARGON2_ID &&
pass == 0 && slice < ARGON2_SYNC_POINTS / 2)) {
uint addr_index = offset % ARGON2_QWORDS_IN_BLOCK;
if (addr_index == 0) {
if (thread == 6) {
++thread_input;
}
next_addresses(thread_input, addr, curr, thread);
}
next_addresses(thread_input, addr, curr, thread);
uint addr_index_x = addr_index % 16;
uint addr_index_y = addr_index / 16;
addr_index = addr_index_y * 16 +
(addr_index_x + (addr_index_y / 2) * 4) % 16;
pseudo_rand_lo = addr->lo[addr_index];
pseudo_rand_hi = addr->hi[addr_index];
} else {
pseudo_rand_lo = prev->lo[0];
pseudo_rand_hi = prev->hi[0];
}
uint addr_index_x = addr_index % 16;
uint addr_index_y = addr_index / 16;
addr_index = addr_index_y * 16 +
(addr_index_x + (addr_index_y / 2) * 4) % 16;
pseudo_rand_lo = addr->lo[addr_index];
pseudo_rand_hi = addr->hi[addr_index];
#else
pseudo_rand_lo = prev->lo[0];
pseudo_rand_hi = prev->hi[0];
#endif
uint ref_lane = pseudo_rand_hi % lanes;
......@@ -539,7 +542,7 @@ __kernel void argon2_kernel_oneshot(
}
barrier(CLK_GLOBAL_MEM_FENCE);
#if ARGON2_TYPE == ARGON2_I
#if ARGON2_TYPE == ARGON2_I || ARGON2_TYPE == ARGON2_ID
if (thread == 2) {
++thread_input;
}
......
......@@ -13,6 +13,7 @@ enum {
enum Type {
ARGON2_D = 0,
ARGON2_I = 1,
ARGON2_ID = 2,
};
enum Version {
......
This diff is collapsed.
......@@ -35,7 +35,7 @@ ProcessingUnit::ProcessingUnit(
kernel.setArg<cl_uint>(3, params->getSegmentBlocks());
} else {
auto localMemSize = (std::size_t)lanes * ARGON2_BLOCK_SIZE;
if (programContext->getArgon2Type() == ARGON2_I) {
if (programContext->getArgon2Type() != ARGON2_D) {
localMemSize *= 3;
} else {
localMemSize *= 2;
......
......@@ -60,7 +60,7 @@ static CommandLineParser<Arguments> buildCmdLineParser()
new ArgumentOption<Arguments>(
[] (Arguments &state, const std::string &type) { state.type = type; },
"type", 't', "Argon2 type (i|d)", "i", "TYPE"),
"type", 't', "Argon2 type (i|d|id)", "i", "TYPE"),
new ArgumentOption<Arguments>(
[] (Arguments &state, const std::string &type) { state.version = type; },
"version", 'v', "Argon2 version (1.0|1.3)", "1.3", "VERSION"),
......@@ -120,6 +120,8 @@ int main(int, const char * const *argv)
type = argon2::ARGON2_I;
} else if (args.type == "d") {
type = argon2::ARGON2_D;
} else if (args.type == "id") {
type = argon2::ARGON2_ID;
} else {
std::cerr << argv[0] << ": Invalid Argon2 type!" << std::endl;
return 1;
......
......@@ -72,9 +72,15 @@ std::size_t runTests(const GlobalContext &global, const Device &device,
Type type, Version version,
const TestCase *casesFrom, const TestCase *casesTo)
{
std::cout << "Running tests for Argon2"
<< (type == ARGON2_I ? "i" : "d")
<< " v" << (version == ARGON2_VERSION_10 ? "1.0" : "1.3")
std::cout << "Running tests for Argon2";
if (type == ARGON2_I) {
std::cout << "i";
} else if (type == ARGON2_D) {
std::cout << "d";
} else if (type == ARGON2_ID) {
std::cout << "id";
}
std::cout << " v" << (version == ARGON2_VERSION_10 ? "1.0" : "1.3")
<< "..." << std::endl;
std::size_t failures = 0;
......@@ -83,7 +89,7 @@ std::size_t runTests(const GlobalContext &global, const Device &device,
const std::array<bool, 2> precomputeOpts = { false, true };
auto precBegin = precomputeOpts.begin();
auto precEnd = precomputeOpts.end();
if (type != ARGON2_I) {
if (type == ARGON2_D) {
precEnd--;
}
for (auto precIt = precBegin; precIt != precEnd; precIt++) {
......@@ -302,6 +308,28 @@ const TestCase CASES_I_13[] = {
},
};
const TestCase CASES_D_10[] = {
{
{
32,
"\x02\x02\x02\x02\x02\x02\x02\x02"
"\x02\x02\x02\x02\x02\x02\x02\x02", 16,
"\x03\x03\x03\x03\x03\x03\x03\x03", 8,
"\x04\x04\x04\x04\x04\x04\x04\x04"
"\x04\x04\x04\x04", 12,
3, 32, 4
},
"\x96\xa9\xd4\xe5\xa1\x73\x40\x92"
"\xc8\x5e\x29\xf4\x10\xa4\x59\x14"
"\xa5\xdd\x1f\x5c\xbf\x08\xb2\x67"
"\x0d\xa6\x8a\x02\x85\xab\xf3\x2b",
"\x01\x01\x01\x01\x01\x01\x01\x01"
"\x01\x01\x01\x01\x01\x01\x01\x01"
"\x01\x01\x01\x01\x01\x01\x01\x01"
"\x01\x01\x01\x01\x01\x01\x01\x01", 32
},
};
const TestCase CASES_D_13[] = {
{
{
......@@ -324,6 +352,50 @@ const TestCase CASES_D_13[] = {
},
};
const TestCase CASES_ID_10[] = {
{
{
32,
"\x02\x02\x02\x02\x02\x02\x02\x02"
"\x02\x02\x02\x02\x02\x02\x02\x02", 16,
"\x03\x03\x03\x03\x03\x03\x03\x03", 8,
"\x04\x04\x04\x04\x04\x04\x04\x04"
"\x04\x04\x04\x04", 12,
3, 32, 4
},
"\xb6\x46\x15\xf0\x77\x89\xb6\x6b"
"\x64\x5b\x67\xee\x9e\xd3\xb3\x77"
"\xae\x35\x0b\x6b\xfc\xbb\x0f\xc9"
"\x51\x41\xea\x8f\x32\x26\x13\xc0",
"\x01\x01\x01\x01\x01\x01\x01\x01"
"\x01\x01\x01\x01\x01\x01\x01\x01"
"\x01\x01\x01\x01\x01\x01\x01\x01"
"\x01\x01\x01\x01\x01\x01\x01\x01", 32
},
};
const TestCase CASES_ID_13[] = {
{
{
32,
"\x02\x02\x02\x02\x02\x02\x02\x02"
"\x02\x02\x02\x02\x02\x02\x02\x02", 16,
"\x03\x03\x03\x03\x03\x03\x03\x03", 8,
"\x04\x04\x04\x04\x04\x04\x04\x04"
"\x04\x04\x04\x04", 12,
3, 32, 4
},
"\x0d\x64\x0d\xf5\x8d\x78\x76\x6c"
"\x08\xc0\x37\xa3\x4a\x8b\x53\xc9"
"\xd0\x1e\xf0\x45\x2d\x75\xb6\x5e"
"\xb5\x25\x20\xe9\x6b\x01\xe6\x59",
"\x01\x01\x01\x01\x01\x01\x01\x01"
"\x01\x01\x01\x01\x01\x01\x01\x01"
"\x01\x01\x01\x01\x01\x01\x01\x01"
"\x01\x01\x01\x01\x01\x01\x01\x01", 32
},
};
#define ARRAY_SIZE(a) (sizeof(a) / sizeof((a)[0]))
#define ARRAY_BEGIN(a) (a)
#define ARRAY_END(a) ((a) + ARRAY_SIZE(a))
......@@ -342,9 +414,18 @@ void runAllTests(std::size_t &failures)
failures += runTests<Device, GlobalContext, ProgramContext, ProcessingUnit>
(global, device, ARGON2_I, ARGON2_VERSION_13,
ARRAY_BEGIN(CASES_I_13), ARRAY_END(CASES_I_13));
failures += runTests<Device, GlobalContext, ProgramContext, ProcessingUnit>
(global, device, ARGON2_D, ARGON2_VERSION_10,
ARRAY_BEGIN(CASES_D_10), ARRAY_END(CASES_D_10));
failures += runTests<Device, GlobalContext, ProgramContext, ProcessingUnit>
(global, device, ARGON2_D, ARGON2_VERSION_13,
ARRAY_BEGIN(CASES_D_13), ARRAY_END(CASES_D_13));
failures += runTests<Device, GlobalContext, ProgramContext, ProcessingUnit>
(global, device, ARGON2_ID, ARGON2_VERSION_10,
ARRAY_BEGIN(CASES_ID_10), ARRAY_END(CASES_ID_10));
failures += runTests<Device, GlobalContext, ProgramContext, ProcessingUnit>
(global, device, ARGON2_ID, ARGON2_VERSION_13,
ARRAY_BEGIN(CASES_ID_13), ARRAY_END(CASES_ID_13));
}
#include "argon2-opencl/processingunit.h"
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment