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

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

【免费下载链接】proxmark3Iceman Fork - Proxmark3项目地址: https://gitcode.com/GitHub_Trending/pr/proxmark3

HiTag2 是汽车远程无钥匙进入系统中曾广泛使用的 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 对，即加密的 nonce（nR）与挑战响应（aR）值，均为十六进制：

./ht2crack5gpu <UID> <nR1> <aR1> <nR2> <aR2>

其中 UID 是采集这两组 nR/aR 时所用的标签的 UID。源码中参数解析见 ht2crack5gpu.c：UID、nR1、nR2 支持可选的0x前缀，且 UID 与 nR 会经rev32(hexreversetouint32(...))做字节序翻转（RFID 传输序与主机序的转换），而两个 aR 直接按 16 进制解析，不做翻转。

构建：OpenCL 框架依赖与平台适配

crack5gpu_disabled的 README 说明构建需要 OpenCL 框架，并在必要时手动编辑 Makefile 调整 INCLUDE 与 LIBS：

make clean make

对照实际的 Makefile 可以看到构建细节：

• 参与编译的源文件除ht2crack5gpu.c外，还包括 common/ht2crackutils.c 与 common/hitagcrypto.c（MYSRCS = 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 -lOpenCL（Makefile 中保留了指向/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 端的分工如下：

1. 解析输入并构造目标比特流。target = ~aR1：由于 aR 是加密后的响应，PRNG 的真实输出（keystream）与~aR1逐位对应；随后bitslice(~target, keystream, 32, true)把这 32 位目标反向位切片为 32 个bitslice_t（每个为 32 位向量，一次承载 32 个并行状态），上传到 GPU 只读缓冲区。

2. 层 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。

3. GPU 端 64 片轮转执行内核。内核以 2 维 NDRange8192 × 1024启动（见 runKernel），每片处理8192 × 1024 / (1<<13) = 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进度。

4. 主机端验证候选。每个 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 32，bitslice_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_bs（2 次 LUT）、f_b_bs（2 次 LUT + 2 次 XOR）、f_c_bs（4 次 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]异或，并与前序结果的位向量做 AND（results1…results8）；只要位向量为 0（if (!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 版固定使用第一个平台的第一个 GPU（clGetDeviceIDs(..., CL_DEVICE_TYPE_GPU, ...)，且无-Werror之外的构建选项），这也是它被标记为 disabled、让位于 5opencl 的原因。

实操步骤小结

结合 README 与源码，完整流程如下：

1. 采集数据：从目标标签的一次正常交互中取得两组十六进制 nR/aR 对及 UID（套件总文档提醒，Proxmark3 的lf hitag sniff当时尚不足以稳定采集所需数据，通常需要借助 RFIDler 等设备）。
2. 进入tools/hitag2crack/crack5gpu_disabled，按本机 OpenCL 环境核对 Makefile 的MYLDLIBS，执行make clean && make。
3. 在该目录下运行./ht2crack5gpu <UID> <nR1> <aR1> <nR2> <aR2>，观察 64 片（slice x/64）的候选数进度。
4. 命中时输出 6 字节（48 位）密钥并退出码 0；穷尽后输出Key not found且退出码 1。
5. 恢复密钥后的下一步（用 Proxmark3 模拟 RWD 或直接复制标签）在 tools/hitag2crack/README.md 的 “Next steps” 一节有说明，仓库中对应命令示例仍标记为 TODO。

需要注意的适用前提：该工具假设目标是标准 HiTag2 协议且 nR 可自由重放（这正是论文中 Attack 1/5 依赖的 nonce 重放前提）；GPU 上编译失败时可借助程序打印的 build log（clGetProgramBuildInfo的输出会完整打印）排查 LOP3/PTX 相关的设备兼容性问题——内核使用了面向 NVIDIA 的lop3.b32内联汇编，在缺少该指令的 OpenCL 设备上可能无法通过编译或构建。

【免费下载链接】proxmark3Iceman Fork - Proxmark3项目地址: https://gitcode.com/GitHub_Trending/pr/proxmark3