fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-05 00:58:05 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 2

etnaviv: isa: Add div opcode

Browse source 
Encoded instruction is taken from blob running the following CL kernel:

 __kernel void simple(__global float *out, __global float *in)
{
    int iGID = get_global_id(0);
    out[iGID] = 4.5f / in[iGID];
}

Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
This commit is contained in:
Christian Gmeiner 2024-03-01 17:52:10 +01:00 • committed by Marge Bot
parent e2a9bc73f5
commit 9c6378abec
2 changed files with 43 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

42
src/etnaviv/isa/etnaviv.xml
View file

@ -540,6 +540,43 @@ SPDX-License-Identifier: MIT
<field name="SRC2_RGROUP" low="124" high="126" type="#reg_group"/>
</bitset>
<bitset name="#instruction-alu-src1-src2" extends="#instruction-alu">
<display>
{INSTR_ALU} {DST:align=18}, void, {SRC1}, {SRC2}
</display>
<!-- SRC0 -->
<pattern pos="43">0</pattern> <!-- SRC0_USE -->
<pattern low="44" high="52">000000000</pattern> <!-- SRC0_REG -->
<pattern low="54" high="61">00000000</pattern> <!-- SRC0_SWIZ -->
<pattern pos="62">0</pattern> <!-- SRC0_NEG -->
<pattern pos="63">0</pattern> <!-- SRC0_ABS -->
<pattern pos="64">0</pattern> <!-- SRC0_AMODE -->
<pattern low="67" high="69">000</pattern> <!-- SRC0_RGROUP -->
<!-- SRC1 -->
<pattern pos="70">1</pattern> <!-- SRC1_USE -->
<field name="SRC1_REG" low="71" high="79" type="uint"/>
<field name="SRC1" low="81" high="90" type="#instruction-src">
<param name="SRC1_REG" as="SRC_REG"/>
<param name="SRC1_AMODE" as="SRC_AMODE"/>
<param name="SRC1_RGROUP" as="SRC_RGROUP"/>
</field>
<field name="SRC1_AMODE" low="91" high="93" type="#reg_addressing_mode"/>
<field name="SRC1_RGROUP" low="96" high="98" type="#reg_group"/>
<!-- SRC2 -->
<pattern pos="99">1</pattern> <!-- SRC2_USE -->
<field name="SRC2_REG" low="100" high="108" type="uint"/>
<field name="SRC2" low="110" high="119" type="#instruction-src">
<param name="SRC2_REG" as="SRC_REG"/>
<param name="SRC2_AMODE" as="SRC_AMODE"/>
<param name="SRC2_RGROUP" as="SRC_RGROUP"/>
</field>
<field name="SRC2_AMODE" low="121" high="123" type="#reg_addressing_mode"/>
<field name="SRC2_RGROUP" low="124" high="126" type="#reg_group"/>
</bitset>
<bitset name="#instruction-alu-src0-src1-src2" extends="#instruction-alu">
<display>
{INSTR_ALU} {DST:align=18}, {SRC0}, {SRC1}, {SRC2}
@ -1187,6 +1224,11 @@ SPDX-License-Identifier: MIT
<pattern pos="80">1</pattern> <!-- OPCODE_BIT6 -->
</bitset>
<bitset name="div" extends="#instruction-alu-src1-src2">
<pattern low="0" high="5">100100</pattern> <!-- OPC -->
<pattern pos="80">1</pattern> <!-- OPCODE_BIT6 -->
</bitset>
<bitset name="atomic_add" extends="#instruction-alu-atomic">
<pattern low="0" high="5">100101</pattern> <!-- OPC -->
<pattern pos="80">1</pattern> <!-- OPCODE_BIT6 -->

1
src/etnaviv/isa/tests/disasm.cpp
View file

@ -147,6 +147,7 @@ INSTANTIATE_TEST_SUITE_P(Opcodes, DisasmTest,
disasm_state{ {0x00801019, 0x15400804, 0x40010000, 0x74000028}, "lshift.s32 t0.x___, t0.yyyy, void, 2\n"},
disasm_state{ {0x0080101a, 0x00001804, 0x40010000, 0x78000018}, "rshift.s32 t0.x___, t1.xxxx, void, 1\n"},
disasm_state{ {0x0080101b, 0x00001804, 0x40010000, 0x00000008}, "rotate.s32 t0.x___, t1.xxxx, void, t0.xxxx\n"},
disasm_state{ {0x03001024, 0x00000005, 0x04098040, 0x0015400f}, "div.rtz t0.zy, void, 4.500000, t0.yyyy\n"},
disasm_state{ {0x01061025, 0x2aa00804, 0xa0010050, 0x7800001f}, "atomic_add.u32 t6._y__, u0.zzzz, 0, 1\n"},
disasm_state{ {0x00801025, 0x2a800884, 0x50010050, 0x0000000f}, "atomic_add.s32 t0.x___, u0.zzzz, 0, t0.xxxx\t; dontcare bits in atomic_add: 00000000000000000000008000000000\n"},
disasm_state{ {0x00821026, 0x2a800884, 0x50010050, 0x0000001f}, "atomic_xchg.s32 t2.x___, u0.zzzz, 0, t1.xxxx\t; dontcare bits in atomic_xchg: 00000000000000000000008000000000\n"},