Proxmark3 的 HiTag2 Attack 5 GPU 版：ht2crack5gpu 的 OpenCL 位切片破解原理与实操

Proxmark3 的 HiTag2 Attack 5 GPU 版ht2crack5gpu 的 OpenCL 位切片破解原理与实操【免费下载链接】proxmark3Iceman Fork - Proxmark3项目地址: https://gitcode.com/GitHub_Trending/pr/proxmark3HiTag2 是汽车远程无钥匙进入系统中曾广泛使用的 RFID 标签加密协议其 PRNG 的线性性使其密钥可被离线破解。本文围绕 Proxmark3 仓库中的 crack5gpu_disabled 文档展开完整讲解ht2crack5gpu这套 HiTag2 Attack 5 的 OpenCL GPU 加速实现如何用仅 2 组加密的nR, aR交互数据恢复 48 位共享密钥其位切片bitslicing内核是如何在 GPU 上并行搜索 PRNG 状态的以及构建、参数调整与运行方法。需要先说明的一点该目录名为crack5gpu_disabled且 tools/hitag2crack/Makefile 的顶层构建目标只包含 crack2/crack3/crack4/crack5 与在SKIPOPENCL未置 1 时的crack5opencl并未引用 crack5gpu 目录。从源码结构看这是已被后继者crack5opencl取代的早期版本保留在仓库中作为历史参考若目标是实际跑通破解建议同时阅读 crack5opencl 文档两者的输入格式完全一致。攻击背景与所需数据crack5gpu属于仓库 tools/hitag2crack 套件中的第五种攻击。该套件的总文档指出Attack 1–4 来自 Verdult/Garcia/Balasch 与 Garcia/Oswald/Kasper/Pavlides 的论文Attack 5 与 5gpu 则基于 FactorIT B.V. 的 HiTag2 Hell 实现移植而来。HiTag2 加密的核心是一个 48 位 PRNG由 48 位反馈移位寄存器非线性函数输入与一个并行 48 位 LFSR加速反馈位计算组成见 hitagcrypto.h 中对hitag2_init/hitag2_nstep接口的定义——初始状态由 48 位共享密钥、32 位标签序列号UID和读取器发送的 32 位初始向量nR确定。Attack 5gpu 与 Attack 5 的攻击逻辑完全相同只是将搜索部分移植到 OpenCL 以利用 GPU 并行度。所需输入非常简单——只要两组 nR/aR 对即加密的 noncenR与挑战响应aR值均为十六进制./ht2crack5gpu UID nR1 aR1 nR2 aR2其中 UID 是采集这两组 nR/aR 时所用的标签的 UID。源码中参数解析见 ht2crack5gpu.cUID、nR1、nR2 支持可选的0x前缀且 UID 与 nR 会经rev32(hexreversetouint32(...))做字节序翻转RFID 传输序与主机序的转换而两个 aR 直接按 16 进制解析不做翻转。构建OpenCL 框架依赖与平台适配crack5gpu_disabled的 README 说明构建需要 OpenCL 框架并在必要时手动编辑 Makefile 调整 INCLUDE 与 LIBSmake clean make对照实际的 Makefile 可以看到构建细节参与编译的源文件除ht2crack5gpu.c外还包括 common/ht2crackutils.c 与 common/hitagcrypto.cMYSRCS ht2crackutils.c hitagcrypto.c即复用了套件中其他攻击的 HiTag 辅助函数头文件搜索路径固定加入../common与../common/OpenCL-Headers。仓库在 crack5gpu_disabled 目录 之外自带了一份 OpenCL 头文件位于 crack5opencl/common/OpenCL-Headers这正是 Makefile 中MYINCLUDES -I ../common/OpenCL-Headers所指的位置可在没有系统 OpenCL SDK 时编译库链接按平台区分Darwin 用-framework OpenCL其他平台默认-L/opt/nvidia/cuda/lib64 -lOpenCLMakefile 中保留了指向/usr/local/cuda-7.5的注释备选路径。如果你的 NVIDIA 驱动 OpenCL 库不在/opt/nvidia/cuda/lib64就需要按 README 提示修改MYLDLIBS主程序构建规则为ht2crack5gpu : $(OBJDIR)/ht2crack5gpu.o $(MYOBJS)最终产物即为ht2crack5gpu可执行文件Mingw 下还会额外加-D_ISOC99_SOURCE以保证%z等格式符可用。运行时还有一个隐含前提内核以源码形式存放程序按相对当前工作目录打开ht2crack5kernel.cl代码中KERNELFILENAME ht2crack5kernel.cl见 ht2crack5gpu.c 与 打开内核源文件 的逻辑因此必须在crack5gpu_disabled目录下运行程序否则会得到Cannot open ht2crack5kernel.cl的错误。主机端流程候选状态预筛与 OpenCL 调度阅读 ht2crack5gpu.c 的main函数主机端CPU与 GPU 端的分工如下解析输入并构造目标比特流。target ~aR1由于 aR 是加密后的响应PRNG 的真实输出keystream与~aR1逐位对应随后bitslice(~target, keystream, 32, true)把这 32 位目标反向位切片为 32 个bitslice_t每个为 32 位向量一次承载 32 个并行状态上传到 GPU 只读缓冲区。层 0 预筛CPU 完成 2^20 状态空间的第一层过滤。HiTag2 的初始状态中最低 2 位从未被真正使用源码注释 “we never actually set or use the lowest 2 bits”且状态的高位结构是确定的0x5806b4a2d16c掩码模板经expand()展开即固定比特0, 1, 1, 1, 1, 1的填充模式。CPU 遍历1 20种 20 位自由位组合i0用查表宏f(state0)计算非线性函数第一个输出位与target 31比对通过者被拆成 3 个 16 位字存入candidates[(1 20) * 3]数组——数组下标乘以 3 是因为每个候选状态占 3 个uint16_t。GPU 端 64 片轮转执行内核。内核以 2 维 NDRange8192 × 1024启动见 runKernel每片处理8192 × 1024 / (113) 2^20 / 2^13个候选主机端循环for (step 0; step 64; step)并以step 13作为cand_base传入恰好覆盖全部2^20 / 2^13 × 64候选空间。每轮把匹配到的状态matches[8192]个uint64_t与matches_found计数读回主机打印slice %3u/64: %5u candidates进度。主机端验证候选。每个 GPU 命中的状态交给 try_state()将状态回滚 2 步rollback(hstate, 2)恢复初始状态从状态中提取低 16 位密钥低半与nR1xk明文 nonce 异或密钥流的部分再用 UID 逐位驱动非线性函数 32 步重算密钥高 32 位得到完整 48 位keyrev用hitag2_init(hstate, keyrev, uid, nR2)以第二组nR 重新初始化执行 32 步后将输出与aR2异或若结果全 1 0xffffffff即密钥正确——打印 6 字节密钥并以 0 退出64 片全部跑完仍未命中则输出Key not found并以 1 退出。这种“GPU 粗筛 CPU 精验”的划分避免了把hitag2_init/rollback这类难以位切片的状态机搬上设备只把最密集的“状态 → 首个 keystream 位”搜索留给 GPU。GPU 内核46 层位切片 LFSR 展开ht2crack5kernel.cl 是整个工具的性能核心。它不逐个模拟状态而是一次模拟 32 个状态MAX_BITSLICES 32bitslice_t为 4 字节向量状态装载。每个工作项用index 3 * (candidate_index_base get_global_id(0))从设备端candidates数组取出 3 个 16 位字拼成 48 位候选状态bitslice()将其 46 个有效位写入state数组其余 2 个未知初始位由第 1 维get_global_id(1)的 1024 个通道逐一假设10 个未知位来自i1的 10 个 bit再叠加第 2 个未知位state[-2 10]等由i2的 32 个组合覆盖从而完整覆盖剩余自由度非线性函数用 LOP3 指令加速。内核通过内联 PTXlop3.b32实现lut3()f_a_bs2 次 LUT、f_b_bs2 次 LUT 2 次 XOR、f_c_bs4 次 LUT 2 次 AND 1 次 XOR分别对应 HiTag2 非线性函数 F 的子表达式LFSR 反馈则用lfsr_bs(i)7 次 LUT 1 次 XOR替代主机端的 16 次异或。文件头注释给出了优化收益非 LUT 版本约 66 操作LUT 版本约 43 操作逐位与 keystream 比对、早退剪枝。内核按 17 个非线性过滤输出filter1…filter17的顺序逐位推进每算出一个 filter 输出就与keystream[n]异或并与前序结果的位向量做 ANDresults1…results8只要位向量为 0if (!resultsX) continue/return;就立即剪枝不再展开后续 3 层未知位i3–i8 对应的循环。这种“每 2–3 个 keystream 位过滤一次”的嵌套结构使绝大多数错误状态在展开 46 步 LFSR 之前就被丢弃命中回传。通过全部 32 个 keystream 位验证后内核用unbitslice()把 48 位状态解包matches[atomic_inc(matches_found)] ...原子写入结果数组。注释特意说明取的是“layer 2”的状态因为最低 2 位可以留在主机端通过反转 LFSR 恢复对应try_state中的rollback(hstate, 2)。与 CPU 版及 5opencl 版的对照与 Attack 5纯 CPU的关系crack5 文档 与crack5gpu的输入格式完全相同./ht2crack5 12345678 71DA20AA 7EFDF3FA 2A4265F9 59653B07。总文档明确 “Attack 5gpu is identical to attack 5, simply the code has been ported to OpenCL to run on GPUs and is therefore much faster than attack 5”而 5gpu 的改动点在 ht2crack5gpu.c 文件头注释 中列出main 接收 UID 与两组 {nR, aR}、用第一组 aR 搜索状态并重建密钥候选、再用第二组 {nR, aR} 验证为适配 OpenCL 缩小了max_bitslices与类型尺寸候选数组由 64 位降为 3 个 16 位字复用其他攻击的 HiTag 辅助函数。与 Attack 5opencl 的关系5opencl 是“基于 5gpu 的优化 OpenCL 版本”支持多 GPU/CPU、设备/平台选择与调度参数见 crack5opencl 文档 的-p/-d/-D/-S/-P/-F/-Q等选项示例命令./ht2crack5opencl -D 2 -Q 2 -p 1,2 -d 1,2,3 2ab12bf2 4B71E49D 6A606453 D79BD94B 16A2255B。5gpu 版固定使用第一个平台的第一个 GPUclGetDeviceIDs(..., CL_DEVICE_TYPE_GPU, ...)且无-Werror之外的构建选项这也是它被标记为 disabled、让位于 5opencl 的原因。实操步骤小结结合 README 与源码完整流程如下采集数据从目标标签的一次正常交互中取得两组十六进制 nR/aR 对及 UID套件总文档提醒Proxmark3 的lf hitag sniff当时尚不足以稳定采集所需数据通常需要借助 RFIDler 等设备。进入tools/hitag2crack/crack5gpu_disabled按本机 OpenCL 环境核对 Makefile 的MYLDLIBS执行make clean make。在该目录下运行./ht2crack5gpu UID nR1 aR1 nR2 aR2观察 64 片slice x/64的候选数进度。命中时输出 6 字节48 位密钥并退出码 0穷尽后输出Key not found且退出码 1。恢复密钥后的下一步用 Proxmark3 模拟 RWD 或直接复制标签在 tools/hitag2crack/README.md 的 “Next steps” 一节有说明仓库中对应命令示例仍标记为 TODO。需要注意的适用前提该工具假设目标是标准 HiTag2 协议且 nR 可自由重放这正是论文中 Attack 1/5 依赖的 nonce 重放前提GPU 上编译失败时可借助程序打印的 build logclGetProgramBuildInfo的输出会完整打印排查 LOP3/PTX 相关的设备兼容性问题——内核使用了面向 NVIDIA 的lop3.b32内联汇编在缺少该指令的 OpenCL 设备上可能无法通过编译或构建。【免费下载链接】proxmark3Iceman Fork - Proxmark3项目地址: https://gitcode.com/GitHub_Trending/pr/proxmark3创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考