github/RRG-Proxmark3
1
0
Fork 0
mirror of https://github.com/RfidResearchGroup/proxmark3.git synced 2025-08-14 18:48:13 -07:00
Code Issues Projects Releases Packages Wiki Activity Actions

style

Browse source
This commit is contained in:
iceman1001 2024-04-22 09:36:42 +02:00
commit 042ba20d58
3 changed files with 37 additions and 21 deletions
Download patch file Download diff file Expand all files Collapse all files

8
tools/hitag2crack/crack5opencl/hitag2.c
View file

@ -135,12 +135,16 @@ void hitag2_init(Hitag_State *pstate, uint64_t sharedkey, uint32_t serialnum, ui
int i;
for (i = 0; i < 16; i++) cur_state = (cur_state >> 1) ^ (uint64_t) hitag2_crypt(cur_state) << 46;
for (i = 0; i < 16; i++) {
cur_state = (cur_state >> 1) ^ (uint64_t) hitag2_crypt(cur_state) << 46;
}
// highest 16 bits of IV XOR Shared Key
cur_state |= (uint64_t) initvector << 47;
for (i = 0; i < 15; i++) cur_state = (cur_state >> 1) ^ (uint64_t) hitag2_crypt(cur_state) << 46;
for (i = 0; i < 15; i++) {
cur_state = (cur_state >> 1) ^ (uint64_t) hitag2_crypt(cur_state) << 46;
}
cur_state ^= (uint64_t) hitag2_crypt(cur_state) << 47;

42
tools/hitag2crack/crack5opencl/ht2crack5opencl_kernel.cl
View file

@ -309,17 +309,20 @@ static uint hitag2_nstep2 (ulong state, ulong lfsr)
return result;
}
inline static int bitn(ulong x, int bit)
{
inline static int bitn(ulong x, int bit) {
const ulong bitmask = 1UL << bit;
return (x & bitmask) ? 1 : 0;
}
static int fnR (ulong x)
{
return (bitn(x, 1) ^ bitn(x, 2) ^ bitn(x, 5) ^ bitn(x, 6) ^ bitn(x, 7) ^
bitn(x, 15) ^ bitn(x, 21) ^ bitn(x, 22) ^ bitn(x, 25) ^ bitn(x, 29) ^ bitn(x, 40) ^
bitn(x, 41) ^ bitn(x, 42) ^ bitn(x, 45) ^ bitn(x, 46) ^ bitn(x, 47));
static int fnR (ulong x) {
return (
bitn(x, 1) ^ bitn(x, 2) ^ bitn(x, 5) ^
bitn(x, 6) ^ bitn(x, 7) ^ bitn(x, 15) ^
bitn(x, 21) ^ bitn(x, 22) ^ bitn(x, 25) ^
bitn(x, 29) ^ bitn(x, 40) ^ bitn(x, 41) ^
bitn(x, 42) ^ bitn(x, 45) ^ bitn(x, 46) ^
bitn(x, 47)
);
}
inline static int fa(unsigned int i) {
@ -330,8 +333,7 @@ inline static int fb(unsigned int i) {
return bitn(0x6671, i);
}
static int fnf (ulong s)
{
static int fnf (ulong s) {
const uint x1 = (bitn(s, 2) << 0) | lut3_0x96( (bitn(s, 3) << 1), (bitn(s, 5) << 2), (bitn(s, 6) << 3));
const uint x2 = (bitn(s, 8) << 0) | lut3_0x96( (bitn(s, 12) << 1), (bitn(s, 14) << 2), (bitn(s, 15) << 3));
const uint x3 = (bitn(s, 17) << 0) | lut3_0x96( (bitn(s, 21) << 1), (bitn(s, 23) << 2), (bitn(s, 26) << 3));
@ -362,16 +364,21 @@ void find_state(const uint candidate_index_base,
{
const size_t gid[2] = { get_global_id(0), get_global_id(1) };
// if (gid[0] == 0) printf("work-item 1,%u\n", gid[1]);
#ifdef HAVE_LOCAL_MEMORY
const size_t lid = get_local_id(0);
const size_t lsize = get_local_size(0);
#endif // HAVE_LOCAL_MEMORY
const uint index = 3 * (candidate_index_base + gid[0]); // dimension 0 should at least keep the execution units saturated - 8k is fine
// dimension 0 should at least keep the execution units saturated - 8k is fine
const uint index = 3 * (candidate_index_base + gid[0]);
const ulong3 c = { candidates[index], candidates[index + 1], candidates[index + 2] };
const ulong3 c = {
candidates[index],
candidates[index + 1],
candidates[index + 2]
};
const ulong candidate = ( c.x << 32 | c.y << 16 | c.z );
@ -379,13 +386,17 @@ void find_state(const uint candidate_index_base,
// store keystream in local memory
__local bitslice_t keystream[32];
for (size_t i = lid; i < 32; i+= lsize) keystream[i] = _keystream[i];
for (size_t i = lid; i < 32; i+= lsize) {
keystream[i] = _keystream[i];
}
#ifdef WITH_HITAG2_FULL
// store uid, aR2, nR1, nR2 in local memory
__local uint checks[4];
for (uint i = lid; i < 4; i+= lsize) checks[i] = _checks[i];
for (uint i = lid; i < 4; i+= lsize) {
checks[i] = _checks[i];
}
#endif
// threads synchronization
@ -437,6 +448,7 @@ void find_state(const uint candidate_index_base,
const bitslice_t filter1 = f_c_bs(filter1_0, filter1_1, filter1_2, filter1_3, filter1_4);
const bitslice_t results1 = filter1 ^ keystream[1];
if (!results1) return;
const bitslice_t filter2_0 = f_a_bs(state[-2 + 4], state[-2 + 5], state[-2 + 7], state[-2 + 8]);

8
tools/hitag2crack/crack5opencl/opencl.c
View file

@ -310,13 +310,13 @@ int discoverDevices(unsigned int profile_selected, uint32_t device_types_selecte
} else {
(*cd_ctx)[platform_idx].device[device_idx].selected = plat_dev_enabled(global_device_id, dev_sel, dev_cnt, (unsigned int) device_type, device_types_selected);
}
global_device_id++;
if ((*cd_ctx)[platform_idx].device[device_idx].selected) {
(*selected_devices_cnt)++;
}
continue;
} else if (info_idx == 5) {
@ -338,7 +338,7 @@ int discoverDevices(unsigned int profile_selected, uint32_t device_types_selecte
}
if ((*cd_ctx)[platform_idx].is_apple) {
if (strncmp((*cd_ctx)[platform_idx].device[device_idx].vendor, "Intel", 5) != 0) {
(*cd_ctx)[platform_idx].device[device_idx].have_local_memory = true;