Displaying 4 results from an estimated 4 matches for "emitgpr".
2018 Apr 19
3
[Bug 106132] New: bar.sync encoding incorrect for GM107
...at nv50_ir_emit_gm107.cpp to understand the instruction encoding.
> [Mesa-dev] [PATCH] gm107/ir: add emission for BAR
> Samuel Pitoiset samuel.pitoiset 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).getFi...
2017 Mar 26
5
[PATCH v5 0/5] nvc0/ir: add support for MAD/FMA PostRALoadPropagation
was "nv50/ir: PostRaConstantFolding improvements" before.
nothing really changed from the last version, just minor things.
Karol Herbst (5):
nv50/ir: restructure and rename postraconstantfolding pass
nv50/ir: implement mad post ra folding for nvc0+
gk110/ir: add LIMM form of mad
gm107/ir: add LIMM form of mad
nv50/ir: also do PostRaLoadPropagation for FMA
2015 May 09
5
[PATCH 1/4] nvc0/ir: avoid jumping to a sched instruction
Signed-off-by: Ilia Mirkin <imirkin at alum.mit.edu>
---
Pretty sure there's nothing wrong with it, but it looks odd in the code.
src/gallium/drivers/nouveau/codegen/nv50_ir_emit_gk110.cpp | 2 ++
src/gallium/drivers/nouveau/codegen/nv50_ir_emit_gm107.cpp | 7 +++++--
src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp | 2 ++
3 files changed, 9 insertions(+), 2 deletions(-)
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
---