search for: asimm

Displaying 9 results from an estimated 9 matches for "asimm".

Did you mean: isimm
2018 Apr 19
3
[Bug 106132] New: bar.sync encoding incorrect for GM107
...mission 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).getFile() == FILE_GPR) { > + emitGPR(0x14, insn->src(1)); > + } else { > + ImmediateVa...
2017 Apr 03
5
[PATCH v2 0/3] nv50/ir: Preapre for running Opts inside a loop
Slowly we are getting to the point, that we miss enough optimization opportunities as the result of our own passes. For this we need to fix AlgebraicOpt to be able to handle mods on sources without creating new issues. The last patch enables looping opts. v2: update commit author Karol Herbst (3): nv50/ir: fix AlgebraicOpt for slcts with mods nv50/ir: handle logops with NOT in AlgebraicOpt
2017 Apr 03
0
[PATCH v2 1/3] nv50/ir: fix AlgebraicOpt for slcts with mods
...ouveau/codegen/nv50_ir_peephole.cpp index 4c92a1efb5..bd60a84998 100644 --- a/src/gallium/drivers/nouveau/codegen/nv50_ir_peephole.cpp +++ b/src/gallium/drivers/nouveau/codegen/nv50_ir_peephole.cpp @@ -1797,10 +1797,10 @@ AlgebraicOpt::handleSLCT(Instruction *slct) if (slct->getSrc(2)->asImm()->compare(slct->asCmp()->setCond, 0.0f)) slct->setSrc(0, slct->getSrc(1)); } else - if (slct->getSrc(0) != slct->getSrc(1)) { + if (slct->getSrc(0) != slct->getSrc(1) || slct->src(0).mod != slct->src(1).mod) return; - } - slct->op =...
2014 Sep 01
0
[PATCH] nv50/ir: avoid creating instructions that can't be emitted
...- a/src/gallium/drivers/nouveau/codegen/nv50_ir_peephole.cpp +++ b/src/gallium/drivers/nouveau/codegen/nv50_ir_peephole.cpp @@ -567,6 +567,10 @@ ConstantFolding::expr(Instruction *i, ImmediateValue src0; if (i->src(0).getImmediate(src0)) expr(i, src0, *i->getSrc(1)->asImm()); + if (i->saturate && !prog->getTarget()->isSatSupported(i)) { + bld.setPosition(i, false); + i->setSrc(1, bld.loadImm(NULL, res.data.u32)); + } } else { i->op = i->saturate ? OP_SAT : OP_MOV; /* SAT handled by unary() */ } -- 1...
2017 Apr 03
3
[PATCH 0/3] nv50/ir: Preapre for running Opts inside a loop
Slowly we are getting to the point, that we miss enough optimization opportunities as the result of our own passes. For this we need to fix AlgebraicOpt to be able to handle mods on sources without creating new issues. The last patch enables looping opts. Karol Herbst (3): nv50/ir: fix AlgebraicOpt for slcts with mods nv50/ir: handle logops with NOT in AlgebraicOpt nv50/ir: run some
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)
2014 May 20
14
[PATCH 00/12] Cherry-pick nv50/nvc0 patches from gallium-nine
I went through the gallium-nine tree and picked out nouveau patches that are general bug-fixes. The first bunch I'd like to also get into 10.2. I've reviewed all of them and they make sense to me, but sending them out for public review as well in case there are any objections. Unless I hear objections, I'd like to push this by Friday. Christoph Bumiller (11): nv50,nvc0: always pull
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 ---
2014 Jan 13
20
[PATCH 00/19] nv50: add sampler2DMS/GP support to get OpenGL 3.2
OK, so there's a bunch of stuff in here. The geometry stuff is based on the work started by Bryan Cain and Christoph Bumiller. Patches 01-12: Add support for geometry shaders and fix related issues Patches 13-14: Make it possible for fb clears to operate on texture attachments with an explicit layer set (as is allowed in gl 3.2). Patches 15-17: Make ARB_texture_multisample work