From 0e1a942421955d270a863b763f192436a53bbd47 Mon Sep 17 00:00:00 2001 From: HD Moore Date: Sun, 5 Jul 2026 12:08:01 -0500 Subject: [PATCH] feat: add RAKP MD5/SHA256 format support --- doc/NEWS | 5 + run/opencl/rakp_md5_kernel.cl | 39 +++ run/opencl/rakp_sha256_kernel.cl | 39 +++ src/opencl_rakp_md5_fmt_plug.c | 384 +++++++++++++++++++++ src/opencl_rakp_sha256_fmt_plug.c | 384 +++++++++++++++++++++ src/rakp_md5_fmt_plug.c | 534 ++++++++++++++++++++++++++++++ src/rakp_sha256_fmt_plug.c | 517 +++++++++++++++++++++++++++++ 7 files changed, 1902 insertions(+) create mode 100644 run/opencl/rakp_md5_kernel.cl create mode 100644 run/opencl/rakp_sha256_kernel.cl create mode 100644 src/opencl_rakp_md5_fmt_plug.c create mode 100644 src/opencl_rakp_sha256_fmt_plug.c create mode 100644 src/rakp_md5_fmt_plug.c create mode 100644 src/rakp_sha256_fmt_plug.c diff --git a/doc/NEWS b/doc/NEWS index a6fda198c55..62399bd7f26 100644 --- a/doc/NEWS +++ b/doc/NEWS @@ -488,6 +488,11 @@ Major changes from 1.9.0-jumbo-1 (May 2019) in this bleeding-edge version: - Add support to keepass2john for the current format of KeePass XML keyfiles. The formats needed no changes. [magnum; 2025] +- Added RAKP-MD5 and RAKP-SHA256 CPU and OpenCL formats for IPMI 2.0 RAKP + (RMCP+) challenge-response hashes in the existing "$rakp$$" + syntax, where digest length selects MD5 (32), SHA1 (40), or SHA256 (64). + [HD Moore; 2026] + Major changes from 1.8.0-jumbo-1 (December 2014) to 1.9.0-jumbo-1 (May 2019): diff --git a/run/opencl/rakp_md5_kernel.cl b/run/opencl/rakp_md5_kernel.cl new file mode 100644 index 00000000000..9d61283eb6f --- /dev/null +++ b/run/opencl/rakp_md5_kernel.cl @@ -0,0 +1,39 @@ +/* + * OpenCL kernel for IPMI 2.0 RAKP (RMCP+) HMAC-MD5 (RAKP auth algorithm 2). + * + * Copyright (c) 2026 HD Moore, released to the general public under the + * following terms: + * Redistribution and use in source and binary forms, with or without + * modification, are permitted. + * + * One work-item per candidate. Keys are packed (key_array + index[gid]); the + * RAKP salt (raw bytes + length) is the HMAC message. Uses the shared + * opencl_hmac_md5 helper so the HMAC construction stays correct and simple. + */ + +#include "opencl_misc.h" +#define HMAC_KEY_TYPE __global const +#include "opencl_hmac_md5.h" + +__kernel void rakp_md5_kernel(__global const uchar *salt, + const uint salt_len, + __global const uint *key_array, + __global const uint *index, + __global uint *digest) +{ + uint gid = get_global_id(0); + uint base = index[gid]; + uint key_len = base & 63; + __global const uchar *key = (__global const uchar *)(key_array + (base >> 6)); + uchar saltbuf[128]; + uint out[4]; + uint i; + + for (i = 0; i < salt_len; i++) + saltbuf[i] = salt[i]; + + hmac_md5(key, key_len, saltbuf, salt_len, out, 16); + + for (i = 0; i < 4; i++) + digest[gid * 4 + i] = out[i]; +} diff --git a/run/opencl/rakp_sha256_kernel.cl b/run/opencl/rakp_sha256_kernel.cl new file mode 100644 index 00000000000..817973f87df --- /dev/null +++ b/run/opencl/rakp_sha256_kernel.cl @@ -0,0 +1,39 @@ +/* + * OpenCL kernel for IPMI 2.0 RAKP (RMCP+) HMAC-SHA256 (RAKP auth algorithm 3). + * + * Copyright (c) 2026 HD Moore, released to the general public under the + * following terms: + * Redistribution and use in source and binary forms, with or without + * modification, are permitted. + * + * One work-item per candidate. Keys are packed (key_array + index[gid]); the + * RAKP salt (raw bytes + length) is the HMAC message. Uses the shared + * opencl_hmac_sha256 helper so the HMAC construction stays correct and simple. + */ + +#include "opencl_misc.h" +#define HMAC_KEY_TYPE __global const +#include "opencl_hmac_sha256.h" + +__kernel void rakp_sha256_kernel(__global const uchar *salt, + const uint salt_len, + __global const uint *key_array, + __global const uint *index, + __global uint *digest) +{ + uint gid = get_global_id(0); + uint base = index[gid]; + uint key_len = base & 63; + __global const uchar *key = (__global const uchar *)(key_array + (base >> 6)); + uchar saltbuf[128]; + uint out[8]; + uint i; + + for (i = 0; i < salt_len; i++) + saltbuf[i] = salt[i]; + + hmac_sha256(key, key_len, saltbuf, salt_len, out, 32); + + for (i = 0; i < 8; i++) + digest[gid * 8 + i] = out[i]; +} diff --git a/src/opencl_rakp_md5_fmt_plug.c b/src/opencl_rakp_md5_fmt_plug.c new file mode 100644 index 00000000000..e4d7e450ed6 --- /dev/null +++ b/src/opencl_rakp_md5_fmt_plug.c @@ -0,0 +1,384 @@ +/* + * OpenCL IPMI 2.0 RAKP (RMCP+) HMAC-MD5 (RAKP auth algorithm 2). + * + * Host based on opencl_rakp_fmt_plug.c (HMAC-SHA1 RAKP, (C) 2013 Harrison Neal, + * magnum, Dhiru Kholia). Copyright (c) 2026 HD Moore. Released under GPLv2 like + * the original. Scalar one-work-item-per-candidate kernel that calls the shared + * hmac_md5 helper; the raw RAKP salt (binary) + length are passed to the GPU. + */ + +#ifdef HAVE_OPENCL + +#if FMT_EXTERNS_H +extern struct fmt_main fmt_opencl_rakp_md5; +#elif FMT_REGISTERS_H +john_register_one(&fmt_opencl_rakp_md5); +#else + +#include +#include + +#include "path.h" +#include "arch.h" +#include "misc.h" +#include "common.h" +#include "formats.h" +#include "johnswap.h" +#include "opencl_common.h" +#include "options.h" + +#define FORMAT_LABEL "RAKP-MD5-opencl" +#define FORMAT_NAME "IPMI 2.0 RAKP (RMCP+)" +#define ALGORITHM_NAME "HMAC-MD5 OpenCL" + +#define BENCHMARK_COMMENT "" +#define BENCHMARK_LENGTH 7 + +#define BLOCK_SIZE 64 +#define SALT_STORAGE_SIZE (BLOCK_SIZE * 2) +#define SALT_SIZE (SALT_STORAGE_SIZE - 9) +#define SALT_MIN_SIZE (SALT_SIZE - BLOCK_SIZE + 1) + +#define PLAINTEXT_LENGTH (BLOCK_SIZE - 1) /* idx & 63 */ +#define BUFFER_SIZE ((PLAINTEXT_LENGTH + 63) / 64 * 64) + +#define BINARY_SIZE 16 + +#define MIN_KEYS_PER_CRYPT 1 +#define MAX_KEYS_PER_CRYPT 1 + +#define FORMAT_TAG "$rakp$" +#define TAG_LENGTH (sizeof(FORMAT_TAG) - 1) + +#define BINARY_ALIGN sizeof(uint32_t) +#define SALT_ALIGN sizeof(uint32_t) + +#define STEP 0 +#define SEED 65536 +#define ROUNDS 5 + +static const char * warn[] = { + "pass xfer: ", ", index xfer: ", ", crypt: ", ", result xfer: " +}; + +typedef struct { + uint32_t length; + unsigned char salt[SALT_STORAGE_SIZE]; +} rakp_salt_t; + +static cl_uint salt_len_arg; + +static cl_mem salt_buffer, keys_buffer, idx_buffer, digest_buffer; + +static int new_keys; +static unsigned int *keys; +static uint32_t *idx; +static unsigned char *digest; +static unsigned int key_idx = 0; +static struct fmt_main *self; + +#include "opencl_autotune.h" + +static struct fmt_tests tests[] = { + {"$rakp$000476044b237ebb4f2a415ce488ebc91996043c8011e26079633b706424119e09dcaad4acf21b10535000000000000000000000000000001404726f6f74$102546a8888f78d420fc48ba84108ee2", "superuser"}, + {"$rakp$000aae1770480e004b31afa4a6e13297b09fc5d60b4a018b4fc50ecf9a8963745bf6a56a6873f50f000000000000000000000000000000001400$ef1f3c6c029975343a4d1195dee130da", "admin"}, + {"$rakp$001108876d5e010001e176e46f686d5e296145b5d6454181da9a9e2f0e9eba82b7c6c1c83dcf3a5a000000000000000000000000000000001400$3ea2e6b997bc34959ef1617fd4be1348", "admin"}, + {NULL} +}; + +static size_t get_task_max_work_group_size() +{ + return autotune_get_task_max_work_group_size(FALSE, 0, crypt_kernel); +} + +static int valid(char *ciphertext, struct fmt_main *self) +{ + char *p, *q = NULL; + int len; + + p = ciphertext; + if (!strncmp(p, FORMAT_TAG, TAG_LENGTH)) + p += TAG_LENGTH; + + q = strrchr(ciphertext, '$'); + if (!q) + return 0; + q = q + 1; + if ((q - p - 1) > SALT_SIZE * 2) + return 0; + if ((q - p - 1) < SALT_MIN_SIZE * 2) + return 0; + + len = strspn(q, HEXCHARS_lc); + if (len != BINARY_SIZE * 2 || len != strlen(q)) + return 0; + if (strspn(p, HEXCHARS_lc) != q - p - 1) + return 0; + + return 1; +} + +static void release_clobj(void); + +static void create_clobj(size_t gws, struct fmt_main *self) +{ + release_clobj(); + + keys = mem_alloc((PLAINTEXT_LENGTH + 1) * gws); + idx = mem_calloc(gws, sizeof(*idx)); + digest = mem_alloc(gws * BINARY_SIZE); + + salt_buffer = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, SALT_STORAGE_SIZE, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating salt_buffer"); + keys_buffer = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, (PLAINTEXT_LENGTH + 1) * gws, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating keys_buffer"); + idx_buffer = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, 4 * gws, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating idx_buffer"); + digest_buffer = clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, BINARY_SIZE * gws, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating digest_buffer"); + + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 0, sizeof(salt_buffer), &salt_buffer), "Error arg 0"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(cl_uint), &salt_len_arg), "Error arg 1"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(keys_buffer), &keys_buffer), "Error arg 2"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 3, sizeof(idx_buffer), &idx_buffer), "Error arg 3"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 4, sizeof(digest_buffer), &digest_buffer), "Error arg 4"); +} + +static void release_clobj(void) +{ + if (keys) { + HANDLE_CLERROR(clReleaseMemObject(digest_buffer), "Release digest_buffer"); + HANDLE_CLERROR(clReleaseMemObject(idx_buffer), "Release idx_buffer"); + HANDLE_CLERROR(clReleaseMemObject(keys_buffer), "Release keys_buffer"); + HANDLE_CLERROR(clReleaseMemObject(salt_buffer), "Release salt_buffer"); + MEM_FREE(digest); + MEM_FREE(idx); + MEM_FREE(keys); + } +} + +static void done(void) +{ + if (program[gpu_id]) { + release_clobj(); + HANDLE_CLERROR(clReleaseKernel(crypt_kernel), "Release kernel"); + HANDLE_CLERROR(clReleaseProgram(program[gpu_id]), "Release program"); + program[gpu_id] = NULL; + } +} + +static void init(struct fmt_main *_self) +{ + self = _self; + opencl_prepare_dev(gpu_id); +} + +static void reset(struct db_main *db) +{ + if (!program[gpu_id]) { + opencl_init("$JOHN/opencl/rakp_md5_kernel.cl", gpu_id, NULL); + crypt_kernel = clCreateKernel(program[gpu_id], "rakp_md5_kernel", &ret_code); + HANDLE_CLERROR(ret_code, "Error creating kernel"); + } + + size_t gws_limit = MIN((1 << 26) * 4 / BUFFER_SIZE, + get_max_mem_alloc_size(gpu_id) / BUFFER_SIZE); + + opencl_init_auto_setup(SEED, 0, NULL, warn, 2, self, + create_clobj, release_clobj, BUFFER_SIZE, gws_limit, db); + autotune_run(self, ROUNDS, gws_limit, 200); +} + +static void clear_keys(void) +{ + key_idx = 0; +} + +static void set_key(char *key, int index) +{ + const unsigned int *key32 = (unsigned int*)key; + int len = strlen(key); + + idx[index] = (key_idx << 6) | len; + while (len > 4) { + keys[key_idx++] = *key32++; + len -= 4; + } + if (len) + keys[key_idx++] = *key32 & (0xffffffffU >> (32 - (len << 3))); + new_keys = 1; +} + +static char *get_key(int index) +{ + static char out[PLAINTEXT_LENGTH + 1]; + int i, len = idx[index] & 63; + char *key = (char*)&keys[idx[index] >> 6]; + + for (i = 0; i < len; i++) + out[i] = key[i]; + out[i] = 0; + return out; +} + +static void *get_binary(char *ciphertext) +{ + static union { + unsigned char c[BINARY_SIZE]; + uint32_t dummy; + } buf; + unsigned char *out = buf.c; + char *p; + int i; + + p = strrchr(ciphertext, '$') + 1; + for (i = 0; i < BINARY_SIZE; i++) { + out[i] = (atoi16[ARCH_INDEX(*p)] << 4) | atoi16[ARCH_INDEX(p[1])]; + p += 2; + } + return out; +} + +static void *get_salt(char *ciphertext) +{ + static rakp_salt_t out; + char *p; + unsigned int i, len; + + memset(&out, 0, sizeof(out)); + p = ciphertext; + if (!strncmp(p, FORMAT_TAG, TAG_LENGTH)) + p += TAG_LENGTH; + len = (strrchr(p, '$') - p) / 2; + for (i = 0; i < len; i++) + out.salt[i] = (atoi16[ARCH_INDEX(p[2 * i])] << 4) | atoi16[ARCH_INDEX(p[2 * i + 1])]; + out.length = len; + return &out; +} + +static void set_salt(void *salt) +{ + rakp_salt_t *s = salt; + + salt_len_arg = s->length; + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(cl_uint), &salt_len_arg), "Error setting salt_len"); + HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], salt_buffer, CL_FALSE, 0, s->length, s->salt, 0, NULL, NULL), + "Error updating salt_buffer"); + HANDLE_CLERROR(clFlush(queue[gpu_id]), "clFlush"); +} + +static int cmp_all(void *binary, int count) +{ + int index; + + for (index = 0; index < count; index++) + if (((uint32_t*)binary)[0] == ((uint32_t*)(digest + index * BINARY_SIZE))[0]) + return 1; + return 0; +} + +static int cmp_one(void *binary, int index) +{ + return !memcmp(binary, digest + index * BINARY_SIZE, BINARY_SIZE); +} + +static int cmp_exact(char *source, int index) +{ + return 1; +} + +static int crypt_all(int *pcount, struct db_salt *salt) +{ + const int count = *pcount; + size_t *lws = local_work_size ? &local_work_size : NULL; + + global_work_size = GET_KPC_MULTIPLE(count, local_work_size); + + if (new_keys && key_idx) { + BENCH_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], keys_buffer, CL_FALSE, 0, 4 * key_idx, keys, 0, NULL, multi_profilingEvent[0]), + "Error updating keys_buffer"); + BENCH_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], idx_buffer, CL_FALSE, 0, 4 * global_work_size, idx, 0, NULL, multi_profilingEvent[1]), + "Error updating idx_buffer"); + new_keys = 0; + } + + BENCH_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1, NULL, &global_work_size, lws, 0, NULL, multi_profilingEvent[2]), + "Error running kernel"); + BENCH_CLERROR(clEnqueueReadBuffer(queue[gpu_id], digest_buffer, CL_TRUE, 0, BINARY_SIZE * global_work_size, digest, 0, NULL, multi_profilingEvent[3]), + "Error reading digest_buffer"); + + return count; +} + +#define OCL_GET_HASH(index) (((uint32_t*)(digest + (index) * BINARY_SIZE))[0]) +static int get_hash_0(int index) { return OCL_GET_HASH(index) & PH_MASK_0; } +static int get_hash_1(int index) { return OCL_GET_HASH(index) & PH_MASK_1; } +static int get_hash_2(int index) { return OCL_GET_HASH(index) & PH_MASK_2; } +static int get_hash_3(int index) { return OCL_GET_HASH(index) & PH_MASK_3; } +static int get_hash_4(int index) { return OCL_GET_HASH(index) & PH_MASK_4; } +static int get_hash_5(int index) { return OCL_GET_HASH(index) & PH_MASK_5; } +static int get_hash_6(int index) { return OCL_GET_HASH(index) & PH_MASK_6; } + +struct fmt_main fmt_opencl_rakp_md5 = { + { + FORMAT_LABEL, + FORMAT_NAME, + ALGORITHM_NAME, + BENCHMARK_COMMENT, + BENCHMARK_LENGTH, + 0, + PLAINTEXT_LENGTH, + BINARY_SIZE, + BINARY_ALIGN, + sizeof(rakp_salt_t), + SALT_ALIGN, + MIN_KEYS_PER_CRYPT, + MAX_KEYS_PER_CRYPT, + FMT_CASE | FMT_8_BIT, + { NULL }, + { FORMAT_TAG }, + tests + }, { + init, + done, + reset, + fmt_default_prepare, + valid, + fmt_default_split, + get_binary, + get_salt, + { NULL }, + fmt_default_source, + { + fmt_default_binary_hash_0, + fmt_default_binary_hash_1, + fmt_default_binary_hash_2, + fmt_default_binary_hash_3, + fmt_default_binary_hash_4, + fmt_default_binary_hash_5, + fmt_default_binary_hash_6 + }, + fmt_default_salt_hash, + NULL, + set_salt, + set_key, + get_key, + clear_keys, + crypt_all, + { + get_hash_0, + get_hash_1, + get_hash_2, + get_hash_3, + get_hash_4, + get_hash_5, + get_hash_6 + }, + cmp_all, + cmp_one, + cmp_exact + } +}; + +#endif /* plugin stanza */ +#endif /* HAVE_OPENCL */ diff --git a/src/opencl_rakp_sha256_fmt_plug.c b/src/opencl_rakp_sha256_fmt_plug.c new file mode 100644 index 00000000000..06daa3ef09d --- /dev/null +++ b/src/opencl_rakp_sha256_fmt_plug.c @@ -0,0 +1,384 @@ +/* + * OpenCL IPMI 2.0 RAKP (RMCP+) HMAC-SHA256 (RAKP auth algorithm 3). + * + * Host based on opencl_rakp_fmt_plug.c (HMAC-SHA1 RAKP, (C) 2013 Harrison Neal, + * magnum, Dhiru Kholia). Copyright (c) 2026 HD Moore. Released under GPLv2 like + * the original. Scalar one-work-item-per-candidate kernel that calls the shared + * hmac_sha256 helper; the raw RAKP salt (binary) + length are passed to the GPU. + */ + +#ifdef HAVE_OPENCL + +#if FMT_EXTERNS_H +extern struct fmt_main fmt_opencl_rakp_sha256; +#elif FMT_REGISTERS_H +john_register_one(&fmt_opencl_rakp_sha256); +#else + +#include +#include + +#include "path.h" +#include "arch.h" +#include "misc.h" +#include "common.h" +#include "formats.h" +#include "johnswap.h" +#include "opencl_common.h" +#include "options.h" + +#define FORMAT_LABEL "RAKP-SHA256-opencl" +#define FORMAT_NAME "IPMI 2.0 RAKP (RMCP+)" +#define ALGORITHM_NAME "HMAC-SHA256 OpenCL" + +#define BENCHMARK_COMMENT "" +#define BENCHMARK_LENGTH 7 + +#define BLOCK_SIZE 64 +#define SALT_STORAGE_SIZE (BLOCK_SIZE * 2) +#define SALT_SIZE (SALT_STORAGE_SIZE - 9) +#define SALT_MIN_SIZE (SALT_SIZE - BLOCK_SIZE + 1) + +#define PLAINTEXT_LENGTH (BLOCK_SIZE - 1) /* idx & 63 */ +#define BUFFER_SIZE ((PLAINTEXT_LENGTH + 63) / 64 * 64) + +#define BINARY_SIZE 32 + +#define MIN_KEYS_PER_CRYPT 1 +#define MAX_KEYS_PER_CRYPT 1 + +#define FORMAT_TAG "$rakp$" +#define TAG_LENGTH (sizeof(FORMAT_TAG) - 1) + +#define BINARY_ALIGN sizeof(uint32_t) +#define SALT_ALIGN sizeof(uint32_t) + +#define STEP 0 +#define SEED 65536 +#define ROUNDS 5 + +static const char * warn[] = { + "pass xfer: ", ", index xfer: ", ", crypt: ", ", result xfer: " +}; + +typedef struct { + uint32_t length; + unsigned char salt[SALT_STORAGE_SIZE]; +} rakp_salt_t; + +static cl_uint salt_len_arg; + +static cl_mem salt_buffer, keys_buffer, idx_buffer, digest_buffer; + +static int new_keys; +static unsigned int *keys; +static uint32_t *idx; +static unsigned char *digest; +static unsigned int key_idx = 0; +static struct fmt_main *self; + +#include "opencl_autotune.h" + +static struct fmt_tests tests[] = { + {"$rakp$1eaf72c001407792e65364f6100997e2946b8d9c2a57a4a8574b99694ca321f30401769c58cfbcd09cc2c43a2c74de030010debf6022a771140474657374$669c87a9a6871f77134f3159c9f421a571723f5493fdeec173309b362ee8b259", "1234"}, + {"$rakp$41987b483ceaad20ffa718e58a99c6ae7ade09383b49f50a70c44febde6e2c7665896457e8b583f10102030405060708090a0b0c0d0e0f101404726f6f74$6a7ad41787afe8831a759c98d65b12f0f52726bb1b020819c29a695376410997", "0penBmc"}, + {"$rakp$4b85d9c83828f0ce031267bc45d72f928f50185296bf7b8fb62eacbe91d2667ee63b5a507f79c758000000000000d7030010debf80adac6b140561646d696e$374bb6a0cc22dc88087fd3638f73584905f7c69fd8ed57decbe93e71c688d107", "admin"}, + {NULL} +}; + +static size_t get_task_max_work_group_size() +{ + return autotune_get_task_max_work_group_size(FALSE, 0, crypt_kernel); +} + +static int valid(char *ciphertext, struct fmt_main *self) +{ + char *p, *q = NULL; + int len; + + p = ciphertext; + if (!strncmp(p, FORMAT_TAG, TAG_LENGTH)) + p += TAG_LENGTH; + + q = strrchr(ciphertext, '$'); + if (!q) + return 0; + q = q + 1; + if ((q - p - 1) > SALT_SIZE * 2) + return 0; + if ((q - p - 1) < SALT_MIN_SIZE * 2) + return 0; + + len = strspn(q, HEXCHARS_lc); + if (len != BINARY_SIZE * 2 || len != strlen(q)) + return 0; + if (strspn(p, HEXCHARS_lc) != q - p - 1) + return 0; + + return 1; +} + +static void release_clobj(void); + +static void create_clobj(size_t gws, struct fmt_main *self) +{ + release_clobj(); + + keys = mem_alloc((PLAINTEXT_LENGTH + 1) * gws); + idx = mem_calloc(gws, sizeof(*idx)); + digest = mem_alloc(gws * BINARY_SIZE); + + salt_buffer = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, SALT_STORAGE_SIZE, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating salt_buffer"); + keys_buffer = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, (PLAINTEXT_LENGTH + 1) * gws, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating keys_buffer"); + idx_buffer = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, 4 * gws, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating idx_buffer"); + digest_buffer = clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, BINARY_SIZE * gws, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating digest_buffer"); + + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 0, sizeof(salt_buffer), &salt_buffer), "Error arg 0"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(cl_uint), &salt_len_arg), "Error arg 1"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(keys_buffer), &keys_buffer), "Error arg 2"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 3, sizeof(idx_buffer), &idx_buffer), "Error arg 3"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 4, sizeof(digest_buffer), &digest_buffer), "Error arg 4"); +} + +static void release_clobj(void) +{ + if (keys) { + HANDLE_CLERROR(clReleaseMemObject(digest_buffer), "Release digest_buffer"); + HANDLE_CLERROR(clReleaseMemObject(idx_buffer), "Release idx_buffer"); + HANDLE_CLERROR(clReleaseMemObject(keys_buffer), "Release keys_buffer"); + HANDLE_CLERROR(clReleaseMemObject(salt_buffer), "Release salt_buffer"); + MEM_FREE(digest); + MEM_FREE(idx); + MEM_FREE(keys); + } +} + +static void done(void) +{ + if (program[gpu_id]) { + release_clobj(); + HANDLE_CLERROR(clReleaseKernel(crypt_kernel), "Release kernel"); + HANDLE_CLERROR(clReleaseProgram(program[gpu_id]), "Release program"); + program[gpu_id] = NULL; + } +} + +static void init(struct fmt_main *_self) +{ + self = _self; + opencl_prepare_dev(gpu_id); +} + +static void reset(struct db_main *db) +{ + if (!program[gpu_id]) { + opencl_init("$JOHN/opencl/rakp_sha256_kernel.cl", gpu_id, NULL); + crypt_kernel = clCreateKernel(program[gpu_id], "rakp_sha256_kernel", &ret_code); + HANDLE_CLERROR(ret_code, "Error creating kernel"); + } + + size_t gws_limit = MIN((1 << 26) * 4 / BUFFER_SIZE, + get_max_mem_alloc_size(gpu_id) / BUFFER_SIZE); + + opencl_init_auto_setup(SEED, 0, NULL, warn, 2, self, + create_clobj, release_clobj, BUFFER_SIZE, gws_limit, db); + autotune_run(self, ROUNDS, gws_limit, 200); +} + +static void clear_keys(void) +{ + key_idx = 0; +} + +static void set_key(char *key, int index) +{ + const unsigned int *key32 = (unsigned int*)key; + int len = strlen(key); + + idx[index] = (key_idx << 6) | len; + while (len > 4) { + keys[key_idx++] = *key32++; + len -= 4; + } + if (len) + keys[key_idx++] = *key32 & (0xffffffffU >> (32 - (len << 3))); + new_keys = 1; +} + +static char *get_key(int index) +{ + static char out[PLAINTEXT_LENGTH + 1]; + int i, len = idx[index] & 63; + char *key = (char*)&keys[idx[index] >> 6]; + + for (i = 0; i < len; i++) + out[i] = key[i]; + out[i] = 0; + return out; +} + +static void *get_binary(char *ciphertext) +{ + static union { + unsigned char c[BINARY_SIZE]; + uint32_t dummy; + } buf; + unsigned char *out = buf.c; + char *p; + int i; + + p = strrchr(ciphertext, '$') + 1; + for (i = 0; i < BINARY_SIZE; i++) { + out[i] = (atoi16[ARCH_INDEX(*p)] << 4) | atoi16[ARCH_INDEX(p[1])]; + p += 2; + } + return out; +} + +static void *get_salt(char *ciphertext) +{ + static rakp_salt_t out; + char *p; + unsigned int i, len; + + memset(&out, 0, sizeof(out)); + p = ciphertext; + if (!strncmp(p, FORMAT_TAG, TAG_LENGTH)) + p += TAG_LENGTH; + len = (strrchr(p, '$') - p) / 2; + for (i = 0; i < len; i++) + out.salt[i] = (atoi16[ARCH_INDEX(p[2 * i])] << 4) | atoi16[ARCH_INDEX(p[2 * i + 1])]; + out.length = len; + return &out; +} + +static void set_salt(void *salt) +{ + rakp_salt_t *s = salt; + + salt_len_arg = s->length; + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(cl_uint), &salt_len_arg), "Error setting salt_len"); + HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], salt_buffer, CL_FALSE, 0, s->length, s->salt, 0, NULL, NULL), + "Error updating salt_buffer"); + HANDLE_CLERROR(clFlush(queue[gpu_id]), "clFlush"); +} + +static int cmp_all(void *binary, int count) +{ + int index; + + for (index = 0; index < count; index++) + if (((uint32_t*)binary)[0] == ((uint32_t*)(digest + index * BINARY_SIZE))[0]) + return 1; + return 0; +} + +static int cmp_one(void *binary, int index) +{ + return !memcmp(binary, digest + index * BINARY_SIZE, BINARY_SIZE); +} + +static int cmp_exact(char *source, int index) +{ + return 1; +} + +static int crypt_all(int *pcount, struct db_salt *salt) +{ + const int count = *pcount; + size_t *lws = local_work_size ? &local_work_size : NULL; + + global_work_size = GET_KPC_MULTIPLE(count, local_work_size); + + if (new_keys && key_idx) { + BENCH_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], keys_buffer, CL_FALSE, 0, 4 * key_idx, keys, 0, NULL, multi_profilingEvent[0]), + "Error updating keys_buffer"); + BENCH_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], idx_buffer, CL_FALSE, 0, 4 * global_work_size, idx, 0, NULL, multi_profilingEvent[1]), + "Error updating idx_buffer"); + new_keys = 0; + } + + BENCH_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1, NULL, &global_work_size, lws, 0, NULL, multi_profilingEvent[2]), + "Error running kernel"); + BENCH_CLERROR(clEnqueueReadBuffer(queue[gpu_id], digest_buffer, CL_TRUE, 0, BINARY_SIZE * global_work_size, digest, 0, NULL, multi_profilingEvent[3]), + "Error reading digest_buffer"); + + return count; +} + +#define OCL_GET_HASH(index) (((uint32_t*)(digest + (index) * BINARY_SIZE))[0]) +static int get_hash_0(int index) { return OCL_GET_HASH(index) & PH_MASK_0; } +static int get_hash_1(int index) { return OCL_GET_HASH(index) & PH_MASK_1; } +static int get_hash_2(int index) { return OCL_GET_HASH(index) & PH_MASK_2; } +static int get_hash_3(int index) { return OCL_GET_HASH(index) & PH_MASK_3; } +static int get_hash_4(int index) { return OCL_GET_HASH(index) & PH_MASK_4; } +static int get_hash_5(int index) { return OCL_GET_HASH(index) & PH_MASK_5; } +static int get_hash_6(int index) { return OCL_GET_HASH(index) & PH_MASK_6; } + +struct fmt_main fmt_opencl_rakp_sha256 = { + { + FORMAT_LABEL, + FORMAT_NAME, + ALGORITHM_NAME, + BENCHMARK_COMMENT, + BENCHMARK_LENGTH, + 0, + PLAINTEXT_LENGTH, + BINARY_SIZE, + BINARY_ALIGN, + sizeof(rakp_salt_t), + SALT_ALIGN, + MIN_KEYS_PER_CRYPT, + MAX_KEYS_PER_CRYPT, + FMT_CASE | FMT_8_BIT, + { NULL }, + { FORMAT_TAG }, + tests + }, { + init, + done, + reset, + fmt_default_prepare, + valid, + fmt_default_split, + get_binary, + get_salt, + { NULL }, + fmt_default_source, + { + fmt_default_binary_hash_0, + fmt_default_binary_hash_1, + fmt_default_binary_hash_2, + fmt_default_binary_hash_3, + fmt_default_binary_hash_4, + fmt_default_binary_hash_5, + fmt_default_binary_hash_6 + }, + fmt_default_salt_hash, + NULL, + set_salt, + set_key, + get_key, + clear_keys, + crypt_all, + { + get_hash_0, + get_hash_1, + get_hash_2, + get_hash_3, + get_hash_4, + get_hash_5, + get_hash_6 + }, + cmp_all, + cmp_one, + cmp_exact + } +}; + +#endif /* plugin stanza */ +#endif /* HAVE_OPENCL */ diff --git a/src/rakp_md5_fmt_plug.c b/src/rakp_md5_fmt_plug.c new file mode 100644 index 00000000000..8f24e78e9ef --- /dev/null +++ b/src/rakp_md5_fmt_plug.c @@ -0,0 +1,534 @@ +/* + * IPMI 2.0 RAKP (RMCP+) with HMAC-MD5 key-exchange authentication code + * (RAKP auth algorithm 2). Same $rakp$$ hash format -- the + * 32-hex (16-byte) HMAC length selects this vs RAKP (40) / RAKP-SHA256 (64). + * + * Derived from rakp_fmt_plug.c (HMAC-SHA1 RAKP, (c) 2013 magnum) and the SIMD + * HMAC-MD5 engine of hmacMD5_fmt_plug.c ((c) 2010 bartavelle, (c) 2011-2015 + * magnum). Released to the general public under the same terms: redistribution + * and use in source and binary forms, with or without modification, permitted. + * + * The key difference from generic HMAC-MD5: RAKP salts are BINARY (contain NUL + * bytes), so the salt is hex-encoded in the hash and loaded by known length -- + * never null-terminated -- in both the SIMD and non-SIMD paths. + */ + +#if FMT_EXTERNS_H +extern struct fmt_main fmt_rakp_md5; +#elif FMT_REGISTERS_H +john_register_one(&fmt_rakp_md5); +#else + +#include + +#ifdef _OPENMP +#include +#endif + +#include "arch.h" +#include "misc.h" +#include "common.h" +#include "formats.h" +#include "md5.h" +#include "aligned.h" +#include "johnswap.h" +#include "simd-intrinsics.h" + +#define FORMAT_LABEL "RAKP-MD5" +#define FORMAT_NAME "IPMI 2.0 RAKP (RMCP+)" + +#ifdef SIMD_COEF_32 +#define MD5_N (SIMD_PARA_MD5 * SIMD_COEF_32) +#endif +#define ALGORITHM_NAME "HMAC-MD5 " MD5_ALGORITHM_NAME +#define BENCHMARK_COMMENT "" +#define BENCHMARK_LENGTH 7 +#define PLAINTEXT_LENGTH MAX_PLAINTEXT_LENGTH +#define PAD_SIZE 64 +#define PAD_SIZE_W (PAD_SIZE / 4) +#define BINARY_SIZE 16 +#define BINARY_ALIGN sizeof(uint32_t) +#ifdef SIMD_COEF_32 +#define SALT_LIMBS 2 /* RAKP salt <= ~74 bytes -> 2 limbs */ +#define SALT_LENGTH (SALT_LIMBS * PAD_SIZE - 9) +#define SALT_ALIGN MEM_ALIGN_SIMD +#else +#define SALT_LENGTH (2 * PAD_SIZE) +#define SALT_ALIGN MEM_ALIGN_NONE +#endif +#define SALT_MIN_SIZE (PAD_SIZE - 8) +#define SALT_MAX_SIZE (2 * PAD_SIZE - 8 - 1) +#define FORMAT_TAG "$rakp$" +#define TAG_LENGTH (sizeof(FORMAT_TAG) - 1) + +#ifndef OMP_SCALE +#define OMP_SCALE 2 +#endif + +#ifdef SIMD_COEF_32 +#define MIN_KEYS_PER_CRYPT MD5_N +#define MAX_KEYS_PER_CRYPT (MD5_N * 128) +#if ARCH_LITTLE_ENDIAN==1 +#define GETPOS(i, index) ((index & (SIMD_COEF_32 - 1)) * 4 + ((i&63) & (0xffffffff - 3)) * SIMD_COEF_32 + ((i&63) & 3) + (unsigned int)index/SIMD_COEF_32 * PAD_SIZE * SIMD_COEF_32) +#else +#define GETPOS(i, index) ((index & (SIMD_COEF_32 - 1)) * 4 + ((i&63) & (0xffffffff - 3)) * SIMD_COEF_32 + (3-((i&63)&3)) + (unsigned int)index/SIMD_COEF_32 * PAD_SIZE * SIMD_COEF_32) +#endif +#else +#define MIN_KEYS_PER_CRYPT 1 +#define MAX_KEYS_PER_CRYPT 256 +#endif + +static struct fmt_tests tests[] = { + {"$rakp$000476044b237ebb4f2a415ce488ebc91996043c8011e26079633b706424119e09dcaad4acf21b10535000000000000000000000000000001404726f6f74$102546a8888f78d420fc48ba84108ee2", "superuser"}, + {"$rakp$000aae1770480e004b31afa4a6e13297b09fc5d60b4a018b4fc50ecf9a8963745bf6a56a6873f50f000000000000000000000000000000001400$ef1f3c6c029975343a4d1195dee130da", "admin"}, + {"$rakp$001108876d5e010001e176e46f686d5e296145b5d6454181da9a9e2f0e9eba82b7c6c1c83dcf3a5a000000000000000000000000000000001400$3ea2e6b997bc34959ef1617fd4be1348", "admin"}, + {NULL} +}; + +typedef struct { + int length; + unsigned char salt[SALT_LENGTH + 1]; +} rakp_salt; + +#ifdef SIMD_COEF_32 +static unsigned char *crypt_key; +static unsigned char *ipad, *prep_ipad; +static unsigned char *opad, *prep_opad; +typedef struct cur_salt_t { + unsigned char salt[SALT_LIMBS][PAD_SIZE * MAX_KEYS_PER_CRYPT]; + int salt_len; +} cur_salt_t; +static cur_salt_t *cur_salt; +static int bufsize; +#define SALT_SIZE sizeof(cur_salt_t) +#else +static rakp_salt cur_salt; +static uint32_t (*crypt_key)[BINARY_SIZE / sizeof(uint32_t)]; +static unsigned char (*ipad)[PAD_SIZE]; +static unsigned char (*opad)[PAD_SIZE]; +static MD5_CTX *ipad_ctx; +static MD5_CTX *opad_ctx; +#define SALT_SIZE sizeof(cur_salt) +#endif +static char (*saved_plain)[PLAINTEXT_LENGTH + 1]; +static int new_keys; + +#ifdef SIMD_COEF_32 +static void clear_keys(void) +{ + memset(ipad, 0x36, bufsize); + memset(opad, 0x5C, bufsize); +} +#endif + +static void init(struct fmt_main *self) +{ +#ifdef SIMD_COEF_32 + int i; +#endif + omp_autotune(self, OMP_SCALE); + +#ifdef SIMD_COEF_32 + bufsize = sizeof(*opad) * self->params.max_keys_per_crypt * PAD_SIZE; + crypt_key = mem_calloc_align(1, bufsize, MEM_ALIGN_SIMD); + ipad = mem_calloc_align(1, bufsize, MEM_ALIGN_SIMD); + opad = mem_calloc_align(1, bufsize, MEM_ALIGN_SIMD); + prep_ipad = mem_calloc_align(self->params.max_keys_per_crypt, BINARY_SIZE, MEM_ALIGN_SIMD); + prep_opad = mem_calloc_align(self->params.max_keys_per_crypt, BINARY_SIZE, MEM_ALIGN_SIMD); + for (i = 0; i < self->params.max_keys_per_crypt; ++i) { + crypt_key[GETPOS(BINARY_SIZE, i)] = 0x80; + ((unsigned int*)crypt_key)[14 * SIMD_COEF_32 + (i&(SIMD_COEF_32-1)) + (i/SIMD_COEF_32) * PAD_SIZE_W * SIMD_COEF_32] = (BINARY_SIZE + PAD_SIZE) << 3; + } + clear_keys(); +#else + crypt_key = mem_calloc(self->params.max_keys_per_crypt, sizeof(*crypt_key)); + ipad = mem_calloc(self->params.max_keys_per_crypt, sizeof(*ipad)); + opad = mem_calloc(self->params.max_keys_per_crypt, sizeof(*opad)); + ipad_ctx = mem_calloc(self->params.max_keys_per_crypt, sizeof(*ipad_ctx)); + opad_ctx = mem_calloc(self->params.max_keys_per_crypt, sizeof(*opad_ctx)); +#endif + saved_plain = mem_calloc(self->params.max_keys_per_crypt, sizeof(*saved_plain)); +} + +static void done(void) +{ + MEM_FREE(saved_plain); +#ifdef SIMD_COEF_32 + MEM_FREE(prep_opad); + MEM_FREE(prep_ipad); +#else + MEM_FREE(opad_ctx); + MEM_FREE(ipad_ctx); +#endif + MEM_FREE(opad); + MEM_FREE(ipad); + MEM_FREE(crypt_key); +} + +static int valid(char *ciphertext, struct fmt_main *self) +{ + char *p, *q; + int len; + + p = ciphertext; + if (!strncmp(p, FORMAT_TAG, TAG_LENGTH)) + p += TAG_LENGTH; + + q = strrchr(ciphertext, '$'); + if (!q) + return 0; + q = q + 1; + if ((q - p - 1) > SALT_MAX_SIZE * 2) + return 0; + if ((q - p - 1) < SALT_MIN_SIZE * 2) + return 0; + + len = strspn(q, HEXCHARS_lc); + if (len != BINARY_SIZE * 2 || len != strlen(q)) + return 0; + if (strspn(p, HEXCHARS_lc) != q - p - 1) + return 0; + + return 1; +} + +static void set_salt(void *salt) +{ +#ifdef SIMD_COEF_32 + cur_salt = salt; +#else + memcpy(&cur_salt, salt, SALT_SIZE); +#endif +} + +static void set_key(char *key, int index) +{ + int len; +#ifdef SIMD_COEF_32 +#if ARCH_LITTLE_ENDIAN==1 + uint32_t *ipadp = (uint32_t*)&ipad[GETPOS(0, index)]; + uint32_t *opadp = (uint32_t*)&opad[GETPOS(0, index)]; +#else + uint32_t *ipadp = (uint32_t*)&ipad[GETPOS(3, index)]; + uint32_t *opadp = (uint32_t*)&opad[GETPOS(3, index)]; +#endif + const uint32_t *keyp = (uint32_t*)key; + unsigned int temp; + + len = strlen(key); + memcpy(saved_plain[index], key, len); + saved_plain[index][len] = 0; + + if (len > PAD_SIZE) { + unsigned char k0[BINARY_SIZE]; + MD5_CTX ctx; + int i; + + MD5_Init(&ctx); + MD5_Update(&ctx, key, len); + MD5_Final(k0, &ctx); + + keyp = (unsigned int*)k0; + for (i = 0; i < BINARY_SIZE / 4; i++, ipadp += SIMD_COEF_32, opadp += SIMD_COEF_32) { +#if ARCH_LITTLE_ENDIAN==1 + temp = *keyp++; +#else + temp = JOHNSWAP(*keyp++); +#endif + *ipadp ^= temp; + *opadp ^= temp; + } + } else { +#if ARCH_LITTLE_ENDIAN==1 + while((unsigned char)(temp = *keyp++)) { + if (!(temp & 0xff00) || !(temp & 0xff0000)) { + *ipadp ^= (unsigned short)temp; + *opadp ^= (unsigned short)temp; + break; + } + *ipadp ^= temp; + *opadp ^= temp; + if (!(temp & 0xff000000)) + break; + ipadp += SIMD_COEF_32; + opadp += SIMD_COEF_32; + } +#else + while((temp = *keyp++) & 0xff000000) { + if (!(temp & 0xff0000) || !(temp & 0xff00)) { + *ipadp ^= (unsigned short)JOHNSWAP(temp); + *opadp ^= (unsigned short)JOHNSWAP(temp); + break; + } + *ipadp ^= JOHNSWAP(temp); + *opadp ^= JOHNSWAP(temp); + if (!(temp & 0xff)) + break; + ipadp += SIMD_COEF_32; + opadp += SIMD_COEF_32; + } +#endif + } +#else + int i; + + len = strlen(key); + memcpy(saved_plain[index], key, len); + saved_plain[index][len] = 0; + + memset(ipad[index], 0x36, PAD_SIZE); + memset(opad[index], 0x5C, PAD_SIZE); + + if (len > PAD_SIZE) { + MD5_CTX ctx; + unsigned char k0[BINARY_SIZE]; + + MD5_Init(&ctx); + MD5_Update(&ctx, key, len); + MD5_Final(k0, &ctx); + + len = BINARY_SIZE; + for (i = 0; i < len; i++) { + ipad[index][i] ^= k0[i]; + opad[index][i] ^= k0[i]; + } + } else + for (i = 0; i < len; i++) { + ipad[index][i] ^= key[i]; + opad[index][i] ^= key[i]; + } +#endif + new_keys = 1; +} + +static char *get_key(int index) +{ + return saved_plain[index]; +} + +static int cmp_all(void *binary, int count) +{ +#ifdef SIMD_COEF_32 + unsigned int x, y; + + for (y = 0; y < (unsigned int)(count + SIMD_COEF_32 - 1) / SIMD_COEF_32; y++) + for (x = 0; x < SIMD_COEF_32; x++) + if (((uint32_t*)binary)[0] == ((uint32_t*)crypt_key)[x + y * SIMD_COEF_32 * PAD_SIZE_W]) + return 1; + return 0; +#else + int index; + + for (index = 0; index < count; index++) + if (((uint32_t*)binary)[0] == crypt_key[index][0]) + return 1; + return 0; +#endif +} + +static int cmp_one(void *binary, int index) +{ +#ifdef SIMD_COEF_32 + int i; + + for (i = 0; i < (BINARY_SIZE/4); i++) + if (((uint32_t*)binary)[i] != ((uint32_t*)crypt_key)[i * SIMD_COEF_32 + (index&(SIMD_COEF_32-1)) + (unsigned int)index/SIMD_COEF_32 * PAD_SIZE_W * SIMD_COEF_32]) + return 0; + return 1; +#else + return !memcmp(binary, crypt_key[index], BINARY_SIZE); +#endif +} + +static int cmp_exact(char *source, int index) +{ + return 1; +} + +static int crypt_all(int *pcount, struct db_salt *salt) +{ + const int count = *pcount; + int index; + +#if _OPENMP +#pragma omp parallel for +#endif + for (index = 0; index < count; index += MIN_KEYS_PER_CRYPT) { +#ifdef SIMD_COEF_32 + int i; + + if (new_keys) { + SIMDmd5body(&ipad[index * PAD_SIZE], + (unsigned int*)&prep_ipad[index * BINARY_SIZE], NULL, SSEi_MIXED_IN); + SIMDmd5body(&opad[index * PAD_SIZE], + (unsigned int*)&prep_opad[index * BINARY_SIZE], NULL, SSEi_MIXED_IN); + } + SIMDmd5body(cur_salt->salt[0], + (unsigned int*)&crypt_key[index * PAD_SIZE], + (unsigned int*)&prep_ipad[index * BINARY_SIZE], + SSEi_MIXED_IN|SSEi_RELOAD|SSEi_OUTPUT_AS_INP_FMT); + for (i = 1; i <= (cur_salt->salt_len + 8) / PAD_SIZE; i++) + SIMDmd5body(cur_salt->salt[i], + (unsigned int*)&crypt_key[index * PAD_SIZE], + (unsigned int*)&crypt_key[index * PAD_SIZE], + SSEi_MIXED_IN|SSEi_RELOAD_INP_FMT|SSEi_OUTPUT_AS_INP_FMT); + SIMDmd5body(&crypt_key[index * PAD_SIZE], + (unsigned int*)&crypt_key[index * PAD_SIZE], + (unsigned int*)&prep_opad[index * BINARY_SIZE], + SSEi_MIXED_IN|SSEi_RELOAD|SSEi_OUTPUT_AS_INP_FMT); +#else + MD5_CTX ctx; + + if (new_keys) { + MD5_Init(&ipad_ctx[index]); + MD5_Update(&ipad_ctx[index], ipad[index], PAD_SIZE); + MD5_Init(&opad_ctx[index]); + MD5_Update(&opad_ctx[index], opad[index], PAD_SIZE); + } + + memcpy(&ctx, &ipad_ctx[index], sizeof(ctx)); + MD5_Update(&ctx, cur_salt.salt, cur_salt.length); + MD5_Final((unsigned char*) crypt_key[index], &ctx); + + memcpy(&ctx, &opad_ctx[index], sizeof(ctx)); + MD5_Update(&ctx, crypt_key[index], BINARY_SIZE); + MD5_Final((unsigned char*) crypt_key[index], &ctx); +#endif + } + new_keys = 0; + + return count; +} + +static void *get_binary(char *ciphertext) +{ + static union { + unsigned char c[BINARY_SIZE]; + uint32_t dummy; + } buf; + unsigned char *out = buf.c; + char *p; + int i; + + p = strrchr(ciphertext, '$') + 1; + for (i = 0; i < BINARY_SIZE; i++) { + out[i] = (atoi16[ARCH_INDEX(*p)] << 4) | atoi16[ARCH_INDEX(p[1])]; + p += 2; + } +#if !ARCH_LITTLE_ENDIAN && defined(SIMD_COEF_32) + alter_endianity(out, BINARY_SIZE); +#endif + return out; +} + +static void *get_salt(char *ciphertext) +{ + unsigned char salt[SALT_LENGTH + 1]; + unsigned int i, len; +#ifdef SIMD_COEF_32 + unsigned int j; + static JTR_ALIGN(MEM_ALIGN_SIMD) cur_salt_t cs; +#else + static rakp_salt out; +#endif + + if (!strncmp(ciphertext, FORMAT_TAG, TAG_LENGTH)) + ciphertext += TAG_LENGTH; + + len = (strrchr(ciphertext, '$') - ciphertext) / 2; + memset(salt, 0, sizeof(salt)); + for (i = 0; i < len; i++) + salt[i] = (atoi16[ARCH_INDEX(ciphertext[2 * i])] << 4) | + atoi16[ARCH_INDEX(ciphertext[2 * i + 1])]; + +#ifdef SIMD_COEF_32 + memset(&cs, 0, sizeof(cs)); + for (i = 0; i < len; i++) + for (j = 0; j < MIN_KEYS_PER_CRYPT; ++j) + cs.salt[i / PAD_SIZE][GETPOS(i, j)] = salt[i]; + cs.salt_len = len; + for (j = 0; j < MIN_KEYS_PER_CRYPT; ++j) { + cs.salt[len / PAD_SIZE][GETPOS(len, j)] = 0x80; + ((unsigned int*)cs.salt[(len + 8) / PAD_SIZE])[14 * SIMD_COEF_32 + (j&(SIMD_COEF_32-1)) + j/SIMD_COEF_32 * PAD_SIZE_W * SIMD_COEF_32] = (len + PAD_SIZE) << 3; + } + return &cs; +#else + memset(&out, 0, sizeof(out)); + out.length = len; + memcpy(out.salt, salt, len); + return &out; +#endif +} + +/* + * crypt_key is kept in SIMD "input format" (SSEi_OUTPUT_AS_INP_FMT), so the + * per-hash stride is PAD_SIZE_W (16) words, not BINARY_SIZE/4. + */ +#define COMMON_GET_HASH_SIMD32 PAD_SIZE_W +#define COMMON_GET_HASH_VAR crypt_key +#include "common-get-hash.h" + +struct fmt_main fmt_rakp_md5 = { + { + FORMAT_LABEL, + FORMAT_NAME, + ALGORITHM_NAME, + BENCHMARK_COMMENT, + BENCHMARK_LENGTH, + 0, + PLAINTEXT_LENGTH, + BINARY_SIZE, + BINARY_ALIGN, + SALT_SIZE, + SALT_ALIGN, + MIN_KEYS_PER_CRYPT, + MAX_KEYS_PER_CRYPT, + /* RAKP salts are bounded (<= ~119 bytes) so FMT_HUGE_INPUT is not needed; + * omitting it enables the per-hash binary-hash table, which matters when + * loading large RAKP hash sets. */ + FMT_CASE | FMT_8_BIT | FMT_OMP, + { NULL }, + { NULL }, + tests + }, { + init, + done, + fmt_default_reset, + fmt_default_prepare, + valid, + fmt_default_split, + get_binary, + get_salt, + { NULL }, + fmt_default_source, + { + fmt_default_binary_hash_0, + fmt_default_binary_hash_1, + fmt_default_binary_hash_2, + fmt_default_binary_hash_3, + fmt_default_binary_hash_4, + fmt_default_binary_hash_5, + fmt_default_binary_hash_6 + }, + fmt_default_salt_hash, + NULL, + set_salt, + set_key, + get_key, +#ifdef SIMD_COEF_32 + clear_keys, +#else + fmt_default_clear_keys, +#endif + crypt_all, + { +#define COMMON_GET_HASH_LINK +#include "common-get-hash.h" + }, + cmp_all, + cmp_one, + cmp_exact + } +}; + +#endif /* plugin stanza */ diff --git a/src/rakp_sha256_fmt_plug.c b/src/rakp_sha256_fmt_plug.c new file mode 100644 index 00000000000..7faf8f250eb --- /dev/null +++ b/src/rakp_sha256_fmt_plug.c @@ -0,0 +1,517 @@ +/* + * IPMI 2.0 RAKP (RMCP+) with HMAC-SHA256 key-exchange authentication code + * (RAKP auth algorithm 3). Same $rakp$$ hash format -- the + * 64-hex (32-byte) HMAC length selects this vs RAKP (40) / RAKP-MD5 (32). + * + * Derived from rakp_fmt_plug.c (HMAC-SHA1 RAKP, (c) 2013 magnum) and the SIMD + * HMAC-SHA256 engine of hmacSHA256_fmt_plug.c ((c) 2012 magnum, SIMD 2015 JimF). + * Released to the general public under the same terms: redistribution and use + * in source and binary forms, with or without modification, permitted. + * + * RAKP salts are BINARY (contain NUL bytes), so the salt is hex-encoded in the + * hash and loaded by known length -- never null-terminated -- in both paths. + */ + +#if FMT_EXTERNS_H +extern struct fmt_main fmt_rakp_sha256; +#elif FMT_REGISTERS_H +john_register_one(&fmt_rakp_sha256); +#else + +#include + +#ifdef _OPENMP +#include +#endif + +#include "arch.h" +#include "sha2.h" +#include "misc.h" +#include "common.h" +#include "formats.h" +#include "johnswap.h" +#include "simd-intrinsics.h" + +#define FORMAT_LABEL "RAKP-SHA256" +#define FORMAT_NAME "IPMI 2.0 RAKP (RMCP+)" +#define ALGORITHM_NAME "HMAC-SHA256 " SHA256_ALGORITHM_NAME +#define BENCHMARK_COMMENT "" +#define BENCHMARK_LENGTH 7 +#define PLAINTEXT_LENGTH MAX_PLAINTEXT_LENGTH +#define PAD_SIZE 64 +#define PAD_SIZE_W (PAD_SIZE / 4) +#define BINARY_SIZE (256 / 8) +#define BINARY_ALIGN 4 +#ifdef SIMD_COEF_32 +#define SALT_LIMBS 2 /* RAKP salt <= ~74 bytes -> 2 limbs */ +#define SALT_LENGTH (SALT_LIMBS * PAD_SIZE - 9) +#define SALT_ALIGN MEM_ALIGN_SIMD +#else +#define SALT_LENGTH (2 * PAD_SIZE) +#define SALT_ALIGN MEM_ALIGN_NONE +#endif +#define SALT_MIN_SIZE (PAD_SIZE - 8) +#define SALT_MAX_SIZE (2 * PAD_SIZE - 8 - 1) +#define FORMAT_TAG "$rakp$" +#define TAG_LENGTH (sizeof(FORMAT_TAG) - 1) + +#ifndef OMP_SCALE +#define OMP_SCALE 4 +#endif + +#ifdef SIMD_COEF_32 +#define MIN_KEYS_PER_CRYPT (SIMD_COEF_32 * SIMD_PARA_SHA256) +#define MAX_KEYS_PER_CRYPT (SIMD_COEF_32 * SIMD_PARA_SHA256 * 64) +#if ARCH_LITTLE_ENDIAN==1 +#define GETPOS(i, index) ((index & (SIMD_COEF_32 - 1)) * 4 + ((i&63) & (0xffffffff - 3)) * SIMD_COEF_32 + (3 - ((i&63) & 3)) + (unsigned int)index/SIMD_COEF_32 * PAD_SIZE * SIMD_COEF_32) +#else +#define GETPOS(i, index) ((index & (SIMD_COEF_32 - 1)) * 4 + ((i&63) & (0xffffffff - 3)) * SIMD_COEF_32 + ((i&63) & 3) + (unsigned int)index/SIMD_COEF_32 * PAD_SIZE * SIMD_COEF_32) +#endif +#else +#define MIN_KEYS_PER_CRYPT 1 +#define MAX_KEYS_PER_CRYPT 128 +#endif + +static struct fmt_tests tests[] = { + {"$rakp$1eaf72c001407792e65364f6100997e2946b8d9c2a57a4a8574b99694ca321f30401769c58cfbcd09cc2c43a2c74de030010debf6022a771140474657374$669c87a9a6871f77134f3159c9f421a571723f5493fdeec173309b362ee8b259", "1234"}, + {"$rakp$41987b483ceaad20ffa718e58a99c6ae7ade09383b49f50a70c44febde6e2c7665896457e8b583f10102030405060708090a0b0c0d0e0f101404726f6f74$6a7ad41787afe8831a759c98d65b12f0f52726bb1b020819c29a695376410997", "0penBmc"}, + {"$rakp$4b85d9c83828f0ce031267bc45d72f928f50185296bf7b8fb62eacbe91d2667ee63b5a507f79c758000000000000d7030010debf80adac6b140561646d696e$374bb6a0cc22dc88087fd3638f73584905f7c69fd8ed57decbe93e71c688d107", "admin"}, + {NULL} +}; + +typedef struct { + int length; + unsigned char salt[SALT_LENGTH + 1]; +} rakp_salt; + +#ifdef SIMD_COEF_32 +static unsigned char *crypt_key; +static unsigned char *ipad, *prep_ipad; +static unsigned char *opad, *prep_opad; +typedef struct cur_salt_t { + unsigned char salt[SALT_LIMBS][PAD_SIZE * MAX_KEYS_PER_CRYPT]; + int salt_len; +} cur_salt_t; +static cur_salt_t *cur_salt; +static int bufsize; +#define SALT_SIZE sizeof(cur_salt_t) +#else +static rakp_salt cur_salt; +static uint32_t (*crypt_key)[BINARY_SIZE / sizeof(uint32_t)]; +static unsigned char (*ipad)[PAD_SIZE]; +static unsigned char (*opad)[PAD_SIZE]; +static SHA256_CTX *ipad_ctx; +static SHA256_CTX *opad_ctx; +#define SALT_SIZE sizeof(cur_salt) +#endif +static char (*saved_plain)[PLAINTEXT_LENGTH + 1]; +static int new_keys; + +#ifdef SIMD_COEF_32 +static void clear_keys(void) +{ + memset(ipad, 0x36, bufsize); + memset(opad, 0x5C, bufsize); +} +#endif + +static void init(struct fmt_main *self) +{ +#ifdef SIMD_COEF_32 + int i; +#endif + omp_autotune(self, OMP_SCALE); + +#ifdef SIMD_COEF_32 + bufsize = sizeof(*opad) * self->params.max_keys_per_crypt * PAD_SIZE; + crypt_key = mem_calloc_align(1, bufsize, MEM_ALIGN_SIMD); + ipad = mem_calloc_align(1, bufsize, MEM_ALIGN_SIMD); + opad = mem_calloc_align(1, bufsize, MEM_ALIGN_SIMD); + prep_ipad = mem_calloc_align(self->params.max_keys_per_crypt, BINARY_SIZE, MEM_ALIGN_SIMD); + prep_opad = mem_calloc_align(self->params.max_keys_per_crypt, BINARY_SIZE, MEM_ALIGN_SIMD); + for (i = 0; i < self->params.max_keys_per_crypt; ++i) { + crypt_key[GETPOS(BINARY_SIZE, i)] = 0x80; + ((unsigned int*)crypt_key)[15 * SIMD_COEF_32 + (i&(SIMD_COEF_32-1)) + (i/SIMD_COEF_32) * PAD_SIZE_W * SIMD_COEF_32] = (BINARY_SIZE + PAD_SIZE) << 3; + } + clear_keys(); +#else + crypt_key = mem_calloc(self->params.max_keys_per_crypt, sizeof(*crypt_key)); + ipad = mem_calloc(self->params.max_keys_per_crypt, sizeof(*ipad)); + opad = mem_calloc(self->params.max_keys_per_crypt, sizeof(*opad)); + ipad_ctx = mem_calloc(self->params.max_keys_per_crypt, sizeof(*ipad_ctx)); + opad_ctx = mem_calloc(self->params.max_keys_per_crypt, sizeof(*opad_ctx)); +#endif + saved_plain = mem_calloc(self->params.max_keys_per_crypt, sizeof(*saved_plain)); +} + +static void done(void) +{ + MEM_FREE(saved_plain); +#ifdef SIMD_COEF_32 + MEM_FREE(prep_opad); + MEM_FREE(prep_ipad); +#else + MEM_FREE(opad_ctx); + MEM_FREE(ipad_ctx); +#endif + MEM_FREE(opad); + MEM_FREE(ipad); + MEM_FREE(crypt_key); +} + +static int valid(char *ciphertext, struct fmt_main *self) +{ + char *p, *q; + int len; + + p = ciphertext; + if (!strncmp(p, FORMAT_TAG, TAG_LENGTH)) + p += TAG_LENGTH; + + q = strrchr(ciphertext, '$'); + if (!q) + return 0; + q = q + 1; + if ((q - p - 1) > SALT_MAX_SIZE * 2) + return 0; + if ((q - p - 1) < SALT_MIN_SIZE * 2) + return 0; + + len = strspn(q, HEXCHARS_lc); + if (len != BINARY_SIZE * 2 || len != strlen(q)) + return 0; + if (strspn(p, HEXCHARS_lc) != q - p - 1) + return 0; + + return 1; +} + +static void set_salt(void *salt) +{ +#ifdef SIMD_COEF_32 + cur_salt = salt; +#else + memcpy(&cur_salt, salt, SALT_SIZE); +#endif +} + +static void set_key(char *key, int index) +{ + int len; +#ifdef SIMD_COEF_32 +#if ARCH_LITTLE_ENDIAN==1 + uint32_t *ipadp = (uint32_t*)&ipad[GETPOS(3, index)]; + uint32_t *opadp = (uint32_t*)&opad[GETPOS(3, index)]; +#else + uint32_t *ipadp = (uint32_t*)&ipad[GETPOS(0, index)]; + uint32_t *opadp = (uint32_t*)&opad[GETPOS(0, index)]; +#endif + const uint32_t *keyp = (uint32_t*)key; + unsigned int temp; + + len = strlen(key); + memcpy(saved_plain[index], key, len); + saved_plain[index][len] = 0; + + if (len > PAD_SIZE) { + unsigned char k0[BINARY_SIZE]; + SHA256_CTX ctx; + int i; + + SHA256_Init(&ctx); + SHA256_Update(&ctx, key, len); + SHA256_Final(k0, &ctx); + + keyp = (unsigned int*)k0; + for (i = 0; i < BINARY_SIZE / 4; i++, ipadp += SIMD_COEF_32, opadp += SIMD_COEF_32) { +#if ARCH_LITTLE_ENDIAN==1 + temp = JOHNSWAP(*keyp++); +#else + temp = *keyp++; +#endif + *ipadp ^= temp; + *opadp ^= temp; + } + } else +#if ARCH_LITTLE_ENDIAN==1 + while(((temp = JOHNSWAP(*keyp++)) & 0xff000000)) { +#else + while(((temp = *keyp++) & 0xff000000)) { +#endif + if (!(temp & 0x00ff0000) || !(temp & 0x0000ff00)) { +#if ARCH_LITTLE_ENDIAN==1 + ((unsigned short*)ipadp)[1] ^= (unsigned short)(temp >> 16); + ((unsigned short*)opadp)[1] ^= (unsigned short)(temp >> 16); +#else + ((unsigned short*)ipadp)[0] ^= (unsigned short)(temp >> 16); + ((unsigned short*)opadp)[0] ^= (unsigned short)(temp >> 16); +#endif + break; + } + *ipadp ^= temp; + *opadp ^= temp; + if (!(temp & 0x000000ff)) + break; + ipadp += SIMD_COEF_32; + opadp += SIMD_COEF_32; + } +#else + int i; + + len = strlen(key); + memcpy(saved_plain[index], key, len); + saved_plain[index][len] = 0; + + memset(ipad[index], 0x36, PAD_SIZE); + memset(opad[index], 0x5C, PAD_SIZE); + + if (len > PAD_SIZE) { + SHA256_CTX ctx; + unsigned char k0[BINARY_SIZE]; + + SHA256_Init(&ctx); + SHA256_Update(&ctx, key, len); + SHA256_Final(k0, &ctx); + + len = BINARY_SIZE; + for (i = 0; i < len; i++) { + ipad[index][i] ^= k0[i]; + opad[index][i] ^= k0[i]; + } + } else + for (i = 0; i < len; i++) { + ipad[index][i] ^= key[i]; + opad[index][i] ^= key[i]; + } +#endif + new_keys = 1; +} + +static char *get_key(int index) +{ + return saved_plain[index]; +} + +static int cmp_all(void *binary, int count) +{ + unsigned int index; + +#ifdef SIMD_COEF_32 + for (index = 0; index < (unsigned int)count; index++) + if (((uint32_t*)binary)[0] == ((uint32_t*)crypt_key)[(index&(SIMD_COEF_32-1)) + index/SIMD_COEF_32 * PAD_SIZE_W * SIMD_COEF_32]) + return 1; + return 0; +#else + for (index = 0; index < (unsigned int)count; index++) + if (((uint32_t*)binary)[0] == crypt_key[index][0]) + return 1; + return 0; +#endif +} + +static int cmp_one(void *binary, int index) +{ +#ifdef SIMD_COEF_32 + int i; + + for (i = 0; i < (BINARY_SIZE/4); i++) + if (((uint32_t*)binary)[i] != ((uint32_t*)crypt_key)[i * SIMD_COEF_32 + (index&(SIMD_COEF_32-1)) + (unsigned int)index/SIMD_COEF_32 * PAD_SIZE_W * SIMD_COEF_32]) + return 0; + return 1; +#else + return !memcmp(binary, crypt_key[index], BINARY_SIZE); +#endif +} + +static int cmp_exact(char *source, int index) +{ + return 1; +} + +static int crypt_all(int *pcount, struct db_salt *salt) +{ + const int count = *pcount; + int index; + +#ifdef _OPENMP +#pragma omp parallel for +#endif + for (index = 0; index < count; index += MIN_KEYS_PER_CRYPT) { +#ifdef SIMD_COEF_32 + unsigned int i; + + if (new_keys) { + SIMDSHA256body(&ipad[index * PAD_SIZE], + (unsigned int*)&prep_ipad[index * BINARY_SIZE], NULL, SSEi_MIXED_IN); + SIMDSHA256body(&opad[index * PAD_SIZE], + (unsigned int*)&prep_opad[index * BINARY_SIZE], NULL, SSEi_MIXED_IN); + } + SIMDSHA256body(cur_salt->salt[0], + (unsigned int*)&crypt_key[index * PAD_SIZE], + (unsigned int*)&prep_ipad[index * BINARY_SIZE], + SSEi_MIXED_IN|SSEi_RELOAD|SSEi_OUTPUT_AS_INP_FMT); + for (i = 1; i <= (cur_salt->salt_len + 8) / PAD_SIZE; i++) + SIMDSHA256body(cur_salt->salt[i], + (unsigned int*)&crypt_key[index * PAD_SIZE], + (unsigned int*)&crypt_key[index * PAD_SIZE], + SSEi_MIXED_IN|SSEi_RELOAD_INP_FMT|SSEi_OUTPUT_AS_INP_FMT); + SIMDSHA256body(&crypt_key[index * PAD_SIZE], + (unsigned int*)&crypt_key[index * PAD_SIZE], + (unsigned int*)&prep_opad[index * BINARY_SIZE], + SSEi_MIXED_IN|SSEi_RELOAD|SSEi_OUTPUT_AS_INP_FMT); +#else + SHA256_CTX ctx; + + if (new_keys) { + SHA256_Init(&ipad_ctx[index]); + SHA256_Update(&ipad_ctx[index], ipad[index], PAD_SIZE); + SHA256_Init(&opad_ctx[index]); + SHA256_Update(&opad_ctx[index], opad[index], PAD_SIZE); + } + + memcpy(&ctx, &ipad_ctx[index], sizeof(ctx)); + SHA256_Update(&ctx, cur_salt.salt, cur_salt.length); + SHA256_Final((unsigned char*) crypt_key[index], &ctx); + + memcpy(&ctx, &opad_ctx[index], sizeof(ctx)); + SHA256_Update(&ctx, crypt_key[index], BINARY_SIZE); + SHA256_Final((unsigned char*) crypt_key[index], &ctx); +#endif + } + new_keys = 0; + + return count; +} + +static void *get_binary(char *ciphertext) +{ + static union { + unsigned char c[BINARY_SIZE]; + uint32_t dummy; + } buf; + unsigned char *out = buf.c; + char *p; + int i; + + p = strrchr(ciphertext, '$') + 1; + for (i = 0; i < BINARY_SIZE; i++) { + out[i] = (atoi16[ARCH_INDEX(*p)] << 4) | atoi16[ARCH_INDEX(p[1])]; + p += 2; + } +#if defined(SIMD_COEF_32) && ARCH_LITTLE_ENDIAN==1 + alter_endianity(out, BINARY_SIZE); +#endif + return out; +} + +static void *get_salt(char *ciphertext) +{ + unsigned char salt[SALT_LENGTH + 1]; + unsigned int i, len; +#ifdef SIMD_COEF_32 + unsigned int j; + static JTR_ALIGN(MEM_ALIGN_SIMD) cur_salt_t cs; +#else + static rakp_salt out; +#endif + + if (!strncmp(ciphertext, FORMAT_TAG, TAG_LENGTH)) + ciphertext += TAG_LENGTH; + + len = (strrchr(ciphertext, '$') - ciphertext) / 2; + memset(salt, 0, sizeof(salt)); + for (i = 0; i < len; i++) + salt[i] = (atoi16[ARCH_INDEX(ciphertext[2 * i])] << 4) | + atoi16[ARCH_INDEX(ciphertext[2 * i + 1])]; + +#ifdef SIMD_COEF_32 + memset(&cs, 0, sizeof(cs)); + for (i = 0; i < len; i++) + for (j = 0; j < MIN_KEYS_PER_CRYPT; ++j) + cs.salt[i / PAD_SIZE][GETPOS(i, j)] = salt[i]; + cs.salt_len = len; + for (j = 0; j < MIN_KEYS_PER_CRYPT; ++j) { + cs.salt[len / PAD_SIZE][GETPOS(len, j)] = 0x80; + ((unsigned int*)cs.salt[(len + 8) / PAD_SIZE])[15 * SIMD_COEF_32 + (j&(SIMD_COEF_32-1)) + j/SIMD_COEF_32 * PAD_SIZE_W * SIMD_COEF_32] = (len + PAD_SIZE) << 3; + } + return &cs; +#else + memset(&out, 0, sizeof(out)); + out.length = len; + memcpy(out.salt, salt, len); + return &out; +#endif +} + +/* + * crypt_key is kept in SIMD "input format" (SSEi_OUTPUT_AS_INP_FMT), so the + * per-hash stride is PAD_SIZE_W (16) words, not BINARY_SIZE/4. + */ +#define COMMON_GET_HASH_SIMD32 PAD_SIZE_W +#define COMMON_GET_HASH_VAR crypt_key +#include "common-get-hash.h" + +struct fmt_main fmt_rakp_sha256 = { + { + FORMAT_LABEL, + FORMAT_NAME, + ALGORITHM_NAME, + BENCHMARK_COMMENT, + BENCHMARK_LENGTH, + 0, + PLAINTEXT_LENGTH, + BINARY_SIZE, + BINARY_ALIGN, + SALT_SIZE, + SALT_ALIGN, + MIN_KEYS_PER_CRYPT, + MAX_KEYS_PER_CRYPT, + /* RAKP salts are bounded (<= ~119 bytes) so FMT_HUGE_INPUT is not needed; + * omitting it enables the per-hash binary-hash table, which matters when + * loading large RAKP hash sets. */ + FMT_CASE | FMT_8_BIT | FMT_OMP, + { NULL }, + { NULL }, + tests + }, { + init, + done, + fmt_default_reset, + fmt_default_prepare, + valid, + fmt_default_split, + get_binary, + get_salt, + { NULL }, + fmt_default_source, + { + fmt_default_binary_hash_0, + fmt_default_binary_hash_1, + fmt_default_binary_hash_2, + fmt_default_binary_hash_3, + fmt_default_binary_hash_4, + fmt_default_binary_hash_5, + fmt_default_binary_hash_6 + }, + fmt_default_salt_hash, + NULL, + set_salt, + set_key, + get_key, +#ifdef SIMD_COEF_32 + clear_keys, +#else + fmt_default_clear_keys, +#endif + crypt_all, + { +#define COMMON_GET_HASH_LINK +#include "common-get-hash.h" + }, + cmp_all, + cmp_one, + cmp_exact + } +}; + +#endif /* plugin stanza */