Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions doc/NEWS
Original file line number Diff line number Diff line change
Expand Up @@ -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$<salt>$<hmac>"
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):

Expand Down
39 changes: 39 additions & 0 deletions run/opencl/rakp_md5_kernel.cl
Original file line number Diff line number Diff line change
@@ -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];
}
39 changes: 39 additions & 0 deletions run/opencl/rakp_sha256_kernel.cl
Original file line number Diff line number Diff line change
@@ -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];
}
Loading