我正在尝试在CUDA C设备代码中实现一些AES函数.在尝试实现左侧按字节旋转运算符时,我感到不安的是看到没有原生的SIMD内向.所以我开始了一个天真的实现,但……它是巨大的,虽然我还没有尝试过,但由于昂贵的拆包/打包,它不会很快……所以,有什么意思吗每字节位旋转操作至少有些效率？
如果您不想看看,这是代码.
__inline__ __device__ uint32_t per_byte_bit_left_rotate(uint32_t input,uint8_t amount) {
return ((((input & 0xFF) >> 0) << amount) | (((input & 0xFF) >> 0) >> 7) & ~0x100) << 0 |
((((input & 0xFF00) >> 8) << amount) | ((input & 0xFF00 >> 8) >> 7) & ~0x100) << 8 |
((((input & 0xFF0000) >> 16) << amount) | ((input & 0xFF0000 >> 16) >> 7) & ~0x100) << 16 |
((((input & 0xFF000000) >> 24) << amount) | ((input & 0xFF000000 >> 24) >> 7) & ~0x100) << 24; } // The XORs are for clearing the old 7th bit who is getting pushed to the next byte of the intermediate int
解决方法
左移和右移整个输入,然后对一个AND中的所有4个字节的那些掩码使所有跨越字节边界的位归零.我认为数量始终是AES中的编译时常量,因此您不必担心动态生成掩码的运行时成本.让编译器去做吧. (IDK CUDA,但这似乎与为正常C写入具有32位整数的SWAR bit-hack相同的问题)
这是基于通常的(x << count) | (x >> (32-count)) rotate idiom,具有遮蔽和不同的右移计数,使其成为单独的8位旋转.
inline
uint32_t per_byte_bit_left_rotate(uint32_t input,unsigned amount)
{
// With constant amount,the left/right masks are constants
uint32_t rmask = 0xFF >> ((8 - amount) & 7);
rmask = (rmask<<24 | rmask<<16 | rmask<<8 | rmask);
uint32_t lmask = ~rmask;
uint32_t lshift = input << amount;
lshift &= lmask;
if (amount == 1) { // special case left-shift by 1 using an in-lane add instead of shift&mask
lshift = __vadd4(input,input);
}
uint32_t rshift = input >> ((8 - amount) & 7);
rshift &= rmask;
uint32_t rotated = lshift | rshift;
return rotated;
}
在移位之前单向屏蔽输入可能更有效,并且在移位之后屏蔽输出((在& lmask中)<< amount |((>>(8-amount))& rmask),用不同的lmask). NVidia硬件有序超标量,shifts have limited throughput.这样做更有可能作为两个独立的移位掩码对执行.
(这并不试图避免数量> = 32的C UB.参见Best practices for circular shift (rotate) operations in C++.在这种情况下,我认为改为lshift = input<<<(amount& 7)就可以了. 为了测试这是否有效编译,我查看了x00-64的clang -O3 asm output,数量不变. Godbolt编译器资源管理器具有各种体系结构的编译器(但不是CUDA),因此如果您可以比x86更容易地阅读这些asm语言,请单击该链接并翻转到ARM,MIPS或PowerPC.
uint32_t rol7(uint32_t a) {
return per_byte_bit_left_rotate(a,7);
}
mov eax,edi
shl eax,7
shr edi
and eax,-2139062144 # 0x80808080
and edi,2139062143 # 0x7F7F7F7F
lea eax,[rdi + rax] # ADD = OR when no bits intersect
ret
完美,正是我所希望的.
几个测试用例：
uint32_t test_rol() {
return per_byte_bit_left_rotate(0x02ffff04,0);
}
// yup,returns the input with count=0
// return 0x2FFFF04
uint32_t test2_rol() {
return per_byte_bit_left_rotate(0x02f73804,4);
}
// yup,swaps nibbles
// return 0x207F8340
这与使用x86 SSE2 / AVX2进行8位移位需要做同样的事情,因为硬件支持的最小移位粒度是16位.
