Displaying 6 results from an estimated 6 matches for "emitfield".
2015 Aug 19
5
[PATCH 1/2] nvc0/ir: detect AND/SHR pairs and convert into EXTBF
Some shaders appear to extract bits using shift/and combos. Detect
(some) of those and convert to EXTBF instead.
Signed-off-by: Ilia Mirkin <imirkin at alum.mit.edu>
---
.../drivers/nouveau/codegen/nv50_ir_peephole.cpp | 66 +++++++++++++++-------
1 file changed, 46 insertions(+), 20 deletions(-)
diff --git a/src/gallium/drivers/nouveau/codegen/nv50_ir_peephole.cpp
2018 Apr 19
3
[Bug 106132] New: bar.sync encoding incorrect for GM107
...toiset at gmail.com
> Tue Mar 1 17:44:42 UTC 2016
>
> + // barrier id
> + if (insn->src(0).getFile() == FILE_GPR) {
> + emitGPR(0x08, insn->src(0));
> + } else {
> + ImmediateValue *imm = insn->getSrc(0)->asImm();
> + assert(imm);
> + emitField(0x08, 8, imm->reg.data.u32);
> + emitField(0x2b, 1, 1);
> + }
> +
> + // thread count
> + if (insn->src(1).getFile() == FILE_GPR) {
> + emitGPR(0x14, insn->src(1));
> + } else {
> + ImmediateValue *imm = insn->getSrc(0)->asImm();
This s...
2015 May 09
5
[PATCH 1/4] nvc0/ir: avoid jumping to a sched instruction
...107::emitBRA()
emitCond5(0x00, CC_TR);
if (!insn->srcExists(0) || insn->src(0).getFile() != FILE_MEMORY_CONST) {
+ int32_t pos = insn->target.bb->binPos;
+ if (writeIssueDelays && !(pos & 0x1f))
+ pos += 8;
if (!insn->absolute)
- emitField(0x14, 24, insn->target.bb->binPos - (codeSize + 8));
+ emitField(0x14, 24, pos - (codeSize + 8));
else
- emitField(0x14, 32, insn->target.bb->binPos);
+ emitField(0x14, 32, pos);
} else {
emitCBUF (0x24, gpr, 20, 16, 0, insn->src(0));...
2015 Feb 20
10
[PATCH 01/11] nvc0/ir: add emission of dadd/dmul/dmad opcodes, fix minmax
Signed-off-by: Ilia Mirkin <imirkin at alum.mit.edu>
---
.../drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp | 66 +++++++++++++++++++++-
1 file changed, 63 insertions(+), 3 deletions(-)
diff --git a/src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp b/src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp
index dfb093c..e38a3b8 100644
---
2019 Oct 14
1
[PATCH] gm107/ir: fix loading z offset for layered 3d image bindings
...oid emitS2R();
void emitCS2R();
@@ -690,6 +692,31 @@ CodeEmitterGM107::emitRAM()
* predicate/cc
******************************************************************************/
+void
+CodeEmitterGM107::emitPSETP()
+{
+
+ emitInsn(0x50900000);
+
+ switch (insn->op) {
+ case OP_AND: emitField(0x18, 3, 0); break;
+ case OP_OR: emitField(0x18, 3, 1); break;
+ case OP_XOR: emitField(0x18, 3, 2); break;
+ default:
+ assert(!"unexpected operation");
+ break;
+ }
+
+ // emitINV (0x2a);
+ emitPRED(0x27); // TODO: support 3-arg
+ emitINV (0x20, insn->src(1)...
2015 Nov 05
7
[PATCH mesa 0/5] nouveau: codegen: Make use of double immediates
Hi All,
This series implements using double immediates in the nouveau codegen code.
This turns the following (nvc0) code:
1: mov u32 $r2 0x00000000 (8)
2: mov u32 $r3 0x3fe00000 (8)
3: add f64 $r0d $r0d $r2d (8)
Into:
1: add f64 $r0d $r0d 0.500000 (8)
This has been tested with the 2 double shader tests which I just send to
the piglet list. On a gk208 (gk110 / SM35)