https://bugs.freedesktop.org/show_bug.cgi?id=106132

            Bug ID: 106132
           Summary: bar.sync encoding incorrect for GM107
           Product: Mesa
           Version: git
          Hardware: All
                OS: All
            Status: NEW
          Severity: normal
          Priority: medium
         Component: Drivers/DRI/nouveau
          Assignee: nouveau-PD4FTy7X32lNgt0PjOBp9y5qC8QIuHrW@public.gmane.org
          Reporter: vriestj-Re5JQEeQqe8AvxtiuMwx3w@public.gmane.org
        QA Contact: nouveau-PD4FTy7X32lNgt0PjOBp9y5qC8QIuHrW@public.gmane.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->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 should probably be using getSrc(1) ?

> +      assert(imm);
> +      emitField(0x14, 12, imm->reg.data.u32);
> +      emitField(0x2c, 1, 1);
> +   }

From these examples:
...
                                x        x
BAR.SYNC 0x0;          /* 0xf0a81b8000070000 */
BAR.SYNC 0xf;          /* 0xf0a81b8000070f00 */
BAR.SYNC R0;           /* 0xf0a80b8000070000 */
BAR.SYNC R2;           /* 0xf0a80b8000070200 */
...
I derive these offsets:
- barrier id           : 0x08
- barrier id imm vs reg: 0x2c

and from these examples:
...
                                 x    x
BAR.SYNC 0x0;          /* 0xf0a81b8000070000 */
BAR.SYNC 0x0, R2;      /* 0xf0a8138000270000 */
...
I derive these offsets:
- thread count           : 0x14
- thread count imm vs reg: 0x2b


But when looking at the code snippet above, the roles seem reversed: 0x2b is
used for barrier id, and 0x2c is used for the thread count.

-- 
You are receiving this mail because:
You are the QA Contact for the bug.
You are the assignee for the bug.