<html>
<head>
<base href="https://bugs.freedesktop.org/">
</head>
<body><table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Bug ID</th>
<td><a class="bz_bug_link
bz_status_NEW "
title="NEW - bar.sync encoding incorrect for GM107"
href="https://bugs.freedesktop.org/show_bug.cgi?id=106132">106132</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>bar.sync encoding incorrect for GM107
</td>
</tr>
<tr>
<th>Product</th>
<td>Mesa
</td>
</tr>
<tr>
<th>Version</th>
<td>git
</td>
</tr>
<tr>
<th>Hardware</th>
<td>All
</td>
</tr>
<tr>
<th>OS</th>
<td>All
</td>
</tr>
<tr>
<th>Status</th>
<td>NEW
</td>
</tr>
<tr>
<th>Severity</th>
<td>normal
</td>
</tr>
<tr>
<th>Priority</th>
<td>medium
</td>
</tr>
<tr>
<th>Component</th>
<td>Drivers/DRI/nouveau
</td>
</tr>
<tr>
<th>Assignee</th>
<td>nouveau@lists.freedesktop.org
</td>
</tr>
<tr>
<th>Reporter</th>
<td>vriestj@gmail.com
</td>
</tr>
<tr>
<th>QA Contact</th>
<td>nouveau@lists.freedesktop.org
</td>
</tr></table>
<p>
<div>
<pre>[ Quoted text copied from
<a href="https://lists.freedesktop.org/archives/mesa-dev/2016-March/108926.html">https://lists.freedesktop.org/archives/mesa-dev/2016-March/108926.html</a> ]
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.
<span class="quote">> [Mesa-dev] [PATCH] gm107/ir: add emission for BAR
> Samuel Pitoiset samuel.pitoiset at gmail.com
> Tue Mar 1 17:44:42 UTC 2016</span >
>
<span class="quote">> + // 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();</span >
This should probably be using getSrc(1) ?
<span class="quote">> + assert(imm);
> + emitField(0x14, 12, imm->reg.data.u32);
> + emitField(0x2c, 1, 1);
> + }</span >
>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.</pre>
</div>
</p>
<hr>
<span>You are receiving this mail because:</span>
<ul>
<li>You are the QA Contact for the bug.</li>
<li>You are the assignee for the bug.</li>
</ul>
</body>
</html>