yggm/libs/keymanip.cu
2025-03-15 15:54:52 +05:00

48 lines
1.8 KiB
Plaintext

#include <keymanip.cuh>
__device__ ds64 ktos(const unsigned char* key) noexcept {
ds64 str;
const char* hexDigits = "0123456789abcdef";
#pragma unroll 32
for (unsigned char i = 0; i < 32; i++) {
str.data[2 * i] = hexDigits[key[i] >> 4];
str.data[2 * i + 1] = hexDigits[key[i] & 0x0F];
}
str.data[65] = '\0';
return str;
}
__device__ ds46 getAddr(const unsigned char rawAddr[16]) noexcept {
ds46 addrStr;
const char* hexDigits = "0123456789abcdef";
unsigned pos = 0;
#pragma unroll 8
for (unsigned char group = 0; group < 8; group++) {
int idx = group * 2;
addrStr.data[pos++] = hexDigits[rawAddr[idx] >> 4];
addrStr.data[pos++] = hexDigits[rawAddr[idx] & 0x0F];
addrStr.data[pos++] = hexDigits[rawAddr[idx + 1] >> 4];
addrStr.data[pos++] = hexDigits[rawAddr[idx + 1] & 0x0F];
if (group < 7) { addrStr.data[pos++] = ':'; }
}
addrStr.data[pos] = '\0';
return addrStr;
}
__device__ void getRawAddress(int lErase, Key32& InvertedPublicKey, Addr16& rawAddr) noexcept {
lErase++;
const int bitsToShift = lErase & 7;
const int start = lErase >> 3;
if (bitsToShift) {
#pragma unroll
for (int i = start; i < start + 15; i++) {
InvertedPublicKey[i] = static_cast<unsigned char>((InvertedPublicKey[i] << bitsToShift) | (InvertedPublicKey[i + 1] >> (8 - bitsToShift)));
}
}
rawAddr[0] = 0x02;
rawAddr[1] = static_cast<unsigned char>(lErase - 1);
memcpy(&rawAddr[2], &InvertedPublicKey[start], 14);
}
__device__ void invertKey(const unsigned char* key, unsigned char* inverted) {
for (int i = 0; i < 32; i += 4) {
uchar4 k = *(reinterpret_cast<const uchar4*>(&key[i]));
*(reinterpret_cast<uchar4*>(&inverted[i])) = make_uchar4(k.x ^ 0xFF, k.y ^ 0xFF, k.z ^ 0xFF, k.w ^ 0xFF);
}
}