首先呢这份代码是这学期选修课的大作业，前后也花了点心思，于是还是贴到这儿，顺便给博客除除草 ww
作业的题目有两大类，一类是使用常见的对称/非对称加密算法（DES/AES/RC5/RSA/Blowfish/blabla）实现一个可用的加解密工具，另一个则是实现一个 HASH 程序，类似于 md5sum 或者是 sha256sum 这种。我选择了做 AES 主要是有几个原因：
- 老师说了，杂凑算法于加解密算法比较起来，实现难度低，因此基础成绩就会相对低一点（然而我觉得两个都蛮简单的啊）。
- 在上课之前就已经接触过 AES 这个算法了（归功于 Shadowsocks 啦），用了这么久的东西，也算是对它的实现比较感兴趣的（然而怎么想都知道自己写的破代码性能会被 OpenSSL 吊着打）。
当然按照老师的说法，简单的实现只是最基本的要求，所以我这里打算做三分实现，分别是 CPU 单线程，CPU 多线程和 GPGPU。其实这里面核心代码都是完全一样的，无非就是分组加密和密钥更新上有点区别，但是听起来就高大上了很多（没毛病）。
AES 本身一次只能使用 128bit 的密钥对 128bit 的数据进行加密，所以在实际中必须使用分组加密算法来完成大文件的处理。往简单里讲，分组加密不过是把数据分成等长的分片，然后每次对一个分片进行加密，然后在将加密后的分片组合成文件。解密过程与之类似。这种最简单粗暴地分组加密被称为电子密码本（ECB）模式。然而实际上，由于相同分片数据使用相同的密钥进行加密后，输出必然也是相同的，这样也就会使加密的内容带有特 征，存在潜在的隐患，于是人们提出一些会修改后续分组的密钥的分组加密算法，比如密码分组链接模式（CBC）、密文反馈模式（CFB）和计数器模式（CTR）。CBC 和 CFB 在设计上，下一分片的密钥受到上一分片的影响，也就意味着这个算法必须线性的对密码进行加密，不利于多线程的实现。那么可选的就是 CTR 和 ECB 模式了。二选一的话，选择 ECB 而不是 CTR 主要是因为我懒，ECB 只要一个函数改变传入指针就能跑了（
分区模式敲定之后，就轮到 AES 算法的了解和实现了。整个算法主要由两部分组成：数据变换和密钥拓展。
首先实现一份密钥拓展函数，其算法参考：Rijndael key schedule – Wikipedia，使用上很简单，直接将密钥指针传入即可，返回的是堆上分配的结果，用完记得要 free（（
extern uint8_t round_const[];
uint8_t *key_schedule(const uint8_t *input_key) {
uint8_t *output = (uint8_t *)malloc((10 + 1) * 16 * sizeof(uint8_t));
uint8_t temp[4], *cursor_4_before, *cursor_16_before, *cursor_current;
size_t counter, i, j;
// Copy the first 16 bytes
memcpy(output, input_key, 16 * sizeof(uint8_t));
// Initialize cursors.
cursor_current = output + 16;
cursor_4_before = cursor_current - 4;
cursor_16_before = cursor_current - 16;
// Rock and roll.
for (counter = 0; counter != 10; counter++) {
// RotWord
temp[3] = cursor_4_before[0];
temp[0] = cursor_4_before[1];
temp[1] = cursor_4_before[2];
temp[2] = cursor_4_before[3];
// SubBytes and add
for (i = 0; i != 4; ++i) {
cursor_current[i] = sbox[temp[i]] ^ cursor_16_before[i];
if (i == 0) {
cursor_current[i] ^= round_const[counter];
}
}
// Forward all cursors
cursor_current += 4, cursor_4_before += 4, cursor_16_before += 4;
// Three step remaining is simple.
for (i = 0; i != 3; ++i) {
for (j = 0; j != 4; ++j) {
cursor_current[j] = cursor_4_before[j] ^ cursor_16_before[j];
}
cursor_current += 4, cursor_4_before += 4, cursor_16_before += 4;
}
}
return output;
}
AES 128 会对输入数据进行十轮处理，其中 AddRoundKey 步骤使用的密钥就是由上面函数生成的。
AES 128 的十轮处理中，必然会有的四个步骤：字节替换、行移位、列混淆和密钥相加（XOR）。
字节替换提供了加密法非线性的变换能力，实现起来相当容易：将每一个元素代入 S 盒中即可：
extern uint8_t sbox[];
void sub_bytes(uint8_t *input) {
int i = 15;
do {
input[i] = sbox[input[i]];
} while (--i >= 0);
}
这里补充一下，AES 对输入的明文和密钥的处理有一点反人类，比如明文 {a0, a1, a2, a3, …, a15}，我们要把它看成这样的一个矩阵：
a0 a4 a8 a12
a1 a5 a9 a13
a2 a6 a10 a14
a3 a7 a11 a15
所以上面给出的 OFFSET 宏乍一看比较反人类。。对应的，行移位后的矩阵应该是这样的：
a0 a4 a8 a12
a5 a9 a13 a1
a10 a14 a2 a6
a15 a3 a7 a11
在内存里也就是这样的数组： {a0, a5, a10, a15, a4, a9, a14, a3, …}
剩下两个函数就是列混淆和密钥相加。其中列混淆可以视为 Rijndael 有限域之下的矩阵乘法，详情可以看 Wikipedia，或者是翻翻书。这里我提前打了个表，来加速完成。变量名叫 cache 是本来想实时计算缓存结果的，最后还是干脆直接打全了。。
static inline uint8_t mul2(uint8_t a) {
static uint8_t cache[256] = {
0x00, 0x02, 0x04, 0x06, 0x08, 0x0A, 0x0C, 0x0E, 0x10, 0x12, 0x14, 0x16, 0x18, 0x1A, 0x1C, 0x1E,
0x20, 0x22, 0x24, 0x26, 0x28, 0x2A, 0x2C, 0x2E, 0x30, 0x32, 0x34, 0x36, 0x38, 0x3A, 0x3C, 0x3E,
0x40, 0x42, 0x44, 0x46, 0x48, 0x4A, 0x4C, 0x4E, 0x50, 0x52, 0x54, 0x56, 0x58, 0x5A, 0x5C, 0x5E,
0x60, 0x62, 0x64, 0x66, 0x68, 0x6A, 0x6C, 0x6E, 0x70, 0x72, 0x74, 0x76, 0x78, 0x7A, 0x7C, 0x7E,
0x80, 0x82, 0x84, 0x86, 0x88, 0x8A, 0x8C, 0x8E, 0x90, 0x92, 0x94, 0x96, 0x98, 0x9A, 0x9C, 0x9E,
0xA0, 0xA2, 0xA4, 0xA6, 0xA8, 0xAA, 0xAC, 0xAE, 0xB0, 0xB2, 0xB4, 0xB6, 0xB8, 0xBA, 0xBC, 0xBE,
0xC0, 0xC2, 0xC4, 0xC6, 0xC8, 0xCA, 0xCC, 0xCE, 0xD0, 0xD2, 0xD4, 0xD6, 0xD8, 0xDA, 0xDC, 0xDE,
0xE0, 0xE2, 0xE4, 0xE6, 0xE8, 0xEA, 0xEC, 0xEE, 0xF0, 0xF2, 0xF4, 0xF6, 0xF8, 0xFA, 0xFC, 0xFE,
0x1B, 0x19, 0x1F, 0x1D, 0x13, 0x11, 0x17, 0x15, 0x0B, 0x09, 0x0F, 0x0D, 0x03, 0x01, 0x07, 0x05,
0x3B, 0x39, 0x3F, 0x3D, 0x33, 0x31, 0x37, 0x35, 0x2B, 0x29, 0x2F, 0x2D, 0x23, 0x21, 0x27, 0x25,
0x5B, 0x59, 0x5F, 0x5D, 0x53, 0x51, 0x57, 0x55, 0x4B, 0x49, 0x4F, 0x4D, 0x43, 0x41, 0x47, 0x45,
0x7B, 0x79, 0x7F, 0x7D, 0x73, 0x71, 0x77, 0x75, 0x6B, 0x69, 0x6F, 0x6D, 0x63, 0x61, 0x67, 0x65,
0x9B, 0x99, 0x9F, 0x9D, 0x93, 0x91, 0x97, 0x95, 0x8B, 0x89, 0x8F, 0x8D, 0x83, 0x81, 0x87, 0x85,
0xBB, 0xB9, 0xBF, 0xBD, 0xB3, 0xB1, 0xB7, 0xB5, 0xAB, 0xA9, 0xAF, 0xAD, 0xA3, 0xA1, 0xA7, 0xA5,
0xDB, 0xD9, 0xDF, 0xDD, 0xD3, 0xD1, 0xD7, 0xD5, 0xCB, 0xC9, 0xCF, 0xCD, 0xC3, 0xC1, 0xC7, 0xC5,
0xFB, 0xF9, 0xFF, 0xFD, 0xF3, 0xF1, 0xF7, 0xF5, 0xEB, 0xE9, 0xEF, 0xED, 0xE3, 0xE1, 0xE7, 0xE5,
};
return cache[a];
}
void mix_columns(uint8_t *input) {
size_t i;
uint8_t temp[16], tmp;
memcpy(temp, input, 16);
for (i = 0; i != 16; i += 4) {
tmp = temp[i] ^ temp[i + 1] ^ temp[i + 2] ^ temp[i + 3];
input[i + 0] = mul2(temp[i + 0] ^ temp[i + 1]) ^ temp[i + 0] ^ tmp;
input[i + 1] = mul2(temp[i + 1] ^ temp[i + 2]) ^ temp[i + 1] ^ tmp;
input[i + 2] = mul2(temp[i + 2] ^ temp[i + 3]) ^ temp[i + 2] ^ tmp;
input[i + 3] = mul2(temp[i + 3] ^ temp[i + 0]) ^ temp[i + 3] ^ tmp;
}
}
最后是密钥相加。太简单了对应位置 Xor 就行了，甚至不用单独写个函数。将这以上四部分组合起来之后，只要写一个函数对明文进行十轮加密，函数大概这样：
void aes_128_single_block(const uint8_t *input, uint8_t *round_keys, uint8_t *output) {
size_t i, j;
uint8_t *result, *cursor_round_key;
// initialize cursors.
result = output;
cursor_round_key = round_keys;
// Copy plain text.
for (i = 0; i != 16; ++i) {
result[i] = input[i] ^ cursor_round_key[i];
}
cursor_round_key += 16;
// Run rounds excluding the last round.
for (i = 0; i != 9; ++i) {
sub_bytes(result);
shift_rows(result);
mix_columns(result);
// Add round
for (j = 0; j != 16; ++j) {
result[j] = result[j] ^ cursor_round_key[j];
}
cursor_round_key += 16;
}
// Now the last round.
sub_bytes(result);
shift_rows(result);
for (i = 0; i != 16; ++i) {
result[i] = result[i] ^ cursor_round_key[i];
}
}
其中 round_keys 就是最开头给的函数生成的带拓展的密钥。
完成以上的加密相关代码后，只要小幅改动部分代码就能实现密文的解密了。比如 AES 里四个步骤的对应函数，除了轮密钥相加以外，其他三个函数都需要编写对应的反函数（姑且这么叫吧）。至于轮密钥相加以外，本质工作就是将轮密钥和体（state）进行异或，根据异或的性质，两次异或直接能就是本身，所以没有再去折腾了。
void aes::inverse_sub_bytes(byte *input) {
int i = 15;
do {
input[i] = inverse_sbox[input[i]];
} while (--i >= 0);
}
void inverse_shift_rows(byte *input) {
int index, inner_index;
#define OFFSET(i, j) ((j) * 4 + (i))
for (index = 0; index != 4; ++index) {
byte temp;
inner_index = index;
while (inner_index-- > 0) {
temp = input[OFFSET(index, 3)];
input[OFFSET(index, 3)] = input[OFFSET(index, 2)];
input[OFFSET(index, 2)] = input[OFFSET(index, 1)];
input[OFFSET(index, 1)] = input[OFFSET(index, 0)];
input[OFFSET(index, 0)] = temp;
}
}
#undef OFFSET
}
void aes::inverse_mix_columns(byte *input) {
size_t i;
byte temp[16], tmp;
memcpy(temp, input, 16);
for (i = 0; i != 16; i += 4) {
tmp = temp[i] ^ temp[i + 1] ^ temp[i + 2] ^ temp[i + 3];
input[i + 0] = mul2(temp[i + 0] ^ temp[i + 1]) ^ temp[i + 0] ^ tmp;
input[i + 1] = mul2(temp[i + 1] ^ temp[i + 2]) ^ temp[i + 1] ^ tmp;
input[i + 2] = mul2(temp[i + 2] ^ temp[i + 3]) ^ temp[i + 2] ^ tmp;
input[i + 3] = mul2(temp[i + 3] ^ temp[i + 0]) ^ temp[i + 3] ^ tmp;
byte u, v;
u = mul2(mul2(temp[i + 0] ^ temp[i + 2]));
v = mul2(mul2(temp[i + 1] ^ temp[i + 3]));
tmp = mul2(v ^ u);
input[i + 0] ^= tmp ^ u;
input[i + 2] ^= tmp ^ u;
input[i + 1] ^= tmp ^ v;
input[i + 3] ^= tmp ^ v;
}
}
最后的解密过程就是加密过程反着来，写起来也没什么困难：
void aes::aes_decrypt_single_block(const byte *input, byte * const round_keys, byte *output) {
size_t i, j;
byte *result, *cursor_round_key;
// initialize cursors.
result = output;
cursor_round_key = round_keys + 160;
for (i = 0; i != 16; ++i) {
result[i] = input[i];
}
// The last round
for (i = 0; i != 16; ++i) {
result[i] = input[i] ^ cursor_round_key[i];
}
inverse_shift_rows(result);
inverse_sub_bytes(result);
cursor_round_key -= 16;
#ifdef DEBUG
printf("# InvRound %d\n", 10);
print_key(result);
#endif
// Run rounds excluding the last round.
for (i = 0; i != 9; ++i) {
for (j = 0; j != 16; ++j) {
result[j] = result[j] ^ cursor_round_key[j];
}
inverse_mix_columns(result);
inverse_shift_rows(result);
inverse_sub_bytes(result);
cursor_round_key -= 16;
#ifdef DEBUG
printf("# InvRound %d\n", 10 - i);
print_key(result);
#endif
}
// Now the last addRoundKey.
for (i = 0; i != 16; ++i) {
result[i] = result[i] ^ cursor_round_key[i];
}
}
以上就是 AES 的所有核心代码了，剩下的就是如何去调用这些函数。多线程方面，依靠 C++ 的 STL 库能很容易地实现线程的创建和锁控制。首先准备一个函数，将 key_schedule 和 aes 加密/解密函数包裹在一起。再准备一个函数按照指定步长去对输入指针内容进行加密/解密操作，代码很简单：
void process_unit(uint8_t *data_in, uint8_t *data_out, uint8_t *key, bool isEncrypt) {
aes_cpu::byte *rkey = aes_cpu::aes::key_schedule(key);
if (isEncrypt) {
aes_cpu::aes::aes_128_single_block(data_in, rkey, data_out);
}
else {
aes_cpu::aes::aes_decrypt_single_block(data_in, rkey, data_out);
}
delete[] rkey;
}
void thread_entry(uint8_t *data_in, uint8_t *data_out, uint8_t *key, size_t file_size, size_t current, size_t total, bool isEncrypt) {
for (uint32_t i = current; i < file_size / 16; i += total) {
if (i % 10000 == 0 && current == 0) {
printf("process: %.2lf\r", (double)i * 100 / (file_size / 16));
}
process_unit(data_in + 16 * i, data_out + 16 * i, key, isEncrypt);
}
}
现在只需要根据用户传入的线程数，创建若干线程即可完成工作：
bool isEncrypt = true;
std::vector<std::thread> threads(thread_count);
// Construct all threads, and let them go.
for (int i = 0; i != threads.size(); ++i) {
threads[i] = std::thread(thread_entry, data_in, data_out, key, full_block_size, i, threads.size(), isEncrypt);
}
// Sync for all threads.
for (int i = 0; i != threads.size(); ++i) {
threads[i].join();
}
涉及到任意文件加的时候就会遇到一个问题：文件大小不一定是 128 比特的整数倍。这时候就需要使用一些密码学上的填充算法。我简单的看了一下 PKCS#7 下的填充方法，简单说来就是：
- 若文件恰好为 16 字节的整数倍，后面强行补十六位；
- 填充内容为当前大小与填充后大小的差，比如 当前是 126 字节，就补充两个 0x02。对于上方的特例，就补充 16 个 0x16；
这里就不把函数列出来了，一个 realloc 一个 for 就能解决。这样我们的多线程 AES 加解密就已经全部完成了，接下来就是将以上代码迁移到 CUDA 平台。
从整体上看，CUDA 和 CPU 的代码基本一致（废话都是 AES 能不一样有鬼了），主要的区别就在于调用 GPU 是一个类似于 Master 和 Slave 通讯的过程。在完成 PKCS#7 的填充后，我们要把明文，密钥等复制到显存，同时在显存里分配一块缓冲区来保存密文，代码大致如下：
uint8_t *round_keys = key_schedule(keys);
uint8_t *gpu_input = nullptr;
uint8_t *gpu_round_keys = nullptr;
uint8_t *gpu_output = nullptr;
cudaError_t cudaStatus;
#define ASSERT(s) \
if (cudaStatus != cudaSuccess) { \
fprintf(stderr, s"\n"); \
goto Error; \
}
// Select which GPU to work with.
cudaStatus = cudaSetDevice(0);
ASSERT("cudaSetDevice");
// Prepare GPU memory block
cudaStatus = cudaMalloc((void**)&gpu_input, size * sizeof(uint8_t));
ASSERT("cudaMalloc");
cudaStatus = cudaMalloc((void**)&gpu_round_keys, (11 * 16) * sizeof(uint8_t));
ASSERT("cudaMalloc");
cudaStatus = cudaMalloc((void**)&gpu_output, size * sizeof(uint8_t));
ASSERT("cudaMalloc");
// Copy data from memory to GPU memory.
cudaStatus = cudaMemcpy(gpu_input, input, size * sizeof(uint8_t), cudaMemcpyHostToDevice);
ASSERT("cudaMemcpy: gpu_input");
cudaStatus = cudaMemcpy(gpu_round_keys, round_keys, (11 * 16) * sizeof(uint8_t), cudaMemcpyHostToDevice);
ASSERT("cudaMemcpy: gpu_round_keys");
然后就可以启动 GPU 上的函数了：
size_t offset = 0;
const size_t threadsPerBlock = 512;
const size_t blockSize = 16;
while (offset * threadsPerBlock * blockSize * 16 < size) {
if (offset % 100 == 0) {
printf("progress: %lf\r", offset * (double)threadsPerBlock * blockSize * 1600 / size);
}
if (isEncrypt) {
aes_128_single_block <<<blockSize, threadsPerBlock >>>(gpu_input, gpu_round_keys, gpu_output, size / 16, offset, threadsPerBlock * blockSize);
}
else {
aes_decrypt_single_block <<<blockSize, threadsPerBlock >>>(gpu_input, gpu_round_keys, gpu_output, size / 16, offset, threadsPerBlock * blockSize);
}
offset++;
}
最后，把显存上的密文复制回内存，再写入到文件，就完成了整个加密过程：
cudaStatus = cudaGetLastError();
ASSERT("cudaGetLastError");
cudaStatus = cudaDeviceSynchronize();
ASSERT("cudaDeviceSynchronize");
cudaStatus = cudaMemcpy(output, gpu_output, size * sizeof(uint8_t), cudaMemcpyDeviceToHost);
ASSERT("cudaMemcpy");