diff --git a/run/opencl/sha384_kernel.cl b/run/opencl/sha384_kernel.cl new file mode 100644 index 0000000000..2ee4e5aa8d --- /dev/null +++ b/run/opencl/sha384_kernel.cl @@ -0,0 +1,403 @@ +/* + * Developed by Claudio André in 2012 + * + * More information at https://openwall.info/wiki/john/OpenCL-RAWSHA-512 + * + * Copyright (c) 2012-2016 Claudio André + * This program comes with ABSOLUTELY NO WARRANTY; express or implied. + * + * This is free software, and you are welcome to redistribute it + * under certain conditions; as expressed here + * http://www.gnu.org/licenses/gpl-2.0.html + */ + +#include "opencl_rawsha512.h" +#include "opencl_mask_extras.h" +#include "opencl_device_info.h" + +/* + * SHA384 differs from SHA512 only in the initial hash values and the + * truncated output (6 x uint64_t = 48 bytes instead of 8 x 64 = 64). + * The round constants, Sigma/sigma functions, and block logic are identical. + * We override H0-H7 (defined by opencl_sha512.h via opencl_rawsha512.h) + * with the SHA384 initial values here. + */ + +#undef H0 +#undef H1 +#undef H2 +#undef H3 +#undef H4 +#undef H5 +#undef H6 +#undef H7 + +#define H0 0xcbbb9d5dc1059ed8UL +#define H1 0x629a292a367cd507UL +#define H2 0x9159015a3070dd17UL +#define H3 0x152fecd8f70e5939UL +#define H4 0x67332667ffc00b31UL +#define H5 0x8eb44a8768581511UL +#define H6 0xdb0c2e0d64f98fa7UL +#define H7 0x47b5481dbefa4fa4UL + +#ifndef UNROLL_LOOP + /// *** UNROLL *** + ///AMD: sometimes a bad thing(?). + ///NVIDIA: GTX 570 don't allow full unroll. + #if amd_vliw4(DEVICE_INFO) || amd_vliw5(DEVICE_INFO) + #define UNROLL_LOOP 133128 + #elif amd_gcn(DEVICE_INFO) && DEV_VER_MAJOR < 2500 + #define UNROLL_LOOP 132098 + #elif amd_gcn(DEVICE_INFO) && DEV_VER_MAJOR >= 2500 + #define UNROLL_LOOP 34079748 + #elif (nvidia_sm_2x(DEVICE_INFO) || nvidia_sm_3x(DEVICE_INFO)) + #define UNROLL_LOOP 132098 + #elif nvidia_sm_5x(DEVICE_INFO) + #define UNROLL_LOOP 33686536 + #elif nvidia_sm_5plus(DEVICE_INFO) + #define UNROLL_LOOP 132104 + #elif gpu_intel(DEVICE_INFO) + #define UNROLL_LOOP 262658 + #else + #define UNROLL_LOOP 0 + #endif +#endif + +#if (UNROLL_LOOP & (1 << 25)) + #define VECTOR_USAGE 1 +#endif + +#if gpu_amd(DEVICE_INFO) + #define USE_LOCAL 1 +#endif + +INLINE void _memcpy( uint32_t * dest, + __global const uint32_t * src, + const uint32_t len) { + + for (uint i = 0; i < len; i += 4) + *dest++ = *src++; +} + +INLINE void any_hash_cracked( + const uint32_t iter, //which candidates_number is this one + volatile __global uint32_t * const hash_id, //information about how recover the cracked password + const uint64_t * const hash, //the hash calculated by this kernel + __global const uint32_t * const bitmap) { + + uint32_t bit_mask_x, bit_mask_y, found; + + SPREAD_64(hash[0], hash[1], BITMAP_SIZE_MINUS1, bit_mask_x, bit_mask_y) + + if (bitmap[bit_mask_x >> 5] & (1U << (bit_mask_x & 31))) { + + if (bitmap[bit_mask_y >> 5] & (1U << (bit_mask_y & 31))) { + //A possible crack have been found. + found = atomic_inc(&hash_id[0]); + + { + //Save (the probably) hashed key metadata. + uint32_t base = get_global_id(0); + + hash_id[1 + 3 * found] = base; + hash_id[2 + 3 * found] = iter; + hash_id[3 + 3 * found] = (uint32_t) hash[0]; + } + } + } +} + +/* + * Reverse 4 rounds more in SHA-256 and SHA-512 OpenCL kernels + * + * 1. To avoid data transfers from GPU to CPU, I'm using a Bloom filter; + * 2. It can produce false positives; + * 3. The best approach is to discard as many candidates as possible on GPU + * and avoid costly transfers; + * 4. If I reverse rounds I will have less data to compare; + * 5. So I'll face more false positives; + * 6. That will impact performance. + * + * That said, for a future version, based on how many hashes were loaded to + * crack, we can: + * -> use a reversed sha_block() version when running a session with only a few + * keys. + * + * I tested it using a session like this one: + * Loaded 4000002 password hashes with no different salts (Raw-SHA256-opencl [SHA256 OpenCL]) + * + * Using mask on a "Titan X Maxwell", JtR will hash (and discard) more than + * 110 millions keys per crypt_all() call. + * + * 1. Running JtR as it is now: + * False positives per crypt_all() [1]: 1515: 0,0014% + * + * 2. When I reverse steps (less bytes to use on filtering, more data transfers): + * False positives per crypt_all() [1]: 208320: 0,0939% + * Since data transfers GPU->CPU are slow: + * => Result is a 300000Kp/s penalty. + * + * [1] Of course, one could be a crack. + */ +INLINE void sha384_block( const uint64_t * const buffer, + const uint32_t total, uint64_t * const H) { + + uint64_t a = H0; + uint64_t b = H1; + uint64_t c = H2; + uint64_t d = H3; + uint64_t e = H4; + uint64_t f = H5; + uint64_t g = H6; + uint64_t h = H7; + uint64_t t; + uint64_t w[16]; //#define w buffer + +#ifdef VECTOR_USAGE + ulong16 w_vector; + w_vector = vload16(0, buffer); + w_vector = SWAP64_V(w_vector); + vstore16(w_vector, 0, w); +#else + #pragma unroll + for (uint i = 0U; i < 15U; i++) + w[i] = SWAP64(buffer[i]); +#endif + w[15] = (total * 8UL); + +#if (UNROLL_LOOP & (1 << 1)) + #pragma unroll 1 +#elif (UNROLL_LOOP & (1 << 2)) + #pragma unroll 4 +#elif (UNROLL_LOOP & (1 << 3)) + #pragma unroll +#endif + for (uint i = 0U; i < 16U; i++) { + t = k[i] + w[i] + h + Sigma1(e) + Ch(e, f, g); + + h = g; + g = f; + f = e; + e = d + t; + t = t + Maj(a, b, c) + Sigma0(a); + d = c; + c = b; + b = a; + a = t; + } + +#if (UNROLL_LOOP & (1 << 9)) + #pragma unroll 1 +#elif (UNROLL_LOOP & (1 << 10)) + #pragma unroll 16 +#elif (UNROLL_LOOP & (1 << 11)) + #pragma unroll +#endif + for (uint i = 16U; i < 80U; i++) { + w[i & 15] = sigma1(w[(i - 2) & 15]) + sigma0(w[(i - 15) & 15]) + w[(i - 16) & 15] + w[(i - 7) & 15]; + t = k[i] + w[i & 15] + h + Sigma1(e) + Ch(e, f, g); + + h = g; + g = f; + f = e; + e = d + t; + t = t + Maj(a, b, c) + Sigma0(a); + d = c; + c = b; + b = a; + a = t; + } + + /* Put checksum in context given as argument. */ + H[0] = (a + H0); + H[1] = (b + H1); + H[2] = (c + H2); + H[3] = (d + H3); + H[4] = (e + H4); + H[5] = (f + H5); + H[6] = (g + H6); + H[7] = (h + H7); +} + +/* ***************** +- index, //keys offset and length +- int_key_loc, //the position of the mask to apply +- int_keys, //mask to be applied +- candidates_number, //the number of candidates by mask mode +- hash_id, //information about how recover the cracked password +- bitmap, //bitmap containing all to crack hashes +***************** */ +__kernel +void kernel_plaintext_raw( + MAYBE_CONSTANT sha512_salt * salt, + __global const uint32_t * __restrict keys_buffer, + __global const uint32_t * const __restrict index, + __global const uint32_t * const __restrict int_key_loc, + __global const uint32_t * const __restrict int_keys, + const uint32_t candidate_id, + __global uint32_t * const __restrict computed_total, + __global uint64_t * const __restrict computed_w) { + + //Compute buffers (on CPU and NVIDIA, better private) + uint64_t w[16]; + size_t gid = get_global_id(0); + +#ifdef USE_LOCAL + __local uint32_t _ltotal[512]; + #define total _ltotal[get_local_id(0)] +#else + uint32_t _ltotal; + #define total _ltotal +#endif + + { + //Get the position and length of the target key. + uint32_t base = index[gid]; + total = base & 63; + + //Ajust keys to it start position. + keys_buffer += (base >> 6); + } + //- Differences ------------------------------- + #define W_OFFSET 0 + + //Clear the buffer. + #pragma unroll + for (uint i = 0; i < 15; i++) + w[i] = 0; + + //Get password. + _memcpy((uint32_t *) w, keys_buffer, total); + //--------------------------------------------- + + //Prepare buffer. + CLEAR_BUFFER_64_SINGLE(w, total); + APPEND_SINGLE(w, 0x80UL, total); + +#ifdef GPU_MASK_MODE + //Mask Mode: keys generation/finalization. + MASK_KEYS_GENERATION(candidate_id) +#endif + + //save computed w[] + computed_total[gid] = total; + + #pragma unroll + for (uint i = 0; i < 15; i++) + computed_w[gid * 16 + i] = w[i]; +} +#undef W_OFFSET + +__kernel +void kernel_crypt( + const uint32_t candidate_id, + volatile __global uint32_t * const __restrict hash_id, + __global uint32_t * const __restrict bitmap, + __global const uint32_t * __restrict computed_total, + __global const uint64_t * __restrict computed_w) { + + //Compute buffers (on CPU and NVIDIA, better private) + uint64_t w[16]; + uint64_t H[8]; + size_t gid = get_global_id(0); + +#ifdef USE_LOCAL + __local uint32_t _ltotal[512]; + #define total _ltotal[get_local_id(0)] +#else + uint32_t _ltotal; + #define total _ltotal +#endif + + //Get w[]. + total = computed_total[gid]; + + #pragma unroll + for (uint i = 0; i < 15; i++) + w[i] = computed_w[gid * 16 + i]; + + /* Run the collected hash value through sha384. */ + sha384_block(w, total, H); + + any_hash_cracked(candidate_id, hash_id, H, bitmap); +} + +__kernel +void kernel_crypt_raw( + MAYBE_CONSTANT sha512_salt * salt, + __global const uint32_t * __restrict keys_buffer, + __global const uint32_t * const __restrict index, + __global const uint32_t * const __restrict int_key_loc, + __global const uint32_t * const __restrict int_keys, + const uint32_t candidates_number, + volatile __global uint32_t * const __restrict hash_id, + __global uint32_t * const __restrict bitmap) { + + //Compute buffers (on CPU and NVIDIA, better private) + uint64_t w[16]; + uint64_t H[8]; +#ifdef USE_LOCAL + __local uint32_t _ltotal[512]; + #define total _ltotal[get_local_id(0)] +#else + uint32_t _ltotal; + #define total _ltotal +#endif + + { + //Get position and length of informed key. + uint32_t base = index[get_global_id(0)]; + total = base & 63; + + //Ajust keys to it start position. + keys_buffer += (base >> 6); + } + //- Differences ------------------------------- + #define W_OFFSET 0 + + //Clear the buffer. + #pragma unroll + for (uint i = 0; i < 15; i++) + w[i] = 0; + + //Get password. + _memcpy((uint32_t *) w, keys_buffer, total); + //--------------------------------------------- + + //Prepare buffer. + CLEAR_BUFFER_64_SINGLE(w, total); + APPEND_SINGLE(w, 0x80UL, total); + + { + uint32_t i = 0; + +#ifdef GPU_MASK_MODE + //Handle the GPU mask mode candidates generation. + for (; i < candidates_number; i++) { +#endif + +#ifdef GPU_MASK_MODE + //Mask Mode: keys generation/finalization. + MASK_KEYS_GENERATION(i) +#endif + /* Run the collected hash value through sha384. */ + sha384_block(w, total, H); + + any_hash_cracked(i, hash_id, H, bitmap); +#ifdef GPU_MASK_MODE + } +#endif + } +} +#undef W_OFFSET + +__kernel +void kernel_prepare( + const uint32_t candidates_number, + __global uint32_t * const __restrict hash_id) { + + //Clean bitmap and result buffer + if (get_global_id(0) == 0) + hash_id[0] = 0; +} diff --git a/src/opencl_rawsha384_fmt_plug.c b/src/opencl_rawsha384_fmt_plug.c new file mode 100644 index 0000000000..1643d3fe30 --- /dev/null +++ b/src/opencl_rawsha384_fmt_plug.c @@ -0,0 +1,790 @@ +/* + * Developed by Claudio André in 2012 + * Based on source code provided by Samuele Giovanni Tonon + * + * More information at https://openwall.info/wiki/john/OpenCL-RAWSHA-384 + * + * Copyright (c) 2012-2016 Claudio André + * This program comes with ABSOLUTELY NO WARRANTY; express or implied. + * + * This is free software, and you are welcome to redistribute it + * under certain conditions; as expressed here + * http://www.gnu.org/licenses/gpl-2.0.html + */ + +#ifdef HAVE_OPENCL + +#if FMT_EXTERNS_H +extern struct fmt_main fmt_opencl_rawsha384; +#elif FMT_REGISTERS_H +john_register_one(&fmt_opencl_rawsha384); +#else + +#include + +#include "sha2.h" +#include "johnswap.h" +#include "opencl_common.h" +#include "config.h" +#include "options.h" +#include "../run/opencl/opencl_rawsha512.h" +#include "rawSHA384_common.h" + +#include "mask_ext.h" +#include "../run/opencl/opencl_mask_extras.h" + +#define FORMAT_LABEL "raw-SHA384-opencl" +#define FORMAT_NAME "" + +#define ALGORITHM_NAME "SHA384 OpenCL" + +#define BINARY_SIZE DIGEST_SIZE + +//plaintext: keys to compute the hash function +//saved_idx: offset and length of each plaintext (data is sent using chunks) +static uint32_t *plaintext, *saved_idx; + +static cl_mem salt_buffer; //Salt information. +static cl_mem pass_buffer; //Plaintext buffer. +static cl_mem idx_buffer; //Sizes and offsets buffer. +static cl_kernel prepare_kernel; + +//Pinned buffers +static cl_mem pinned_plaintext, pinned_saved_idx, pinned_int_key_loc; + +//Reference to self +static struct fmt_main *self; + +//Reference to the first element in salt list +static struct db_main *main_db; + +//Device (GPU) buffers +//int_keys: mask to apply +//hash_ids: information about how recover the cracked password +//bitmap: a bitmap memory space. +//int_key_loc: the position of the mask to apply. +static cl_mem buffer_int_keys, buffer_hash_ids, buffer_bitmap, buffer_int_key_loc; + +//Host buffers +//saved_int_key_loc: the position of the mask to apply +//num_loaded_hashes: number of binary hashes transferred/loaded to GPU +//hash_ids: information about how recover the cracked password +static uint32_t *saved_int_key_loc, num_loaded_hashes, *hash_ids, *saved_bitmap; + +//ocl_initialized: a reference counter of the openCL objetcts (expect to be 0 or 1) +static unsigned ocl_initialized = 0; + +// Keeps track of whether we should tune for this reset() call. +static int should_tune; + +//Used to control partial key transfers. +static uint32_t key_idx = 0; +static size_t offset = 0, offset_idx = 0; +static int new_keys; + +static uint32_t bitmap_size, previous_size; + +static void load_hash(); +static char *get_key(int index); +static void build_kernel(); +static void release_kernel(); +static void release_mask_buffers(void); + +//This file contains auto-tuning routine(s). It has to be included after formats definitions. +#include "opencl_autotune.h" + +/* ------- Helper functions ------- */ +static size_t get_task_max_work_group_size() +{ + size_t s; + + s = autotune_get_task_max_work_group_size(FALSE, 0, crypt_kernel); + s = MIN(s, autotune_get_task_max_work_group_size(FALSE, 0, + prepare_kernel)); + return MIN(s, 512); +} + +static uint32_t get_num_loaded_hashes() +{ + uint32_t num_hashes; + struct db_salt *current_salt; + + num_hashes = 0; + current_salt = main_db->salts; + + do + num_hashes += current_salt->count; + while ((current_salt = current_salt->next)); + + return num_hashes; +} + +static uint64_t *crypt_one(int index) { + SHA512_CTX ctx; + static uint64_t hash[DIGEST_SIZE / sizeof(uint64_t)]; + + char * key = get_key(index); + int len = strlen(key); + + SHA384_Init(&ctx); + SHA384_Update(&ctx, key, len); + SHA384_Final((unsigned char *) (hash), &ctx); + + alter_endianity_to_BE64(hash, DIGEST_SIZE / sizeof(uint64_t)); + + return hash; +} + +/* ------- Create and destroy necessary objects ------- */ +static void create_mask_buffers() +{ + release_mask_buffers(); + + saved_bitmap = (uint32_t *) + mem_alloc((bitmap_size / 32 + 1) * sizeof(uint32_t)); + buffer_bitmap = clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, + (bitmap_size / 32 + 1) * sizeof(uint32_t), NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating buffer argument buffer_bitmap"); + + //Set crypt kernel arguments + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 7, sizeof(buffer_bitmap), + (void *)&buffer_bitmap), "Error setting argument 7"); +} + +static void release_mask_buffers() +{ + MEM_FREE(saved_bitmap); + + if (buffer_bitmap) + clReleaseMemObject(buffer_bitmap); + buffer_bitmap = NULL; +} + +static void create_clobj(size_t gws, struct fmt_main *self) +{ + uint32_t hash_id_size; + size_t mask_cand = 1, mask_gws = 1; + + release_clobj(); + + if (mask_int_cand.num_int_cand > 1) { + mask_cand = mask_int_cand.num_int_cand; + mask_gws = gws; + } + + pinned_plaintext = clCreateBuffer(context[gpu_id], + CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, + BUFFER_SIZE * gws, NULL, &ret_code); + HANDLE_CLERROR(ret_code, + "Error creating page-locked memory pinned_plaintext"); + + plaintext = (uint32_t *) clEnqueueMapBuffer(queue[gpu_id], + pinned_plaintext, CL_TRUE, CL_MAP_WRITE, 0, + BUFFER_SIZE * gws, 0, NULL, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error mapping page-locked memory plaintext"); + + pinned_saved_idx = clCreateBuffer(context[gpu_id], + CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, + sizeof(uint32_t) * gws, NULL, &ret_code); + HANDLE_CLERROR(ret_code, + "Error creating page-locked memory pinned_saved_idx"); + + saved_idx = (uint32_t *) clEnqueueMapBuffer(queue[gpu_id], + pinned_saved_idx, CL_TRUE, CL_MAP_WRITE, 0, + sizeof(uint32_t) * gws, 0, NULL, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error mapping page-locked memory saved_idx"); + + // create arguments (buffers) + salt_buffer = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, + sizeof(sha512_salt), NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating salt_buffer out argument"); + + pass_buffer = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, + BUFFER_SIZE * gws, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating buffer argument pass_buffer"); + + idx_buffer = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, + sizeof(uint32_t) * gws, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating buffer argument idx_buffer"); + + hash_id_size = mask_int_cand.num_int_cand * gws; + hash_ids = (uint32_t *) mem_alloc( + hash_id_size * 3 * sizeof(uint32_t) + sizeof(uint32_t)); + buffer_hash_ids = clCreateBuffer(context[gpu_id], CL_MEM_READ_WRITE, + hash_id_size * 3 * sizeof(uint32_t) + sizeof(uint32_t), + NULL, &ret_code); + + HANDLE_CLERROR(ret_code, "Error creating buffer argument buffer_buffer_hash_ids"); + + //Mask mode + pinned_int_key_loc = clCreateBuffer(context[gpu_id], + CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, + sizeof(uint32_t) * mask_gws, NULL, &ret_code); + HANDLE_CLERROR(ret_code, + "Error creating page-locked memory pinned_int_key_loc"); + + saved_int_key_loc = (uint32_t *) clEnqueueMapBuffer(queue[gpu_id], + pinned_int_key_loc, CL_TRUE, CL_MAP_WRITE, 0, + sizeof(uint32_t) * mask_gws, 0, NULL, NULL, &ret_code); + HANDLE_CLERROR(ret_code, + "Error mapping page-locked memory saved_int_key_loc"); + + buffer_int_key_loc = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, + sizeof(uint32_t) * mask_gws, NULL, &ret_code); + HANDLE_CLERROR(ret_code, + "Error creating buffer argument buffer_int_key_loc"); + + buffer_int_keys = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, + 4 * mask_cand, NULL, &ret_code); + HANDLE_CLERROR(ret_code, "Error creating buffer argument buffer_int_keys"); + + //Set prepare kernel arguments + HANDLE_CLERROR(clSetKernelArg(prepare_kernel, 0, sizeof(cl_uint), + (void *)&mask_int_cand.num_int_cand), "Error setting argument 0"); + HANDLE_CLERROR(clSetKernelArg(prepare_kernel, 1, sizeof(buffer_hash_ids), + (void *)&buffer_hash_ids), "Error setting argument 1"); + + //Set kernel arguments + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 0, sizeof(cl_mem), + (void *)&salt_buffer), "Error setting argument 0"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(cl_mem), + (void *)&pass_buffer), "Error setting argument 1"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(cl_mem), + (void *)&idx_buffer), "Error setting argument 2"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 3, sizeof(buffer_int_key_loc), + (void *)&buffer_int_key_loc), "Error setting argument 3"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 4, sizeof(buffer_int_keys), + (void *)&buffer_int_keys), "Error setting argument 4"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 5, sizeof(cl_uint), + (void *)&(mask_int_cand.num_int_cand)), + "Error setting argument 5"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 6, sizeof(buffer_hash_ids), + (void *)&buffer_hash_ids), "Error setting argument 6"); + + //Indicates that the OpenCL objetcs are initialized. + ocl_initialized++; + + //Assure buffers have no "trash data". + memset(plaintext, '\0', BUFFER_SIZE * gws); + memset(saved_idx, '\0', sizeof(uint32_t) * gws); + memset(saved_int_key_loc, '\0', sizeof(uint32_t) * mask_gws); +} + +static void release_clobj() +{ + cl_int ret_code; + + if (ocl_initialized) { + ret_code = clEnqueueUnmapMemObject(queue[gpu_id], pinned_plaintext, + plaintext, 0, NULL, NULL); + HANDLE_CLERROR(ret_code, "Error Unmapping keys"); + ret_code = clEnqueueUnmapMemObject(queue[gpu_id], pinned_saved_idx, + saved_idx, 0, NULL, NULL); + HANDLE_CLERROR(ret_code, "Error Unmapping indexes"); + ret_code = clEnqueueUnmapMemObject(queue[gpu_id], pinned_int_key_loc, + saved_int_key_loc, 0, NULL, NULL); + HANDLE_CLERROR(ret_code, "Error Unmapping key locations"); + + ret_code = clReleaseMemObject(salt_buffer); + HANDLE_CLERROR(ret_code, "Error Releasing salt_buffer"); + ret_code = clReleaseMemObject(pass_buffer); + HANDLE_CLERROR(ret_code, "Error Releasing pass_buffer"); + ret_code = clReleaseMemObject(idx_buffer); + HANDLE_CLERROR(ret_code, "Error Releasing idx_buffer"); + + MEM_FREE(hash_ids); + clReleaseMemObject(buffer_hash_ids); + HANDLE_CLERROR(ret_code, "Error Releasing buffer_hash_ids"); + + ret_code = clReleaseMemObject(buffer_int_key_loc); + HANDLE_CLERROR(ret_code, "Error Releasing buffer_int_key_loc"); + ret_code = clReleaseMemObject(buffer_int_keys); + HANDLE_CLERROR(ret_code, "Error Releasing buffer_int_keys"); + ret_code = clReleaseMemObject(pinned_plaintext); + HANDLE_CLERROR(ret_code, "Error Releasing pinned_plaintext"); + ret_code = clReleaseMemObject(pinned_saved_idx); + HANDLE_CLERROR(ret_code, "Error Releasing pinned_saved_idx"); + ret_code = clReleaseMemObject(pinned_int_key_loc); + HANDLE_CLERROR(ret_code, "Error Releasing pinned_int_key_loc"); + + ocl_initialized = 0; + HANDLE_CLERROR(clFinish(queue[gpu_id]), "Error releasing memory"); + } +} + +/* ------- Reset functions ------- */ +static void tune(struct db_main *db) +{ + char *tmp_value; + size_t gws_limit; + int autotune_limit = 500; + + if ((tmp_value = getenv("_GPU_AUTOTUNE_LIMIT"))) + autotune_limit = atoi(tmp_value); + + // Auto-tune / Benckmark / Self-test. + gws_limit = MIN((0xf << 22) * 4 / BUFFER_SIZE, + get_max_mem_alloc_size(gpu_id) / BUFFER_SIZE); + + if (options.flags & FLG_MASK_CHK) + gws_limit = MIN(gws_limit, + get_max_mem_alloc_size(gpu_id) / + (mask_int_cand.num_int_cand * 3 * sizeof(uint32_t))); + + //Initialize openCL tuning (library) for this format. + opencl_init_auto_setup(SEED, 0, NULL, + warn, 3, self, create_clobj, release_clobj, + 2 * BUFFER_SIZE, gws_limit, db); + + //Auto tune execution from shared/included code. + autotune_run(self, 1, gws_limit, autotune_limit); +} + +static void reset(struct db_main *db) +{ + offset = 0; + offset_idx = 0; + key_idx = 0; + + main_db = db; + num_loaded_hashes = get_num_loaded_hashes(); + + //Adjust kernel parameters and rebuild (if necessary). + build_kernel(); + + tune(db); + + hash_ids[0] = 0; + load_hash(); +} + +/* ------- Key functions ------- */ +static void clear_keys(void) +{ + offset = 0; + offset_idx = 0; + key_idx = 0; +} + +static void set_key(char *_key, int index) +{ + + const uint32_t *key = (uint32_t *) _key; + int len = strlen(_key); + + saved_idx[index] = (key_idx << 6) | len; + + do { + plaintext[key_idx++] = *key++; + len -= 4; + } while (len > 4); + + if (len > 0) + plaintext[key_idx++] = *key; + + //Mask Mode ranges setup + if (mask_int_cand.num_int_cand > 1) { + int i; + + saved_int_key_loc[index] = 0; + + for (i = 0; i < MASK_FMT_INT_PLHDR; i++) { + + if (mask_skip_ranges[i] != -1) { + saved_int_key_loc[index] |= + ((mask_int_cand.int_cpu_mask_ctx-> + ranges[mask_skip_ranges[i]].offset + + mask_int_cand.int_cpu_mask_ctx-> + ranges[mask_skip_ranges[i]].pos) & 0xff) + << (i << 3); + } else + saved_int_key_loc[index] |= 0x80 << (i << 3); + } + } + //Batch transfers to GPU. + if ((index % TRANSFER_SIZE) == 0 && (index > 0)) { + HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], pass_buffer, + CL_FALSE, sizeof(uint32_t) * offset, + sizeof(uint32_t) * TRANSFER_SIZE, + plaintext + offset, 0, NULL, NULL), + "failed in clEnqueueWriteBuffer pass_buffer"); + HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], idx_buffer, + CL_FALSE, sizeof(uint32_t) * offset, + sizeof(uint32_t) * TRANSFER_SIZE, + saved_idx + offset, 0, NULL, NULL), + "failed in clEnqueueWriteBuffer idx_buffer"); + + HANDLE_CLERROR(clFlush(queue[gpu_id]), "failed in clFlush"); + offset += TRANSFER_SIZE; + offset_idx = key_idx; + } + new_keys = 1; +} + +static char *get_key(int index) +{ + static char *ret; + int int_index, t, i; + + if (!ret) + ret = mem_alloc_tiny(PLAINTEXT_LENGTH + 1, MEM_ALIGN_WORD); + + //Mask Mode plaintext recovery + if (hash_ids == NULL || hash_ids[0] == 0 || index > hash_ids[0]) { + t = index; + int_index = 0; + + } else { + t = hash_ids[1 + 3 * index]; + int_index = hash_ids[2 + 3 * index]; + } + + //Mask Mode plaintext recovery. + if (t >= global_work_size) + t = 0; + + memcpy(ret, ((char *)&plaintext[saved_idx[t] >> 6]), PLAINTEXT_LENGTH); + ret[saved_idx[t] & 63] = '\0'; + + if (saved_idx[t] & 63 && + mask_skip_ranges && mask_int_cand.num_int_cand > 1) { + for (i = 0; i < MASK_FMT_INT_PLHDR && mask_skip_ranges[i] != -1; i++) + ret[(saved_int_key_loc[t] & (0xff << (i * 8))) >> (i * 8)] = + mask_int_cand.int_cand[int_index].x[i]; + } + + return ret; +} + +/* ------- Initialization ------- */ +static void build_kernel() +{ + static int num_int_cand; + + char *task = "$JOHN/opencl/sha384_kernel.cl"; + char opt[MAX_OCLINFO_STRING_LEN]; + + bitmap_size = get_bitmap_size_bits(num_loaded_hashes, gpu_id); + + if (previous_size != bitmap_size || num_int_cand != mask_int_cand.num_int_cand) { + previous_size = bitmap_size; + num_int_cand = mask_int_cand.num_int_cand; + + release_kernel(); + + snprintf(opt, sizeof(opt), "-DBITMAP_SIZE_MINUS1=%u", bitmap_size - 1U); + + if (mask_int_cand.num_int_cand > 1) + strncat(opt, " -DGPU_MASK_MODE", 64U); + + opencl_build_kernel(task, gpu_id, opt, 0); + + // create kernel(s) to execute + prepare_kernel = clCreateKernel(program[gpu_id], "kernel_prepare", + &ret_code); + HANDLE_CLERROR(ret_code, + "Error creating kernel_prepare. Double-check kernel name?"); + + crypt_kernel = clCreateKernel(program[gpu_id], + "kernel_crypt_raw", &ret_code); + HANDLE_CLERROR(ret_code, + "Error creating kernel. Double-check kernel name?"); + } + //Allocate bit array and pass its size to OpenCL. + create_mask_buffers(); +} + +static void release_kernel() +{ + if (program[gpu_id]) { + HANDLE_CLERROR(clReleaseKernel(crypt_kernel), "Release kernel"); + HANDLE_CLERROR(clReleaseKernel(prepare_kernel), "Release kernel"); + HANDLE_CLERROR(clReleaseProgram(program[gpu_id]), "Release Program"); + + program[gpu_id] = NULL; + } +} + +static void init_raw(struct fmt_main *_self) +{ + char *tmp_value; + + self = _self; + opencl_prepare_dev(gpu_id); + mask_int_cand_target = opencl_speed_index(gpu_id) / 300; + previous_size = 0; + + if ((tmp_value = getenv("_GPU_MASK_CAND"))) + mask_int_cand_target = atoi(tmp_value); +} + +static void done(void) +{ + if (program[gpu_id]) { + release_clobj(); + release_kernel(); + release_mask_buffers(); + } + should_tune = 0; + ocl_initialized = 0; +} + +static void prepare_bit_array() +{ + uint64_t *binary; + struct db_password *pw; + struct db_salt *current_salt; + + current_salt = main_db->salts; +#ifdef DEBUG + fprintf(stderr, "Clear bitmap array\n"); +#endif + memset(saved_bitmap, '\0', (bitmap_size / 8 + 1)); + + do { + pw = current_salt->list; + + do { + unsigned int bit_mask_x, bit_mask_y; + binary = (uint64_t *) pw->binary; + // Skip cracked. + if (binary) { + SPREAD_64(binary[0], binary[1], (bitmap_size - 1U), + bit_mask_x, bit_mask_y) +#ifdef DEBUG + if (saved_bitmap[bit_mask_x >> 5] & (1U << (bit_mask_x & 31)) && + saved_bitmap[bit_mask_y >> 5] & (1U << (bit_mask_y & 31))) + fprintf(stderr, "Collision: %u %08x %08x %08x %08x\n", + num_loaded_hashes, (unsigned int) binary[0], + bit_mask_x, bit_mask_y, + saved_bitmap[bit_mask_x >> 5]); +#endif + saved_bitmap[bit_mask_x >> 5] |= (1U << (bit_mask_x & 31)); + saved_bitmap[bit_mask_y >> 5] |= (1U << (bit_mask_y & 31)); + } + } while ((pw = pw->next)); + + } while ((current_salt = current_salt->next)); +} + +/* ------- Send hashes to crack (binary) to GPU ------- */ +static void load_hash() +{ + num_loaded_hashes = get_num_loaded_hashes(); + + prepare_bit_array(); + + HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], buffer_bitmap, CL_TRUE, 0, + (bitmap_size / 32 + 1) * sizeof(uint32_t), + saved_bitmap, 0, NULL, NULL), + "failed in clEnqueueWriteBuffer buffer_bitmap"); + + HANDLE_CLERROR(clFinish(queue[gpu_id]), "failed in clFinish"); +} + +/* ------- Crypt function ------- */ +static int crypt_all(int *pcount, struct db_salt *_salt) +{ + const int count = *pcount; + size_t gws; + size_t *lws = local_work_size ? &local_work_size : NULL; + + gws = GET_NEXT_MULTIPLE(count, local_work_size); + + //Check if any password was cracked and reload (if necessary) + if (num_loaded_hashes != get_num_loaded_hashes()) + load_hash(); + + BENCH_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], prepare_kernel, 1, + NULL, &gws, lws, 0, NULL, multi_profilingEvent[0]), + "failed in clEnqueueNDRangeKernel I"); + + //Send data to device. + if (new_keys && key_idx > offset) + BENCH_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], pass_buffer, + CL_FALSE, sizeof(uint32_t) * offset, + sizeof(uint32_t) * (key_idx - offset), plaintext + offset, 0, + NULL, multi_profilingEvent[1]), + "failed in clEnqueueWriteBuffer pass_buffer"); + + BENCH_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], idx_buffer, CL_FALSE, + sizeof(uint32_t) * offset, + sizeof(uint32_t) * (gws - offset), + saved_idx + offset, 0, NULL, multi_profilingEvent[2]), + "failed in clEnqueueWriteBuffer idx_buffer"); + + if (new_keys && mask_int_cand.num_int_cand > 1) { + BENCH_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], buffer_int_key_loc, + CL_FALSE, 0, 4 * gws, saved_int_key_loc, 0, NULL, + multi_profilingEvent[5]), + "failed in clEnqueueWriteBuffer buffer_int_key_loc"); + + BENCH_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], buffer_int_keys, + CL_FALSE, 0, 4 * mask_int_cand.num_int_cand, + mask_int_cand.int_cand, 0, NULL, multi_profilingEvent[6]), + "failed in clEnqueueWriteBuffer buffer_int_keys"); + } + //Enqueue the kernel + BENCH_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1, NULL, + &gws, lws, 0, NULL, multi_profilingEvent[3]), + "failed in clEnqueueNDRangeKernel"); + + //Possible cracked hashes + BENCH_CLERROR(clEnqueueReadBuffer(queue[gpu_id], buffer_hash_ids, CL_FALSE, + 0, sizeof(uint32_t), hash_ids, + 0, NULL, multi_profilingEvent[4]), + "failed in reading data back buffer_hash_ids"); + + //Do the work + BENCH_CLERROR(clFinish(queue[gpu_id]), "failed in clFinish"); + new_keys = 0; + +#ifdef DEBUG + if (hash_ids[0]) + fprintf(stderr, "Some checks are going to be done on CPU: %u: %1.4f%%\n", hash_ids[0], + ((double) hash_ids[0]) / (global_work_size * mask_int_cand.num_int_cand) * 100); +#endif + if (hash_ids[0] > global_work_size * mask_int_cand.num_int_cand) { + fprintf(stderr, "Error, crypt_all() kernel: %u.\n", hash_ids[0]); + error(); + } + + if (hash_ids[0]) { + BENCH_CLERROR(clEnqueueReadBuffer(queue[gpu_id], buffer_hash_ids, CL_FALSE, + 0, (hash_ids[0] * 3 * sizeof(uint32_t) + sizeof(uint32_t)), hash_ids, + 0, NULL, NULL), + "failed in reading data back buffer_hash_ids"); + + //Do the work + BENCH_CLERROR(clFinish(queue[gpu_id]), "failed in clFinish"); + } + *pcount *= mask_int_cand.num_int_cand; + return hash_ids[0]; +} + +/* ------- Compare functions ------- */ +static int cmp_all(void *binary, int count) +{ + return (count > 0); +} + +static int cmp_one(void *binary, int index) +{ + return (hash_ids[3 + 3 * index] == ((uint32_t *) binary)[0]); +} + +static int cmp_exact_raw(char *source, int index) +{ + uint64_t *binary; + uint64_t *full_hash; + +#ifdef DEBUG + fprintf(stderr, "Stressing CPU\n"); +#endif + binary = (uint64_t *) sha384_common_binary_BE(source); + + full_hash = crypt_one(index); + return !memcmp(binary, (void *) full_hash, BINARY_SIZE); +} + +//Get Hash functions group. +static int get_hash_0(int index) +{ + return hash_ids[3 + 3 * index] & PH_MASK_0; +} + +static int get_hash_1(int index) +{ + return hash_ids[3 + 3 * index] & PH_MASK_1; +} + +static int get_hash_2(int index) +{ + return hash_ids[3 + 3 * index] & PH_MASK_2; +} + +static int get_hash_3(int index) +{ + return hash_ids[3 + 3 * index] & PH_MASK_3; +} + +static int get_hash_4(int index) +{ + return hash_ids[3 + 3 * index] & PH_MASK_4; +} + +static int get_hash_5(int index) +{ + return hash_ids[3 + 3 * index] & PH_MASK_5; +} + +static int get_hash_6(int index) +{ + return hash_ids[3 + 3 * index] & PH_MASK_6; +} + +/* ------- Format structure ------- */ +struct fmt_main fmt_opencl_rawsha384 = { + { + FORMAT_LABEL, + FORMAT_NAME, + ALGORITHM_NAME, + BENCHMARK_COMMENT, + BENCHMARK_LENGTH, + 0, + PLAINTEXT_LENGTH, + BINARY_SIZE, + BINARY_ALIGN, + SALT_SIZE_RAW, + SALT_ALIGN, + MIN_KEYS_PER_CRYPT, + MAX_KEYS_PER_CRYPT, + FMT_CASE | FMT_8_BIT | FMT_SPLIT_UNIFIES_CASE | FMT_MASK, + {NULL}, + {FORMAT_TAG}, + sha384_common_tests_rawsha384 + }, { + init_raw, + done, + reset, + fmt_default_prepare, + sha384_common_valid, + sha384_common_split, + sha384_common_binary_BE, + fmt_default_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, + fmt_default_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_raw + } +}; + +#endif /* plugin stanza */ + +#endif /* HAVE_OPENCL */ diff --git a/src/rawSHA384_common.h b/src/rawSHA384_common.h new file mode 100644 index 0000000000..a05518bd6c --- /dev/null +++ b/src/rawSHA384_common.h @@ -0,0 +1,38 @@ +/* + * This file is part of John the Ripper password cracker, + * Copyright (c) 2012 magnum + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted. + * + * There's ABSOLUTELY NO WARRANTY, express or implied. + */ + +#ifndef _COMMON_RAWSHA384_H +#define _COMMON_RAWSHA384_H + +/* ------ Contains (at least) prepare(), valid() and split() ------ */ + +#define DIGEST_SIZE 48 +#define BINARY_ALIGN sizeof(uint64_t) + +#define BENCHMARK_COMMENT "" +#define BENCHMARK_LENGTH 7 + +#define FORMAT_TAG "$SHA384$" + +#define TAG_LENGTH (sizeof(FORMAT_TAG) - 1) + +#define CIPHERTEXT_LENGTH 96 + +int sha384_common_valid(char *ciphertext, struct fmt_main *self); + +void * sha384_common_binary(char *ciphertext); +void * sha384_common_binary_BE(char *ciphertext); +void * sha384_common_binary_rev(char *ciphertext); + +char * sha384_common_split(char *ciphertext, int index, struct fmt_main *self); + +extern struct fmt_tests sha384_common_tests_rawsha384[]; + +#endif diff --git a/src/rawSHA384_common_plug.c b/src/rawSHA384_common_plug.c new file mode 100644 index 0000000000..a3f1a49bf3 --- /dev/null +++ b/src/rawSHA384_common_plug.c @@ -0,0 +1,139 @@ +/* + * This file is part of John the Ripper password cracker, + * Copyright (c) 2012 magnum + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted. + * + * There's ABSOLUTELY NO WARRANTY, express or implied. + */ + +#include "formats.h" +#include "johnswap.h" +#include "base64_convert.h" +#include "rawSHA384_common.h" + +struct fmt_tests sha384_common_tests_rawsha384[] = { + {"a8b64babd0aca91a59bdbb7761b421d4f2bb38280d3a75ba0f21f2bebc45583d446c598660c94ce680c47d19c30783a7", "password"}, + {FORMAT_TAG "a8b64babd0aca91a59bdbb7761b421d4f2bb38280d3a75ba0f21f2bebc45583d446c598660c94ce680c47d19c30783a7", "password"}, + {"8cafed2235386cc5855e75f0d34f103ccc183912e5f02446b77c66539f776e4bf2bf87339b4518a7cb1c2441c568b0f8", "12345678"}, + {FORMAT_TAG "8cafed2235386cc5855e75f0d34f103ccc183912e5f02446b77c66539f776e4bf2bf87339b4518a7cb1c2441c568b0f8", "12345678"}, + {"38b060a751ac96384cd9327eb1b1e36a21fdb71114be07434c0cc7bf63f6e1da274edebfe76f65fbd51ad2f14898b95b", ""}, + {FORMAT_TAG "38b060a751ac96384cd9327eb1b1e36a21fdb71114be07434c0cc7bf63f6e1da274edebfe76f65fbd51ad2f14898b95b", ""}, + {"94e75dd8e1f16d7df761d76c021ad98c283791008b98368e891f411fc5aa1a83ef289e348abdecf5e1ba6971604a0cb0", "UPPERCASE"}, + {"47f05d367b0c32e438fb63e6cf4a5f35c2aa2f90dc7543f8a41a0f95ce8a40a313ab5cf36134a2068c4c969cb50db776", "1"}, + {"1e237288d39d815abc653befcab0eb70966558a5bbc10a24739c116ed2f615be31e81670f02af48fe3cf5112f0fa03e8", "12"}, + {"9a0a82f0c0cf31470d7affede3406cc9aa8410671520b727044eda15b4c25532a9b5cd8aaf9cec4919d76255b6bfb00f", "123"}, + {"504f008c8fcf8b2ed5dfcde752fc5464ab8ba064215d9c5b5fc486af3d9ab8c81b14785180d2ad7cee1ab792ad44798c", "1234"}, + {"0fa76955abfa9dafd83facca8343a92aa09497f98101086611b0bfa95dbc0dcc661d62e9568a5a032ba81960f3e55d4a", "12345"}, + {"0a989ebc4a77b56a6e2bb7b19d995d185ce44090c13e2984b7ecc6d446d4b61ea9991b76a4c2f04b1b4d244841449454", "123456"}, + {"826227b9dfb593ae4ddbd3f5b7e24b6cb92e342c951cce56546fa68a2e56557b5ebac824a5e778438a7f35c985dfe082", "1234567"}, + {NULL} +}; + +static uint64_t H[8] = { + 0xcbbb9d5dc1059ed8LL, + 0x629a292a367cd507LL, + 0x9159015a3070dd17LL, + 0x152fecd8f70e5939LL, + 0x67332667ffc00b31LL, + 0x8eb44a8768581511LL, + 0xdb0c2e0d64f98fa7LL, + 0x47b5481dbefa4fa4LL +}; + +/* ------- Check if the ciphertext if a valid SHA384 hash ------- */ +int sha384_common_valid(char *ciphertext, struct fmt_main *self) +{ + char *p, *q; + + p = ciphertext; + if (!strncmp(p, FORMAT_TAG, TAG_LENGTH)) + p += TAG_LENGTH; + + q = p; + while (atoi16[ARCH_INDEX(*q)] != 0x7F) + q++; + return !*q && q - p == CIPHERTEXT_LENGTH; +} + +/* ------- Binary ------- */ +void *sha384_common_binary(char *ciphertext) +{ + static unsigned char * out; + char *p; + int i; + + if (!out) out = mem_calloc_tiny(DIGEST_SIZE, BINARY_ALIGN); + + p = ciphertext + TAG_LENGTH; + for (i = 0; i < DIGEST_SIZE; i++) { + out[i] = + (atoi16[ARCH_INDEX(*p)] << 4) | + atoi16[ARCH_INDEX(p[1])]; + p += 2; + } + +#if defined(SIMD_COEF_64) && ARCH_LITTLE_ENDIAN==1 + alter_endianity_to_BE64(out, DIGEST_SIZE/8); +#endif + return out; +} + +void *sha384_common_binary_BE(char *ciphertext) +{ + static unsigned char * out; + char *p; + int i; + + if (!out) out = mem_calloc_tiny(DIGEST_SIZE, BINARY_ALIGN); + + p = ciphertext + TAG_LENGTH; + for (i = 0; i < DIGEST_SIZE; i++) { + out[i] = + (atoi16[ARCH_INDEX(*p)] << 4) | + atoi16[ARCH_INDEX(p[1])]; + p += 2; + } + alter_endianity_to_BE64(out, DIGEST_SIZE/8); + return out; +} + +void *sha384_common_binary_rev(char *ciphertext) +{ + static union { + unsigned char out[DIGEST_SIZE]; + uint64_t x; + } x; + unsigned char *out = x.out; + char *p; + int i; + uint64_t *b; + + p = ciphertext + TAG_LENGTH; + for (i = 0; i < DIGEST_SIZE; i++) { + out[i] = + (atoi16[ARCH_INDEX(*p)] << 4) | + atoi16[ARCH_INDEX(p[1])]; + p += 2; + } + b = (uint64_t*)out; + for (i = 0; i < 8; i++) { + uint64_t t = JOHNSWAP64(b[i])-H[i]; + b[i] = JOHNSWAP64(t); + } + return out; +} + +/* ------- Split ------- */ +char * sha384_common_split(char *ciphertext, int index, struct fmt_main *self) +{ + static char out[TAG_LENGTH + CIPHERTEXT_LENGTH + 1]; + + if (!strncmp(ciphertext, FORMAT_TAG, TAG_LENGTH)) + ciphertext += TAG_LENGTH; + + memcpy(out, FORMAT_TAG, TAG_LENGTH); + memcpylwr(out + TAG_LENGTH, ciphertext, CIPHERTEXT_LENGTH + 1); + return out; +}