Nouveau - Apr 2018 - [Bug 106132] New: bar.sync encoding incorrect for GM107

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 at lists.freedesktop.org
          Reporter: vriestj at gmail.com
        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->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.
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<https://lists.freedesktop.org/archives/nouveau/attachments/20180419/393119cf/attachment.html>

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

--- Comment #1 from Tom de Vries <vriestj at gmail.com> ---
Previously reported at
https://lists.freedesktop.org/archives/mesa-dev/2018-April/192621.html .

-- 
You are receiving this mail because:
You are the assignee for the bug.
You are the QA Contact for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<https://lists.freedesktop.org/archives/nouveau/attachments/20180419/92d2132c/attachment.html>

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

--- Comment #2 from Tom de Vries <vriestj at gmail.com> ---
Created attachment 138920
  --> https://bugs.freedesktop.org/attachment.cgi?id=138920&action=edit
Tentative patch

-- 
You are receiving this mail because:
You are the assignee for the bug.
You are the QA Contact for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<https://lists.freedesktop.org/archives/nouveau/attachments/20180419/1cc50a3c/attachment.html>

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

GitLab Migration User <gitlab-migration at fdo.invalid> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
         Resolution|---                         |MOVED
             Status|NEW                         |RESOLVED

--- Comment #3 from GitLab Migration User <gitlab-migration at
fdo.invalid> ---
-- GitLab Migration Automatic Message --

This bug has been migrated to freedesktop.org's GitLab instance and has been
closed from further activity.

You can subscribe and participate further through the new bug through this link
to our GitLab instance: https://gitlab.freedesktop.org/mesa/mesa/issues/1156.

-- 
You are receiving this mail because:
You are the assignee for the bug.
You are the QA Contact for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<https://lists.freedesktop.org/archives/nouveau/attachments/20190918/a404e524/attachment.html>