search for: nv50_ir_emit_gm107

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 ---

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(-) diff --git a/src/gallium/drivers/nouveau/codegen/nv50_ir_emit_gk110.cpp b/src/gallium/drivers/nouveau/codegen/nv50_ir_emit_gk110.cpp index 6bb9620..28081fa 100644 --...

...ore work > for e.g. atomic ops. > > Signed-off-by: Hans de Goede <hdegoede at redhat.com> > --- > src/gallium/drivers/nouveau/codegen/nv50_ir.h | 1 + > .../drivers/nouveau/codegen/nv50_ir_emit_gk110.cpp | 31 +++++++++++++++++----- > .../drivers/nouveau/codegen/nv50_ir_emit_gm107.cpp | 5 +++- > .../drivers/nouveau/codegen/nv50_ir_emit_nv50.cpp | 10 ++++--- > .../drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp | 26 +++++++++++++----- > .../drivers/nouveau/codegen/nv50_ir_from_tgsi.cpp | 14 +++++++--- > .../drivers/nouveau/codegen/nv50_ir_peephole.cpp...

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)

...processing (!), and the surface conversion function is seriously hacked up. But splitting it up is harder, since a lot of information has to flow from stage to stage, like when to do what kind of access, and cloning the surface op is much easier in the coord processing stage. .../nouveau/codegen/nv50_ir_emit_gm107.cpp | 34 ++- .../nouveau/codegen/nv50_ir_lowering_nvc0.cpp | 206 +++++++++++++----- .../nouveau/codegen/nv50_ir_lowering_nvc0.h | 4 +- src/gallium/drivers/nouveau/nvc0/nvc0_tex.c | 10 +- 4 files changed, 201 insertions(+), 53 deletions(-) diff --git a/src/gallium/drivers/nouveau/cod...

...ar load and stores and likely needs more work for e.g. atomic ops. Signed-off-by: Hans de Goede <hdegoede at redhat.com> --- src/gallium/drivers/nouveau/codegen/nv50_ir.h | 1 + .../drivers/nouveau/codegen/nv50_ir_emit_gk110.cpp | 31 +++++++++++++++++----- .../drivers/nouveau/codegen/nv50_ir_emit_gm107.cpp | 5 +++- .../drivers/nouveau/codegen/nv50_ir_emit_nv50.cpp | 10 ++++--- .../drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp | 26 +++++++++++++----- .../drivers/nouveau/codegen/nv50_ir_from_tgsi.cpp | 14 +++++++--- .../drivers/nouveau/codegen/nv50_ir_peephole.cpp | 5 +++- .../drivers/n...

...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 .../drivers/nouveau/codegen/nv50_ir_emit_gk110.cpp | 50 ++++--- .../drivers/nouveau/codegen/nv50_ir_emit_gm107.cpp | 34 +++-- .../drivers/nouveau/codegen/nv50_ir_peephole.cpp | 162 +++++++++++++-------- src/gallium/drivers/nouveau/codegen/nv50_ir_ra.cpp | 2 +- 4 files changed, 164 insertions(+), 84 deletions(-) -- 2.12.0

tgsi_default_instruction_memory / tgsi_build_instruction_memory were returning uninitialized memory for tgsi_instruction_memory.Texture and tgsi_instruction_memory.Format. Note 0 means not set, and thus is a correct default initializer for these. Fixes: 3243b6fc97 ("tgsi: add Texture and Format to tgsi_instruction_memory") Cc: Nicolai Hähnle <nicolai.haehnle at amd.com>

...global memory. > > Signed-off-by: Hans de Goede <hdegoede at redhat.com> > --- > src/gallium/drivers/nouveau/codegen/nv50_ir.h | 2 +- > src/gallium/drivers/nouveau/codegen/nv50_ir_emit_gk110.cpp | 10 > +++++----- > src/gallium/drivers/nouveau/codegen/nv50_ir_emit_gm107.cpp | 6 +++--- > src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nv50.cpp | 10 > +++++----- > src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp | 12 > ++++++------ > src/gallium/drivers/nouveau/codegen/nv50_ir_from_tgsi.cpp | 8 ++++---- > .../drivers/nouv...

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

...tomic ops. >> >> Signed-off-by: Hans de Goede <hdegoede at redhat.com> >> --- >> src/gallium/drivers/nouveau/codegen/nv50_ir.h | 1 + >> .../drivers/nouveau/codegen/nv50_ir_emit_gk110.cpp | 31 +++++++++++++++++----- >> .../drivers/nouveau/codegen/nv50_ir_emit_gm107.cpp | 5 +++- >> .../drivers/nouveau/codegen/nv50_ir_emit_nv50.cpp | 10 ++++--- >> .../drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp | 26 +++++++++++++----- >> .../drivers/nouveau/codegen/nv50_ir_from_tgsi.cpp | 14 +++++++--- >> .../drivers/nouveau/codegen/nv50_ir...

...of FILE_MEMORY_GLOBAL for regular global memory. Signed-off-by: Hans de Goede <hdegoede at redhat.com> --- src/gallium/drivers/nouveau/codegen/nv50_ir.h | 2 +- src/gallium/drivers/nouveau/codegen/nv50_ir_emit_gk110.cpp | 10 +++++----- src/gallium/drivers/nouveau/codegen/nv50_ir_emit_gm107.cpp | 6 +++--- src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nv50.cpp | 10 +++++----- src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp | 12 ++++++------ src/gallium/drivers/nouveau/codegen/nv50_ir_from_tgsi.cpp | 8 ++++---- .../drivers/nouveau/codegen/nv50_ir_lowering_n...

...t; Signed-off-by: Hans de Goede <hdegoede at redhat.com> >> --- >> src/gallium/drivers/nouveau/codegen/nv50_ir.h | 2 +- >> src/gallium/drivers/nouveau/codegen/nv50_ir_emit_gk110.cpp | 10 >> +++++----- >> src/gallium/drivers/nouveau/codegen/nv50_ir_emit_gm107.cpp | 6 +++--- >> src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nv50.cpp | 10 >> +++++----- >> src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp | 12 >> ++++++------ >> src/gallium/drivers/nouveau/codegen/nv50_ir_from_tgsi.cpp | 8 ++++----...

So I'm trying to track an special value in IPA instruction generation. https://cgit.freedesktop.org/mesa/mesa/tree/src/gallium/drivers/nouveau/codegen/nv50_ir_emit_gm107.cpp#L2561 Register on 0x14 (20) is set to some source on "insn->op == OP_PINTERP" I have found while emulation that such value can be set sometimes to FragCoord.w, I don't however know what that value is and how to represent it on glsl. Do you guys know where does that value come...

Hello, I got some doubts on how texture modes work on TEX, TEXS, TLD4, etc instructions. I got: DC, AOFFI, NDV, NODEP, MZ, PTP modes as well as LZ Mode. How does this work or change the behavior of the texture instruction. So far of those I know AOFFI defines an Offset but I'm on blanks for the rest. -------------- next part -------------- An HTML attachment was scrubbed... URL:

...QA Contact: nouveau at lists.freedesktop.org [ Quoted text copied from https://lists.freedesktop.org/archives/mesa-dev/2016-March/108926.html ] I've been playing around with bar.sync in ptx, JIT-compiling it to GM107 (my quadro m1200 card), and disassembling with cuobjdump -sass. I looked 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-&gt...