【问题标题】:When is the (default-variant) PTX instruction `prmt` useful?(默认变体)PTX 指令“prmt”何时有用?
【发布时间】:2023-05-12 18:02:02
【问题描述】:

PTX 有一个prmt instruction 有许多变体。这个问题是关于默认的,如果格式化为 C/C++ 函数,它看起来像这样:

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

这就是它的作用(改编自官方文档):

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

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

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

【问题讨论】:

  • this 可能感兴趣
  • @RobertCrovella:但是提取是您可以通过 PRMT、IIANM 的非默认模式实现的。默认模式有这个奇怪的符号复制选项。我确信有某种用例激发了这些特定的语义。也许与图形有关?
  • 其实我的意思是另一个帖子,但我不能立即找到它:如果我没记错的话,我在几年前发布了一些字节原子操作的示例代码,可能在Nvidia论坛而不是这里。我使用了带有字节选择器的默认 prmt 模式。我也偶尔在其他地方用过。
  • @njuffa:令人惊讶的是,__byte_perm() 没有使用默认模式 :-(

标签: cuda operators nvidia ptx


【解决方案1】:

PTX 指令prmt 公开了机器指令PRMT 的功能。当没有指定特殊模式.f4e, .b4e, .rc8, .ecl, .ecr, .rc16 时,使用prmt 指令的默认模式。

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

另一个子模式是特殊,它不是复制整个源字节,而是跨目标字节复制指定源字节的最高有效位。为此,需要将选择器字段的 msb 设置为 1。程序员必须使用 PTX 内联汇编来访问此功能。

我没有设计 GPU 硬件,因此无法说明为什么要实现该子模式。当每个字节的 msb 用作需要将其转换为整个字节的掩码的布尔值时,它通常很有用。这对于 32 位寄存器中的逐字节处理通常很有用。请注意,CUDA 包含许多用于此类处理的设备函数内部函数,反汇编将确认prmt 默认模式的 msb 复制子模式已用于其中的许多。

一个完整的示例,模拟paddsb 操作(带符号饱和的逐字节加法),如下所示。请注意 prmtmasked_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;
}

【讨论】:

  • 这看起来很巴洛克......这真的比简单的版本快吗? 8-|
  • @einpoklum 试一试,计时,让我们知道哪个版本更快。从内存来看,在 Pascal 级 GPU 上,my_paddsb 的速度是功能等效的直接解决方案的两倍,但是您可能正在考虑一种与我尝试过的不同的直接解决方案。简单地调用__vaddss4() 对于设备代码当然也很简单,但我试图为主机和设备提供一个完全可移植的代码。