lop3.b32命令的含义研究
作用
- 对输入的三个数进行任意的逻辑计算
语法
lop3.b32 d, a, b, c, immLut;
lop3.BoolOp.b32 d|p, a, b, c, immLut, q;
lop3.b32 _, a, b, c, immLut; // 执行位操作,但不保存结果
.BoolOp = { .or , .and };
描述
- 对输入进行按位的逻辑运算（针对每一位上做一次再把结果结合起来），并将结果存入a, b, c, d 中。
lop3.b32指令可以选择性地根据目标操作数和谓词(predicate)执行布尔运算，来计算最终的结果。- 这里的”谓词”指的是
<lut>参数，它定义了一个查找表，用于确定最终的位操作结果。
- 这里的”谓词”指的是
- 逻辑运算由查找表定义，对于 3 个输入，该查找表可以表示为操作数指定的 8 位值，可以取 0 到 255 之间的值，从而允许对输入进行多达 256 个不同的逻辑运算。
- 我们可以根据预先确定的三个常量确定immLut的值：
- 我们通过想对a, b, c进行的逻辑运算提前对ta, tb, tc 进行一次得到immLut。
ta = 0xF0;
tb = 0xCC;
tc = 0xAA;
immLut = F(ta, tb, tc);
//Examples
If F = (a & b & c);
immLut = 0xF0 & 0xCC & 0xAA = 0x80
If F = (a | b | c);
immLut = 0xF0 | 0xCC | 0xAA = 0xFE
If F = (a & b & ~c);
immLut = 0xF0 & 0xCC & (~0xAA) = 0x40
If F = ((a & b | c) ^ a);
immLut = (0xF0 & 0xCC | 0xAA) ^ 0xF0 = 0x1A
应用
template <int lut>
__device__ inline uint32_t lop3(uint32_t a, uint32_t b, uint32_t c) {
uint32_t res;
asm volatile(
"lop3.b32 %0, %1, %2, %3, %4;\n"
: "=r"(res) : "r"(a), "r"(b), "r"(c), "n"(lut)
);
return res;
}
static constexpr uint32_t LO = 0x000f000f;
static constexpr uint32_t HI = 0x00f000f0;
static constexpr uint32_t EX = 0x64006400;
// Guarantee that the `(a & b) | c` operations are LOP3s.
uint32_t t0 = lop3<(0xf0 & 0xcc) | 0xaa>(q, LO, EX);
//(q & LO) | EX
uint32_t t1 = lop3<(0xf0 & 0xcc) | 0xaa>(q, HI, EX);
//(q & HI) | EX
文档信息
- 本文作者：Yizhuo Zheng
- 本文链接：https://amark071.github.io/learnai//2024/08/16/ptx-isa-lop3-b32/
- 版权声明：自由转载-非商用-非衍生-保持署名（创意共享3.0许可证）