在 CUDA 中使用 SIMD 实现位循环运算符

Posted 2023-02-16

【中文标题】在 CUDA 中使用 SIMD 实现位循环运算符

【英文标题】：Implementation of bit rotate operators using SIMD in CUDA

【发布时间】：2017-08-27 00:02:24

【问题描述】：

我知道 *** 不是用来向其他人询问代码的，但请允许我说一下。

我正在尝试在 CUDA C++ 设备代码中实现一些 AES 函数。在尝试实现左字节旋转运算符时，我很不安地看到没有本机 SIMD intrisic。所以我开始了一个幼稚的实现，但是....它很大，虽然我还没有尝试过，但它不会很快，因为解包/打包成本很高......所以，有没有办法做至少有点高效的每字节位循环操作？

如果你不想看，这里是代码。

__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

【问题讨论】：

我不认为你的 XOR 工作。如果一开始是零，它将设置该位。你需要一个 AND 来无条件地清除一点。

哦，是的，没听懂，我要编辑它。正如我所说，这甚至还没有编译，只是大约

我认为你需要& ~0x100。或者只是& 0xFF，因为您将所有内容都转移到低字节。 AND 与 0x100 会将所有 except 归零。

我想我的编辑速度太快了。对不起。

【参考方案1】：

CUDA 有一个 __byte_perm() 内在函数，它直接映射到机器代码 (SASS) 级别的 PRMT 指令，这是一个按字节排列的指令。它可用于有效地提取和合并字节。为了实现逐字节左旋转，我们可以将每个字节加倍，将字节对移动所需的数量，然后提取并合并字节对的四个高字节。

对于逐字节旋转，我们只需要移位量的最低三位，因为s 的旋转与s mod 8 的旋转相同。为了提高效率，最好避免包含少于 32 位的整数类型，因为 C++ 语义要求在表达式中使用之前，将窄于 int 的整数类型扩大到 int。这可能而且确实会在包括 GPU 在内的许多架构上产生转换开销。

PRMT 指令的吞吐量取决于架构，因此使用__byte_perm() 可能会导致代码比使用another answer 中演示的经典SIMD-in-a-register 方法更快或更慢，因此，请务必在部署之前根据您的用例环境进行基准测试。

#include <stdio.h>
#include <stdint.h>
#include <stdlib.h>

__device__ uint32_t per_byte_bit_left_rotate (uint32_t input, uint32_t amount)

     uint32_t l = __byte_perm (input, 0, 0x1100) << (amount & 7);
     uint32_t h = __byte_perm (input, 0, 0x3322) << (amount & 7);
     return __byte_perm (l, h, 0x7531);


__global__ void rotl_kernel (uint32_t input, uint32_t amount, uint32_t *res)

    *res = per_byte_bit_left_rotate (input, amount);


uint32_t ref_per_byte_bit_left_rotate (uint32_t input, uint32_t amount)

   int s = amount & 7;
   uint8_t b0 = (input >>  0) & 0xff;
   uint8_t b1 = (input >>  8) & 0xff;
   uint8_t b2 = (input >> 16) & 0xff;
   uint8_t b3 = (input >> 24) & 0xff;
   b0 = s ? ((b0 << s) | (b0 >> (8 - s))) : b0;
   b1 = s ? ((b1 << s) | (b1 >> (8 - s))) : b1;
   b2 = s ? ((b2 << s) | (b2 >> (8 - s))) : b2;
   b3 = s ? ((b3 << s) | (b3 >> (8 - s))) : b3;
   return (b3 << 24) | (b2 << 16) | (b1 << 8) | (b0 << 0);


// Fixes via: Greg Rose, KISS: A Bit Too Simple. http://eprint.iacr.org/2011/007
static unsigned int z=362436069,w=521288629,jsr=362436069,jcong=123456789;
#define znew (z=36969*(z&0xffff)+(z>>16))
#define wnew (w=18000*(w&0xffff)+(w>>16))
#define MWC  ((znew<<16)+wnew)
#define SHR3 (jsr^=(jsr<<13),jsr^=(jsr>>17),jsr^=(jsr<<5)) /* 2^32-1 */
#define CONG (jcong=69069*jcong+13579)                     /* 2^32 */
#define KISS ((MWC^CONG)+SHR3)

// Macro to catch CUDA errors in CUDA runtime calls
#define CUDA_SAFE_CALL(call)                                          \
do                                                                   \
    cudaError_t err = call;                                           \
    if (cudaSuccess != err)                                          \
        fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\
                 __FILE__, __LINE__, cudaGetErrorString(err) );       \
        exit(EXIT_FAILURE);                                           \
                                                                     \
 while (0)

// Macro to catch CUDA errors in kernel launches
#define CHECK_LAUNCH_ERROR()                                          \
do                                                                   \
    /* Check synchronous errors, i.e. pre-launch */                   \
    cudaError_t err = cudaGetLastError();                             \
    if (cudaSuccess != err)                                          \
        fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\
                 __FILE__, __LINE__, cudaGetErrorString(err) );       \
        exit(EXIT_FAILURE);                                           \
                                                                     \
    /* Check asynchronous errors, i.e. kernel failed (ULF) */         \
    err = cudaThreadSynchronize();                                    \
    if (cudaSuccess != err)                                          \
        fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\
                 __FILE__, __LINE__, cudaGetErrorString( err) );      \
        exit(EXIT_FAILURE);                                           \
                                                                     \
 while (0)

int main (void)

    uint32_t arg, ref, res = 0, *res_d = 0;
    uint32_t shft;

    CUDA_SAFE_CALL (cudaMalloc ((void**)&res_d, sizeof(*res_d)));
    for (int i = 0; i < 100000; i++) 
        arg  = KISS;
        shft = KISS;
        ref = ref_per_byte_bit_left_rotate (arg, shft);
        rotl_kernel <<<1,1>>>(arg, shft, res_d);
        CHECK_LAUNCH_ERROR();
        CUDA_SAFE_CALL (cudaMemcpy (&res, res_d, sizeof (res), 
                                    cudaMemcpyDeviceToHost));
        if (res != ref) 
            printf ("!!!! arg=%08x shft=%d  res=%08x  ref=%08x\n", 
                    arg, shft, res, ref);
        
    
    CUDA_SAFE_CALL (cudaFree (res_d));
    CUDA_SAFE_CALL (cudaDeviceSynchronize());
    return EXIT_SUCCESS;

【讨论】：

嗯...我没想到 __byte_perm() 可以用来选择一些字节，我总是发现在某种意义上说 intrisic 没用（这可能是它的速度不一致的原因，因为如果我猜的话，图形不能使用很多）。感谢闪电我的道路。但是，作为 uint8_t 的 bX 变量，它们不是首先转换为 char，然后在返回行中重新转换为 int 吗？ AND 可以有效地掩盖一些数据，不是吗？ （假设有更多的注册压力）

@Sachiko.Shinozaki b? 变量仅出现在 reference 实现中，该实现是为功能而非性能而编写的。当然，参考函数的其他实现也是可能的。 __byte_perm() 对于移位/旋转 8 位的倍数、交换字节、跨 32 位字的 splat 字节等很有用。根据我的观察，CUDA 编译器有时会生成 PRMT 指令，但不经常生成。我不做图形，所以不能说它在那里的实用性。

我们只能猜测编译器“理解”__byte_perm() 的逻辑并根据架构/管道使用情况使用 PTX/SASS ..？英伟达没有透露那种东西，所以我们只能猜测。

@Sachiko.Shinozaki __byte_perm() 翻译成PRMT，它是一个intrinsic。如果您对此持怀疑态度，您可以随时通过cuobjdump -dump-sass 仔细检查。当我说 CUDA 编译器有时会使用PRMT 时，我指的是那些not 源自__byte_perm() 的显式我们的实例，例如8 点换班。如果不清楚，请见谅。

【参考方案2】：

所有元素的旋转计数都相同，对吧？

将整个输入向左和向右移动，然后将那些带有掩码的那些跨越字节边界的所有位都归零，将所有 4 个字节放在一个 AND 中。我认为amount 始终是 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;

在移位前以一种方式屏蔽输入，并在移位后屏蔽输出（(in&lmask)<<amount | ((in>>(8-amount))&rmask)，使用不同的 lmask）可能会更有效。 NVidia 硬件是有序超标量，shifts have limited throughput。这样做更有可能作为两个独立的 shift+m​​ask 对来执行。

（这并没有试图避免数量>=32 的 C++ UB。请参阅 Best practices for circular shift (rotate) operations in C++。在这种情况下，我认为更改为 lshift = input << (amount & 7) 可以解决问题。

为了测试这个编译是否有效，我查看了 x86-64 的 clang -O3 asm output 和常量 amount。 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 位。

【讨论】：

其实，我的问题的“CUDA”部分是因为CUDA支持一些SIMD intrisics，但没有它也没关系，所有基本操作都在SASS中定义（使用的对象级指令集由 CUDA 设备）。 >= 32 位 UB 并不重要，因为它在任何地方都没有调用 AES

@Sachiko.Shinozaki：啊，我明白了。 docs.nvidia.com/cuda/cuda-math-api/…。半字或字节元素之类的东西在 32 位整数内的通道中添加。对于左移 1，您可以使用 __vadd4 而不是 shift/mask。 (x+x = x<<1)

我实际上发现自己对这个非常傻眼，但无论如何，错误是存在的，所以你可以学习它们。我非常尊重，不知道 CUDA 却挖掘数学知识来帮助我。无限感谢。感谢您向我介绍 SWAR 的世界。虽然我做了一些基本的汇编，但我并没有真正参考汇编世界来学习更高级的习语。

@Sachiko.Shinozaki：回答问题是了解我一直想要了解的新事物的好方法 :) 我很好奇 CUDA 支持什么样的事物。跨度>

@Sachiko.Shinozaki：有关 SWAR / bithack 的更多内容，请参阅 graphics.stanford.edu/~seander/bithacks.html。有一些简洁的东西，但其中很多是现在用内在函数做得更好的东西（例如 popcnt 或 tzcnt）。或一个字中零字节的 SIMD 内在函数。尽管如此，很高兴看到如何通过屏蔽和移位来实现popcnt，这样您就可以进行 16 个不相互进位的 1 位加法，然后是 8 个 2 位加法等。另见 hackersdelight.org