path: root/envydis/g80.c

/*
 * Copyright (C) 2009-2011 Marcin Kościelnicki <koriakin@0x04.net>
 * All Rights Reserved.
 *
 * Permission is hereby granted, free of charge, to any person obtaining a
 * copy of this software and associated documentation files (the "Software"),
 * to deal in the Software without restriction, including without limitation
 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
 * and/or sell copies of the Software, and to permit persons to whom the
 * Software is furnished to do so, subject to the following conditions:
 *
 * The above copyright notice and this permission notice (including the next
 * paragraph) shall be included in all copies or substantial portions of the
 * Software.
 *
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
 * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
 * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
 * OTHER DEALINGS IN THE SOFTWARE.
 */

#include "dis-intern.h"

/*
 * Table of Contents
 *  0. Table of Contents
 *  1. Instructions
 *  2. Differences from PTX
 *  3. General instruction format
 *  4. Misc. hack alerts
 *  5. Color scheme
 *  6. Table format
 *  7. Code target field
 *  8. Immediate field
 *  9. Misc number fields
 * 10. Register fields
 * 11. Memory fields
 * 12. Short instructions
 * 13. Immediate instructions
 * 14. Long instructions
 * 15. Non-predicated long instructions
 */

/*
 * Instructions, from PTX manual:
 *   
 *   1. Integer Arithmetic
 *    - sad.64		TODO	emulated
 *    - div		TODO	emulated
 *    - rem		TODO	emulated
 *    - abs		desc...
 *    - neg		desc...
 *   2. Floating-Point
 *    - add		desc...
 *    - sub		desc...
 *    - mul		desc...
 *    - fma		XXX
 *    - mad		XXX
 *    - div.approxf32	desc...
 *    - div.full.f32	TODO	emulated
 *    - div.f64		TODO	emulated
 *    - abs		desc...
 *    - neg		desc...
 *    - min		desc...
 *    - max		desc...
 *    - rcp.f32		desc...
 *    - rcp.f64		TODO	emulated
 *    - sqrt.f32	desc...
 *    - sqrt.f64	TODO	emulated
 *    - rsqrt.f32	desc...
 *    - rsqrt.f64	TODO	emulated
 *    - sin		XXX
 *    - cos		XXX
 *    - lg2		desc
 *    - ex2		XXX
 *   3. Comparison and selection
 *    - selp		desc...
 *    - slct		desc...
 *   5. Data Movement and Conversion
 *    - mov		desc...
 *    - ld		desc...
 *    - st		desc...
 *    - cvt		XXX cleanup, rounding modes
 *   6. Texture
 *    - tex		desc...
 *   7. Control Flow
 *    - { }		done
 *    - @		done
 *    - bra		desc...
 *    - call		XXX WTF?
 *    - ret		desc...
 *    - exit		desc...
 *   8. Parallel Synchronization and Communication
 *    - bar		desc...
 *    - membar.cta	desc...
 *    - membar.gl	sweet mother of God, what?
 *    - atom		desc...
 *    - red		desc...
 *    - vote		desc...
 *   9. Miscellaneous
 *    - trap		desc...
 *    - brkpt		desc...
 *    - pmevent		desc... XXX
 *  10. Undocumented/implicit stuff
 *    - join		XXX what does it do?
 *    - joinat		desc...
 *    - interp		XXX Needs figuring out.
 *    - discard		XXX what happens after you use it?
 *    - emit/reset	XXX can do combo?
 *
 */

/*
 * ISA differences from what's published about PTX:
 *
 * Instructions that don't really exist and are emulated:
 *  - sqrt.f32 -> rsqrt, rcp
 *  - neg and abs on f32 -> cvt
 *  - cvt.sat to f64 is implemented in software.
 *  - cvt with floating point can only convert up/down one size level, you
 *    need a chain of floats to go f32->u8 and others.
 *  - no 64-bit loads/stores except in l[] and g[].
 *  - vote.uni is emulated by vote.all
 *  - membar.cta is actually a nop that just forces instruction ordering
 *    before and after it.
 *  - div.approx.f32 doesn't exist and is replaced by mul by rcp
 *  - selp and slct. Emulated with set and predicated mov's.
 * Additional instructions and other stuff:
 *  - a lot of instructions can accept c?[], s[] directly
 *  - cvt can optionally abs and neg stuff before use. the sequence of
 *    operations is: abs, neg, convert+round, sat
 * Instructions that behave otherwise than you'd expect:
 *  - sin, cos, ex2 are split into two parts: function and prefunction. you
 *    need to run prefunction on input, and the real function on output from
 *    prefunction. no idea yet what the prefunctions are. presin is same as
 *    precos. presin could be mod-pi, what about preex2?
 *  - cvt to integer types always saturates. if you need chopping, just ignore
 *    the high part [or upconvert from a low part of register to full]
 *  - cvt.sat to f32 doesn't convert NaNs to 0. max.f32 with 0 to finish it off
 *  - cvt.rpi.f64 may or may not work on 1.25, add a specific check for it
 *  - same for rmi and -1.25.
 * Misc:
 *  - explicit join points
 *  - $a registers exist
 *  - there are actually 16 const spaces and global spaces.
 * Incl. graphics stuff:
 *  - v[] space exists
 *  - interp instruction
 *  - discard instruction
 *  - $o registers exist
 *
 * XXX this section is probably useless, remove it once proper ISA
 * documentation exists.
 */

/*
 * Misc. hack alerts:
 *  - 0x2c-0x2d is read twice in addc instructions: it selects $c to use for
 *    a carry flag, and selects $c to use for conditional execution. Printing
 *    it depends on zeroing happening after decoding the instruction.
 *    Same thing also happens in mov-from-$c instruction.
 *  - If instruction has both s[] and c?[] memory operands and sets the AREG
 *    field to non-0, specified $a only applies to the first one. This depends
 *    on bits in *a and *b getting zeroed out by function printing out the first
 *    argument. For added fun, a[] behaves differently.
 *  - long instruction always need to be aligned to even word bounduary.
 *    However, I disabled checking for this case, since renouveau dumps contain
 *    shader code starting at random positions, and the check is mostly useless
 *    anyway.
 */

#define F_SM11		0x01
#define F_SM12		0x02
#define F_FP64		0x04
#define F_D3D10_1	0x08

#define VP 1
#define GP 2
#define FP 4
#define CP 8

/*
 * Code target field
 *
 * This field represents a code address and is used for branching and the
 * likes. Target is counted in 32-bit words from the start point of code
 * segment. Low part of target is in 0/11-26, high part is in 1/14-19. Code
 * addresses max out at 0xfffffc.
 */

static struct rbitfield ctargoff = { { 0xb, 16, 0x2e, 6 }, RBF_UNSIGNED, 2 };
#define BTARG atombtarg, &ctargoff
#define CTARG atomctarg, &ctargoff

/*
 * Immediate fields
 *
 *  - IMM: Used by all immediate insns [tabi table] for last argument. Full 32-bit.
 *  - PM: PM event id for pmevent
 *  - BAR: barrier id for bar.sync
 *  - OFFS: an offset to add in add $a, num, $a, which is basically a LEA
 *  - SHCNT: used in shift-by-immediate insns for shift amount
 *  - HSHCNT: used in $a shift-by-immediate insn for shift amount
 */

static struct rbitfield immbf = { { 0x10, 6, 0x22, 26 }, .wrapok = 1 };
static struct rbitfield offoff = { { 9, 16 }, .wrapok = 1 };
#define IMM atomrimm, &immbf
#define OFFS atomrimm, &offoff

static struct bitfield pmoff = { 0xa, 4 };
static struct bitfield baroff = { 0x15, 4 };
static struct bitfield barmaskoff = { 9, 12 };
static struct bitfield shcntoff = { 0x10, 7 };
static struct bitfield hshcntoff = { 0x10, 4 };
static struct bitfield toffxoff = { { 0x38, 4 }, BF_SIGNED };
static struct bitfield toffyoff = { { 0x34, 4 }, BF_SIGNED };
static struct bitfield toffzoff = { { 0x30, 4 }, BF_SIGNED };
#define PM atomimm, &pmoff
#define BAR atomimm, &baroff
#define BARMASK atomimm, &barmaskoff
#define SHCNT atomimm, &shcntoff
#define HSHCNT atomimm, &hshcntoff
#define TOFFX atomimm, &toffxoff
#define TOFFY atomimm, &toffyoff
#define TOFFZ atomimm, &toffzoff

/*
 * Ignored fields
 *
 * Used in cases when some bitfield doesn't do anything, but nvidia's blob
 * sets it due to stupi^H^H^H^H^Hdesign decisions.
 *  - IGNCE: $c write enable
 *  - IGNPRED: predicates, for instructions that don't use them.
 *
 * The whole point of this op is to kill spurious red when disassembling
 * ptxas output and blob shaders. Don't include all random unused fields here.
 */

static int ignce[] = { 0x26, 1 };
static int ignpred[] = { 0x27, 7 };
#define IGNCE atomign, ignce
#define IGNPRED atomign, ignpred

/*
 * Register fields
 *
 * There are four locations of register fields for $r:
 *  - DST: 2 and up, used for destination reg of most insns, but for source
 *         in store instructions. Insn seems to get ignored if this field is
 *         > number of allocated regs, so don't use that for storing 0s.
 *  - SRC: 9 and up, used as a source register in all types of insns.
 *  - SRC2: 0x10 and up, used as a source register in short and long insns.
 *  - SRC3: 0x2e and up, used as s source register in long insns.
 *
 * These fields can be unused, or used as part of opcode or modifiers if insn
 * doesn't need all available fields.
 *
 * All fields are either 6-bit [S] or 7-bit [L]. Short insns and non-mov
 * immediate insns have 6-bit fields and use the high bit as opcode/mods,
 * while mov immediate and long insns use full 7-bit fields.
 *
 * If encoded registers are 32-bit [no modifier letter] or 64-bit [D], the
 * field contains register number directly. For 16-bit halves of registers [H],
 * lowest bit selects low [0] or high [1] half of given register, and higher
 * bits contain the number of register. For 64-bit registers, the number needs
 * to be even.
 *
 * DST field seems to sometimes store $o registers instead of $r. If it can,
 * using $o instead of $r is indicated by bit 0x23 being set. LLDST means
 * either $o or $r in DST. It's unknown yet in what instructions it can be
 * used [probably all non-tex non-access-g[] longs?]
 *
 * ADST is a smaller version of DST field used for storing $a destination in
 * insns that allow one.
 *
 * COND is $c register used as input, both for predicating and other purposes,
 * like addition with carry or mov from $c.
 *
 * CDST is $c register used for output. It's set according to insn result for
 * most insns that allow it. Writing to $c in this field needs to be enabled
 * by setting bit 0x26. MCDST is an op that reflects that and is used for insns
 * that can function both with or without $c output.
 * 
 * AREG is a field storing $a register used in insn for memory addressing,
 * or as a normal source register sometimes. It's especially ugly, because
 * it's spread out across 0x1a-0x1b [low part] and 0x22 [high part], with the
 * high part assumed 0 if insn is short or immediate. Also, if two memory
 * accesses are used in the same insn, this field applies only to the first.
 *
 */

static struct sreg areg_sr[] = {
	{ 0, 0, SR_ZERO },
	{ -1 },
};
static struct sreg sreg_sr[] = {
	{ 0, "physid" },
	{ 1, "clock" },
	{ 3, "vstride" },
	{ 4, "pm0" },
	{ 5, "pm1" },
	{ 6, "pm2" },
	{ 7, "pm3" },
	{ 8, "sampleid", .fmask = F_D3D10_1 },
	{ -1 },
};
static struct bitfield sdst_bf = { 2, 6 };
static struct bitfield ldst_bf = { 2, 7 };
static struct bitfield ssrc_bf = { 9, 6 };
static struct bitfield lsrc_bf = { 9, 7 };
static struct bitfield ssrc2_bf = { 0x10, 6 };
static struct bitfield lsrc2_bf = { 0x10, 7 };
static struct bitfield lsrc3_bf = { 0x2e, 7 };
static struct bitfield sareg_bf = { 0x1a, 2 };
static struct bitfield lareg_bf = { { 0x1a, 2, 0x22, 1 } };
static struct bitfield adst_bf = { 2, 3 };
static struct bitfield cond_bf = { 0x2c, 2 };
static struct bitfield c0_bf = { 0, 0 };
static struct bitfield cdst_bf = { 0x24, 2 };
static struct bitfield tex_bf = { 9, 8 };
static struct bitfield samp_bf = { 0x11, 5 };
static struct bitfield sreg_bf = { 0x2e, 4 };

static struct reg sdst_r = { &sdst_bf, "r" };
static struct reg ldst_r = { &ldst_bf, "r" };
static struct reg shdst_r = { &sdst_bf, "r", .hilo = 1 };
static struct reg lhdst_r = { &ldst_bf, "r", .hilo = 1 };
static struct reg lddst_r = { &ldst_bf, "r", "d" };
static struct reg lqdst_r = { &ldst_bf, "r", "q" };
static struct reg ssrc_r = { &ssrc_bf, "r" };
static struct reg lsrc_r = { &lsrc_bf, "r" };
static struct reg shsrc_r = { &ssrc_bf, "r", .hilo = 1 };
static struct reg lhsrc_r = { &lsrc_bf, "r", .hilo = 1 };
static struct reg ldsrc_r = { &lsrc_bf, "r", "d" };
static struct reg ssrc2_r = { &ssrc2_bf, "r" };
static struct reg lsrc2_r = { &lsrc2_bf, "r" };
static struct reg shsrc2_r = { &ssrc2_bf, "r", .hilo = 1 };
static struct reg lhsrc2_r = { &lsrc2_bf, "r", .hilo = 1 };
static struct reg ldsrc2_r = { &lsrc2_bf, "r", "d" };
static struct reg lsrc3_r = { &lsrc3_bf, "r" };
static struct reg lhsrc3_r = { &lsrc3_bf, "r", .hilo = 1 };
static struct reg ldsrc3_r = { &lsrc3_bf, "r", "d" };
static struct reg sareg_r = { &sareg_bf, "a", .specials = areg_sr, .cool = 1 }; // for mem operands only
static struct reg lareg_r = { &lareg_bf, "a", .specials = areg_sr, .cool = 1 };
static struct reg adst_r = { &adst_bf, "a", .specials = areg_sr, .cool = 1 };
static struct reg cond_r = { &cond_bf, "c", .cool = 1 };
static struct reg c0_r = { &c0_bf, "c", .cool = 1 };
static struct reg cdst_r = { &cdst_bf, "c", .cool = 1 };
static struct reg tex_r = { &tex_bf, "t", .cool = 1 };
static struct reg samp_r = { &samp_bf, "s", .cool = 1 };
static struct reg sreg_r = { &sreg_bf, "sr", .specials = sreg_sr, .always_special = 1 };
#define SDST atomreg, &sdst_r
#define LDST atomreg, &ldst_r
#define SHDST atomreg, &shdst_r
#define LHDST atomreg, &lhdst_r
#define LDDST atomreg, &lddst_r
#define LQDST atomreg, &lqdst_r
#define SSRC atomreg, &ssrc_r
#define LSRC atomreg, &lsrc_r
#define SHSRC atomreg, &shsrc_r
#define LHSRC atomreg, &lhsrc_r
#define LDSRC atomreg, &ldsrc_r
#define SSRC2 atomreg, &ssrc2_r
#define LSRC2 atomreg, &lsrc2_r
#define SHSRC2 atomreg, &shsrc2_r
#define LHSRC2 atomreg, &lhsrc2_r
#define LDSRC2 atomreg, &ldsrc2_r
#define LSRC3 atomreg, &lsrc3_r
#define LHSRC3 atomreg, &lhsrc3_r
#define LDSRC3 atomreg, &ldsrc3_r
#define LAREG atomreg, &lareg_r
#define ADST atomreg, &adst_r
#define COND atomreg, &cond_r
#define C0 atomreg, &c0_r
#define CDST atomreg, &cdst_r
#define TEX atomreg, &tex_r
#define SAMP atomreg, &samp_r
#define SREG atomreg, &sreg_r

static struct bitfield tdst_cnt = { .addend = 4 };
static struct bitfield ltdst_mask = { { 0x19, 2, 0x2e, 2 } };
static struct bitfield tsrc_cnt = { { 0x16, 2 }, .addend = 1 };
static struct vec ltdst_v = { "r", &ldst_bf, &tdst_cnt, &ltdst_mask };
static struct vec stdst_v = { "r", &sdst_bf, &tdst_cnt, 0 };
static struct vec ltsrc_v = { "r", &ldst_bf, &tsrc_cnt, 0 };
static struct vec stsrc_v = { "r", &sdst_bf, &tsrc_cnt, 0 };
#define LTDST atomvec, &ltdst_v
#define STDST atomvec, &stdst_v
#define LTSRC atomvec, &ltsrc_v
#define STSRC atomvec, &stsrc_v

#define MCDST T(mcdst)
F1(mcdst, 0x26, CDST)

/*
 * Memory fields
 *
 * Memory accesses on g80 are generally of the form x[$a+offset]. $a register
 * is taken from AREG and applies only to first memory operand actually used
 * in a given insn. Offset is taken from a field in insn and maybe multiplied
 * by operand size before use.
 *
 * Fields for reading shared/attribute spaces:
 *  S*SHARED: s/a space, multiplied, shares field with S*SRC.
 *    Used in many short/immediate insns instead of S*SRC when 0/24 set.
 *  L*SHARED, s/a space, multiplied, shares field with L*SRC.
 *    Used in many long insns instead of L*SRC if 1/21 set.
 * High 2 bits of both of these fields are some sort of subspace specifier.
 * Known ones include:
 *  For 8-bit operands: 00 - s[] space.
 *  For 16-bit operands: 01 - s[] space for unsigned operands, 10 - s[] space
 *    for signed operands. I have no idea what's the difference.
 *  For 32-bit operands: 00 - a[] space, 11 - s[] space.
 * Low 4 or 5 bits are offset that gets multiplied by access size before use.
 *
 * Fields for reading const spaces:
 *  S*CONST: c* space, multiplied, shares field with S*SRC2.
 *    Used in many short insns instead of S*SRC2 if 0/23 set.
 *  L*CONST2: c* space, multiplied, shares field with L*SRC2.
 *    Used in many long insns instead of L*SRC2 if 1/23 set.
 *  L*CONST3: c* space, multiplied, shares field with L*SRC3.
 *    Used in many long insns instead of L*SRC3 if 1/24 set.
 *  F*CONST: c* space, multiplied, uses 15-bit offset. Seen only on special
 *    movs from c* spaces.
 * L*CONST* and F*CONST have constant space number in 1/22-25. S*CONST uses
 * only 5 bits of its 6-bit field, presumably 0/21 selects c1 or c0, but needs
 * to be checked.
 *
 * For some reason ptxas hasn't been seen emitting both S*CONST and S*SHARED
 * nor L*CONST2 and L*CONST3 in the same insn. But L*SHARED with L*CONST[23]
 * works ok.
 *
 * Fields for storing to shared space:
 *  F*SHARED: s space, multiplied. Used only in mov to s[].
 *
 * Fields for local space:
 *  LOCAL: l space, not multiplied. Used only in movs to/from l[].
 *
 * g[] space is totally different: its only addressing mode is g[$r], where
 * $r is taken from LSRC operand.
 */

// BF, offset shift, 'l', flags, const space num BF. flags: 1 supports $a, 2 supports full 3-bit $a, 4 supports autoincrement

static struct rbitfield fmem8_imm = { { 9, 16 }, RBF_UNSIGNED, 0, .wrapok = 1 };
static struct rbitfield fmem16_imm = { { 9, 15 }, RBF_UNSIGNED, 1, .wrapok = 1 };
static struct rbitfield fmem32_imm = { { 9, 14 }, RBF_UNSIGNED, 2, .wrapok = 1 };

static struct rbitfield ssmem8_imm = { { 9, 4 }, RBF_UNSIGNED, 0 };
static struct rbitfield ssmem16_imm = { { 9, 4 }, RBF_UNSIGNED, 1 };
static struct rbitfield ssmem32_imm = { { 9, 4 }, RBF_UNSIGNED, 2 };
static struct rbitfield lsmem8_imm = { { 9, 5 }, RBF_UNSIGNED, 0 };
static struct rbitfield lsmem16_imm = { { 9, 5 }, RBF_UNSIGNED, 1 };
static struct rbitfield lsmem32_imm = { { 9, 5 }, RBF_UNSIGNED, 2 };
static struct rbitfield ssmem8pi_imm = { { 9, 4 }, RBF_SIGNED, 0 };
static struct rbitfield ssmem16pi_imm = { { 9, 4 }, RBF_SIGNED, 1 };
static struct rbitfield ssmem32pi_imm = { { 9, 4 }, RBF_SIGNED, 2 };
static struct rbitfield lsmem8pi_imm = { { 9, 5 }, RBF_SIGNED, 0 };
static struct rbitfield lsmem16pi_imm = { { 9, 5 }, RBF_SIGNED, 1 };
static struct rbitfield lsmem32pi_imm = { { 9, 5 }, RBF_SIGNED, 2 };
static struct mem ssmem8_m = { "s", 0, &sareg_r, &ssmem8_imm };
static struct mem ssmem16_m = { "s", 0, &sareg_r, &ssmem16_imm };
static struct mem ssmem32_m = { "s", 0, &sareg_r, &ssmem32_imm };
static struct mem lsmem8_m = { "s", 0, &lareg_r, &lsmem8_imm };
static struct mem lsmem16_m = { "s", 0, &lareg_r, &lsmem16_imm };
static struct mem lsmem32_m = { "s", 0, &lareg_r, &lsmem32_imm };
static struct mem ssmem8pi_m = { "s", 0, &sareg_r, &ssmem8pi_imm, .postincr = 1 };
static struct mem ssmem16pi_m = { "s", 0, &sareg_r, &ssmem16pi_imm, .postincr = 1 };
static struct mem ssmem32pi_m = { "s", 0, &sareg_r, &ssmem32pi_imm, .postincr = 1 };
static struct mem lsmem8pi_m = { "s", 0, &lareg_r, &lsmem8pi_imm, .postincr = 1 };
static struct mem lsmem16pi_m = { "s", 0, &lareg_r, &lsmem16pi_imm, .postincr = 1 };
static struct mem lsmem32pi_m = { "s", 0, &lareg_r, &lsmem32pi_imm, .postincr = 1 };
static struct mem fsmem8_m = { "s", 0, &lareg_r, &fmem8_imm };
static struct mem fsmem16_m = { "s", 0, &lareg_r, &fmem16_imm };
static struct mem fsmem32_m = { "s", 0, &lareg_r, &fmem32_imm };
static struct mem fsmem8pi_m = { "s", 0, &lareg_r, &fmem8_imm, .postincr = 1 };
static struct mem fsmem16pi_m = { "s", 0, &lareg_r, &fmem16_imm, .postincr = 1 };
static struct mem fsmem32pi_m = { "s", 0, &lareg_r, &fmem32_imm, .postincr = 1 };
#define SSHARED8 atommem, &ssmem8_m
#define SSHARED16 atommem, &ssmem16_m
#define SSHARED32 atommem, &ssmem32_m
#define LSHARED8 atommem, &lsmem8_m
#define LSHARED16 atommem, &lsmem16_m
#define LSHARED32 atommem, &lsmem32_m
#define SSHARED8PI atommem, &ssmem8pi_m
#define SSHARED16PI atommem, &ssmem16pi_m
#define SSHARED32PI atommem, &ssmem32pi_m
#define LSHARED8PI atommem, &lsmem8pi_m
#define LSHARED16PI atommem, &lsmem16pi_m
#define LSHARED32PI atommem, &lsmem32pi_m
#define FSHARED8 atommem, &fsmem8_m
#define FSHARED16 atommem, &fsmem16_m
#define FSHARED32 atommem, &fsmem32_m
#define FSHARED8PI atommem, &fsmem8pi_m
#define FSHARED16PI atommem, &fsmem16pi_m
#define FSHARED32PI atommem, &fsmem32pi_m
F(fs8, 0x19, FSHARED8, FSHARED8PI);
F(fs16, 0x19, FSHARED16, FSHARED16PI);
F(fs32, 0x19, FSHARED32, FSHARED32PI);

// XXX: try postincr on remaining types
static struct bitfield scmem_idx = { 0x15, 1 };
static struct bitfield lcmem_idx = { 0x36, 4 };
static struct rbitfield scmem16_imm = { { 0x10, 5 }, RBF_UNSIGNED, 1 };
static struct rbitfield scmem32_imm = { { 0x10, 5 }, RBF_UNSIGNED, 2 };
static struct rbitfield scmem16pi_imm = { { 0x10, 5 }, RBF_SIGNED, 1 };
static struct rbitfield scmem32pi_imm = { { 0x10, 5 }, RBF_SIGNED, 2 };
static struct rbitfield l2cmem16_imm = { { 0x10, 7 }, RBF_UNSIGNED, 1 };
static struct rbitfield l2cmem32_imm = { { 0x10, 7 }, RBF_UNSIGNED, 2 };
static struct rbitfield l3cmem16_imm = { { 0x2e, 7 }, RBF_UNSIGNED, 1 };
static struct rbitfield l3cmem32_imm = { { 0x2e, 7 }, RBF_UNSIGNED, 2 };
static struct rbitfield l2cmem16pi_imm = { { 0x10, 7 }, RBF_SIGNED, 1 };
static struct rbitfield l2cmem32pi_imm = { { 0x10, 7 }, RBF_SIGNED, 2 };
static struct rbitfield l3cmem16pi_imm = { { 0x2e, 7 }, RBF_SIGNED, 1 };
static struct rbitfield l3cmem32pi_imm = { { 0x2e, 7 }, RBF_SIGNED, 2 };
static struct mem scmem16_m = { "c", &scmem_idx, &sareg_r, &scmem16_imm };
static struct mem scmem32_m = { "c", &scmem_idx, &sareg_r, &scmem32_imm };
static struct mem scmem16pi_m = { "c", &scmem_idx, &sareg_r, &scmem16pi_imm, .postincr = 1 };
static struct mem scmem32pi_m = { "c", &scmem_idx, &sareg_r, &scmem32pi_imm, .postincr = 1 };
static struct mem l2cmem16_m = { "c", &lcmem_idx, &lareg_r, &l2cmem16_imm };
static struct mem l2cmem32_m = { "c", &lcmem_idx, &lareg_r, &l2cmem32_imm };
static struct mem l3cmem16_m = { "c", &lcmem_idx, &lareg_r, &l3cmem16_imm };
static struct mem l3cmem32_m = { "c", &lcmem_idx, &lareg_r, &l3cmem32_imm };
static struct mem l2cmem16pi_m = { "c", &lcmem_idx, &lareg_r, &l2cmem16pi_imm, .postincr = 1 };
static struct mem l2cmem32pi_m = { "c", &lcmem_idx, &lareg_r, &l2cmem32pi_imm, .postincr = 1 };
static struct mem l3cmem16pi_m = { "c", &lcmem_idx, &lareg_r, &l3cmem16pi_imm, .postincr = 1 };
static struct mem l3cmem32pi_m = { "c", &lcmem_idx, &lareg_r, &l3cmem32pi_imm, .postincr = 1 };
static struct mem l2cmem16na_m = { "c", &lcmem_idx, 0, &l2cmem16_imm };
static struct mem l2cmem32na_m = { "c", &lcmem_idx, 0, &l2cmem32_imm };
static struct mem l3cmem16na_m = { "c", &lcmem_idx, 0, &l3cmem16_imm };
static struct mem l3cmem32na_m = { "c", &lcmem_idx, 0, &l3cmem32_imm };
static struct mem fcmem8_m = { "c", &lcmem_idx, &lareg_r, &fmem8_imm };
static struct mem fcmem16_m = { "c", &lcmem_idx, &lareg_r, &fmem16_imm };
static struct mem fcmem32_m = { "c", &lcmem_idx, &lareg_r, &fmem32_imm };
static struct mem fcmem8pi_m = { "c", &lcmem_idx, &lareg_r, &fmem8_imm, .postincr = 1 };
static struct mem fcmem16pi_m = { "c", &lcmem_idx, &lareg_r, &fmem16_imm, .postincr = 1 };
static struct mem fcmem32pi_m = { "c", &lcmem_idx, &lareg_r, &fmem32_imm, .postincr = 1 };
#define SCONST16 atommem, &scmem16_m
#define SCONST32 atommem, &scmem32_m
#define SCONST16PI atommem, &scmem16pi_m
#define SCONST32PI atommem, &scmem32pi_m
#define L2CONST16 atommem, &l2cmem16_m
#define L2CONST32 atommem, &l2cmem32_m
#define L3CONST16 atommem, &l3cmem16_m
#define L3CONST32 atommem, &l3cmem32_m
#define L2CONST16PI atommem, &l2cmem16pi_m
#define L2CONST32PI atommem, &l2cmem32pi_m
#define L3CONST16PI atommem, &l3cmem16pi_m
#define L3CONST32PI atommem, &l3cmem32pi_m
#define L2CONST16NA atommem, &l2cmem16na_m
#define L2CONST32NA atommem, &l2cmem32na_m
#define L3CONST16NA atommem, &l3cmem16na_m
#define L3CONST32NA atommem, &l3cmem32na_m
#define FCONST8 atommem, &fcmem8_m
#define FCONST16 atommem, &fcmem16_m
#define FCONST32 atommem, &fcmem32_m
#define FCONST8PI atommem, &fcmem8pi_m
#define FCONST16PI atommem, &fcmem16pi_m
#define FCONST32PI atommem, &fcmem32pi_m

static struct mem lmem_m = { "l", 0, &lareg_r, &fmem8_imm };
static struct mem lmempi_m = { "l", 0, &lareg_r, &fmem8_imm, .postincr = 1 };
#define LOCAL atommem, &lmem_m
#define LOCALPI atommem, &lmempi_m
F(local, 0x19, LOCAL, LOCALPI);

// XXX: try for postincr
static struct rbitfield samem_imm = { { 9, 6 }, RBF_UNSIGNED, 2 };
static struct rbitfield lamem_imm = { { 9, 7 }, RBF_UNSIGNED, 2 };
static struct mem samem_m = { "a", 0, 0, &samem_imm };
static struct mem lamem_m = { "a", 0, 0, &lamem_imm };
static struct mem famem_m = { "a", 0, &lareg_r, &lamem_imm };
static struct mem spmem_m = { "p", 0, &sareg_r, &samem_imm };
static struct mem lpmem_m = { "p", 0, &lareg_r, &lamem_imm };
#define SATTR atommem, &samem_m
#define LATTR atommem, &lamem_m
#define FATTR atommem, &famem_m
#define SPRIM atommem, &spmem_m
#define LPRIM atommem, &lpmem_m

// XXX: try for postincr
static struct rbitfield vmem_imm = { { 0x10, 8 }, RBF_UNSIGNED, 2 };
static struct mem svmem_m = { "v", 0, &sareg_r, &vmem_imm };
static struct mem lvmem_m = { "v", 0, &lareg_r, &vmem_imm };
#define SVAR atommem, &svmem_m
#define LVAR atommem, &lvmem_m

static struct bitfield global_idx = { 0x10, 4 };
static struct bitfield global2_idx = { 0x17, 4 };
static struct mem global_m = { "g", &global_idx, &lsrc_r };
static struct mem global2_m = { "g", &global2_idx, &lsrc_r };
#define GLOBAL atommem, &global_m
#define GLOBAL2 atommem, &global2_m

static struct rbitfield omem16_imm = { { 2, 7 }, .shr = 1 };
static struct rbitfield omem32_imm = { { 2, 7 }, .shr = 2 };
static struct mem omem16_m = { "o", 0, 0, &omem16_imm };
static struct mem omem32_m = { "o", 0, 0, &omem32_imm };
static struct mem fomem32_m = { "o", 0, &lareg_r, &fmem32_imm };
#define OUT16 atommem, &omem16_m
#define OUT32 atommem, &omem32_m
#define FOUT32 atommem, &fomem32_m

static struct insn tabss[] = {
	{ 0x01800000, 0x01800000, SPRIM, .ptype = GP },	// XXX check
	{ 0x01000000, 0x03806000, N("u8"), SSHARED8, .ptype = CP },
	{ 0x01002000, 0x03806000, N("u16"), SSHARED16, .ptype = CP },
	{ 0x01004000, 0x03806000, N("s16"), SSHARED16, .ptype = CP },
	{ 0x01006000, 0x03806000, N("b32"), SSHARED32, .ptype = CP },
	{ 0x03000000, 0x03806000, N("u8"), SSHARED8PI, .ptype = CP },
	{ 0x03002000, 0x03806000, N("u16"), SSHARED16PI, .ptype = CP },
	{ 0x03004000, 0x03806000, N("s16"), SSHARED16PI, .ptype = CP },
	{ 0x03006000, 0x03806000, N("b32"), SSHARED32PI, .ptype = CP },
	{ 0x01000000, 0x01800000, SATTR, .ptype = VP|GP },
	{ 0, 0, OOPS },
};

static struct insn tabrealls[] = {
	{ 0x00000000, 0x0200c000, N("u8"), LSHARED8, .ptype = CP },
	{ 0x00004000, 0x0200c000, N("u16"), LSHARED16, .ptype = CP },
	{ 0x00008000, 0x0200c000, N("s16"), LSHARED16, .ptype = CP },
	{ 0x0000c000, 0x0200c000, N("b32"), LSHARED32, .ptype = CP },
	{ 0x02000000, 0x0200c000, N("u8"), LSHARED8PI, .ptype = CP },
	{ 0x02004000, 0x0200c000, N("u16"), LSHARED16PI, .ptype = CP },
	{ 0x02008000, 0x0200c000, N("s16"), LSHARED16PI, .ptype = CP },
	{ 0x0200c000, 0x0200c000, N("b32"), LSHARED32PI, .ptype = CP },
	{ 0x00000000, 0x00000000, LATTR, .ptype = VP|GP },
	{ 0, 0, OOPS },
};

static struct insn tabaddop[] = {
	{ 0x00000000, 0x10400000, N("add") },
	{ 0x00400000, 0x10400000, N("sub") },
	{ 0x10000000, 0x10400000, N("subr") },
	{ 0x10400000, 0x10400000, N("addc") },
	{ 0, 0, OOPS },
};

/*
 * Short instructions
 */

// various stuff available for filling the misc bits.

F1(sm1sat, 8, N("sat"))
F1V(sm1satg200, F_SM12, 8, N("sat"))
F(sm1us16, 8, N("u16"), N("s16"))
F(sm1us32, 8, N("u32"), N("s32"))
F1(sm1high, 8, N("high"))
F(sm1tex, 8, N("all"), N("live"))

F1(sm2neg, 0xf, N("neg"))
F1(sm2abs, 0xf, N("abs"))
F(sm2us16, 0xf, N("u16"), N("s16"))
F(sm2us24, 0xf, N("u24"), N("s24"))

F1(sm3neg, 0x16, N("neg"))
F1(sm3not, 0x16, N("not"))

F1(ssfu, 0x11, N("sfu"))

static struct insn tabssh[] = {
	{ 0x00000000, 0x01000000, SHSRC },
	{ 0x01000000, 0x01000000, T(ss) },
	{ 0, 0, OOPS }
};

static struct insn tabssw[] = {
	{ 0x00000000, 0x01000000, SSRC },
	{ 0x01000000, 0x01000000, T(ss) },
	{ 0, 0, OOPS }
};

static struct insn tabsch[] = {
	{ 0x00000000, 0x01800000, SHSRC2 },
	{ 0x00800000, 0x03800000, SCONST16 },
	{ 0x02800000, 0x03800000, SCONST16PI },
	{ 0x01000000, 0x01800000, SHSRC2 },
	{ 0x01800000, 0x01800000, SHSRC2 },
	{ 0, 0, OOPS }
};

static struct insn tabscw[] = {
	{ 0x00000000, 0x01800000, SSRC2 },
	{ 0x00800000, 0x03800000, SCONST32 },
	{ 0x02800000, 0x03800000, SCONST32PI },
	{ 0x01000000, 0x01800000, SSRC2 },
	{ 0x01800000, 0x01800000, SSRC2 },
	{ 0, 0, OOPS }
};

static struct insn tabaddc0[] = {
	{ 0x00000000, 0x10400000 },
	{ 0x00400000, 0x10400000 },
	{ 0x10000000, 0x10400000 },
	{ 0x10400000, 0x10400000, C0 },
	{ 0, 0, OOPS }
};

static struct insn tabs[] = {
	{ 0x10000000, 0xf0008002, N("mov"), T(ssfu), N("b16"), SHDST, T(ssh) },
	{ 0x10008000, 0xf0008002, N("mov"), T(ssfu), N("b32"), SDST, T(ssw) },

	{ 0x20000000, 0xe0008002, T(addop), T(sm1sat), N("b16"), SHDST, T(ssh), T(sch), T(addc0) },
	{ 0x20008000, 0xe0008002, T(addop), T(sm1sat), N("b32"), SDST, T(ssw), T(scw), T(addc0) },

	{ 0x40000000, 0xf0400002, N("mul"), SDST, T(sm2us16), T(ssh), T(sm1us16), T(sch) },
	{ 0x40400000, 0xf0400002, N("mul"), SDST, T(sm1high), T(sm2us24), T(ssw), T(scw) },

	{ 0x50000000, 0xf0008002, N("sad"), SDST, T(sm1us16), SHSRC, SHSRC2, SDST },
	{ 0x50008000, 0xf0008002, N("sad"), SDST, T(sm1us32), SSRC, SSRC2, SDST },

	{ 0x60000000, 0xe0008102, T(addop), SDST, SESTART, N("mul"), N("u16"), T(ssh), T(sch), SEEND, SDST, T(addc0) },
	{ 0x60000100, 0xe0008102, T(addop), SDST, SESTART, N("mul"), N("s16"), T(ssh), T(sch), SEEND, SDST, T(addc0) },
	{ 0x60008000, 0xe0008102, T(addop), N("sat"), SDST, SESTART, N("mul"), N("s16"), T(ssh), T(sch), SEEND, SDST, T(addc0) },
	{ 0x60008100, 0xe0008102, T(addop), SDST, SESTART, N("mul"), N("u24"), T(ssw), T(scw), SEEND, SDST, T(addc0) },

	// desc VVV
	{ 0x80000000, 0xf3000102, N("interp"), SDST, SVAR, .ptype = FP },		// most likely something is wrong with this.
	{ 0x81000000, 0xf3000102, N("interp"), SDST, N("cent"), SVAR, .ptype = FP },
	{ 0x82000000, 0xf3000102, N("interp"), SDST, SVAR, SSRC, .ptype = FP },
	{ 0x83000000, 0xf3000102, N("interp"), SDST, N("cent"), SVAR, SSRC, .ptype = FP },
	{ 0x80000100, 0xf3000102, N("interp"), SDST, N("flat"), SVAR, .ptype = FP },

	{ 0x90000000, 0xf0000002, N("rcp"), N("f32"), SDST, T(sm3neg), T(sm2abs), SSRC },

	{ 0xb0000000, 0xf0000002, N("add"), T(sm1sat), N("f32"), SDST, T(sm2neg), T(ssw), T(sm3neg), T(scw) },

	{ 0xc0000000, 0xf0000002, N("mul"), T(sm1satg200), N("f32"), SDST, T(sm2neg), T(ssw), T(sm3neg), T(scw) },

	{ 0xe0000000, 0xf0000002, N("add"), T(sm1sat), N("f32"), SDST, T(sm2neg), SESTART, N("mul"), T(ssw), T(scw), SEEND, T(sm3neg), SDST },

	{ 0xf0000000, 0xf1000002, N("texauto"), T(sm1tex), STDST, TEX, SAMP, STSRC },
	{ 0xf1000000, 0xf1000002, N("texfetch"), T(sm1tex), STDST, TEX, SAMP, STSRC},
	// desc ^^^

	{ 0, 2, OOPS, SDST, T(ssw), T(scw) },

	{ 0x90000002, 0xf0000002, N("trap") },
	{ 0xb0000002, 0xf0000002, N("brkpt"), .fmask = F_SM11 },

	{ 0, 0, OOPS }
};

/*
 * Immediate instructions
 */

static struct insn tabis[] = {
	{ 0x01000000, 0x01000000, SPRIM, .ptype = GP },	// XXX check
	{ 0x01000000, 0x03006000, N("u8"), SSHARED8, .ptype = CP },
	{ 0x01002000, 0x03006000, N("u16"), SSHARED16, .ptype = CP },
	{ 0x01004000, 0x03006000, N("s16"), SSHARED16, .ptype = CP },
	{ 0x01006000, 0x03006000, N("b32"), SSHARED32, .ptype = CP },
	{ 0x03000000, 0x03006000, N("u8"), SSHARED8PI, .ptype = CP },
	{ 0x03002000, 0x03006000, N("u16"), SSHARED16PI, .ptype = CP },
	{ 0x03004000, 0x03006000, N("s16"), SSHARED16PI, .ptype = CP },
	{ 0x03006000, 0x03006000, N("b32"), SSHARED32PI, .ptype = CP },
	{ 0, 0, OOPS },
};

static struct insn tabish[] = {
	{ 0x00000000, 0x01000000, SHSRC },
	{ 0x01000000, 0x01000000, T(is) },
	{ 0, 0, OOPS }
};

static struct insn tabisw[] = {
	{ 0x00000000, 0x01000000, SSRC },
	{ 0x01000000, 0x01000000, T(is) },
	{ 0, 0, OOPS }
};

static struct insn tabi[] = {
	{ 0x10000000, 0xf0008000, N("mov"), N("b16"), LHDST, IMM },
	{ 0x10008000, 0xf0008000, N("mov"), N("b32"), LDST, IMM },	// yes. LDST. special case.

	{ 0x20000000, 0xe0008000, T(addop), T(sm1sat), N("b16"), SHDST, T(ish), IMM, T(addc0) },
	{ 0x20008000, 0xe0008000, T(addop), T(sm1sat), N("b32"), SDST, T(isw), IMM, T(addc0) },

	{ 0x40000000, 0xf0400000, N("mul"), SDST, T(sm2us16), T(ish), T(sm1us16), IMM },
	{ 0x40400000, 0xf0400000, N("mul"), SDST, T(sm1high), T(sm2us24), T(isw), IMM },

	{ 0x60000000, 0xe0008100, T(addop), SDST, SESTART, N("mul"), N("u16"), T(ish), IMM, SEEND, SDST, T(addc0) },
	{ 0x60000100, 0xe0008100, T(addop), SDST, SESTART, N("mul"), N("s16"), T(ish), IMM, SEEND, SDST, T(addc0) },
	{ 0x60008000, 0xe0008100, T(addop), N("sat"), SDST, SESTART, N("mul"), N("s16"), T(ish), IMM, SEEND, SDST, T(addc0) },
	{ 0x60008100, 0xe0008100, T(addop), SDST, SESTART, N("mul"), N("u24"), T(isw), IMM, SEEND, SDST, T(addc0) },

	// desc VVV
	{ 0xb0000000, 0xf0000000, N("add"), T(sm1sat), N("f32"), SDST, T(sm2neg), T(isw), T(sm3neg), IMM },

	{ 0xc0000000, 0xf0000000, N("mul"), T(sm1satg200), N("f32"), SDST, T(sm2neg), T(isw), T(sm3neg), IMM },
	// desc ^^^

	{ 0xd0000000, 0xf0008100, N("and"), N("b32"), SDST, T(sm3not), T(isw), IMM },
	{ 0xd0000100, 0xf0008100, N("or"), N("b32"), SDST, T(sm3not), T(isw), IMM },
	{ 0xd0008000, 0xf0008100, N("xor"), N("b32"), SDST, T(sm3not), T(isw), IMM },
	{ 0xd0008100, 0xf0008100, N("mov2"), N("b32"), SDST, T(sm3not), T(isw), IMM },

	// desc VVV
	{ 0xe0000000, 0xf0000000, N("add"), T(sm1sat), N("f32"), SDST, T(sm2neg), SESTART, N("mul"), T(isw), IMM, SEEND, T(sm3neg), SDST },
	// desc ^^^

	{ 0, 0, OOPS, SDST, T(isw), IMM },
};

/*
 * Long instructions
 */

// A few helper tables for usual operand types.

static struct insn tablsh[] = {
	{ 0x0000000000000000ull, 0x0020000001800000ull, LHSRC },
	{ 0x0000000000800000ull, 0x0020000001800000ull, LHSRC },
	{ 0x0000000001000000ull, 0x0020000001800000ull, LHSRC },
	{ 0x0000000001800000ull, 0x0020000001800000ull, LPRIM, .ptype = GP },	// XXX check
	{ 0x0020000000000000ull, 0x0020000001800000ull, T(realls) },
	{ 0x0020000000800000ull, 0x0020000001800000ull, T(realls) },
	{ 0x0020000001000000ull, 0x0020000001800000ull, T(realls) },
	{ 0x0020000001800000ull, 0x0020000001800000ull, LPRIM, .ptype = GP },	// XXX check
	{ 0, 0, OOPS },
};

static struct insn tablsw[] = {
	{ 0x0000000000000000ull, 0x0020000001800000ull, LSRC },
	{ 0x0000000000800000ull, 0x0020000001800000ull, LSRC },
	{ 0x0000000001000000ull, 0x0020000001800000ull, LSRC },
	{ 0x0000000001800000ull, 0x0020000001800000ull, LPRIM, .ptype = GP },	// XXX check
	{ 0x0020000000000000ull, 0x0020000001800000ull, T(realls) },
	{ 0x0020000000800000ull, 0x0020000001800000ull, T(realls) },
	{ 0x0020000001000000ull, 0x0020000001800000ull, T(realls) },
	{ 0x0020000001800000ull, 0x0020000001800000ull, LPRIM, .ptype = GP },	// XXX check
	{ 0, 0, OOPS },
};

static struct insn tablc2h[] = {
	{ 0x0000000000000000ull, 0x0000000001800000ull, LHSRC2 },
	{ 0x0000000000800000ull, 0x0020000003800000ull, L2CONST16, .ptype = CP },
	{ 0x0000000002800000ull, 0x0020000003800000ull, L2CONST16PI, .ptype = CP },
	{ 0x0020000000800000ull, 0x0020000001800000ull, L2CONST16NA, .ptype = CP },
	{ 0x0000000000800000ull, 0x0000000003800000ull, L2CONST16, .ptype = VP|GP|FP },
	{ 0x0000000002800000ull, 0x0000000003800000ull, L2CONST16PI, .ptype = VP|GP|FP },
	{ 0x0000000001000000ull, 0x0000000001800000ull, LHSRC2 },
	{ 0x0000000001800000ull, 0x0020000001800000ull, L2CONST16NA, .ptype = GP },
	{ 0x0020000001800000ull, 0x0020000001800000ull, LHSRC2 },
	{ 0, 0, OOPS }
};

static struct insn tablc2w[] = {
	{ 0x0000000000000000ull, 0x0000000001800000ull, LSRC2 },
	{ 0x0000000000800000ull, 0x0020000003800000ull, L2CONST32, .ptype = CP },
	{ 0x0000000002800000ull, 0x0020000003800000ull, L2CONST32, .ptype = CP },
	{ 0x0020000000800000ull, 0x0020000001800000ull, L2CONST32NA, .ptype = CP },
	{ 0x0000000000800000ull, 0x0000000003800000ull, L2CONST32, .ptype = VP|GP|FP },
	{ 0x0000000002800000ull, 0x0000000003800000ull, L2CONST32PI, .ptype = VP|GP|FP },
	{ 0x0000000001000000ull, 0x0000000001800000ull, LSRC2 },
	{ 0x0000000001800000ull, 0x0020000001800000ull, L2CONST32NA, .ptype = GP },
	{ 0x0020000001800000ull, 0x0020000001800000ull, LSRC2 },
	{ 0, 0, OOPS },
};

static struct insn tablc3h[] = {
	{ 0x0000000000000000ull, 0x0000000001800000ull, LHSRC3 },
	{ 0x0000000000800000ull, 0x0000000001800000ull, LHSRC3 },
	{ 0x0000000001000000ull, 0x0020000003800000ull, L3CONST16, .ptype = CP },
	{ 0x0000000003000000ull, 0x0020000003800000ull, L3CONST16PI, .ptype = CP },
	{ 0x0020000001000000ull, 0x0020000001800000ull, L3CONST16NA, .ptype = CP },
	{ 0x0000000001000000ull, 0x0000000003800000ull, L3CONST16, .ptype = VP|GP|FP },
	{ 0x0000000003000000ull, 0x0000000003800000ull, L3CONST16PI, .ptype = VP|GP|FP },
	{ 0x0000000001800000ull, 0x0020000001800000ull, L3CONST16NA, .ptype = GP },
	{ 0x0020000001800000ull, 0x0020000001800000ull, LHSRC3 },
	{ 0, 0, OOPS },
};

static struct insn tablc3w[] = {
	{ 0x0000000000000000ull, 0x0000000001800000ull, LSRC3 },
	{ 0x0000000000800000ull, 0x0000000001800000ull, LSRC3 },
	{ 0x0000000001000000ull, 0x0020000003800000ull, L3CONST32, .ptype = CP },
	{ 0x0000000003000000ull, 0x0020000003800000ull, L3CONST32PI, .ptype = CP },
	{ 0x0020000001000000ull, 0x0020000001800000ull, L3CONST32NA, .ptype = CP },
	{ 0x0000000001000000ull, 0x0000000003800000ull, L3CONST32, .ptype = VP|GP|FP },
	{ 0x0000000003000000ull, 0x0000000003800000ull, L3CONST32PI, .ptype = VP|GP|FP },
	{ 0x0000000001800000ull, 0x0020000001800000ull, L3CONST16NA, .ptype = GP },
	{ 0x0020000001800000ull, 0x0020000001800000ull, LSRC3 },
	{ 0, 0, OOPS },
};

static struct insn tablldst[] = {
	{ 0x0000000000000000ull, 0x0000000800000000ull, LDST },
	{ 0x00000008000001fcull, 0x00000008000001fcull, DISCARD },
	{ 0x0000000800000000ull, 0x0000000800000000ull, OUT32 },
	{ 0, 0, OOPS },
};
#define LLDST T(lldst)

static struct insn tabllhdst[] = {
	{ 0x0000000000000000ull, 0x0000000800000000ull, LHDST },
	{ 0x00000008000001fcull, 0x00000008000001fcull, DISCARD },
	{ 0x0000000800000000ull, 0x0000000800000000ull, OUT16 },
	{ 0, 0, OOPS },
};
#define LLHDST T(llhdst)

F(shcnt, 0x34, T(lc2w), SHCNT)
F(hshcnt, 0x34, T(lc2h), SHCNT)

// various stuff available for filling the misc bits.

// stolen from SRC3
F(s30us16, 0x2e, N("u16"), N("s16"))
F1(s30high, 0x2e, N("high"))
F(s31us16, 0x2f, N("u16"), N("s16"))
F(s31us24, 0x2f, N("u24"), N("s24"))
F1(s32not, 0x30, N("not"))
F1(s33not, 0x31, N("not"))
F1(s35sat, 0x33, N("sat"))
F1(s35abs, 0x33, N("abs"))
F1V(s36satg200, F_SM12, 0x34, N("sat"))
F1(s36abs, 0x34, N("abs"))
// the actual misc field
F1(m1neg, 0x3a, N("neg"))
F1(m2neg, 0x3b, N("neg"))
F1(m2sat, 0x3b, N("sat"))
F(lm2us16, 0x3b, N("u16"), N("s16"))
F(lm2us32, 0x3b, N("u32"), N("s32"))
// stolen from opcode field.
F1(o0neg, 0x3d, N("neg"))
F1(o0sat, 0x3d, N("sat"))
F1(lsfu, 0x39, N("sfu"))

static struct insn tablfm1[] = {
	{ 0, 0, T(m1neg), T(s36abs) }
};

static struct insn tablfm2[] = {
	{ 0, 0, T(m2neg), T(s35abs) }
};

static struct insn tabcvtmod[] = {
	{ 0, 0, T(o0neg), T(s36abs) },
};

// for g[] and l[] insns.
static struct insn tabldstm[] = { // we seriously need to unify these, if possible. I wonder if atomics will work with smaller sizes.
	{ 0x0000000000000000ull, 0x00e0000000000000ull, N("u8") },
	{ 0x0020000000000000ull, 0x00e0000000000000ull, N("s8") },
	{ 0x0040000000000000ull, 0x00e0000000000000ull, N("u16") },
	{ 0x0060000000000000ull, 0x00e0000000000000ull, N("s16") },
	{ 0x0080000000000000ull, 0x00e0000000000000ull, N("b64") },
	{ 0x00a0000000000000ull, 0x00e0000000000000ull, N("b128") },
	{ 0x00c0000000000000ull, 0x00e0000000000000ull, N("b32") },
	{ 0, 0, OOPS }
};

static struct insn tabraadd[] = {
	{ 0x0080000000000000ull, 0x00e0000000000000ull, N("u64"), .fmask = F_SM12 },
	{ 0x00c0000000000000ull, 0x00e0000000000000ull, N("u32") },
	{ 0x00e0000000000000ull, 0x00e0000000000000ull, N("s32") },
	{ 0, 0, OOPS }
};

static struct insn tabrab[] = {
	{ 0x0080000000000000ull, 0x00e0000000000000ull, N("b64"), .fmask = F_SM12 },
	{ 0x00c0000000000000ull, 0x00e0000000000000ull, N("b32") },
	{ 0, 0, OOPS }
};

static struct insn tabramm[] = {
	{ 0x00c0000000000000ull, 0x00e0000000000000ull, N("u32") },
	{ 0x00e0000000000000ull, 0x00e0000000000000ull, N("s32") },
	{ 0, 0, OOPS }
};

static struct insn tabrab32[] = {
	{ 0x00c0000000000000ull, 0x00e0000000000000ull, N("b32") },
	{ 0, 0, OOPS }
};

static struct insn tabrau32[] = {
	{ 0x00c0000000000000ull, 0x00e0000000000000ull, N("u32") },
	{ 0, 0, OOPS }
};

static struct insn tabldsto[] = {	// hack alert: reads the bitfield second time.
	{ 0x0000000000000000ull, 0x00e0000000000000ull, LDST },
	{ 0x0020000000000000ull, 0x00e0000000000000ull, LDST },
	{ 0x0040000000000000ull, 0x00e0000000000000ull, LDST },
	{ 0x0060000000000000ull, 0x00e0000000000000ull, LDST },
	{ 0x0080000000000000ull, 0x00e0000000000000ull, LDDST },
	{ 0x00a0000000000000ull, 0x00e0000000000000ull, LQDST },
	{ 0x00c0000000000000ull, 0x00e0000000000000ull, LDST },
	{ 0x00e0000000000000ull, 0x00e0000000000000ull, LDST },
	{ 0, 0, OOPS },
};

static struct insn tabldsts2[] = {
	{ 0x0080000000000000ull, 0x00e0000000000000ull, LDSRC2 },
	{ 0x00c0000000000000ull, 0x00e0000000000000ull, LSRC2 },
	{ 0x00e0000000000000ull, 0x00e0000000000000ull, LSRC2 },
	{ 0, 0, OOPS },
};

static struct insn tabldsts3[] = {
	{ 0x0080000000000000ull, 0x00e0000000000000ull, LDSRC3 },
	{ 0x00c0000000000000ull, 0x00e0000000000000ull, LSRC3 },
	{ 0x00e0000000000000ull, 0x00e0000000000000ull, LSRC3 },
	{ 0, 0, OOPS },
};

static struct insn tabmldsts3[] = {
	{ 0x0000000000000000ull, 0x0000003800000000ull },
	{ 0x0000000800000000ull, 0x0000003c00000000ull, T(ldsts3) },
	{ 0x0000000c00000000ull, 0x0000003c00000000ull },
	{ 0x0000001000000000ull, 0x0000003000000000ull },
	{ 0x0000002000000000ull, 0x0000002000000000ull },
	{ 0, 0, OOPS },
};

static struct insn tabredm[] = {
	{ 0x0000000000000000ull, 0x0000003c00000000ull, N("add"), T(raadd), },
	{ 0x0000001000000000ull, 0x0000003c00000000ull, N("inc"), T(rau32), },
	{ 0x0000001400000000ull, 0x0000003c00000000ull, N("dec"), T(rau32), },
	{ 0x0000001800000000ull, 0x0000003c00000000ull, N("max"), T(ramm), },
	{ 0x0000001c00000000ull, 0x0000003c00000000ull, N("min"), T(ramm), },
	{ 0x0000002800000000ull, 0x0000003c00000000ull, N("and"), T(rab32), },
	{ 0x0000002c00000000ull, 0x0000003c00000000ull, N("or"), T(rab32), },
	{ 0x0000003000000000ull, 0x0000003c00000000ull, N("xor"), T(rab32), },
	{ 0, 0, OOPS },
};

static struct insn tabatomm[] = {
	{ 0x0000000400000000ull, 0x0000003c00000000ull, N("exch"), T(rab), },
	{ 0x0000000800000000ull, 0x0000003c00000000ull, N("cas"), T(rab), },
	{ 0, 0, N("ld"), T(redm) },
};

// rounding modes
static struct insn tabcvtrnd[] = {
	{ 0x0000000000000000ull, 0x0006000000000000ull, N("rn") },
	{ 0x0002000000000000ull, 0x0006000000000000ull, N("rm") },
	{ 0x0004000000000000ull, 0x0006000000000000ull, N("rp") },
	{ 0x0006000000000000ull, 0x0006000000000000ull, N("rz") },
	{ 0, 0, OOPS }
};

static struct insn tabcvtrint[] = {
	{ 0x0000000000000000ull, 0x0006000000000000ull, N("rni") },
	{ 0x0002000000000000ull, 0x0006000000000000ull, N("rmi") },
	{ 0x0004000000000000ull, 0x0006000000000000ull, N("rpi") },
	{ 0x0006000000000000ull, 0x0006000000000000ull, N("rzi") },
	{ 0, 0, OOPS }
};

static struct insn tabaf64r[] = {
	{ 0x00000000, 0x00030000, N("rn") },
	{ 0x00010000, 0x00030000, N("rm") },
	{ 0x00020000, 0x00030000, N("rp") },
	{ 0x00030000, 0x00030000, N("rz") },
	{ 0, 0, OOPS },
};

static struct insn tabaf32r[] = {
	{ 0x00000000, 0x00030000, N("rn") },
	{ 0x00030000, 0x00030000, N("rz") },
	{ 0, 0, OOPS }
};

static struct insn tabmf32r[] = {
	{ 0x0000000000000000ull, 0x0000c00000000000ull, N("rn") },
	{ 0x0000c00000000000ull, 0x0000c00000000000ull, N("rz") },
	{ 0, 0, OOPS }
};

static struct insn tabmad64r[] = {
	{ 0x0000000000000000ull, 0x00c0000000000000ull, N("rn") },
	{ 0x0040000000000000ull, 0x00c0000000000000ull, N("rm") },
	{ 0x0080000000000000ull, 0x00c0000000000000ull, N("rp") },
	{ 0x00c0000000000000ull, 0x00c0000000000000ull, N("rz") },
	{ 0, 0, OOPS }
};

// ops for set
static struct insn tabseti[] = {
	{ 0x0000000000000000ull, 0x0001c00000000000ull, N("never") },
	{ 0x0000400000000000ull, 0x0001c00000000000ull, N("l") },
	{ 0x0000800000000000ull, 0x0001c00000000000ull, N("e") },
	{ 0x0000c00000000000ull, 0x0001c00000000000ull, N("le") },
	{ 0x0001000000000000ull, 0x0001c00000000000ull, N("g") },
	{ 0x0001400000000000ull, 0x0001c00000000000ull, N("lg") },
	{ 0x0001800000000000ull, 0x0001c00000000000ull, N("ge") },
	{ 0x0001c00000000000ull, 0x0001c00000000000ull, N("always") },
	{ 0, 0, OOPS }
};

static struct insn tabsetf[] = {
	{ 0x0000000000000000ull, 0x0003c00000000000ull, N("never") },
	{ 0x0000400000000000ull, 0x0003c00000000000ull, N("l") },
	{ 0x0000800000000000ull, 0x0003c00000000000ull, N("e") },
	{ 0x0000c00000000000ull, 0x0003c00000000000ull, N("le") },
	{ 0x0001000000000000ull, 0x0003c00000000000ull, N("g") },
	{ 0x0001400000000000ull, 0x0003c00000000000ull, N("lg") },
	{ 0x0001800000000000ull, 0x0003c00000000000ull, N("ge") },
	{ 0x0001c00000000000ull, 0x0003c00000000000ull, N("lge") },
	{ 0x0002000000000000ull, 0x0003c00000000000ull, N("u") },
	{ 0x0002400000000000ull, 0x0003c00000000000ull, N("lu") },
	{ 0x0002800000000000ull, 0x0003c00000000000ull, N("eu") },
	{ 0x0002c00000000000ull, 0x0003c00000000000ull, N("leu") },
	{ 0x0003000000000000ull, 0x0003c00000000000ull, N("gu") },
	{ 0x0003400000000000ull, 0x0003c00000000000ull, N("lgu") },
	{ 0x0003800000000000ull, 0x0003c00000000000ull, N("geu") },
	{ 0x0003c00000000000ull, 0x0003c00000000000ull, N("always") },
	{ 0, 0, OOPS }
};

// for cvt
static struct insn tabcvtiisrc[] ={
	{ 0x0000000000000000ull, 0x0001c00000000000ull, N("u16"), T(lsh) },
	{ 0x0000400000000000ull, 0x0001c00000000000ull, N("u32"), T(lsw) },
	{ 0x0000800000000000ull, 0x0001c00000000000ull, N("u8"), T(lsh) },
	{ 0x0000c00000000000ull, 0x0001c00000000000ull, N("u8"), LSRC },	// what about mem?
	{ 0x0001000000000000ull, 0x0001c00000000000ull, N("s16"), T(lsh) },
	{ 0x0001400000000000ull, 0x0001c00000000000ull, N("s32"), T(lsw) },
	{ 0x0001800000000000ull, 0x0001c00000000000ull, N("s8"), T(lsh) },
	{ 0x0001c00000000000ull, 0x0001c00000000000ull, N("s8"), LSRC },	// what about mem?
	{ 0, 0, OOPS }
};


// for tex
F1(dtex, 0x23, N("deriv")) // suspected to enable implicit derivatives on non-FPs.
F(ltex, 0x22, N("all"), N("live"))

static struct insn tabtexf[] = {
	{ 0, 0, T(ltex), T(dtex) },
};

// for mov
static struct insn tablane[] = {
	{ 0x0000000000000000ull, 0x0003c00000000000ull, SESTART, N("lnone"), SEEND, },
	{ 0x0000400000000000ull, 0x0003c00000000000ull, SESTART, N("l0"), SEEND, },
	{ 0x0000800000000000ull, 0x0003c00000000000ull, SESTART, N("l1"), SEEND, },
	{ 0x0000c00000000000ull, 0x0003c00000000000ull, SESTART, N("l01"), SEEND, },
	{ 0x0001000000000000ull, 0x0003c00000000000ull, SESTART, N("l2"), SEEND, },
	{ 0x0001400000000000ull, 0x0003c00000000000ull, SESTART, N("l02"), SEEND, },
	{ 0x0001800000000000ull, 0x0003c00000000000ull, SESTART, N("l12"), SEEND, },
	{ 0x0001c00000000000ull, 0x0003c00000000000ull, SESTART, N("l012"), SEEND, },
	{ 0x0002000000000000ull, 0x0003c00000000000ull, SESTART, N("l3"), SEEND, },
	{ 0x0002400000000000ull, 0x0003c00000000000ull, SESTART, N("l03"), SEEND, },
	{ 0x0002800000000000ull, 0x0003c00000000000ull, SESTART, N("l13"), SEEND, },
	{ 0x0002c00000000000ull, 0x0003c00000000000ull, SESTART, N("l013"), SEEND, },
	{ 0x0003000000000000ull, 0x0003c00000000000ull, SESTART, N("l23"), SEEND, },
	{ 0x0003400000000000ull, 0x0003c00000000000ull, SESTART, N("l023"), SEEND, },
	{ 0x0003800000000000ull, 0x0003c00000000000ull, SESTART, N("l123"), SEEND, },
	{ 0x0003c00000000000ull, 0x0003c00000000000ull },
	{ 0, 0, OOPS },
};

// for mov from c[]
static struct insn tabfcon[] = {
	{ 0x0000000000000000ull, 0x0000c00002000000ull, N("u8"), FCONST8 },
	{ 0x0000400000000000ull, 0x0000c00002000000ull, N("u16"), FCONST16 },
	{ 0x0000800000000000ull, 0x0000c00002000000ull, N("s16"), FCONST16 },
	{ 0x0000c00000000000ull, 0x0000c00002000000ull, N("b32"), FCONST32 },
	{ 0x0000000002000000ull, 0x0000c00002000000ull, N("u8"), FCONST8PI },
	{ 0x0000400002000000ull, 0x0000c00002000000ull, N("u16"), FCONST16PI },
	{ 0x0000800002000000ull, 0x0000c00002000000ull, N("s16"), FCONST16PI },
	{ 0x0000c00002000000ull, 0x0000c00002000000ull, N("b32"), FCONST32PI },
	{ 0, 0, OOPS }
};

// for quadop
static struct insn tabqs1[] = {
	{ 0x00000000, 0x00070000, N("l0") },
	{ 0x00010000, 0x00070000, N("l1") },
	{ 0x00020000, 0x00070000, N("l2") },
	{ 0x00030000, 0x00070000, N("l3") },
	{ 0x00040000, 0x00050000, N("dx") },
	{ 0x00050000, 0x00050000, N("dy") },
	{ 0, 0, OOPS },
};

static struct insn tabqop0[] = {
	{ 0x0000000000000000ull, 0x0c00000000000000ull, N("add") },
	{ 0x0400000000000000ull, 0x0c00000000000000ull, N("subr") },
	{ 0x0800000000000000ull, 0x0c00000000000000ull, N("sub") },
	{ 0x0c00000000000000ull, 0x0c00000000000000ull, N("mov2") },
	{ 0, 0, OOPS },
};

static struct insn tabqop1[] = {
	{ 0x0000000000000000ull, 0x0300000000000000ull, N("add") },
	{ 0x0100000000000000ull, 0x0300000000000000ull, N("subr") },
	{ 0x0200000000000000ull, 0x0300000000000000ull, N("sub") },
	{ 0x0300000000000000ull, 0x0300000000000000ull, N("mov2") },
	{ 0, 0, OOPS },
};

static struct insn tabqop2[] = {
	{ 0x0000000000000000ull, 0x00c0000000000000ull, N("add") },
	{ 0x0040000000000000ull, 0x00c0000000000000ull, N("subr") },
	{ 0x0080000000000000ull, 0x00c0000000000000ull, N("sub") },
	{ 0x00c0000000000000ull, 0x00c0000000000000ull, N("mov2") },
	{ 0, 0, OOPS },
};

static struct insn tabqop3[] = {
	{ 0x00000000, 0x00300000, N("add") },
	{ 0x00100000, 0x00300000, N("subr") },
	{ 0x00200000, 0x00300000, N("sub") },
	{ 0x00300000, 0x00300000, N("mov2") },
	{ 0, 0, OOPS },
};

static struct insn tablogop[] = {
	{ 0x0000000000000000ull, 0x0000c00000000000ull, N("and") },
	{ 0x0000400000000000ull, 0x0000c00000000000ull, N("or") },
	{ 0x0000800000000000ull, 0x0000c00000000000ull, N("xor") },
	{ 0x0000c00000000000ull, 0x0000c00000000000ull, N("mov2") },
	{ 0, 0, OOPS },
};

static struct insn tabaddcond[] = {
	{ 0x00000000, 0x10400000 },
	{ 0x00400000, 0x10400000 },
	{ 0x10000000, 0x10400000 },
	{ 0x10400000, 0x10400000, COND },
	{ 0, 0, OOPS }
};

static struct insn tabaddop2[] = {
	{ 0x0000000000000000ull, 0x0c00000000000000ull, N("add") },
	{ 0x0400000000000000ull, 0x0c00000000000000ull, N("sub") },
	{ 0x0800000000000000ull, 0x0c00000000000000ull, N("subr") },
	{ 0x0c00000000000000ull, 0x0c00000000000000ull, N("addc") },
	{ 0, 0, OOPS },
};

static struct insn tabaddcond2[] = {
	{ 0x0000000000000000ull, 0x0c00000000000000ull },
	{ 0x0400000000000000ull, 0x0c00000000000000ull },
	{ 0x0800000000000000ull, 0x0c00000000000000ull },
	{ 0x0c00000000000000ull, 0x0c00000000000000ull, COND },
	{ 0, 0, OOPS }
};

F(sstreg, 0x35, LHSRC3, LSRC3)
F1V(unlock, F_SM12, 0x37, N("unlock"))
F(csldreg, 0x3a, LLHDST, LLDST)

static struct insn tabl[] = {
	// 0
	// desc VVV
	{ 0x0420000000000000ull, 0xe4200000f0000000ull, T(lane), N("ld"), N("b32"), LLDST, FATTR, .ptype = VP|GP },
	// desc ^^^
	{ 0x2000000000000000ull, 0xe0000000f0000000ull, N("mov"), LDST, COND },
	{ 0x4000000000000000ull, 0xe0000000f0000000ull, N("mov"), LDST, LAREG },
	{ 0x6000000000000000ull, 0xe0000000f0000000ull, N("mov"), LDST, SREG },
	{ 0x80c0000000000000ull, 0xe0c00000f0000000ull, N("st"), N("b32"), FOUT32, LSRC3 },
	{ 0xa000000000000000ull, 0xe0000000f0000000ull, N("mov"), CDST, LSRC, IGNCE },
	{ 0xc000000000000000ull, 0xe0000000f0000000ull, N("shl"), ADST, T(lsw), HSHCNT },
	
	// desc VVV
	{ 0xe040000000000000ull, 0xe0400000f0000000ull, N("st"), N("b8"), T(fs8), T(sstreg), .ptype = CP },
	{ 0xe000000000000000ull, 0xe4400000f0000000ull, N("st"), N("b16"), T(fs16), T(sstreg), .ptype = CP },
	{ 0xe400000000000000ull, 0xe4400000f0000000ull, N("st"), T(unlock), N("b32"), T(fs32), T(sstreg), .ptype = CP },
	// desc ^^^

	// 1
	{ 0x0000000010000000ull, 0xe4000000f0000000ull, T(lane), N("mov"), T(lsfu), N("b16"), LLHDST, T(lsh) },
	{ 0x0400000010000000ull, 0xe4000000f0000000ull, T(lane), N("mov"), T(lsfu), N("b32"), LLDST, T(lsw) },

	// desc VVV
	{ 0x2000000010000000ull, 0xe0000000f0000000ull, N("ld"), T(csldreg), T(fcon) },

	{ 0x4000000010000000ull, 0xe000c000f0000000ull, N("ld"), T(csldreg), N("u8"), T(fs8), .fmask = F_SM11, .ptype = CP },
	{ 0x4000400010000000ull, 0xe000c000f0000000ull, N("ld"), T(csldreg), N("u16"), T(fs16), .fmask = F_SM11, .ptype = CP },
	{ 0x4000800010000000ull, 0xe000c000f0000000ull, N("ld"), T(csldreg), N("s16"), T(fs16), .fmask = F_SM11, .ptype = CP },
	{ 0x4000c00010000000ull, 0xe080c000f0000000ull, N("ld"), T(csldreg), N("b32"), T(fs32), .fmask = F_SM11, .ptype = CP },
	{ 0x4080c00010000000ull, 0xe080c000f0000000ull, N("ld"), N("lock"), CDST, T(csldreg), N("b32"), T(fs32), IGNCE, .fmask = F_SM12, .ptype = CP },

	{ 0x6000000010000200ull, 0xe0000000f0000600ull, N("vote"), N("any"), CDST, IGNCE, .fmask = F_SM12 },
	{ 0x6000000010000400ull, 0xe0000000f0000600ull, N("vote"), N("all"), CDST, IGNCE, .fmask = F_SM12 },
	// desc ^^^

	// 2 and 3
	{ 0x0000000020000000ull, 0xe4000000e0000000ull, T(addop), N("b16"), T(m2sat), MCDST, LLHDST, T(lsh), T(lc3h), T(addcond) },
	{ 0x0400000020000000ull, 0xe4000000e0000000ull, T(addop), N("b32"), T(m2sat), MCDST, LLDST, T(lsw), T(lc3w), T(addcond) },

	{ 0x6000000030000000ull, 0xe4000000f0000000ull, N("set"), MCDST, LLHDST, T(seti), T(lm2us16), T(lsh), T(lc2h) },
	{ 0x6400000030000000ull, 0xe4000000f0000000ull, N("set"), MCDST, LLDST, T(seti), T(lm2us32), T(lsw), T(lc2w) },

	{ 0x8000000030000000ull, 0xe4000000f0000000ull, N("max"), T(lm2us16), MCDST, LLHDST, T(lsh), T(lc2h) },
	{ 0x8400000030000000ull, 0xe4000000f0000000ull, N("max"), T(lm2us32), MCDST, LLDST, T(lsw), T(lc2w) },

	{ 0xa000000030000000ull, 0xe4000000f0000000ull, N("min"), T(lm2us16), MCDST, LLHDST, T(lsh), T(lc2h) },
	{ 0xa400000030000000ull, 0xe4000000f0000000ull, N("min"), T(lm2us32), MCDST, LLDST, T(lsw), T(lc2w) },

	{ 0xc000000030000000ull, 0xe4000000f0000000ull, N("shl"), N("b16"), MCDST, LLHDST, T(lsh), T(hshcnt) },
	{ 0xc400000030000000ull, 0xe4000000f0000000ull, N("shl"), N("b32"), MCDST, LLDST, T(lsw), T(shcnt) },

	{ 0xe000000030000000ull, 0xe4000000f0000000ull, N("shr"), T(lm2us16), MCDST, LLHDST, T(lsh), T(hshcnt) },
	{ 0xe400000030000000ull, 0xe4000000f0000000ull, N("shr"), T(lm2us32), MCDST, LLDST, T(lsw), T(shcnt) },

	// 4
	{ 0x0000000040000000ull, 0xe0010000f0000000ull, N("mul"), MCDST, LLDST, T(s31us16), T(lsh), T(s30us16), T(lc2h) },
	{ 0x0001000040000000ull, 0xe0010000f0000000ull, N("mul"), MCDST, LLDST, T(s30high), T(s31us24), T(lsw), T(lc2w) },

	// 5
	{ 0x0000000050000000ull, 0xe4000000f0000000ull, N("sad"), MCDST, LLDST, T(lm2us16), T(lsh), T(lc2h), T(lc3w) },
	{ 0x0400000050000000ull, 0xe4000000f0000000ull, N("sad"), MCDST, LLDST, T(lm2us32), T(lsw), T(lc2w), T(lc3w) },

	// 6 and 7
	{ 0x0000000060000000ull, 0xe0000000f0000000ull, T(addop2), MCDST, LLDST, SESTART, N("mul"), N("u16"), T(lsh), T(lc2h), SEEND, T(lc3w), T(addcond2) },
	{ 0x2000000060000000ull, 0xe0000000f0000000ull, T(addop2), MCDST, LLDST, SESTART, N("mul"), N("s16"), T(lsh), T(lc2h), SEEND, T(lc3w), T(addcond2) },
	{ 0x4000000060000000ull, 0xe0000000f0000000ull, T(addop2), N("sat"), MCDST, LLDST, SESTART, N("mul"), N("s16"), T(lsh), T(lc2h), SEEND, T(lc3w), T(addcond2) },
	{ 0x6000000060000000ull, 0xe0000000f0000000ull, T(addop2), MCDST, LLDST, SESTART, N("mul"), N("u24"), T(lsw), T(lc2w), SEEND, T(lc3w), T(addcond2) },
	{ 0x8000000060000000ull, 0xe0000000f0000000ull, T(addop2), MCDST, LLDST, SESTART, N("mul"), N("s24"), T(lsw), T(lc2w), SEEND, T(lc3w), T(addcond2) },
	{ 0xa000000060000000ull, 0xe0000000f0000000ull, T(addop2), N("sat"), MCDST, LLDST, SESTART, N("mul"), N("s24"), T(lsw), T(lc2w), SEEND, T(lc3w), T(addcond2) },
	{ 0xc000000060000000ull, 0xe0000000f0000000ull, T(addop2), MCDST, LLDST, SESTART, N("mul"), N("high"), N("u24"), T(lsw), T(lc2w), SEEND, T(lc3w), T(addcond2) },
	{ 0xe000000060000000ull, 0xe0000000f0000000ull, T(addop2), MCDST, LLDST, SESTART, N("mul"), N("high"), N("s24"), T(lsw), T(lc2w), SEEND, T(lc3w), T(addcond2) },
	{ 0x0000000070000000ull, 0xe0000000f0000000ull, T(addop2), N("sat"), MCDST, LLDST, SESTART, N("mul"), N("high"), N("s24"), T(lsw), T(lc2w), SEEND, T(lc3w), T(addcond2) },

	// desc VVV
	// 8
	{ 0x0000000080000000ull, 0x00070000f0000000ull, N("interp"), LDST, LVAR, .ptype = FP },
	{ 0x0001000080000000ull, 0x00070000f0000000ull, N("interp"), LDST, N("cent"), LVAR, .ptype = FP },
	{ 0x0002000080000000ull, 0x00070000f0000000ull, N("interp"), LDST, LVAR, T(m1neg), LSRC, .ptype = FP },
	{ 0x0003000080000000ull, 0x00070000f0000000ull, N("interp"), LDST, N("cent"), LVAR, T(m1neg), LSRC, .ptype = FP },
	{ 0x0004000080000000ull, 0x00070000f0000000ull, N("interp"), LDST, N("flat"), LVAR, .ptype = FP },

	// 9
	{ 0x0000000090000000ull, 0xe0000000f0000000ull, N("rcp"), N("f32"), LLDST, T(lfm1), LSRC },
	{ 0x4000000090000000ull, 0xe0000000f0000000ull, N("rsqrt"), N("f32"), LLDST, T(lfm1), LSRC },
	{ 0x6000000090000000ull, 0xe0000000f0000000ull, N("lg2"), N("f32"), LLDST, T(lfm1), LSRC },
	{ 0x8000000090000000ull, 0xe0000000f0000000ull, N("sin"), N("f32"), LLDST, LSRC },
	{ 0xa000000090000000ull, 0xe0000000f0000000ull, N("cos"), N("f32"), LLDST, LSRC },
	{ 0xc000000090000000ull, 0xe0000000f0000000ull, N("ex2"), N("f32"), T(m2sat), LLDST, LSRC },

	// a
	{ 0xc0000000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), N("f16"), MCDST, LLHDST, N("f16"), T(lsh) },
	{ 0xc8000000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), T(cvtrint), N("f16"), MCDST, LLHDST, N("f16"), T(lsh) },
	{ 0xc0004000a0000000ull, 0xc4404000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), T(cvtrnd), N("f16"), MCDST, LLHDST, N("f32"), T(lsw) },
	{ 0xc4004000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), N("f32"), MCDST, LLDST, N("f32"), T(lsw) },
	{ 0xcc004000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), T(cvtrint), N("f32"), MCDST, LLDST, N("f32"), T(lsw) },
	{ 0xc4000000a0000000ull, 0xc4404000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), N("f32"), MCDST, LLDST, N("f16"), T(lsh) },

	{ 0xc0404000a0000000ull, 0xc4404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrnd), N("f32"), MCDST, LLDST, N("f64"), LDSRC, .fmask = F_FP64 },
	{ 0xc4404000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), N("f64"), MCDST, LDDST, N("f64"), LDSRC, .fmask = F_FP64 },
	{ 0xcc404000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("f64"), MCDST, LDDST, N("f64"), LDSRC, .fmask = F_FP64 },
	{ 0xc4400000a0000000ull, 0xc4404000f0000000ull, N("cvt"), T(cvtmod), N("f64"), MCDST, LDDST, N("f32"), T(lsw), .fmask = F_FP64 },

	{ 0x80000000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("u16"), MCDST, LLHDST, N("f16"), T(lsh) },
	{ 0x80004000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("u16"), MCDST, LLHDST, N("f32"), T(lsw) },
	{ 0x88000000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("s16"), MCDST, LLHDST, N("f16"), T(lsh) },
	{ 0x88004000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("s16"), MCDST, LLHDST, N("f32"), T(lsw) },
	{ 0x84000000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("u32"), MCDST, LLDST, N("f16"), T(lsh) },
	{ 0x84004000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("u32"), MCDST, LLDST, N("f32"), T(lsw) },
	{ 0x8c000000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("s32"), MCDST, LLDST, N("f16"), T(lsh) },
	{ 0x8c004000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("s32"), MCDST, LLDST, N("f32"), T(lsw) },

	{ 0x80404000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("u32"), MCDST, LLDST, N("f64"), LDSRC, .fmask = F_FP64 },
	{ 0x88404000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("s32"), MCDST, LLDST, N("f64"), LDSRC, .fmask = F_FP64 },
	{ 0x84400000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("u64"), MCDST, LDDST, N("f32"), T(lsw), .fmask = F_FP64 },
	{ 0x84404000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("u64"), MCDST, LDDST, N("f64"), LDSRC, .fmask = F_FP64 },
	{ 0x8c400000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("s64"), MCDST, LDDST, N("f32"), T(lsw), .fmask = F_FP64 },
	{ 0x8c404000a0000000ull, 0xcc404000f0000000ull, N("cvt"), T(cvtmod), T(cvtrint), N("s64"), MCDST, LDDST, N("f64"), LDSRC, .fmask = F_FP64 },

	{ 0x40000000a0000000ull, 0xc4400000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), T(cvtrnd), N("f16"), MCDST, LLHDST, T(cvtiisrc) },
	{ 0x44000000a0000000ull, 0xc4400000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), T(cvtrnd), N("f32"), MCDST, LLDST, T(cvtiisrc) },

	{ 0x44400000a0000000ull, 0xc4414000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), T(cvtrnd), N("f64"), MCDST, LDDST, N("u32"), T(lsw), .fmask = F_FP64 },
	{ 0x44410000a0000000ull, 0xc4414000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), T(cvtrnd), N("f64"), MCDST, LDDST, N("s32"), T(lsw), .fmask = F_FP64 },
	{ 0x40404000a0000000ull, 0xc4414000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), T(cvtrnd), N("f32"), MCDST, LLDST, N("u64"), LDSRC, .fmask = F_FP64 },
	{ 0x40414000a0000000ull, 0xc4414000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), T(cvtrnd), N("f32"), MCDST, LLDST, N("s64"), LDSRC, .fmask = F_FP64 },
	{ 0x44404000a0000000ull, 0xc4414000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), T(cvtrnd), N("f64"), MCDST, LDDST, N("u64"), LDSRC, .fmask = F_FP64 },
	{ 0x44414000a0000000ull, 0xc4414000f0000000ull, N("cvt"), T(cvtmod), T(s35sat), T(cvtrnd), N("f64"), MCDST, LDDST, N("s64"), LDSRC, .fmask = F_FP64 },

	{ 0x00000000a0000000ull, 0xcc080000f0000000ull, N("cvt"), T(cvtmod), N("u16"), MCDST, LLHDST, T(cvtiisrc) },
	{ 0x00080000a0000000ull, 0xcc080000f0000000ull, N("cvt"), T(cvtmod), N("u8"), MCDST, LLHDST, T(cvtiisrc) },
	{ 0x04000000a0000000ull, 0xcc080000f0000000ull, N("cvt"), T(cvtmod), N("u32"), MCDST, LLDST, T(cvtiisrc) },
	{ 0x04080000a0000000ull, 0xcc080000f0000000ull, N("cvt"), T(cvtmod), N("u8"), MCDST, LLDST, T(cvtiisrc) },
	{ 0x08000000a0000000ull, 0xcc080000f0000000ull, N("cvt"), T(cvtmod), N("s16"), MCDST, LLHDST, T(cvtiisrc) },
	{ 0x08080000a0000000ull, 0xcc080000f0000000ull, N("cvt"), T(cvtmod), N("s8"), MCDST, LLHDST, T(cvtiisrc) },
	{ 0x0c000000a0000000ull, 0xcc080000f0000000ull, N("cvt"), T(cvtmod), N("s32"), MCDST, LLDST, T(cvtiisrc) },
	{ 0x0c080000a0000000ull, 0xcc080000f0000000ull, N("cvt"), T(cvtmod), N("s8"), MCDST, LLDST, T(cvtiisrc) },


	// b
	{ 0x00000000b0000000ull, 0xc0000000f0000000ull, N("add"), T(o0sat), T(af32r), N("f32"), MCDST, LLDST, T(m1neg), T(lsw), T(m2neg), T(lc3w) },

	{ 0x60000000b0000000ull, 0xe0000000f0000000ull, N("set"), MCDST, LLDST, T(setf), N("f32"), T(lfm1), T(lsw), T(lfm2), T(lc2w) },

	{ 0x80000000b0000000ull, 0xe0000000f0000000ull, N("max"), N("f32"), MCDST, LLDST, T(lfm1), T(lsw), T(lfm2), T(lc2w) },

	{ 0xa0000000b0000000ull, 0xe0000000f0000000ull, N("min"), N("f32"), MCDST, LLDST, T(lfm1), T(lsw), T(lfm2), T(lc2w) },

	{ 0xc0000000b0000000ull, 0xe0004000f0000000ull, N("presin"), N("f32"), LLDST, T(lfm1), T(lsw) },
	{ 0xc0004000b0000000ull, 0xe0004000f0000000ull, N("preex2"), N("f32"), LLDST, T(lfm1), T(lsw) },
	/* preex2 converts float to fixed point, results:
	 * 0-0x3fffffff: 7.23 fixed-point number
	 * 0x40000000: +nan
	 * 0x40800000: +infinity
	 * flip bit 0x80000000 in any of the above for negative numbers.
	 * presin divides by pi/2, mods with 4 [or with 2*pi, pre-div], then does preex2
	 */

	// c
	{ 0x00000000c0000000ull, 0xe0000000f0000000ull, N("mul"), T(s36satg200), T(mf32r), N("f32"), MCDST, LLDST, T(m1neg), T(lsw), T(m2neg), T(lc2w) },

	{ 0x40000000c0000000ull, 0xc0000000f0000000ull, N("slct"), N("b32"), MCDST, LLDST, T(lsw), T(lc2w), N("f32"), T(o0neg), T(lc3w) },

	{ 0x80000000c0000000ull, 0xe0000000f0000000ull, N("quadop"), N("f32"), T(qop0), T(qop1), T(qop2), T(qop3), MCDST, LLDST, T(qs1), LSRC, LSRC3 },
	// desc ^^^

	// d
	{ 0x00000000d0000000ull, 0xe4000000f0000000ull, T(logop), N("b16"), MCDST, LLHDST, T(s32not), T(lsh), T(s33not), T(lc2h) },
	{ 0x04000000d0000000ull, 0xe4000000f0000000ull, T(logop), N("b32"), MCDST, LLDST, T(s32not), T(lsw), T(s33not), T(lc2w) },

	{ 0x20000000d0000000ull, 0xe0000000f0000000ull, N("add"), ADST, LAREG, OFFS },

	// desc VVV
	{ 0x40000000d0000000ull, 0xe0000000f0000000ull, N("ld"), T(ldstm), T(ldsto), T(local) },
	{ 0x60000000d0000000ull, 0xe0000000f0000000ull, N("st"), T(ldstm), T(local), T(ldsto) },

	{ 0x80000000d0000000ull, 0xe0000000f0000000ull, N("ld"), T(ldstm), T(ldsto), GLOBAL, .ptype = CP },
	{ 0xa0000000d0000000ull, 0xe0000000f0000000ull, N("st"), T(ldstm), GLOBAL, T(ldsto), .ptype = CP },
	{ 0xc0000000d0000000ull, 0xe0000000f0000000ull, T(redm), GLOBAL, T(ldsto), .fmask = F_SM11, .ptype = CP },
	{ 0xe0000000d0000000ull, 0xe0000000f0000000ull, T(atomm), T(ldsto), GLOBAL2, T(ldsts2), T(mldsts3), .fmask = F_SM11, .ptype = CP },

	// e
	{ 0x00000000e0000000ull, 0xc0000000f0000000ull, N("add"), T(o0sat), N("f32"), MCDST, LLDST, T(m1neg), SESTART, N("mul"), T(lsw), T(lc2w), SEEND, T(m2neg), T(lc3w) },	// multiply, round, add, round. meh.

	{ 0x40000000e0000000ull, 0xe0000000f0000000ull, N("fma"), T(mad64r), N("f64"), MCDST, LDDST, T(m1neg), LDSRC, LDSRC2, T(m2neg), LDSRC3, .fmask = F_FP64 },	// *fused* multiply-add, no intermediate rounding :)
	{ 0x60000000e0000000ull, 0xe0000000f0000000ull, N("add"), T(af64r), N("f64"), MCDST, LDDST, T(m1neg), LDSRC, T(m2neg), LDSRC3, .fmask = F_FP64 },
	{ 0x80000000e0000000ull, 0xe0000000f0000000ull, N("mul"), T(cvtrnd), N("f64"), MCDST, LDDST, T(m1neg), LDSRC, LDSRC2, .fmask = F_FP64 },
	{ 0xa0000000e0000000ull, 0xe0000000f0000000ull, N("min"), N("f64"), MCDST, LDDST, T(lfm1), LDSRC, T(lfm2), LDSRC2, .fmask = F_FP64 },
	{ 0xc0000000e0000000ull, 0xe0000000f0000000ull, N("max"), N("f64"), MCDST, LDDST, T(lfm1), LDSRC, T(lfm2), LDSRC2, .fmask = F_FP64 },
	{ 0xe0000000e0000000ull, 0xe0000000f0000000ull, N("set"), MCDST, LLDST, T(setf), N("f64"), T(lfm1), LDSRC, T(lfm2), LDSRC2, .fmask = F_FP64 },

	// f
	// order of inputs: x, y, z, index, dref, bias/lod. index is integer, others float.
	{ 0x00000000f0000000ull, 0xf0000000f9000000ull, N("texauto"), T(texf), LTDST, TEX, SAMP, LTSRC, TOFFX, TOFFY, TOFFZ },
	{ 0x00000000f8000000ull, 0xf0000000f9000000ull, N("texauto"), N("cube"), T(texf), LTDST, TEX, SAMP, LTSRC },

	// takes integer inputs.
	{ 0x00000000f1000000ull, 0xf0000000f1000000ull, N("texfetch"), T(texf), LTDST, TEX, SAMP, LTSRC, TOFFX, TOFFY, TOFFZ },

	// bias needs to be same for everything, or else.
	{ 0x20000000f0000000ull, 0xf0000000f8000000ull, N("texbias"), T(texf), LTDST, TEX, SAMP, LTSRC, TOFFX, TOFFY, TOFFZ },
	{ 0x20000000f8000000ull, 0xf0000000f8000000ull, N("texbias"), N("cube"), T(texf), LTDST, TEX, SAMP, LTSRC },

	// lod needs to be same for everything, or else.
	{ 0x40000000f0000000ull, 0xf0000000f8000000ull, N("texlod"), T(texf), LTDST, TEX, SAMP, LTSRC, TOFFX, TOFFY, TOFFZ },
	{ 0x40000000f8000000ull, 0xf0000000f8000000ull, N("texlod"), N("cube"), T(texf), LTDST, TEX, SAMP, LTSRC },

	// integer input and output. in: LOD, out: size.x, size.y, size.z
	{ 0x60000000f0000000ull, 0xf00f0000f0000000ull, N("texsize"), T(texf), LTDST, TEX, SAMP, LDST },

	// input: 3 normalized cube coords [float], layer [int]; output: equivalent x, y, combined layer coords to pass to non-cube tex variants.
	{ 0x60010000f8000000ull, 0xf00f0000f8000000ull, N("texprep"), N("cube"), T(texf), LTDST, TEX, SAMP, LTSRC, .fmask = F_D3D10_1 },

	// returned values are FIXED-point with 6 fractional bits
	{ 0x60020000f0000000ull, 0xf00f0000f8000000ull, N("texquerylod"), T(texf), LTDST, TEX, SAMP, LTSRC, .fmask = F_D3D10_1 },
	{ 0x60020000f8000000ull, 0xf00f0000f8000000ull, N("texquerylod"), N("cube"), T(texf), LTDST, TEX, SAMP, LTSRC, .fmask = F_D3D10_1 },

	// in: float coords + CSAA mask from ZETA
	{ 0x80000000f0000000ull, 0xf0000000f1000000ull, N("texcsaa"), T(texf), LTDST, TEX, SAMP, LTSRC, TOFFX, TOFFY, TOFFZ },

	{ 0x80000000f1000000ull, 0xf0000000f9000000ull, N("texgather"), T(texf), LTDST, TEX, SAMP, LTSRC, TOFFX, TOFFY, TOFFZ, .fmask = F_D3D10_1 },
	{ 0x80000000f9000000ull, 0xf0000000f9000000ull, N("texgather"), N("cube"), T(texf), LTDST, TEX, SAMP, LTSRC, .fmask = F_D3D10_1 },

	{ 0xc0000000f0000200ull, 0xe0000000f0000600ull, N("emit"), .ptype = GP },
	{ 0xc0000000f0000400ull, 0xe0000000f0000600ull, N("restart"), .ptype = GP },

	{ 0xe0000000f0000000ull, 0xe0000004f0000000ull, N("nop") },
	{ 0xe0000004f0000000ull, 0xe0000004f0000000ull, N("pmevent"), PM },

	// desc ^^^
	
	// try to print out *some* info.
	{ 0, 0, OOPS, MCDST, LLDST, T(lsw), T(lc2w), T(lc3w) },

};

/*
 * Predicates
 */
static struct insn tabp[] = {
	{ 0x0000000000000000ull, 0x00000f8000000000ull, SESTART, N("never"), SEEND },
	{ 0x0000008000000000ull, 0x00000f8000000000ull, SESTART, N("l"), COND, SEEND },
	{ 0x0000010000000000ull, 0x00000f8000000000ull, SESTART, N("e"), COND, SEEND },
	{ 0x0000018000000000ull, 0x00000f8000000000ull, SESTART, N("le"), COND, SEEND },
	{ 0x0000020000000000ull, 0x00000f8000000000ull, SESTART, N("g"), COND, SEEND },
	{ 0x0000028000000000ull, 0x00000f8000000000ull, SESTART, N("lg"), COND, SEEND },
	{ 0x0000030000000000ull, 0x00000f8000000000ull, SESTART, N("ge"), COND, SEEND },
	{ 0x0000038000000000ull, 0x00000f8000000000ull, SESTART, N("lge"), COND, SEEND },
	{ 0x0000040000000000ull, 0x00000f8000000000ull, SESTART, N("u"), COND, SEEND },
	{ 0x0000048000000000ull, 0x00000f8000000000ull, SESTART, N("lu"), COND, SEEND },
	{ 0x0000050000000000ull, 0x00000f8000000000ull, SESTART, N("eu"), COND, SEEND },
	{ 0x0000058000000000ull, 0x00000f8000000000ull, SESTART, N("leu"), COND, SEEND },
	{ 0x0000060000000000ull, 0x00000f8000000000ull, SESTART, N("gu"), COND, SEEND },
	{ 0x0000068000000000ull, 0x00000f8000000000ull, SESTART, N("lgu"), COND, SEEND },
	{ 0x0000070000000000ull, 0x00000f8000000000ull, SESTART, N("geu"), COND, SEEND },
	{ 0x0000078000000000ull, 0x00000f8000000000ull },
	{ 0x0000080000000000ull, 0x00000f8000000000ull, SESTART, N("o"), COND, SEEND },
	{ 0x0000088000000000ull, 0x00000f8000000000ull, SESTART, N("c"), COND, SEEND },
	{ 0x0000090000000000ull, 0x00000f8000000000ull, SESTART, N("a"), COND, SEEND },
	{ 0x0000098000000000ull, 0x00000f8000000000ull, SESTART, N("s"), COND, SEEND },
	{ 0x00000e0000000000ull, 0x00000f8000000000ull, SESTART, N("ns"), COND, SEEND },
	{ 0x00000e8000000000ull, 0x00000f8000000000ull, SESTART, N("na"), COND, SEEND },
	{ 0x00000f0000000000ull, 0x00000f8000000000ull, SESTART, N("nc"), COND, SEEND },
	{ 0x00000f8000000000ull, 0x00000f8000000000ull, SESTART, N("no"), COND, SEEND },
	{ 0, 0, OOPS },
};

F1(lim, 0x26, N("lim")) // if set, call is limited, and will be ignored if X limited calls already active on the stack. X is settable by CALL_LIMIT_LOG method.
F1(barinc, 0x19, N("inc"))
F1(barwait, 0x1a, N("wait"))

static struct insn tabc[] = {
	{ 0x0000000000000000ull, 0x00000000f0000000ull, T(p), N("discard"), .ptype = FP },
	{ 0x0000000010000000ull, 0x00000000f0000000ull, T(p), N("bra"), BTARG },
	{ 0x0000000020000000ull, 0x00000000f0000000ull, IGNPRED, N("call"), T(lim), CTARG },
	{ 0x0000000030000000ull, 0x00000000f0000000ull, T(p), N("ret") },
	{ 0x0000000040000000ull, 0x00000000f0000000ull, IGNPRED, N("breakaddr"), BTARG },
	{ 0x0000000050000000ull, 0x00000000f0000000ull, T(p), N("break") },
	{ 0x0000000060000000ull, 0x00000000f0000000ull, IGNPRED, N("quadon") },
	{ 0x0000000070000000ull, 0x00000000f0000000ull, IGNPRED, N("quadpop") },
	{ 0x0000000080000000ull, 0x00004000f0000000ull, IGNPRED, N("bar"), T(barinc), T(barwait), BAR, BARMASK },
	{ 0x0000400080000000ull, 0x00004000f0000000ull, IGNPRED, N("bar"), T(barinc), T(barwait), BAR, N("all") },
	{ 0x0000000090000000ull, 0x00000000f0000000ull, IGNPRED, N("trap") },
	{ 0x00000000a0000000ull, 0x00000000f0000000ull, IGNPRED, N("joinat"), BTARG },
	{ 0x00000000b0000000ull, 0x00000000f0000000ull, T(p), N("brkpt"), .fmask = F_SM11 },
	{ 0x00000000c0000000ull, 0x00000000f0000000ull, IGNPRED, N("bra"), T(lim), FCONST8, .fmask = F_D3D10_1 }, /* XXX: check variants */
	{ 0x00000000d0000000ull, 0x00000000f0000000ull, IGNPRED, N("preret"), T(lim), BTARG, .fmask = F_D3D10_1 },
	{ 0, 0, T(p), OOPS, BTARG },
};

static struct insn tabroot[] = {
	{         0x00000000ull,         0x00000001ull, OP4B, T(s) },
	{ 0x0000000000000003ull, 0x0000000000000003ull, OP8B, T(c) },
	{ 0x0000000000000001ull, 0x0000000300000003ull, OP8B, T(p), T(l) },
	{ 0x0000000100000001ull, 0x0000000300000003ull, OP8B, N("exit"), T(p), T(l) },
	{ 0x0000000200000001ull, 0x0000000300000003ull, OP8B, N("join"), T(p), T(l) },
	{ 0x0000000000000003ull, 0x0000000000000003ull, OP8B, N("long"), T(c) },
	{ 0x0000000000000001ull, 0x0000000300000003ull, OP8B, N("long"), T(p), T(l) },
	{ 0x0000000100000001ull, 0x0000000300000003ull, OP8B, N("long"), N("exit"), T(p), T(l) },
	{ 0x0000000200000001ull, 0x0000000300000003ull, OP8B, N("long"), N("join"), T(p), T(l) },
	{ 0x0000000300000001ull, 0x0000000300000003ull, OP8B, T(i) },
	{ 0, 0, OOPS },
};

static void g80_prep(struct disisa *isa) {
	isa->vardata = vardata_new("g80isa");
	int f_sm11 = vardata_add_feature(isa->vardata, "sm11", "SM1.1 new opcodes");
	int f_sm12 = vardata_add_feature(isa->vardata, "sm12", "SM1.2 new opcodes");
	int f_fp64 = vardata_add_feature(isa->vardata, "fp64", "64-bit floating point");
	int f_d3d10_1 = vardata_add_feature(isa->vardata, "d3d10_1", "Direct3D 10.1 new features");
	if (f_sm11 == -1 || f_sm12 == -1 || f_fp64 == -1 || f_d3d10_1 == -1)
		abort();
	vardata_feature_imply(isa->vardata, f_sm12, f_sm11);
	int vs_chipset = vardata_add_varset(isa->vardata, "chipset", "GPU chipset");
	if (vs_chipset == -1)
		abort();
	int v_g80 = vardata_add_variant(isa->vardata, "g80", "G80", vs_chipset);
	int v_g84 = vardata_add_variant(isa->vardata, "g84", "G84:G200", vs_chipset);
	int v_g200 = vardata_add_variant(isa->vardata, "g200", "G200", vs_chipset);
	int v_mcp77 = vardata_add_variant(isa->vardata, "mcp77", "MCP77, MCP79", vs_chipset);
	int v_gt215 = vardata_add_variant(isa->vardata, "gt215", "GT215+", vs_chipset);
	if (v_g80 == -1 || v_g84 == -1 || v_g200 == -1 || v_mcp77 == -1 || v_gt215 == -1)
		abort();
	vardata_variant_feature(isa->vardata, v_g84, f_sm11);
	vardata_variant_feature(isa->vardata, v_g200, f_sm12);
	vardata_variant_feature(isa->vardata, v_mcp77, f_sm12);
	vardata_variant_feature(isa->vardata, v_gt215, f_sm12);
	vardata_variant_feature(isa->vardata, v_g200, f_fp64);
	vardata_variant_feature(isa->vardata, v_gt215, f_d3d10_1);
	int ms_ptype = vardata_add_modeset(isa->vardata, "ptype", "Program type");
	if (ms_ptype == -1)
		abort();
	int m_vp = vardata_add_mode(isa->vardata, "vp", "vertex program", ms_ptype);
	int m_gp = vardata_add_mode(isa->vardata, "gp", "geometry program", ms_ptype);
	int m_fp = vardata_add_mode(isa->vardata, "fp", "fragment program", ms_ptype);
	int m_cp = vardata_add_mode(isa->vardata, "cp", "compute program", ms_ptype);
	if (m_vp == -1 || m_gp == -1 || m_fp == -1 || m_cp == -1)
		abort();
	if (vardata_validate(isa->vardata))
		abort();
}

struct disisa g80_isa_s = {
	tabroot,
	8,
	4,
	1,
	.i_need_g80as_hack = 1,
	.prep = g80_prep,
};