cuda 逆向

写在前面

CUDA (Compute Unified Device Architecture) 是由 NVIDIA 开发的一种并行计算平台和编程模型；它允许软件开发者和研究人员利用 NVIDIA GPU 的强大并行处理能力来加速通用计算任务，而不仅仅局限于图形渲染
cuda 逆向本身不难，一般是使用 cuobjump 来获得 ptx 汇编 ，根据汇编理出加密逻辑最后进行解密即可

PTX 基础知识

寄存器：
%p<N>：谓词寄存器（1 位，用于条件执行）
%rs<N>：16 位寄存器（例如 %rs1）
%r<N>：32 位寄存器（例如 %r1）
%rd<N>：64 位寄存器（例如 %rd1）
%f<N>：32 位浮点寄存器（例如 %f1）
特殊寄存器：
%tid.x、%tid.y、%tid.z：线程在其线程块内的 ID
%ntid.x、%ntid.y、%ntid.z：线程块的维度（每个块的线程数）
%ctaid.x、%ctaid.y、%ctaid.z ：线程块在网格内的 ID
%nctaid.x、%nctaid.y、%nctaid.z：网格的维度（网格中的块数）
内存空间：
.const：常量内存（只读，缓存）
.global：全局内存（主要的 GPU 内存，所有线程都可访问）
.param：参数内存（用于内核函数的参数）
常见指令：
ld：从内存加载数据
st：将数据存储到内存
mov：在寄存器之间移动值
add、sub、mul、shl、shr：加、减、乘、左移、右移
and、or、xor：按位与、按位或、异或
setp：根据比较结果设置谓词寄存器
bra：分支（无条件跳转） (jmp)
@%pX bra：条件分支（如果谓词 %pX 为真，则跳转）
ret：从内核返回
bar.sync：同步线程块内的所有线程
cvta.to.global.u64：将地址转换为全局内存地址
cvt：数据类型转换
fma.rn.f32：融合乘加（float32，四舍五入到最近的）
mad.lo.s32：乘加（有符号 32 位，结果的低 32 位） (a * b + c)
mul.wide.u32：乘法（无符号 32 位，生成 64 位结果）

ACTF 2025 Deeptx

附件先扔 ida，看到 main

主要步骤:打开并验证 bmp 图像 -> 读取调色板和像素数据 -> 将 S-box、T-box 和 motion 表复制到 GPU 常量内存 -> 调用三层 cuda 处理(Layer1、Layer2、Layer3) -> 保存处理后的图像

使用 cuobjdump.exe,cuobjdump --dump-ptx quiz， dump 出汇编

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
.const .align 1 .b8 cuda_sbox[256]; 
.const .align 1 .b8 cuda_tbox[256]; 
.const .align 4 .b8 cuda_motion[1024]; 

.visible .entry _Z6Layer1PhS_(  
.param .u64 _Z6Layer1PhS__param_0,
.param .u64 _Z6Layer1PhS__param_1 
)
{
.reg .pred %p<6>;
.reg .b16 %rs<2>;
.reg .f32 %f<12>;
.reg .b32 %r<23>;
.reg .b64 %rd<15>;

ld.param.u64 %rd5, [_Z6Layer1PhS__param_0];
ld.param.u64 %rd6, [_Z6Layer1PhS__param_1];
mov.u32 %r1, %tid.x;
setp.lt.u32 %p1, %r1, 241;
mov.u32 %r2, %ctaid.x;
setp.lt.u32 %p2, %r2, 241;
and.pred %p3, %p1, %p2;
@%p3 bra $L__BB0_2;
bra.uni $L__BB0_1;

$L__BB0_2:
mov.u32 %r3, %ntid.x;
cvta.to.global.u64 %rd1, %rd5;
mov.f32 %f10, 0f00000000;
mov.u32 %r11, 0;
mov.u64 %rd8, cuda_motion;
mov.u32 %r20, %r11;

$L__BB0_3:
.pragma "nounroll";
add.s32 %r13, %r20, %r2;
shl.b32 %r14, %r20, 4;
mov.u32 %r15, 240;
sub.s32 %r16, %r15, %r14;
mad.lo.s32 %r21, %r13, %r3, %r1;
mul.wide.u32 %rd7, %r16, 4;
add.s64 %rd14, %rd8, %rd7;
mov.u32 %r22, %r11;

$L__BB0_4:
.pragma "nounroll";
cvt.u64.u32 %rd9, %r21;
add.s64 %rd10, %rd1, %rd9;
ld.global.u8 %rs1, [%rd10];
cvt.rn.f32.u16 %f7, %rs1;
ld.const.f32 %f8, [%rd14];
fma.rn.f32 %f10, %f8, %f7, %f10;
add.s32 %r21, %r21, 1;
add.s64 %rd14, %rd14, 4;
add.s32 %r22, %r22, 1;
setp.ne.s32 %p4, %r22, 16;
@%p4 bra $L__BB0_4;

add.s32 %r20, %r20, 1;
setp.lt.u32 %p5, %r20, 16;
@%p5 bra $L__BB0_3;
bra.uni $L__BB0_6;

$L__BB0_1:
mov.f32 %f10, 0f00000000;

$L__BB0_6:
cvt.rzi.u32.f32 %r17, %f10;
mov.u32 %r18, %ntid.x;
mad.lo.s32 %r19, %r2, %r18, %r1;
cvt.u64.u32 %rd11, %r19;
cvta.to.global.u64 %rd12, %rd6;
add.s64 %rd13, %rd12, %rd11;
st.global.u8 [%rd13], %r17;
ret;

}
//
.visible .entry _Z6Layer2PhS_(
.param .u64 _Z6Layer2PhS__param_0,
.param .u64 _Z6Layer2PhS__param_1
)
{
.reg .b16 %rs<2>;
.reg .b32 %r<8>;
.reg .b64 %rd<14>;


ld.param.u64 %rd1, [_Z6Layer2PhS__param_0];
ld.param.u64 %rd2, [_Z6Layer2PhS__param_1];
cvta.to.global.u64 %rd3, %rd2;
cvta.to.global.u64 %rd4, %rd1;
mov.u32 %r1, %ctaid.x;
mov.u32 %r2, %ntid.x;
mov.u32 %r3, %tid.x;
mad.lo.s32 %r4, %r1, %r2, %r3;
cvt.u64.u32 %rd5, %r4;
add.s64 %rd6, %rd4, %rd5;
ld.global.u8 %rs1, [%rd6];
cvt.u64.u32 %rd7, %r3;
mov.u64 %rd8, cuda_sbox;
add.s64 %rd9, %rd8, %rd7;
ld.const.u8 %r5, [%rd9];
cvt.u64.u32 %rd10, %r1;
add.s64 %rd11, %rd8, %rd10;
ld.const.u8 %r6, [%rd11];
mad.lo.s32 %r7, %r2, %r5, %r6;
cvt.u64.u32 %rd12, %r7;
add.s64 %rd13, %rd3, %rd12;
st.global.u8 [%rd13], %rs1;
ret;

}
//
.visible .entry _Z6Layer3PhS_(
.param .u64 _Z6Layer3PhS__param_0,
.param .u64 _Z6Layer3PhS__param_1
)
{
.reg .pred %p<5>;
.reg .b16 %rs<33>;
.reg .b32 %r<52>;
.reg .b64 %rd<24>;


ld.param.u64 %rd6, [_Z6Layer3PhS__param_0];
ld.param.u64 %rd5, [_Z6Layer3PhS__param_1];
mov.u32 %r21, %ntid.x;
mov.u32 %r1, %ctaid.x;
mul.lo.s32 %r49, %r1, %r21;
mov.u32 %r3, %tid.x;
add.s32 %r22, %r49, %r3;
cvt.u64.u32 %rd1, %r22;
cvta.to.global.u64 %rd2, %rd6;
add.s64 %rd3, %rd2, %rd1;
cvt.u16.u32 %rs8, %r3;
cvt.u16.u32 %rs9, %r1;
or.b16 %rs10, %rs9, %rs8;
ld.global.u8 %rs11, [%rd3];
xor.b16 %rs12, %rs11, %rs10;
st.global.u8 [%rd3], %rs12;
bar.sync 0;
and.b32 %r23, %r3, 7;
setp.ne.s32 %p1, %r23, 0;
@%p1 bra $L__BB2_4;

ld.global.u32 %r47, [%rd3+4];
ld.global.u32 %r48, [%rd3];
mov.u32 %r46, 1786956040;
mov.u32 %r45, 0;

$L__BB2_2:
.pragma "nounroll";
shl.b32 %r26, %r48, 4;
add.s32 %r27, %r26, 1386807340;
shr.u32 %r28, %r48, 5;
add.s32 %r29, %r28, 2007053320;
xor.b32 %r30, %r29, %r27;
add.s32 %r31, %r48, %r46;
xor.b32 %r32, %r30, %r31;
add.s32 %r47, %r32, %r47;
shl.b32 %r33, %r47, 4;
add.s32 %r34, %r33, 621668851;
add.s32 %r35, %r46, %r47;
xor.b32 %r36, %r34, %r35;
shr.u32 %r37, %r47, 5;
add.s32 %r38, %r37, -862448841;
xor.b32 %r39, %r36, %r38;
sub.s32 %r48, %r48, %r39;
add.s32 %r46, %r46, -1708609273;
add.s32 %r45, %r45, 1;
setp.ne.s32 %p2, %r45, 3238567;
@%p2 bra $L__BB2_2;

st.global.u32 [%rd3], %r48;
st.global.u32 [%rd3+4], %r47;

$L__BB2_4:
bar.sync 0;
and.b16 %rs16, %rs9, %rs8;
ld.global.u8 %rs17, [%rd3];
xor.b16 %rs18, %rs17, %rs16;
st.global.u8 [%rd3], %rs18;
bar.sync 0;
cvt.u64.u32 %rd7, %r3;
mov.u64 %rd8, cuda_sbox;
add.s64 %rd9, %rd8, %rd7;
ld.const.u8 %rs31, [%rd9];
cvta.to.global.u64 %rd4, %rd5;
mov.u16 %rs32, 0;
mov.u32 %r50, 0;
mov.u64 %rd14, cuda_tbox;

$L__BB2_5:
.pragma "nounroll";
cvt.u64.u32 %rd10, %r49;
add.s64 %rd11, %rd2, %rd10;
cvt.u64.u16 %rd12, %rs31;
and.b64 %rd13, %rd12, 255;
add.s64 %rd15, %rd14, %rd13;
ld.const.u8 %rs19, [%rd15];
ld.global.u8 %rs20, [%rd11];
mul.lo.s16 %rs21, %rs19, %rs20;
add.s16 %rs32, %rs21, %rs32;
mul.lo.s16 %rs22, %rs31, 5;
add.s16 %rs31, %rs22, 17;
add.s32 %r49, %r49, 1;
add.s32 %r50, %r50, 1;
setp.ne.s32 %p3, %r50, 256;
@%p3 bra $L__BB2_5;

xor.b32 %r18, %r1, %r3;
mov.u32 %r51, 8;

$L__BB2_7:
.pragma "nounroll";
shl.b16 %rs23, %rs32, 3;
and.b16 %rs24, %rs32, 224;
shr.u16 %rs25, %rs24, 5;
or.b16 %rs26, %rs25, %rs23;
cvt.u32.u16 %r42, %rs26;
mad.lo.s32 %r43, %r42, 13, %r18;
and.b32 %r44, %r51, 255;
cvt.u64.u32 %rd16, %r44;
add.s64 %rd18, %rd14, %rd16;
cvt.u16.u32 %rs27, %r43;
ld.const.u8 %rs28, [%rd18];
xor.b16 %rs29, %rs28, %rs27;
cvt.u64.u16 %rd19, %rs29;
and.b64 %rd20, %rd19, 255;
add.s64 %rd22, %rd8, %rd20;
ld.const.u8 %rs32, [%rd22];
add.s32 %r51, %r51, 1;
setp.ne.s32 %p4, %r51, 4137823;
@%p4 bra $L__BB2_7;

add.s64 %rd23, %rd4, %rd1;
st.global.u8 [%rd23], %rs32;
ret;

}

Layer1

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
.visible .entry _Z6Layer1PhS_(
.param .u64 _Z6Layer1PhS__param_0,
.param .u64 _Z6Layer1PhS__param_1
)
{
.reg .pred %p<6>;
.reg .b16 %rs<2>;
.reg .f32 %f<12>;
.reg .b32 %r<23>;
.reg .b64 %rd<15>;

// 加载内核参数（指向全局内存的指针）
ld.param.u64 %rd5, [_Z6Layer1PhS__param_0]; // %rd5 = input_ptr (第一个参数)
ld.param.u64 %rd6, [_Z6Layer1PhS__param_1]; // %rd6 = output_ptr (第二个参数)

// 获取线程和块 ID
mov.u32 %r1, %tid.x;   // %r1 = 当前线程在其块内的 x 轴 ID
mov.u32 %r2, %ctaid.x; // %r2 = 当前块在网格内的 x 轴 ID

// 基于线程和块 ID 的条件执行
setp.lt.u32 %p1, %r1, 241; // %p1 = 如果 %tid.x < 241，则为真
setp.lt.u32 %p2, %r2, 241; // %p2 = 如果 %ctaid.x < 241，则为真
and.pred %p3, %p1, %p2;    // %p3 = 如果 (%tid.x < 241) 且 (%ctaid.x < 241)，则为真

@%p3 bra $L__BB0_2; // 如果 %p3 为真，则跳转到 $L__BB0_2 (主计算部分)
bra.uni $L__BB0_1;   // 否则，无条件跳转到 $L__BB0_1 (跳过计算)

$L__BB0_2: // 主计算路径（如果线程/块 ID 在范围内）
mov.u32 %r3, %ntid.x;     // %r3 = 线程块在 x 轴的线程数
cvta.to.global.u64 %rd1, %rd5; // 将 input_ptr 转换为全局内存地址。%rd1 现在指向输入数据。
mov.f32 %f10, 0f00000000; // %f10 = 0.0 (浮点求和的累加器)
mov.u32 %r11, 0;          // %r11 = 循环计数器初始化
mov.u64 %rd8, cuda_motion; // %rd8 = cuda_motion 常量数组的基地址
mov.u32 %r20, %r11;       // %r20 = 外层循环计数器（初始化为 0）

$L__BB0_3: // 外层循环（迭代 16 次）
.pragma "nounroll"; // 编译器提示：不要展开此循环。
add.s32 %r13, %r20, %r2;    // %r13 = outer_loop_counter + block_id.x
shl.b32 %r14, %r20, 4;      // %r14 = outer_loop_counter * 16 (每个浮点数 4 字节，每个内层循环 16 个元素)
mov.u32 %r15, 240;          // %r15 = 240 (常量)
sub.s32 %r16, %r15, %r14;   // %r16 = 240 - (outer_loop_counter * 16)
mad.lo.s32 %r21, %r13, %r3, %r1; // %r21 = (outer_loop_counter + block_id.x) * num_threads_per_block + thread_id.x
                               // 这计算了访问输入数据的全局索引。
mul.wide.u32 %rd7, %r16, 4; // %rd7 = (240 - (outer_loop_counter * 16)) * 4。这计算了 cuda_motion 中的偏移量。
add.s64 %rd14, %rd8, %rd7;  // %rd14 = cuda_motion 的地址 + 计算出的偏移量。这是此内层循环中 cuda_motion 中权重的起始地址。
mov.u32 %r22, %r11;         // %r22 = 内层循环计数器（初始化为 0）

$L__BB0_4: // 内层循环（迭代 16 次）
.pragma "nounroll"; // 编译器提示：不要展开此循环。
cvt.u64.u32 %rd9, %r21;    // 将输入索引转换为 64 位以进行地址计算
add.s64 %rd10, %rd1, %rd9; // %rd10 = input_ptr + input_index。输入字节的地址。
ld.global.u8 %rs1, [%rd10]; // 从全局输入内存加载 1 字节到 %rs1 (16 位寄存器)
cvt.rn.f32.u16 %f7, %rs1;  // 将加载的字节（作为无符号 16 位）转换为 float32 (%f7) (四舍五入到最近的整数)
ld.const.f32 %f8, [%rd14]; // 从 cuda_motion 加载 4 字节（一个 float32）到 %f8 (这是一个权重)
fma.rn.f32 %f10, %f8, %f7, %f10; // 融合乘加： %f10 = (%f8 * %f7) + %f10。累加加权和。
add.s32 %r21, %r21, 1;     // 递增输入索引，用于下一次迭代
add.s64 %rd14, %rd14, 4;   // 递增 cuda_motion 地址 4 字节（浮点数大小）
add.s32 %r22, %r22, 1;     // 递增内层循环计数器
setp.ne.s32 %p4, %r22, 16; // %p4 = 如果 inner_loop_counter != 16，则为真
@%p4 bra $L__BB0_4;        // 如果 %p4 为真，继续内层循环

add.s32 %r20, %r20, 1;     // 递增外层循环计数器
setp.lt.u32 %p5, %r20, 16; // %p5 = 如果 outer_loop_counter < 16，则为真
@%p5 bra $L__BB0_3;        // 如果 %p5 为真，继续外层循环
bra.uni $L__BB0_6;         // 循环完成后，无条件跳转到 $L__BB0_6

$L__BB0_1: // 对于超出范围的线程/块的路径
mov.f32 %f10, 0f00000000; // 将累加器初始化为 0.0 (它们实际上不进行计算)

$L__BB0_6: // 计算或跳过后的通用代码路径
cvt.rzi.u32.f32 %r17, %f10; // 将累加的 float32 值 (%f10) 转换为无符号 32 位整数，向零舍入。
mov.u32 %r18, %ntid.x;      // %r18 = 线程块在 x 轴的线程数
mad.lo.s32 %r19, %r2, %r18, %r1; // %r19 = block_id.x * num_threads_per_block + thread_id.x (全局线程索引)
cvt.u64.u32 %rd11, %r19;    // 将全局线程索引转换为 64 位
cvta.to.global.u64 %rd12, %rd6; // 将 output_ptr 转换为全局内存地址。%rd12 指向输出数据。
add.s64 %rd13, %rd12, %rd11; // %rd13 = output_ptr + global_thread_index。写入输出的地址。
st.global.u8 [%rd13], %r17; // 将 %r17 的低 8 位（结果）存储到全局输出内存。
ret; // 从内核返回。
}

Layer2

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
.visible .entry _Z6Layer2PhS_(
.param .u64 _Z6Layer2PhS__param_0,
.param .u64 _Z6Layer2PhS__param_1
)
{
.reg .b16 %rs<2>;
.reg .b32 %r<8>;
.reg .b64 %rd<14>;

// 加载内核参数（指向全局内存的指针）
ld.param.u64 %rd1, [_Z6Layer2PhS__param_0]; // %rd1 = input_ptr
ld.param.u64 %rd2, [_Z6Layer2PhS__param_1]; // %rd2 = output_ptr

// 将指针转换为全局内存地址
cvta.to.global.u64 %rd3, %rd2; // %rd3 = global_output_ptr
cvta.to.global.u64 %rd4, %rd1; // %rd4 = global_input_ptr

// 获取线程和块 ID 并计算全局线程索引
mov.u32 %r1, %ctaid.x;   // %r1 = block_id.x
mov.u32 %r2, %ntid.x;   // %r2 = num_threads_per_block.x
mov.u32 %r3, %tid.x;    // %r3 = thread_id.x
mad.lo.s32 %r4, %r1, %r2, %r3; // %r4 = block_id.x * num_threads_per_block.x + thread_id.x (全局线程索引)

// 读取输入字节
cvt.u64.u32 %rd5, %r4;    // 将全局线程索引转换为 64 位
add.s64 %rd6, %rd4, %rd5; // %rd6 = global_input_ptr + global_thread_index。要读取的地址。
ld.global.u8 %rs1, [%rd6]; // 从全局输入内存加载 1 字节到 %rs1

// 执行 S 盒查找
cvt.u64.u32 %rd7, %r3;    // 将 thread_id.x 转换为 64 位以作为索引
mov.u64 %rd8, cuda_sbox;  // %rd8 = cuda_sbox 的基地址
add.s64 %rd9, %rd8, %rd7; // %rd9 = cuda_sbox_base + thread_id.x。用于 S 盒查找的地址。
ld.const.u8 %r5, [%rd9];  // 从 cuda_sbox 中以 thread_id.x 为索引加载 1 字节到 %r5

cvt.u64.u32 %rd10, %r1;   // 将 block_id.x 转换为 64 位以作为索引
add.s64 %rd11, %rd8, %rd10; // %rd11 = cuda_sbox_base + block_id.x。用于另一次 S 盒查找的地址。
ld.const.u8 %r6, [%rd11]; // 从 cuda_sbox 中以 block_id.x 为索引加载 1 字节到 %r6

// 计算输出地址偏移
mad.lo.s32 %r7, %r2, %r5, %r6; // %r7 = num_threads_per_block.x * S_box_val(thread_id.x) + S_box_val(block_id.x)
                               // 这计算了写入输出的最终偏移量。

// 写入输出
cvt.u64.u32 %rd12, %r7;   // 将计算出的偏移量转换为 64 位
add.s64 %rd13, %rd3, %rd12; // %rd13 = global_output_ptr + calculated_offset。最终输出地址。
st.global.u8 [%rd13], %rs1; // 将输入字节 (%rs1) 存储到计算出的输出地址。
ret; // 从内核返回。
}

Layer3

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
.visible .entry _Z6Layer3PhS_(
.param .u64 _Z6Layer3PhS__param_0,
.param .u64 _Z6Layer3PhS__param_1
)
{
.reg .pred %p<5>;
.reg .b16 %rs<33>;
.reg .b32 %r<52>;
.reg .b64 %rd<24>;

// 加载内核参数（指向全局内存的指针）
ld.param.u64 %rd6, [_Z6Layer3PhS__param_0]; // %rd6 = input/output_ptr (数据从这里读取并写入这里)
ld.param.u64 %rd5, [_Z6Layer3PhS__param_1]; // %rd5 = another output_ptr (用于最终写入)

// 获取线程和块 ID 并计算全局线程索引
mov.u32 %r21, %ntid.x;  // %r21 = num_threads_per_block.x
mov.u32 %r1, %ctaid.x;  // %r1 = block_id.x
mul.lo.s32 %r49, %r1, %r21; // %r49 = block_id.x * num_threads_per_block.x (块的基偏移量)
mov.u32 %r3, %tid.x;    // %r3 = thread_id.x
add.s32 %r22, %r49, %r3; // %r22 = 全局线程索引

// 计算初始读/写地址
cvt.u64.u32 %rd1, %r22;    // 将全局线程索引转换为 64 位
cvta.to.global.u64 %rd2, %rd6; // 将 input/output_ptr 转换为全局内存地址
add.s64 %rd3, %rd2, %rd1;  // %rd3 = global_input_output_ptr + global_thread_index。要操作的地址。

// 第一次 XOR 操作
cvt.u16.u32 %rs8, %r3;  // 将 thread_id.x 转换为 16 位
cvt.u16.u32 %rs9, %r1;  // 将 block_id.x 转换为 16 位
or.b16 %rs10, %rs9, %rs8; // %rs10 = (block_id.x OR thread_id.x) (16 位)
ld.global.u8 %rs11, [%rd3]; // 从 %rd3 的全局内存加载 1 字节到 %rs11
xor.b16 %rs12, %rs11, %rs10; // %rs12 = loaded_byte XOR (%rs10)
st.global.u8 [%rd3], %rs12; // 将结果存储回相同的全局内存位置 (%rd3)

bar.sync 0; // 同步块内的所有线程。这在读取可能已被块内其他线程写入的数据之前至关重要。

// 条件哈希类计算
and.b32 %r23, %r3, 7;   // %r23 = thread_id.x AND 7 (检查 if thread_id.x % 8 == 0)
setp.ne.s32 %p1, %r23, 0; // %p1 = 如果 (thread_id.x % 8 != 0)，则为真
@%p1 bra $L__BB2_4;     // 如果 %p1 为真，则跳转到 $L__BB2_4 (跳过哈希计算)

// 只有 thread_id.x % 8 == 0 的线程才执行此块
ld.global.u32 %r47, [%rd3+4]; // 从全局内存 %rd3 + 4 加载 4 字节到 %r47 (第二个 32 位字)
ld.global.u32 %r48, [%rd3];   // 从全局内存 %rd3 加载 4 字节到 %r48 (第一个 32 位字)
mov.u32 %r46, 1786956040;     // %r46 = 初始常量 (哈希状态的一部分)
mov.u32 %r45, 0;              // %r45 = 循环计数器（初始化为 0）

$L__BB2_2: // 哈希类循环（迭代 3,238,567 次！）
.pragma "nounroll"; // 编译器提示：不要展开此循环。
shl.b32 %r26, %r48, 4;      // %r26 = %r48 << 4
add.s32 %r27, %r26, 1386807340; // %r27 = (%r48 << 4) + constant1
shr.u32 %r28, %r48, 5;      // %r28 = %r48 >> 5
add.s32 %r29, %r28, 2007053320; // %r29 = (%r48 >> 5) + constant2
xor.b32 %r30, %r29, %r27;   // %r30 = %r29 XOR %r27
add.s32 %r31, %r48, %r46;   // %r31 = %r48 + %r46
xor.b32 %r32, %r30, %r31;   // %r32 = %r30 XOR %r31
add.s32 %r47, %r32, %r47;   // %r47 = %r32 + %r47 (更新第一个状态变量)

shl.b32 %r33, %r47, 4;      // %r33 = %r47 << 4
add.s32 %r34, %r33, 621668851;  // %r34 = (%r47 << 4) + constant3
add.s32 %r35, %r46, %r47;   // %r35 = %r46 + %r47
xor.b32 %r36, %r34, %r35;   // %r36 = %r34 XOR %r35
shr.u32 %r37, %r47, 5;      // %r37 = %r47 >> 5
add.s32 %r38, %r37, -862448841; // %r38 = (%r47 >> 5) + constant4
xor.b32 %r39, %r36, %r38;   // %r39 = %r36 XOR %r38
sub.s32 %r48, %r48, %r39;   // %r48 = %r48 - %r39 (更新第二个状态变量)

add.s32 %r46, %r46, -1708609273; // 更新下一个迭代的常量
add.s32 %r45, %r45, 1;      // 递增循环计数器
setp.ne.s32 %p2, %r45, 3238567; // %p2 = 如果 loop_counter != 3238567，则为真
@%p2 bra $L__BB2_2;         // 如果 %p2 为真，继续哈希循环

st.global.u32 [%rd3], %r48;     // 将更新后的第一个 32 位字存储回全局内存
st.global.u32 [%rd3+4], %r47;   // 将更新后的第二个 32 位字存储回全局内存

$L__BB2_4: // 哈希计算或跳过后的通用路径
bar.sync 0; // 同步块内的所有线程。在下一次读/写之前是必需的。

// 第二次 XOR 操作（与第一次类似）
and.b16 %rs16, %rs9, %rs8; // %rs16 = (block_id.x AND thread_id.x) (16 位)
ld.global.u8 %rs17, [%rd3]; // 从 %rd3 的全局内存加载 1 字节到 %rs17
xor.b16 %rs18, %rs17, %rs16; // %rs18 = loaded_byte XOR (%rs16)
st.global.u8 [%rd3], %rs18; // 将结果存储回相同的全局内存位置 (%rd3)

bar.sync 0; // 同步块内的所有线程。

// T 盒和 S 盒查找与累积乘法
cvt.u64.u32 %rd7, %r3;    // 将 thread_id.x 转换为 64 位以作为 S 盒索引
mov.u64 %rd8, cuda_sbox;  // %rd8 = cuda_sbox 的基地址
add.s64 %rd9, %rd8, %rd7; // %rd9 = cuda_sbox_base + thread_id.x
ld.const.u8 %rs31, [%rd9]; // 从 cuda_sbox 中以 thread_id.x 为索引加载 1 字节到 %rs31 (这是一个运行中的“状态”值)

cvta.to.global.u64 %rd4, %rd5; // 将 output_ptr (第二个参数) 转换为全局内存地址
mov.u16 %rs32, 0;         // %rs32 = 16 位和的累加器（初始化为 0）
mov.u32 %r50, 0;          // %r50 = 循环计数器（初始化为 0）
mov.u64 %rd14, cuda_tbox; // %rd14 = cuda_tbox 的基地址

$L__BB2_5: // T 盒/S 盒乘法循环（迭代 256 次）
.pragma "nounroll"; // 编译器提示：不要展开此循环。
cvt.u64.u32 %rd10, %r49;   // 将 block_base_offset (%r49 = block_id.x * num_threads_per_block.x) 转换为 64 位
add.s64 %rd11, %rd2, %rd10; // %rd11 = global_input_output_ptr + block_base_offset。这似乎在访问跨块的数据。
cvt.u64.u16 %rd12, %rs31;  // 将当前 S 盒状态 (%rs31) 转换为 64 位
and.b64 %rd13, %rd12, 255; // 掩码 %rd12 以获取低 8 位 (T 盒索引的 0-255 范围)
add.s64 %rd15, %rd14, %rd13; // %rd15 = cuda_tbox_base + (S-box_state & 0xFF)。T 盒查找的地址。
ld.const.u8 %rs19, [%rd15]; // 从 cuda_tbox 加载 1 字节到 %rs19 (T 盒值)

ld.global.u8 %rs20, [%rd11]; // 从 %rd11 的全局内存加载 1 字节 (似乎是块共享输入)
mul.lo.s16 %rs21, %rs19, %rs20; // %rs21 = T-box_value * loaded_byte (乘法，低 16 位)
add.s16 %rs32, %rs21, %rs32; // %rs32 = %rs21 + %rs32 (累积乘积)

mul.lo.s16 %rs22, %rs31, 5; // %rs22 = current_sbox_state * 5
add.s16 %rs31, %rs22, 17;   // %rs31 = %rs22 + 17 (更新 S 盒状态以进行下一次迭代)

add.s32 %r49, %r49, 1;     // 递增 %r49 (这意味着读取地址 %rd11 也递增)
add.s32 %r50, %r50, 1;     // 递增循环计数器
setp.ne.s32 %p3, %r50, 256; // %p3 = 如果 loop_counter != 256，则为真
@%p3 bra $L__BB2_5;        // 如果 %p3 为真，继续 T 盒/S 盒循环

// 基于累积值的最终转换
xor.b32 %r18, %r1, %r3; // %r18 = block_id.x XOR thread_id.x
mov.u32 %r51, 8;        // %r51 = 计数器，初始化为 8

$L__BB2_7: // 最终置换/替换循环（迭代 4,137,815 次！）
.pragma "nounroll"; // 编译器提示：不要展开此循环。
shl.b16 %rs23, %rs32, 3;   // %rs23 = accumulated_value << 3
and.b16 %rs24, %rs32, 224; // %rs24 = accumulated_value AND 224 (0b11100000)
shr.u16 %rs25, %rs24, 5;   // %rs25 = %rs24 >> 5
or.b16 %rs26, %rs25, %rs23; // %rs26 = %rs25 OR %rs23 (对 %rs32 的旋转类操作)
cvt.u32.u16 %r42, %rs26;   // 将 %rs26 转换为 32 位
mad.lo.s32 %r43, %r42, 13, %r18; // %r43 = %r42 * 13 + (%r18)
and.b32 %r44, %r51, 255;   // %r44 = counter AND 255 (使用计数器的低 8 位作为索引)
cvt.u64.u32 %rd16, %r44;   // 转换为 64 位
add.s64 %rd18, %rd14, %rd16; // %rd18 = cuda_tbox_base + (counter & 0xFF)。T 盒查找的地址。
cvt.u16.u32 %rs27, %r43;   // 将 %r43 转换为 16 位
ld.const.u8 %rs28, [%rd18]; // 从 cuda_tbox 加载 1 字节到 %rs28
xor.b16 %rs29, %rs28, %rs27; // %rs29 = loaded_T-box_value XOR %rs27

cvt.u64.u16 %rd19, %rs29;  // 将 %rs29 转换为 64 位
and.b64 %rd20, %rd19, 255; // 掩码 %rd19 以获取低 8 位 (S 盒索引的 0-255 范围)
add.s64 %rd22, %rd8, %rd20; // %rd22 = cuda_sbox_base + (%rs29 & 0xFF)。S 盒查找的地址。
ld.const.u8 %rs32, [%rd22]; // 从 cuda_sbox 加载 1 字节到 %rs32 (这是此迭代的最终输出字节)

add.s32 %r51, %r51, 1;     // 递增计数器
setp.ne.s32 %p4, %r51, 4137823; // %p4 = 如果 counter != 4137823，则为真
@%p4 bra $L__BB2_7;        // 如果 %p4 为真，继续最终循环

add.s64 %rd23, %rd4, %rd1; // %rd23 = global_output_ptr (第二个参数) + global_thread_index。最终输出地址。
st.global.u8 [%rd23], %rs32; // 将最终转换后的字节 (%rs32) 存储到全局输出内存。
ret; // 从内核返回。
}

分析推导
Layer1:二维卷积操作

1
2
3
4
ld.param.u64 %rd5, [_Z6Layer1PhS__param_0];  // 输入指针
ld.param.u64 %rd6, [_Z6Layer1PhS__param_1];  // 输出指针
setp.lt.u32 %p1, %r1, 241;  // 检查threadIdx.x < 241
setp.lt.u32 %p2, %r2, 241;  // 检查blockIdx.x < 241

边界检查，只处理前 241x241 的网格

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
mov.f32 %f10, 0f00000000;   // 累加器初始化为0.0
mov.u64 %rd8, cuda_motion;  // 加载motion矩阵地址

$L__BB0_3:  // 外层循环 (i=0-15)
add.s32 %r13, %r20, %r2;    // bid + i
shl.b32 %r14, %r20, 4;      // i*16
sub.s32 %r16, %r15, %r14;   // 240 - i*16 → motion起始索引

$L__BB0_4:  // 内层循环 (j=0-15)
ld.global.u8 %rs1, [%rd10];  // 加载输入字节
cvt.rn.f32.u16 %f7, %rs1;   // 字节转浮点
ld.const.f32 %f8, [%rd14];  // 加载motion权重
fma.rn.f32 %f10, %f8, %f7, %f10;  // 乘加累加
cvt.rzi.u32.f32 %r17, %f10;  // 浮点转整数(截断)
st.global.u8 [%rd13], %r17;  // 存储结果字节

还原：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
def layer1(input_data):
    block_dim = 256
    output = [0] * (256 * 256)
    
    for bid in range(241):  # 块索引限制
        for tid in range(241):  # 线程索引限制
            total = 0.0
            for i in range(16):  # 外层循环
                motion_idx = 240 - i * 16
                for j in range(16):  # 内层循环
                    input_idx = (bid + i) * block_dim + (tid + j)
                    total += cuda_motion[motion_idx + j] * input_data[input_idx]
            
            output[bid * block_dim + tid] = int(total) & 0xFF
    
    return output

Layer2: 置换操作，使用 sbox 对每个像素的位置进行重新排列

1
2
3
4
5
6
7
8
mad.lo.s32 %r4, %r1, %r2, %r3;  // 计算输入索引: blockIdx.x*blockDim.x + threadIdx.x
ld.global.u8 %rs1, [%rd6];       // 加载输入字节

ld.const.u8 %r5, [%rd9];  // sbox[threadIdx.x]
ld.const.u8 %r6, [%rd11]; // sbox[blockIdx.x]

mad.lo.s32 %r7, %r2, %r5, %r6;  // 计算新位置: blockDim.x * sbox[tid] + sbox[bid]
st.global.u8 [%rd13], %rs1;      // 存储到新位置

还原：

1
2
3
4
5
6
7
8
9
10
11
def layer2(input_data):
    block_dim = 256
    output = [0] * (256 * 256)
    
    for bid in range(256):
        for tid in range(256):
            input_idx = bid * block_dim + tid
            new_idx = block_dim * cuda_sbox[tid] + cuda_sbox[bid]
            output[new_idx] = input_data[input_idx]
    
    return output

Layer3: 核心加密操作
1.初次异或

1
2
3
4
or.b16 %rs10, %rs9, %rs8;  // bid | tid
xor.b16 %rs12, %rs11, %rs10;  // buffer ^= (bid | tid)
st.global.u8 [%rd3], %rs12;
bar.sync 0;  // 同步

推导：

1
2
3
4
for bid in range(256):
    for tid in range(256):
        idx = bid * 256 + tid
        buffer[idx] ^= bid | tid

2.类似 TEA 的加密

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
and.b32 %r23, %r3, 7;      // tid & 7
setp.ne.s32 %p1, %r23, 0;  // 检查是否执行加密
@%p1 bra $L__BB2_4;        // 条件跳转

// 加密核心
ld.global.u32 %r47, [%rd3+4];  // 加载高位字
ld.global.u32 %r48, [%rd3];    // 加载低位字
mov.u32 %r46, 1786956040;      // 初始常量

$L__BB2_2:  // 加密循环 (3,238,567轮)
// 更新高位字
shl.b32 %r26, %r48, 4;        // v0 << 4
add.s32 %r27, %r26, 1386807340; 
shr.u32 %r28, %r48, 5;        // v0 >> 5
add.s32 %r29, %r28, 2007053320;
xor.b32 %r30, %r29, %r27;     // (v0>>5 + C2) ^ (v0<<4 + C1)
add.s32 %r31, %r48, %r46;     // v0 + sum
xor.b32 %r32, %r30, %r31;     // 二次异或
add.s32 %r47, %r32, %r47;     // v1 += ...

// 更新低位字
shl.b32 %r33, %r47, 4;        // v1 << 4
add.s32 %r34, %r33, 621668851;
add.s32 %r35, %r46, %r47;     // sum + v1
xor.b32 %r36, %r34, %r35;     // (v1<<4 + C3) ^ (sum+v1)
shr.u32 %r37, %r47, 5;        // v1 >> 5
add.s32 %r38, %r37, -862448841;
xor.b32 %r39, %r36, %r38;     // 二次异或
sub.s32 %r48, %r48, %r39;     // v0 -= ...

// 更新常量
add.s32 %r46, %r46, -1708609273;  // sum += delta

推导：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
or bid in range(256):
    for tid in range(0, 256, 8):
        idx = bid * 256 + tid
        # 读取8字节(2个32位字)
        v0 = int.from_bytes(buffer[idx:idx+4], 'little')
        v1 = int.from_bytes(buffer[idx+4:idx+8], 'little')
        s = 1786956040
        
        for _ in range(3238567):
            # 更新v1
            t1 = ((v0 << 4) + 1386807340) & 0xFFFFFFFF
            t2 = ((v0 >> 5) + 2007053320) & 0xFFFFFFFF
            t3 = (t1 ^ t2) & 0xFFFFFFFF
            v1 = (v1 + (t3 ^ (v0 + s))) & 0xFFFFFFFF
            
            # 更新v0
            t1 = ((v1 << 4) + 621668851) & 0xFFFFFFFF
            t2 = ((v1 >> 5) - 862448841) & 0xFFFFFFFF
            t3 = (t1 ^ (s + v1)) & 0xFFFFFFFF
            v0 = (v0 - (t3 ^ t2)) & 0xFFFFFFFF
            
            # 更新常量
            s = (s - 1708609273) & 0xFFFFFFFF
        
        # 写回结果
        buffer[idx:idx+4] = v0.to_bytes(4, 'little')
        buffer[idx+4:idx+8] = v1.to_bytes(4, 'little')

3.二次异或：

1
2
3
4
and.b16 %rs16, %rs9, %rs8;  // bid & tid
xor.b16 %rs18, %rs17, %rs16;  // buffer ^= (bid & tid)
st.global.u8 [%rd3], %rs18;
bar.sync 0;  // 同步

推导：

1
2
3
for tid in range(256):
    idx = bid * 256 + tid
    buffer[idx] ^= bid & tid

4.T 盒/S 盒 累加

1
2
3
4
5
6
7
8
9
10
11
12
ld.const.u8 %rs31, [%rd9];  // sbox_val = sbox[tid]
mov.u16 %rs32, 0;           // 累加器初始化

$L__BB2_5:  // 累加循环 (256次)
ld.const.u8 %rs19, [%rd15];  // tbox[sbox_val]
ld.global.u8 %rs20, [%rd11]; // 加载输入字节
mul.lo.s16 %rs21, %rs19, %rs20;  // tbox[sbox] * buffer[i]
add.s16 %rs32, %rs21, %rs32;     // 累加

// 更新S盒值
mul.lo.s16 %rs22, %rs31, 5;
add.s16 %rs31, %rs22, 17;   // sbox_val = sbox_val * 5 + 17

推导：

1
2
3
4
5
6
7
8
9
10
11
12
13
for bid in range(256):
    for tid in range(256):
        s_val = cuda_sbox[tid]
        accum = 0
        
        # 处理整个块
        for i in range(256):
            idx = bid * 256 + i
            accum += cuda_tbox[s_val] * buffer[idx]
            s_val = (s_val * 5 + 17) & 0xFF
        
        # 保留16位结果
        accum &= 0xFFFF

5.混淆循环

1
2
3
4
5
6
7
8
9
10
11
12
$L__BB2_7:  // 混淆循环 (4,137,815轮)
// 循环左移3位
shl.b16 %rs23, %rs32, 3;     // 低13位左移3位
and.b16 %rs24, %rs32, 224;   // 0xE0 (高3位)
shr.u16 %rs25, %rs24, 5;     // 高3位右移5位
or.b16 %rs26, %rs25, %rs23;  // 组合 → 循环左移3位

// 非线性变换
mad.lo.s32 %r43, %r42, 13, %r18;  // rotated*13 + (bid^tid)
ld.const.u8 %rs28, [%rd18];  // tbox[i & 0xFF]
xor.b16 %rs29, %rs28, %rs27; // tbox[i] ^ LSB(t)
ld.const.u8 %rs32, [%rd22];  // sbox[result] → 更新累加器

推导：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
for bid in range(256):
    for tid in range(256):
        mixer = bid ^ tid
        
        # 混淆循环 (从8开始)
        for i in range(8, 4137823):
            # 16位循环左移3位
            rotated = ((accum << 3) | (accum >> 13)) & 0xFFFF
            
            # 乘加运算
            t_val = rotated * 13 + mixer
            
            # T盒/S盒变换
            tbox_val = cuda_tbox[i & 0xFF]
            s_idx = tbox_val ^ (t_val & 0xFF)
            accum = cuda_sbox[s_idx]
        
        output[bid * 256 + tid] = accum & 0xFF

完整加密过程

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
def layer1(input_data):
    block_dim = 256
    output = [0] * (256 * 256)
    
    for bid in range(241):
        for tid in range(241):
            total = 0.0
            for i in range(16):
                motion_idx = 240 - i * 16
                for j in range(16):
                    input_idx = (bid + i) * block_dim + (tid + j)
                    total += cuda_motion[motion_idx + j] * input_data[input_idx]
            
            output[bid * block_dim + tid] = int(total) & 0xFF
    
    return output

def layer2(input_data):
    block_dim = 256
    output = [0] * (256 * 256)
    
    for bid in range(256):
        for tid in range(256):
            input_idx = bid * block_dim + tid
            new_idx = block_dim * cuda_sbox[tid] + cuda_sbox[bid]
            output[new_idx] = input_data[input_idx]
    
    return output

def layer3(input_data):
    block_dim = 256
    buffer = bytearray(input_data)  
    output = [0] * (256 * 256)
    

    for bid in range(256):
        for tid in range(256):
            idx = bid * block_dim + tid
            buffer[idx] ^= bid | tid
    

    for bid in range(256):
        for tid in range(0, 256, 8):
            idx = bid * block_dim + tid

            v0 = int.from_bytes(buffer[idx:idx+4], 'little')
            v1 = int.from_bytes(buffer[idx+4:idx+8], 'little')
            s = 1786956040
            
            for _ in range(3238567):

                t1 = ((v0 << 4) + 1386807340) & 0xFFFFFFFF
                t2 = ((v0 >> 5) + 2007053320) & 0xFFFFFFFF
                t3 = (t1 ^ t2) & 0xFFFFFFFF
                v1 = (v1 + (t3 ^ (v0 + s))) & 0xFFFFFFFF
                
                t1 = ((v1 << 4) + 621668851) & 0xFFFFFFFF
                t2 = ((v1 >> 5) - 862448841) & 0xFFFFFFFF
                t3 = (t1 ^ (s + v1)) & 0xFFFFFFFF
                v0 = (v0 - (t3 ^ t2)) & 0xFFFFFFFF
                
                s = (s - 1708609273) & 0xFFFFFFFF
            
            buffer[idx:idx+4] = v0.to_bytes(4, 'little')
            buffer[idx+4:idx+8] = v1.to_bytes(4, 'little')
    
    for bid in range(256):
        for tid in range(256):
            idx = bid * block_dim + tid
            buffer[idx] ^= bid & tid
    
    for bid in range(256):
        for tid in range(256):
            s_val = cuda_sbox[tid]
            accum = 0
            
            for i in range(256):
                idx = bid * block_dim + i
                accum = (accum + cuda_tbox[s_val] * buffer[idx]) & 0xFFFF
                s_val = (s_val * 5 + 17) & 0xFF
    
    for bid in range(256):
        for tid in range(256):
            mixer = bid ^ tid
            
            for i in range(8, 4137823):
                rotated = ((accum << 3) | (accum >> 13)) & 0xFFFF
                
                t_val = rotated * 13 + mixer
                
                tbox_val = cuda_tbox[i & 0xFF]
                s_idx = tbox_val ^ (t_val & 0xFF)
                accum = cuda_sbox[s_idx]

            output[bid * block_dim + tid] = accum & 0xFF
    
    return output

def process_data(input_data):
    l1 = layer1(input_data)
    l2 = layer2(l1)
    l3 = layer3(l2)
    return l3

接下来就是从 Layer3 开始逐层倒着解密
详细参考
棱晶 の wp

CTF
#Reverse #cuda
cuda 逆向
http://example.com/2025/07/21/cuda/
作者
Eleven
发布于
2025年7月21日
许可协议