From 9c6378abeccf32f676770dcf93e559bebef01671 Mon Sep 17 00:00:00 2001
From: Christian Gmeiner <cgmeiner@igalia.com>
Date: Fri, 1 Mar 2024 17:52:10 +0100
Subject: [PATCH] etnaviv: isa: Add div opcode

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>
---
 src/etnaviv/isa/etnaviv.xml      | 42 ++++++++++++++++++++++++++++++++
 src/etnaviv/isa/tests/disasm.cpp |  1 +
 2 files changed, 43 insertions(+)

diff --git a/src/etnaviv/isa/etnaviv.xml b/src/etnaviv/isa/etnaviv.xml
index 5841f9aa06d..16389fab574 100644
--- a/src/etnaviv/isa/etnaviv.xml
+++ b/src/etnaviv/isa/etnaviv.xml
@@ -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 -->
diff --git a/src/etnaviv/isa/tests/disasm.cpp b/src/etnaviv/isa/tests/disasm.cpp
index df3c0cfb93d..a1755cd2290 100644
--- a/src/etnaviv/isa/tests/disasm.cpp
+++ b/src/etnaviv/isa/tests/disasm.cpp
@@ -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"},