(默认变体) PTX指令‘PTX’什么时候有用？

Question

PTX有一个包含许多变体的prmt instruction。这个问题涉及默认的问题，如果格式化为C/C++函数，则如下所示：

uint32_t prmt(uint32_t a, uint32_t b, uint32_t byte_selectors);

这就是它所做的(改编自官方文档)：

在泛型形式(未指定模式)中，byte_selectors由四个4位选择值组成.两个源参数a和b中的字节编号为0到7：{b，a} ={b7，b6，b5，b4}，{b3，b2，b1，b0}}。对于函数输出中的每个字节，定义了一个4位选择值.

选择值的3 lsb指定应将8个源字节中的哪个移动到目标位置。msb定义是否应该复制字节值，或者是否应该在目标位置的所有8位上复制符号(字节的msb)(字节值的符号扩展)；msb=0表示复制文字值；msb=1表示复制符号。

我的问题是:这种操作什么时候有用？什么样的计算可以利用它呢？

Answer 1

PTX指令prmt公开机器指令PRMT的功能。当未指定任何特殊模式( prmt )时，将使用.f4e, .b4e, .rc8, .ecl, .ecr, .rc16指令的默认模式。

默认模式有两个每字节子模式，由8个源字节中每个4位选择器字段中最重要的位控制。常用的子模式是使选择器字段的msb为零，这意味着从指定的源字节逐字复制目标字节。这种子模式是通过设备函数内部的__byte_perm()公开的，通常用于提取、插入和置换字节或执行8倍的位移位。在this answer中可以看到示例用法。

另一种子模式是特殊的，因为它不是复制整个源字节，而是在目标字节中复制指定源字节中最重要的位。为此，需要将选择器字段的msb设置为1。程序员必须使用PTX内联程序集来访问此功能。

我没有设计GPU硬件，所以无法说明为什么会实现这种子模式.当每个字节的msb充当一个布尔值，需要转换为整个字节的掩码时，它通常是有用的。这对于32位寄存器内的字节处理通常是有用的.请注意，CUDA包含了许多用于此类处理的设备功能本质，反汇编将确认prmt默认模式的msb复制子模式用于其中的许多。

下面是一个完全工作的示例，即paddsb操作的仿真(按字节计算的有符号饱和加法)。注意在prmt中使用masked_sign_to_byte_mask()和msb复制。

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

#if (__CUDACC__)
#define __HOST__ __host__
#define __DEVICE__ __device__
#else // __CUDACC__
#define __HOST__
#define __DEVICE__
#endif // __CUDACC__

#define MSB_MASK (0x80808080U)  // mask for msb of each byte

// r = (a ^ b) & ~c
__HOST__ __DEVICE__ uint32_t lop3_14 (uint32_t a, uint32_t b, uint32_t c)
{
    uint32_t r;
#if (__CUDA_ARCH__ >= 500)
    asm ("lop3.b32 %0,%1,%2,%3,0x14;\n\t" : "=r"(r) : "r"(a), "r"(b), "r"(c));
#else // __CUDA_ARCH__
    r = (a ^ b) & ~c;
#endif // __CUDA_ARCH__
    return r;
}

// r = (a ^ b) & c
__HOST__ __DEVICE__ uint32_t lop3_28 (uint32_t a, uint32_t b, uint32_t c)
{
    uint32_t r;
#if (__CUDA_ARCH__ >= 500)
    asm ("lop3.b32 %0,%1,%2,%3,0x28;\n\t" : "=r"(r) : "r"(a), "r"(b), "r"(c));
#else // __CUDA_ARCH__
    r = (a ^ b) & c;
#endif // __CUDA_ARCH__
    return r;
}

// r = a ^ (~b & c)
__HOST__ __DEVICE__ uint32_t lop3_d2 (uint32_t a, uint32_t b, uint32_t c)
{
    uint32_t r;
#if (__CUDA_ARCH__ >= 500)
    asm ("lop3.b32 %0,%1,%2,%3,0xd2;\n\t" : "=r"(r) : "r"(a), "r"(b), "r"(c));
#else // __CUDA_ARCH__
    r = a ^ (~b & c);
#endif // __CUDA_ARCH__ 
    return r;
}

// r = (a & c) | (b & ~c)
__HOST__ __DEVICE__ uint32_t lop3_f4 (uint32_t a, uint32_t b, uint32_t c)
{
    uint32_t r;
#if (__CUDA_ARCH__ >= 500)
    asm ("lop3.b32 %0,%1,%2,%3,0xf4;\n\t" : "=r"(r) : "r"(a), "r"(b), "r"(c));
#else // __CUDA_ARCH__
    r = (a & c) | (b & ~c);
#endif // __CUDA_ARCH__
    return r;
} 

__HOST__ __DEVICE__ uint32_t masked_sign_to_byte_mask (uint32_t a)
{
#if (__CUDA_ARCH__ >= 200)
    asm ("prmt.b32 %0,%0,0,0xba98;" : "+r"(a)); // convert MSBs to masks
#else
    a = a & MSB_MASK;
    a = a + a - (a >> 7); // extend MSBs to full byte to create mask
#endif
    return a;
}

__HOST__ __DEVICE__ uint32_t masked_select (uint32_t a, uint32_t b, uint32_t m)
{
#if (__CUDA_ARCH__ >= 500) 
    return lop3_f4 (a, b, m);
#elif 0
    return (((a)&(m))|((b)&(~(m))));
#else
    return((((a)^(b))&(m))^(b));
#endif
}

/* 
   my_paddsb() performs byte-wise addition with signed saturation. In the 
   case of overflow, positive results are clamped at 127, while negative 
   results are clamped at -128.
*/
__HOST__ __DEVICE__ uint32_t my_paddsb (uint32_t a, uint32_t b)
{
    uint32_t sum, res, ofl, sga, msk;
    res = (a & ~MSB_MASK) + (b & ~MSB_MASK);
    sum = a ^ b;
    ofl = lop3_14 (res, a, sum); // ofl = (res ^ a) & ~sum
    sga = masked_sign_to_byte_mask (a);  // sign(a)-mask
    msk = masked_sign_to_byte_mask (ofl);// overflow-mask
    res = lop3_d2 (res, ~MSB_MASK, sum); // res = res ^ (MSB_MASK & sum)
    sga = lop3_28 (sga, ~MSB_MASK, msk); // sga = (sga ^ ~MSB_MASK) & msk
    res = masked_select (sga, res, msk); // res = (sga & msk) | (res & ~msk)
    return res;
}

__global__ void kernel (uint32_t a, uint32_t b)
{
    printf ("GPU: %08x\n", my_paddsb (a, b));
}

int main (void)
{
    uint32_t a = 0x12ef70a0;
    uint32_t b = 0x34cd6090;
    kernel<<<1,1>>>(a, b);
    cudaDeviceSynchronize();
    printf ("CPU: %08x\n", my_paddsb (a, b));
    return EXIT_SUCCESS;
}