c – 在CUDA中使用SIMD实现位旋转运算符

发布时间：2020-12-16 03:27:05 所属栏目：百科来源：网络整理

导读：我知道StackOverflow不是为了向其他人询问代码,而是让我说话. 我正在尝试在CUDA C设备代码中实现一些AES函数.在尝试实现左侧按字节旋转运算符时,我感到不安的是看到没有原生的SIMD内向.所以我开始了一个天真的实现,但……它是巨大的,虽然我还没有尝试过,但由

我知道StackOverflow不是为了向其他人询问代码,而是让我说话.

我正在尝试在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位.

（编辑：李大同）

【声明】本站内容均来自网络，其相关言论仅代表作者个人观点，不代表本站立场。若无意侵犯到您的权利，请及时与联系站长删除相关内容!