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
操作(带符号饱和的逐字节加法),如下所示。请注意 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;
}