github/tinygrad
1
1
Fork 0
mirror of https://github.com/tinygrad/tinygrad.git synced 2026-01-09 15:08:02 -05:00
Code Issues Packages Projects Releases Wiki Activity
Files
49d1bf93d636f032316f7d6f2203f0ee38374080
tinygrad/extra/assembly/amd/autogen/cdna/__init__.py
George Hotz 49d1bf93d6 assembly/amd: refactor asm.py to be simpler (#13900) 
* assembly/amd: refactor asm.py

* assembly/amd: refactor asm.py to be simpler

* multiple fxns

* fast

* more tests pass

* regen

* stop decode
2025-12-30 13:51:40 -05:00

3480 lines
147 KiB
Python
Raw Blame History

# autogenerated from AMD CDNA3+CDNA4 ISA PDF by dsl.py - do not edit
from enum import IntEnum
from typing import Annotated
from extra.assembly.amd.dsl import bits, BitField, Inst32, Inst64, SGPR, VGPR, TTMP as TTMP, s as s, v as v, ttmp as ttmp, SSrc, Src, SImm, Imm, VDSTYEnc, SGPRField, VGPRField
import functools
class SrcEnum(IntEnum):
S_ADD_U32 = 0
S_SUB_U32 = 1
S_ADD_I32 = 2
S_SUB_I32 = 3
S_ADDC_U32 = 4
S_SUBB_U32 = 5
S_MIN_I32 = 6
FLAT_SCRATCH_LO = 102
FLAT_SCRATCH_HI = 103
XNACK_MASK_LO = 104
XNACK_MASK_HI = 105
VCC_LO = 106
VCC_HI = 107
M0 = 124
EXEC_LO = 126
EXEC_HI = 127
ZERO = 128
DPP8 = 233
DPP8FI = 234
SHARED_BASE = 235
SHARED_LIMIT = 236
PRIVATE_BASE = 237
PRIVATE_LIMIT = 238
RESERVED = 239
POS_HALF = 240
NEG_HALF = 241
POS_ONE = 242
NEG_ONE = 243
POS_TWO = 244
NEG_TWO = 245
POS_FOUR = 246
NEG_FOUR = 247
INV_2PI = 248
DPP16 = 250
VCCZ = 251
EXECZ = 252
SCC = 253
LDS_DIRECT = 254
class DSOp(IntEnum):
DS_ADD_U32 = 0
DS_SUB_U32 = 1
DS_RSUB_U32 = 2
DS_INC_U32 = 3
DS_DEC_U32 = 4
DS_MIN_I32 = 5
DS_MAX_I32 = 6
DS_MIN_U32 = 7
DS_MAX_U32 = 8
DS_AND_B32 = 9
DS_OR_B32 = 10
DS_XOR_B32 = 11
DS_MSKOR_B32 = 12
DS_WRITE_B32 = 13
DS_WRITE2_B32 = 14
DS_WRITE2ST64_B32 = 15
DS_CMPST_B32 = 16
DS_CMPST_F32 = 17
DS_MIN_F32 = 18
DS_MAX_F32 = 19
DS_NOP = 20
DS_ADD_F32 = 21
DS_PK_ADD_F16 = 23
DS_PK_ADD_BF16 = 24
DS_WRITE_ADDTID_B32 = 29
DS_WRITE_B8 = 30
DS_WRITE_B16 = 31
DS_ADD_RTN_U32 = 32
DS_SUB_RTN_U32 = 33
DS_RSUB_RTN_U32 = 34
DS_INC_RTN_U32 = 35
DS_DEC_RTN_U32 = 36
DS_MIN_RTN_I32 = 37
DS_MAX_RTN_I32 = 38
DS_MIN_RTN_U32 = 39
DS_MAX_RTN_U32 = 40
DS_AND_RTN_B32 = 41
DS_OR_RTN_B32 = 42
DS_XOR_RTN_B32 = 43
DS_MSKOR_RTN_B32 = 44
DS_WRXCHG_RTN_B32 = 45
DS_WRXCHG2_RTN_B32 = 46
DS_WRXCHG2ST64_RTN_B32 = 47
DS_CMPST_RTN_B32 = 48
DS_CMPST_RTN_F32 = 49
DS_MIN_RTN_F32 = 50
DS_MAX_RTN_F32 = 51
DS_WRAP_RTN_B32 = 52
DS_ADD_RTN_F32 = 53
DS_READ_B32 = 54
DS_READ2_B32 = 55
DS_READ2ST64_B32 = 56
DS_READ_I8 = 57
DS_READ_U8 = 58
DS_READ_I16 = 59
DS_READ_U16 = 60
DS_SWIZZLE_B32 = 61
DS_PERMUTE_B32 = 62
DS_BPERMUTE_B32 = 63
DS_ADD_U64 = 64
DS_SUB_U64 = 65
DS_RSUB_U64 = 66
DS_INC_U64 = 67
DS_DEC_U64 = 68
DS_MIN_I64 = 69
DS_MAX_I64 = 70
DS_MIN_U64 = 71
DS_MAX_U64 = 72
DS_AND_B64 = 73
DS_OR_B64 = 74
DS_XOR_B64 = 75
DS_MSKOR_B64 = 76
DS_WRITE_B64 = 77
DS_WRITE2_B64 = 78
DS_WRITE2ST64_B64 = 79
DS_CMPST_B64 = 80
DS_CMPST_F64 = 81
DS_MIN_F64 = 82
DS_MAX_F64 = 83
DS_WRITE_B8_D16_HI = 84
DS_WRITE_B16_D16_HI = 85
DS_READ_U8_D16 = 86
DS_READ_U8_D16_HI = 87
DS_READ_I8_D16 = 88
DS_READ_I8_D16_HI = 89
DS_READ_U16_D16 = 90
DS_READ_U16_D16_HI = 91
DS_ADD_F64 = 92
DS_ADD_RTN_U64 = 96
DS_SUB_RTN_U64 = 97
DS_RSUB_RTN_U64 = 98
DS_INC_RTN_U64 = 99
DS_DEC_RTN_U64 = 100
DS_MIN_RTN_I64 = 101
DS_MAX_RTN_I64 = 102
DS_MIN_RTN_U64 = 103
DS_MAX_RTN_U64 = 104
DS_AND_RTN_B64 = 105
DS_OR_RTN_B64 = 106
DS_XOR_RTN_B64 = 107
DS_MSKOR_RTN_B64 = 108
DS_WRXCHG_RTN_B64 = 109
DS_WRXCHG2_RTN_B64 = 110
DS_WRXCHG2ST64_RTN_B64 = 111
DS_CMPST_RTN_B64 = 112
DS_CMPST_RTN_F64 = 113
DS_MIN_RTN_F64 = 114
DS_MAX_RTN_F64 = 115
DS_READ_B64 = 118
DS_READ2_B64 = 119
DS_READ2ST64_B64 = 120
DS_ADD_RTN_F64 = 124
DS_CONDXCHG32_RTN_B64 = 126
DS_GWS_SEMA_RELEASE_ALL = 152
DS_GWS_INIT = 153
DS_GWS_SEMA_V = 154
DS_GWS_SEMA_BR = 155
DS_GWS_SEMA_P = 156
DS_GWS_BARRIER = 157
DS_READ_ADDTID_B32 = 182
DS_PK_ADD_RTN_F16 = 183
DS_PK_ADD_RTN_BF16 = 184
DS_CONSUME = 189
DS_APPEND = 190
DS_WRITE_B96 = 222
DS_WRITE_B128 = 223
DS_READ_B64_TR_B4 = 224
DS_READ_B96_TR_B6 = 225
DS_READ_B64_TR_B8 = 226
DS_READ_B64_TR_B16 = 227
DS_READ_B96 = 254
DS_READ_B128 = 255
CDNA4 = 600
class FLATOp(IntEnum):
FLAT_LOAD_UBYTE = 16
FLAT_LOAD_SBYTE = 17
FLAT_LOAD_USHORT = 18
FLAT_LOAD_SSHORT = 19
FLAT_LOAD_DWORD = 20
FLAT_LOAD_DWORDX2 = 21
FLAT_LOAD_DWORDX3 = 22
FLAT_LOAD_DWORDX4 = 23
FLAT_STORE_BYTE = 24
FLAT_STORE_BYTE_D16_HI = 25
FLAT_STORE_SHORT = 26
FLAT_STORE_SHORT_D16_HI = 27
FLAT_STORE_DWORD = 28
FLAT_STORE_DWORDX2 = 29
FLAT_STORE_DWORDX3 = 30
FLAT_STORE_DWORDX4 = 31
FLAT_LOAD_UBYTE_D16 = 32
FLAT_LOAD_UBYTE_D16_HI = 33
FLAT_LOAD_SBYTE_D16 = 34
FLAT_LOAD_SBYTE_D16_HI = 35
FLAT_LOAD_SHORT_D16 = 36
FLAT_LOAD_SHORT_D16_HI = 37
FLAT_ATOMIC_SWAP = 64
FLAT_ATOMIC_CMPSWAP = 65
FLAT_ATOMIC_ADD = 66
FLAT_ATOMIC_SUB = 67
FLAT_ATOMIC_SMIN = 68
FLAT_ATOMIC_UMIN = 69
FLAT_ATOMIC_SMAX = 70
FLAT_ATOMIC_UMAX = 71
FLAT_ATOMIC_AND = 72
FLAT_ATOMIC_OR = 73
FLAT_ATOMIC_XOR = 74
FLAT_ATOMIC_INC = 75
FLAT_ATOMIC_DEC = 76
FLAT_ATOMIC_ADD_F32 = 77
FLAT_ATOMIC_PK_ADD_F16 = 78
FLAT_ATOMIC_ADD_F64 = 79
FLAT_ATOMIC_MIN_F64 = 80
FLAT_ATOMIC_MAX_F64 = 81
FLAT_ATOMIC_PK_ADD_BF16 = 82
FLAT_ATOMIC_SWAP_X2 = 96
FLAT_ATOMIC_CMPSWAP_X2 = 97
FLAT_ATOMIC_ADD_X2 = 98
FLAT_ATOMIC_SUB_X2 = 99
FLAT_ATOMIC_SMIN_X2 = 100
FLAT_ATOMIC_UMIN_X2 = 101
FLAT_ATOMIC_SMAX_X2 = 102
FLAT_ATOMIC_UMAX_X2 = 103
FLAT_ATOMIC_AND_X2 = 104
FLAT_ATOMIC_OR_X2 = 105
FLAT_ATOMIC_XOR_X2 = 106
FLAT_ATOMIC_INC_X2 = 107
FLAT_ATOMIC_DEC_X2 = 108
CDNA4 = 600
class GLOBALOp(IntEnum):
GLOBAL_LOAD_UBYTE = 16
GLOBAL_LOAD_SBYTE = 17
GLOBAL_LOAD_USHORT = 18
GLOBAL_LOAD_SSHORT = 19
GLOBAL_LOAD_DWORD = 20
GLOBAL_LOAD_DWORDX2 = 21
GLOBAL_LOAD_DWORDX3 = 22
GLOBAL_LOAD_DWORDX4 = 23
GLOBAL_STORE_BYTE = 24
GLOBAL_STORE_BYTE_D16_HI = 25
GLOBAL_STORE_SHORT = 26
GLOBAL_STORE_SHORT_D16_HI = 27
GLOBAL_STORE_DWORD = 28
GLOBAL_STORE_DWORDX2 = 29
GLOBAL_STORE_DWORDX3 = 30
GLOBAL_STORE_DWORDX4 = 31
GLOBAL_LOAD_UBYTE_D16 = 32
GLOBAL_LOAD_UBYTE_D16_HI = 33
GLOBAL_LOAD_SBYTE_D16 = 34
GLOBAL_LOAD_SBYTE_D16_HI = 35
GLOBAL_LOAD_SHORT_D16 = 36
GLOBAL_LOAD_SHORT_D16_HI = 37
GLOBAL_LOAD_LDS_UBYTE = 38
GLOBAL_LOAD_LDS_SBYTE = 39
GLOBAL_LOAD_LDS_USHORT = 40
GLOBAL_LOAD_LDS_SSHORT = 41
GLOBAL_LOAD_LDS_DWORD = 42
GLOBAL_ATOMIC_SWAP = 64
GLOBAL_ATOMIC_CMPSWAP = 65
GLOBAL_ATOMIC_ADD = 66
GLOBAL_ATOMIC_SUB = 67
GLOBAL_ATOMIC_SMIN = 68
GLOBAL_ATOMIC_UMIN = 69
GLOBAL_ATOMIC_SMAX = 70
GLOBAL_ATOMIC_UMAX = 71
GLOBAL_ATOMIC_AND = 72
GLOBAL_ATOMIC_OR = 73
GLOBAL_ATOMIC_XOR = 74
GLOBAL_ATOMIC_INC = 75
GLOBAL_ATOMIC_DEC = 76
GLOBAL_ATOMIC_ADD_F32 = 77
GLOBAL_ATOMIC_PK_ADD_F16 = 78
GLOBAL_ATOMIC_ADD_F64 = 79
GLOBAL_ATOMIC_MIN_F64 = 80
GLOBAL_ATOMIC_MAX_F64 = 81
GLOBAL_ATOMIC_PK_ADD_BF16 = 82
GLOBAL_ATOMIC_SWAP_X2 = 96
GLOBAL_ATOMIC_CMPSWAP_X2 = 97
GLOBAL_ATOMIC_ADD_X2 = 98
GLOBAL_ATOMIC_SUB_X2 = 99
GLOBAL_ATOMIC_SMIN_X2 = 100
GLOBAL_ATOMIC_UMIN_X2 = 101
GLOBAL_ATOMIC_SMAX_X2 = 102
GLOBAL_ATOMIC_UMAX_X2 = 103
GLOBAL_ATOMIC_AND_X2 = 104
GLOBAL_ATOMIC_OR_X2 = 105
GLOBAL_ATOMIC_XOR_X2 = 106
GLOBAL_ATOMIC_INC_X2 = 107
GLOBAL_ATOMIC_DEC_X2 = 108
GLOBAL_LOAD_LDS_DWORDX4 = 125
GLOBAL_LOAD_LDS_DWORDX3 = 126
CDNA4 = 600
class MTBUFOp(IntEnum):
TBUFFER_LOAD_FORMAT_X = 0
TBUFFER_LOAD_FORMAT_XY = 1
TBUFFER_LOAD_FORMAT_XYZ = 2
TBUFFER_LOAD_FORMAT_XYZW = 3
TBUFFER_STORE_FORMAT_X = 4
TBUFFER_STORE_FORMAT_XY = 5
TBUFFER_STORE_FORMAT_XYZ = 6
TBUFFER_STORE_FORMAT_XYZW = 7
TBUFFER_LOAD_FORMAT_D16_X = 8
TBUFFER_LOAD_FORMAT_D16_XY = 9
TBUFFER_LOAD_FORMAT_D16_XYZ = 10
TBUFFER_LOAD_FORMAT_D16_XYZW = 11
TBUFFER_STORE_FORMAT_D16_X = 12
TBUFFER_STORE_FORMAT_D16_XY = 13
TBUFFER_STORE_FORMAT_D16_XYZ = 14
TBUFFER_STORE_FORMAT_D16_XYZW = 15
class MUBUFOp(IntEnum):
BUFFER_LOAD_FORMAT_X = 0
BUFFER_LOAD_FORMAT_XY = 1
BUFFER_LOAD_FORMAT_XYZ = 2
BUFFER_LOAD_FORMAT_XYZW = 3
BUFFER_STORE_FORMAT_X = 4
BUFFER_STORE_FORMAT_XY = 5
BUFFER_STORE_FORMAT_XYZ = 6
BUFFER_STORE_FORMAT_XYZW = 7
BUFFER_LOAD_FORMAT_D16_X = 8
BUFFER_LOAD_FORMAT_D16_XY = 9
BUFFER_LOAD_FORMAT_D16_XYZ = 10
BUFFER_LOAD_FORMAT_D16_XYZW = 11
BUFFER_STORE_FORMAT_D16_X = 12
BUFFER_STORE_FORMAT_D16_XY = 13
BUFFER_STORE_FORMAT_D16_XYZ = 14
BUFFER_STORE_FORMAT_D16_XYZW = 15
BUFFER_LOAD_UBYTE = 16
BUFFER_LOAD_SBYTE = 17
BUFFER_LOAD_USHORT = 18
BUFFER_LOAD_SSHORT = 19
BUFFER_LOAD_DWORD = 20
BUFFER_LOAD_DWORDX2 = 21
BUFFER_LOAD_DWORDX3 = 22
BUFFER_LOAD_DWORDX4 = 23
BUFFER_STORE_BYTE = 24
BUFFER_STORE_BYTE_D16_HI = 25
BUFFER_STORE_SHORT = 26
BUFFER_STORE_SHORT_D16_HI = 27
BUFFER_STORE_DWORD = 28
BUFFER_STORE_DWORDX2 = 29
BUFFER_STORE_DWORDX3 = 30
BUFFER_STORE_DWORDX4 = 31
BUFFER_LOAD_UBYTE_D16 = 32
BUFFER_LOAD_UBYTE_D16_HI = 33
BUFFER_LOAD_SBYTE_D16 = 34
BUFFER_LOAD_SBYTE_D16_HI = 35
BUFFER_LOAD_SHORT_D16 = 36
BUFFER_LOAD_SHORT_D16_HI = 37
BUFFER_LOAD_FORMAT_D16_HI_X = 38
BUFFER_STORE_FORMAT_D16_HI_X = 39
BUFFER_WBL2 = 40
BUFFER_INV = 41
BUFFER_ATOMIC_SWAP = 64
BUFFER_ATOMIC_CMPSWAP = 65
BUFFER_ATOMIC_ADD = 66
BUFFER_ATOMIC_SUB = 67
BUFFER_ATOMIC_SMIN = 68
BUFFER_ATOMIC_UMIN = 69
BUFFER_ATOMIC_SMAX = 70
BUFFER_ATOMIC_UMAX = 71
BUFFER_ATOMIC_AND = 72
BUFFER_ATOMIC_OR = 73
BUFFER_ATOMIC_XOR = 74
BUFFER_ATOMIC_INC = 75
BUFFER_ATOMIC_DEC = 76
BUFFER_ATOMIC_ADD_F32 = 77
BUFFER_ATOMIC_PK_ADD_F16 = 78
BUFFER_ATOMIC_ADD_F64 = 79
BUFFER_ATOMIC_MIN_F64 = 80
BUFFER_ATOMIC_MAX_F64 = 81
BUFFER_ATOMIC_PK_ADD_BF16 = 82
BUFFER_ATOMIC_SWAP_X2 = 96
BUFFER_ATOMIC_CMPSWAP_X2 = 97
BUFFER_ATOMIC_ADD_X2 = 98
BUFFER_ATOMIC_SUB_X2 = 99
BUFFER_ATOMIC_SMIN_X2 = 100
BUFFER_ATOMIC_UMIN_X2 = 101
BUFFER_ATOMIC_SMAX_X2 = 102
BUFFER_ATOMIC_UMAX_X2 = 103
BUFFER_ATOMIC_AND_X2 = 104
BUFFER_ATOMIC_OR_X2 = 105
BUFFER_ATOMIC_XOR_X2 = 106
BUFFER_ATOMIC_INC_X2 = 107
BUFFER_ATOMIC_DEC_X2 = 108
CDNA4 = 600
class SCRATCHOp(IntEnum):
SCRATCH_LOAD_UBYTE = 16
SCRATCH_LOAD_SBYTE = 17
SCRATCH_LOAD_USHORT = 18
SCRATCH_LOAD_SSHORT = 19
SCRATCH_LOAD_DWORD = 20
SCRATCH_LOAD_DWORDX2 = 21
SCRATCH_LOAD_DWORDX3 = 22
SCRATCH_LOAD_DWORDX4 = 23
SCRATCH_STORE_BYTE = 24
SCRATCH_STORE_BYTE_D16_HI = 25
SCRATCH_STORE_SHORT = 26
SCRATCH_STORE_SHORT_D16_HI = 27
SCRATCH_STORE_DWORD = 28
SCRATCH_STORE_DWORDX2 = 29
SCRATCH_STORE_DWORDX3 = 30
SCRATCH_STORE_DWORDX4 = 31
SCRATCH_LOAD_UBYTE_D16 = 32
SCRATCH_LOAD_UBYTE_D16_HI = 33
SCRATCH_LOAD_SBYTE_D16 = 34
SCRATCH_LOAD_SBYTE_D16_HI = 35
SCRATCH_LOAD_SHORT_D16 = 36
SCRATCH_LOAD_SHORT_D16_HI = 37
SCRATCH_LOAD_LDS_UBYTE = 38
SCRATCH_LOAD_LDS_SBYTE = 39
SCRATCH_LOAD_LDS_USHORT = 40
SCRATCH_LOAD_LDS_SSHORT = 41
SCRATCH_LOAD_LDS_DWORD = 42
class SMEMOp(IntEnum):
S_LOAD_DWORD = 0
S_LOAD_DWORDX2 = 1
S_LOAD_DWORDX4 = 2
S_LOAD_DWORDX8 = 3
S_LOAD_DWORDX16 = 4
S_SCRATCH_LOAD_DWORD = 5
S_SCRATCH_LOAD_DWORDX2 = 6
S_SCRATCH_LOAD_DWORDX4 = 7
S_BUFFER_LOAD_DWORD = 8
S_BUFFER_LOAD_DWORDX2 = 9
S_BUFFER_LOAD_DWORDX4 = 10
S_BUFFER_LOAD_DWORDX8 = 11
S_BUFFER_LOAD_DWORDX16 = 12
S_STORE_DWORD = 16
S_STORE_DWORDX2 = 17
S_STORE_DWORDX4 = 18
S_SCRATCH_STORE_DWORD = 21
S_SCRATCH_STORE_DWORDX2 = 22
S_SCRATCH_STORE_DWORDX4 = 23
S_BUFFER_STORE_DWORD = 24
S_BUFFER_STORE_DWORDX2 = 25
S_BUFFER_STORE_DWORDX4 = 26
S_DCACHE_INV = 32
S_DCACHE_WB = 33
S_DCACHE_INV_VOL = 34
S_DCACHE_WB_VOL = 35
S_MEMTIME = 36
S_MEMREALTIME = 37
S_DCACHE_DISCARD = 40
S_DCACHE_DISCARD_X2 = 41
S_BUFFER_ATOMIC_SWAP = 64
S_BUFFER_ATOMIC_CMPSWAP = 65
S_BUFFER_ATOMIC_ADD = 66
S_BUFFER_ATOMIC_SUB = 67
S_BUFFER_ATOMIC_SMIN = 68
S_BUFFER_ATOMIC_UMIN = 69
S_BUFFER_ATOMIC_SMAX = 70
S_BUFFER_ATOMIC_UMAX = 71
S_BUFFER_ATOMIC_AND = 72
S_BUFFER_ATOMIC_OR = 73
S_BUFFER_ATOMIC_XOR = 74
S_BUFFER_ATOMIC_INC = 75
S_BUFFER_ATOMIC_DEC = 76
S_BUFFER_ATOMIC_SWAP_X2 = 96
S_BUFFER_ATOMIC_CMPSWAP_X2 = 97
S_BUFFER_ATOMIC_ADD_X2 = 98
S_BUFFER_ATOMIC_SUB_X2 = 99
S_BUFFER_ATOMIC_SMIN_X2 = 100
S_BUFFER_ATOMIC_UMIN_X2 = 101
S_BUFFER_ATOMIC_SMAX_X2 = 102
S_BUFFER_ATOMIC_UMAX_X2 = 103
S_BUFFER_ATOMIC_AND_X2 = 104
S_BUFFER_ATOMIC_OR_X2 = 105
S_BUFFER_ATOMIC_XOR_X2 = 106
S_BUFFER_ATOMIC_INC_X2 = 107
S_BUFFER_ATOMIC_DEC_X2 = 108
S_ATOMIC_SWAP = 128
S_ATOMIC_CMPSWAP = 129
S_ATOMIC_ADD = 130
S_ATOMIC_SUB = 131
S_ATOMIC_SMIN = 132
S_ATOMIC_UMIN = 133
S_ATOMIC_SMAX = 134
S_ATOMIC_UMAX = 135
S_ATOMIC_AND = 136
S_ATOMIC_OR = 137
S_ATOMIC_XOR = 138
S_ATOMIC_INC = 139
S_ATOMIC_DEC = 140
S_ATOMIC_SWAP_X2 = 160
S_ATOMIC_CMPSWAP_X2 = 161
S_ATOMIC_ADD_X2 = 162
S_ATOMIC_SUB_X2 = 163
S_ATOMIC_SMIN_X2 = 164
S_ATOMIC_UMIN_X2 = 165
S_ATOMIC_SMAX_X2 = 166
S_ATOMIC_UMAX_X2 = 167
S_ATOMIC_AND_X2 = 168
S_ATOMIC_OR_X2 = 169
S_ATOMIC_XOR_X2 = 170
S_ATOMIC_INC_X2 = 171
S_ATOMIC_DEC_X2 = 172
CDNA4 = 600
class SOP1Op(IntEnum):
S_MOV_B32 = 0
S_MOV_B64 = 1
S_CMOV_B32 = 2
S_CMOV_B64 = 3
S_NOT_B32 = 4
S_NOT_B64 = 5
S_WQM_B32 = 6
S_WQM_B64 = 7
S_BREV_B32 = 8
S_BREV_B64 = 9
S_BCNT0_I32_B32 = 10
S_BCNT0_I32_B64 = 11
S_BCNT1_I32_B32 = 12
S_BCNT1_I32_B64 = 13
S_FF0_I32_B32 = 14
S_FF0_I32_B64 = 15
S_FF1_I32_B32 = 16
S_FF1_I32_B64 = 17
S_FLBIT_I32_B32 = 18
S_FLBIT_I32_B64 = 19
S_FLBIT_I32 = 20
S_FLBIT_I32_I64 = 21
S_SEXT_I32_I8 = 22
S_SEXT_I32_I16 = 23
S_BITSET0_B32 = 24
S_BITSET0_B64 = 25
S_BITSET1_B32 = 26
S_BITSET1_B64 = 27
S_GETPC_B64 = 28
S_SETPC_B64 = 29
S_SWAPPC_B64 = 30
S_RFE_B64 = 31
S_AND_SAVEEXEC_B64 = 32
S_OR_SAVEEXEC_B64 = 33
S_XOR_SAVEEXEC_B64 = 34
S_ANDN2_SAVEEXEC_B64 = 35
S_ORN2_SAVEEXEC_B64 = 36
S_NAND_SAVEEXEC_B64 = 37
S_NOR_SAVEEXEC_B64 = 38
S_XNOR_SAVEEXEC_B64 = 39
S_QUADMASK_B32 = 40
S_QUADMASK_B64 = 41
S_MOVRELS_B32 = 42
S_MOVRELS_B64 = 43
S_MOVRELD_B32 = 44
S_MOVRELD_B64 = 45
S_CBRANCH_JOIN = 46
S_ABS_I32 = 48
S_SET_GPR_IDX_IDX = 50
S_ANDN1_SAVEEXEC_B64 = 51
S_ORN1_SAVEEXEC_B64 = 52
S_ANDN1_WREXEC_B64 = 53
S_ANDN2_WREXEC_B64 = 54
S_BITREPLICATE_B64_B32 = 55
CDNA4 = 600
class SOP2Op(IntEnum):
S_ADD_U32 = 0
S_SUB_U32 = 1
S_ADD_I32 = 2
S_SUB_I32 = 3
S_ADDC_U32 = 4
S_SUBB_U32 = 5
S_MIN_I32 = 6
S_MIN_U32 = 7
S_MAX_I32 = 8
S_MAX_U32 = 9
S_CSELECT_B32 = 10
S_CSELECT_B64 = 11
S_AND_B32 = 12
S_AND_B64 = 13
S_OR_B32 = 14
S_OR_B64 = 15
S_XOR_B32 = 16
S_XOR_B64 = 17
S_ANDN2_B32 = 18
S_ANDN2_B64 = 19
S_ORN2_B32 = 20
S_ORN2_B64 = 21
S_NAND_B32 = 22
S_NAND_B64 = 23
S_NOR_B32 = 24
S_NOR_B64 = 25
S_XNOR_B32 = 26
S_XNOR_B64 = 27
S_LSHL_B32 = 28
S_LSHL_B64 = 29
S_LSHR_B32 = 30
S_LSHR_B64 = 31
S_ASHR_I32 = 32
S_ASHR_I64 = 33
S_BFM_B32 = 34
S_BFM_B64 = 35
S_MUL_I32 = 36
S_BFE_U32 = 37
S_BFE_I32 = 38
S_BFE_U64 = 39
S_BFE_I64 = 40
S_CBRANCH_G_FORK = 41
S_ABSDIFF_I32 = 42
S_MUL_HI_U32 = 44
S_MUL_HI_I32 = 45
S_LSHL1_ADD_U32 = 46
S_LSHL2_ADD_U32 = 47
S_LSHL3_ADD_U32 = 48
S_LSHL4_ADD_U32 = 49
S_PACK_LL_B32_B16 = 50
S_PACK_LH_B32_B16 = 51
S_PACK_HH_B32_B16 = 52
CDNA4 = 600
class SOPCOp(IntEnum):
S_CMP_EQ_I32 = 0
S_CMP_LG_I32 = 1
S_CMP_GT_I32 = 2
S_CMP_GE_I32 = 3
S_CMP_LT_I32 = 4
S_CMP_LE_I32 = 5
S_CMP_EQ_U32 = 6
S_CMP_LG_U32 = 7
S_CMP_GT_U32 = 8
S_CMP_GE_U32 = 9
S_CMP_LT_U32 = 10
S_CMP_LE_U32 = 11
S_BITCMP0_B32 = 12
S_BITCMP1_B32 = 13
S_BITCMP0_B64 = 14
S_BITCMP1_B64 = 15
S_SETVSKIP = 16
S_SET_GPR_IDX_ON = 17
S_CMP_EQ_U64 = 18
S_CMP_LG_U64 = 19
CDNA4 = 600
class SOPKOp(IntEnum):
S_MOVK_I32 = 0
S_CMOVK_I32 = 1
S_CMPK_EQ_I32 = 2
S_CMPK_LG_I32 = 3
S_CMPK_GT_I32 = 4
S_CMPK_GE_I32 = 5
S_CMPK_LT_I32 = 6
S_CMPK_LE_I32 = 7
S_CMPK_EQ_U32 = 8
S_CMPK_LG_U32 = 9
S_CMPK_GT_U32 = 10
S_CMPK_GE_U32 = 11
S_CMPK_LT_U32 = 12
S_CMPK_LE_U32 = 13
S_ADDK_I32 = 14
S_MULK_I32 = 15
S_CBRANCH_I_FORK = 16
S_GETREG_B32 = 17
S_SETREG_B32 = 18
S_SETREG_IMM32_B32 = 20
S_CALL_B64 = 21
class SOPPOp(IntEnum):
S_NOP = 0
S_ENDPGM = 1
S_BRANCH = 2
S_WAKEUP = 3
S_CBRANCH_SCC0 = 4
S_CBRANCH_SCC1 = 5
S_CBRANCH_VCCZ = 6
S_CBRANCH_VCCNZ = 7
S_CBRANCH_EXECZ = 8
S_CBRANCH_EXECNZ = 9
S_BARRIER = 10
S_SETKILL = 11
S_WAITCNT = 12
S_SETHALT = 13
S_SLEEP = 14
S_SETPRIO = 15
S_SENDMSG = 16
S_SENDMSGHALT = 17
S_TRAP = 18
S_ICACHE_INV = 19
S_INCPERFLEVEL = 20
S_DECPERFLEVEL = 21
S_TTRACEDATA = 22
S_CBRANCH_CDBGSYS = 23
S_CBRANCH_CDBGUSER = 24
S_CBRANCH_CDBGSYS_OR_USER = 25
S_CBRANCH_CDBGSYS_AND_USER = 26
S_ENDPGM_SAVED = 27
S_SET_GPR_IDX_OFF = 28
S_SET_GPR_IDX_MODE = 29
CDNA4 = 600
class VOP1Op(IntEnum):
V_NOP = 0
V_MOV_B32 = 1
V_READFIRSTLANE_B32 = 2
V_CVT_I32_F64 = 3
V_CVT_F64_I32 = 4
V_CVT_F32_I32 = 5
V_CVT_F32_U32 = 6
V_CVT_U32_F32 = 7
V_CVT_I32_F32 = 8
V_CVT_F16_F32 = 10
V_CVT_F32_F16 = 11
V_CVT_RPI_I32_F32 = 12
V_CVT_FLR_I32_F32 = 13
V_CVT_OFF_F32_I4 = 14
V_CVT_F32_F64 = 15
V_CVT_F64_F32 = 16
V_CVT_F32_UBYTE0 = 17
V_CVT_F32_UBYTE1 = 18
V_CVT_F32_UBYTE2 = 19
V_CVT_F32_UBYTE3 = 20
V_CVT_U32_F64 = 21
V_CVT_F64_U32 = 22
V_TRUNC_F64 = 23
V_CEIL_F64 = 24
V_RNDNE_F64 = 25
V_FLOOR_F64 = 26
V_FRACT_F32 = 27
V_TRUNC_F32 = 28
V_CEIL_F32 = 29
V_RNDNE_F32 = 30
V_FLOOR_F32 = 31
V_EXP_F32 = 32
V_LOG_F32 = 33
V_RCP_F32 = 34
V_RCP_IFLAG_F32 = 35
V_RSQ_F32 = 36
V_RCP_F64 = 37
V_RSQ_F64 = 38
V_SQRT_F32 = 39
V_SQRT_F64 = 40
V_SIN_F32 = 41
V_COS_F32 = 42
V_NOT_B32 = 43
V_BFREV_B32 = 44
V_FFBH_U32 = 45
V_FFBL_B32 = 46
V_FFBH_I32 = 47
V_FREXP_EXP_I32_F64 = 48
V_FREXP_MANT_F64 = 49
V_FRACT_F64 = 50
V_FREXP_EXP_I32_F32 = 51
V_FREXP_MANT_F32 = 52
V_CLREXCP = 53
V_MOV_B64 = 56
V_CVT_F16_U16 = 57
V_CVT_F16_I16 = 58
V_CVT_U16_F16 = 59
V_CVT_I16_F16 = 60
V_RCP_F16 = 61
V_SQRT_F16 = 62
V_RSQ_F16 = 63
V_LOG_F16 = 64
V_EXP_F16 = 65
V_FREXP_MANT_F16 = 66
V_FREXP_EXP_I16_F16 = 67
V_FLOOR_F16 = 68
V_CEIL_F16 = 69
V_TRUNC_F16 = 70
V_RNDNE_F16 = 71
V_FRACT_F16 = 72
V_SIN_F16 = 73
V_COS_F16 = 74
V_CVT_NORM_I16_F16 = 77
V_CVT_NORM_U16_F16 = 78
V_SAT_PK_U8_I16 = 79
V_SWAP_B32 = 81
V_ACCVGPR_MOV_B32 = 82
V_CVT_F32_FP8 = 84
V_CVT_F32_BF8 = 85
V_CVT_PK_F32_FP8 = 86
V_CVT_PK_F32_BF8 = 87
V_PRNG_B32 = 88
V_PERMLANE16_SWAP_B32 = 89
V_PERMLANE32_SWAP_B32 = 90
V_CVT_F32_BF16 = 91
CDNA4 = 600
class VOP2Op(IntEnum):
V_CNDMASK_B32 = 0
V_ADD_F32 = 1
V_SUB_F32 = 2
V_SUBREV_F32 = 3
V_FMAC_F64 = 4
V_MUL_F32 = 5
V_MUL_I32_I24 = 6
V_MUL_HI_I32_I24 = 7
V_MUL_U32_U24 = 8
V_MUL_HI_U32_U24 = 9
V_MIN_F32 = 10
V_MAX_F32 = 11
V_MIN_I32 = 12
V_MAX_I32 = 13
V_MIN_U32 = 14
V_MAX_U32 = 15
V_LSHRREV_B32 = 16
V_ASHRREV_I32 = 17
V_LSHLREV_B32 = 18
V_AND_B32 = 19
V_OR_B32 = 20
V_XOR_B32 = 21
V_DOT2C_F32_BF16 = 22
V_FMAMK_F32 = 23
V_FMAAK_F32 = 24
V_ADD_CO_U32 = 25
V_SUB_CO_U32 = 26
V_SUBREV_CO_U32 = 27
V_ADDC_CO_U32 = 28
V_SUBB_CO_U32 = 29
V_SUBBREV_CO_U32 = 30
V_ADD_F16 = 31
V_SUB_F16 = 32
V_SUBREV_F16 = 33
V_MUL_F16 = 34
V_MAC_F16 = 35
V_MADMK_F16 = 36
V_MADAK_F16 = 37
V_ADD_U16 = 38
V_SUB_U16 = 39
V_SUBREV_U16 = 40
V_MUL_LO_U16 = 41
V_LSHLREV_B16 = 42
V_LSHRREV_B16 = 43
V_ASHRREV_I16 = 44
V_MAX_F16 = 45
V_MIN_F16 = 46
V_MAX_U16 = 47
V_MAX_I16 = 48
V_MIN_U16 = 49
V_MIN_I16 = 50
V_LDEXP_F16 = 51
V_ADD_U32 = 52
V_SUB_U32 = 53
V_SUBREV_U32 = 54
V_DOT2C_F32_F16 = 55
V_DOT2C_I32_I16 = 56
V_DOT4C_I32_I8 = 57
V_DOT8C_I32_I4 = 58
V_FMAC_F32 = 59
V_PK_FMAC_F16 = 60
V_XNOR_B32 = 61
CDNA4 = 600
class VOP3AOp(IntEnum):
V_CMP_CLASS_F32 = 16
V_CMPX_CLASS_F32 = 17
V_CMP_CLASS_F64 = 18
V_CMPX_CLASS_F64 = 19
V_CMP_CLASS_F16 = 20
V_CMPX_CLASS_F16 = 21
V_CMP_F_F16 = 32
V_CMP_LT_F16 = 33
V_CMP_EQ_F16 = 34
V_CMP_LE_F16 = 35
V_CMP_GT_F16 = 36
V_CMP_LG_F16 = 37
V_CMP_GE_F16 = 38
V_CMP_O_F16 = 39
V_CMP_U_F16 = 40
V_CMP_NGE_F16 = 41
V_CMP_NLG_F16 = 42
V_CMP_NGT_F16 = 43
V_CMP_NLE_F16 = 44
V_CMP_NEQ_F16 = 45
V_CMP_NLT_F16 = 46
V_CMP_TRU_F16 = 47
V_CMPX_F_F16 = 48
V_CMPX_LT_F16 = 49
V_CMPX_EQ_F16 = 50
V_CMPX_LE_F16 = 51
V_CMPX_GT_F16 = 52
V_CMPX_LG_F16 = 53
V_CMPX_GE_F16 = 54
V_CMPX_O_F16 = 55
V_CMPX_U_F16 = 56
V_CMPX_NGE_F16 = 57
V_CMPX_NLG_F16 = 58
V_CMPX_NGT_F16 = 59
V_CMPX_NLE_F16 = 60
V_CMPX_NEQ_F16 = 61
V_CMPX_NLT_F16 = 62
V_CMPX_TRU_F16 = 63
V_CMP_F_F32 = 64
V_CMP_LT_F32 = 65
V_CMP_EQ_F32 = 66
V_CMP_LE_F32 = 67
V_CMP_GT_F32 = 68
V_CMP_LG_F32 = 69
V_CMP_GE_F32 = 70
V_CMP_O_F32 = 71
V_CMP_U_F32 = 72
V_CMP_NGE_F32 = 73
V_CMP_NLG_F32 = 74
V_CMP_NGT_F32 = 75
V_CMP_NLE_F32 = 76
V_CMP_NEQ_F32 = 77
V_CMP_NLT_F32 = 78
V_CMP_TRU_F32 = 79
V_CMPX_F_F32 = 80
V_CMPX_LT_F32 = 81
V_CMPX_EQ_F32 = 82
V_CMPX_LE_F32 = 83
V_CMPX_GT_F32 = 84
V_CMPX_LG_F32 = 85
V_CMPX_GE_F32 = 86
V_CMPX_O_F32 = 87
V_CMPX_U_F32 = 88
V_CMPX_NGE_F32 = 89
V_CMPX_NLG_F32 = 90
V_CMPX_NGT_F32 = 91
V_CMPX_NLE_F32 = 92
V_CMPX_NEQ_F32 = 93
V_CMPX_NLT_F32 = 94
V_CMPX_TRU_F32 = 95
V_CMP_F_F64 = 96
V_CMP_LT_F64 = 97
V_CMP_EQ_F64 = 98
V_CMP_LE_F64 = 99
V_CMP_GT_F64 = 100
V_CMP_LG_F64 = 101
V_CMP_GE_F64 = 102
V_CMP_O_F64 = 103
V_CMP_U_F64 = 104
V_CMP_NGE_F64 = 105
V_CMP_NLG_F64 = 106
V_CMP_NGT_F64 = 107
V_CMP_NLE_F64 = 108
V_CMP_NEQ_F64 = 109
V_CMP_NLT_F64 = 110
V_CMP_TRU_F64 = 111
V_CMPX_F_F64 = 112
V_CMPX_LT_F64 = 113
V_CMPX_EQ_F64 = 114
V_CMPX_LE_F64 = 115
V_CMPX_GT_F64 = 116
V_CMPX_LG_F64 = 117
V_CMPX_GE_F64 = 118
V_CMPX_O_F64 = 119
V_CMPX_U_F64 = 120
V_CMPX_NGE_F64 = 121
V_CMPX_NLG_F64 = 122
V_CMPX_NGT_F64 = 123
V_CMPX_NLE_F64 = 124
V_CMPX_NEQ_F64 = 125
V_CMPX_NLT_F64 = 126
V_CMPX_TRU_F64 = 127
V_CMP_F_I16 = 160
V_CMP_LT_I16 = 161
V_CMP_EQ_I16 = 162
V_CMP_LE_I16 = 163
V_CMP_GT_I16 = 164
V_CMP_NE_I16 = 165
V_CMP_GE_I16 = 166
V_CMP_T_I16 = 167
V_CMP_F_U16 = 168
V_CMP_LT_U16 = 169
V_CMP_EQ_U16 = 170
V_CMP_LE_U16 = 171
V_CMP_GT_U16 = 172
V_CMP_NE_U16 = 173
V_CMP_GE_U16 = 174
V_CMP_T_U16 = 175
V_CMPX_F_I16 = 176
V_CMPX_LT_I16 = 177
V_CMPX_EQ_I16 = 178
V_CMPX_LE_I16 = 179
V_CMPX_GT_I16 = 180
V_CMPX_NE_I16 = 181
V_CMPX_GE_I16 = 182
V_CMPX_T_I16 = 183
V_CMPX_F_U16 = 184
V_CMPX_LT_U16 = 185
V_CMPX_EQ_U16 = 186
V_CMPX_LE_U16 = 187
V_CMPX_GT_U16 = 188
V_CMPX_NE_U16 = 189
V_CMPX_GE_U16 = 190
V_CMPX_T_U16 = 191
V_CMP_F_I32 = 192
V_CMP_LT_I32 = 193
V_CMP_EQ_I32 = 194
V_CMP_LE_I32 = 195
V_CMP_GT_I32 = 196
V_CMP_NE_I32 = 197
V_CMP_GE_I32 = 198
V_CMP_T_I32 = 199
V_CMP_F_U32 = 200
V_CMP_LT_U32 = 201
V_CMP_EQ_U32 = 202
V_CMP_LE_U32 = 203
V_CMP_GT_U32 = 204
V_CMP_NE_U32 = 205
V_CMP_GE_U32 = 206
V_CMP_T_U32 = 207
V_CMPX_F_I32 = 208
V_CMPX_LT_I32 = 209
V_CMPX_EQ_I32 = 210
V_CMPX_LE_I32 = 211
V_CMPX_GT_I32 = 212
V_CMPX_NE_I32 = 213
V_CMPX_GE_I32 = 214
V_CMPX_T_I32 = 215
V_CMPX_F_U32 = 216
V_CMPX_LT_U32 = 217
V_CMPX_EQ_U32 = 218
V_CMPX_LE_U32 = 219
V_CMPX_GT_U32 = 220
V_CMPX_NE_U32 = 221
V_CMPX_GE_U32 = 222
V_CMPX_T_U32 = 223
V_CMP_F_I64 = 224
V_CMP_LT_I64 = 225
V_CMP_EQ_I64 = 226
V_CMP_LE_I64 = 227
V_CMP_GT_I64 = 228
V_CMP_NE_I64 = 229
V_CMP_GE_I64 = 230
V_CMP_T_I64 = 231
V_CMP_F_U64 = 232
V_CMP_LT_U64 = 233
V_CMP_EQ_U64 = 234
V_CMP_LE_U64 = 235
V_CMP_GT_U64 = 236
V_CMP_NE_U64 = 237
V_CMP_GE_U64 = 238
V_CMP_T_U64 = 239
V_CMPX_F_I64 = 240
V_CMPX_LT_I64 = 241
V_CMPX_EQ_I64 = 242
V_CMPX_LE_I64 = 243
V_CMPX_GT_I64 = 244
V_CMPX_NE_I64 = 245
V_CMPX_GE_I64 = 246
V_CMPX_T_I64 = 247
V_CMPX_F_U64 = 248
V_CMPX_LT_U64 = 249
V_CMPX_EQ_U64 = 250
V_CMPX_LE_U64 = 251
V_CMPX_GT_U64 = 252
V_CMPX_NE_U64 = 253
V_CMPX_GE_U64 = 254
V_CMPX_T_U64 = 255
V_CNDMASK_B32 = 256
V_ADD_F32 = 257
V_SUB_F32 = 258
V_SUBREV_F32 = 259
V_FMAC_F64 = 260
V_MUL_F32 = 261
V_MUL_I32_I24 = 262
V_MUL_HI_I32_I24 = 263
V_MUL_U32_U24 = 264
V_MUL_HI_U32_U24 = 265
V_MIN_F32 = 266
V_MAX_F32 = 267
V_MIN_I32 = 268
V_MAX_I32 = 269
V_MIN_U32 = 270
V_MAX_U32 = 271
V_LSHRREV_B32 = 272
V_ASHRREV_I32 = 273
V_LSHLREV_B32 = 274
V_AND_B32 = 275
V_OR_B32 = 276
V_XOR_B32 = 277
V_DOT2C_F32_BF16 = 278
V_ADD_F16 = 287
V_SUB_F16 = 288
V_SUBREV_F16 = 289
V_MUL_F16 = 290
V_MAC_F16 = 291
V_ADD_U16 = 294
V_SUB_U16 = 295
V_SUBREV_U16 = 296
V_MUL_LO_U16 = 297
V_LSHLREV_B16 = 298
V_LSHRREV_B16 = 299
V_ASHRREV_I16 = 300
V_MAX_F16 = 301
V_MIN_F16 = 302
V_MAX_U16 = 303
V_MAX_I16 = 304
V_MIN_U16 = 305
V_MIN_I16 = 306
V_LDEXP_F16 = 307
V_ADD_U32 = 308
V_SUB_U32 = 309
V_SUBREV_U32 = 310
V_DOT2C_F32_F16 = 311
V_DOT2C_I32_I16 = 312
V_DOT4C_I32_I8 = 313
V_DOT8C_I32_I4 = 314
V_FMAC_F32 = 315
V_PK_FMAC_F16 = 316
V_XNOR_B32 = 317
V_NOP = 384
V_MOV_B32 = 385
V_READFIRSTLANE_B32 = 386
V_CVT_I32_F64 = 387
V_CVT_F64_I32 = 388
V_CVT_F32_I32 = 389
V_CVT_F32_U32 = 390
V_CVT_U32_F32 = 391
V_CVT_I32_F32 = 392
V_CVT_F16_F32 = 394
V_CVT_F32_F16 = 395
V_CVT_RPI_I32_F32 = 396
V_CVT_FLR_I32_F32 = 397
V_CVT_OFF_F32_I4 = 398
V_CVT_F32_F64 = 399
V_CVT_F64_F32 = 400
V_CVT_F32_UBYTE0 = 401
V_CVT_F32_UBYTE1 = 402
V_CVT_F32_UBYTE2 = 403
V_CVT_F32_UBYTE3 = 404
V_CVT_U32_F64 = 405
V_CVT_F64_U32 = 406
V_TRUNC_F64 = 407
V_CEIL_F64 = 408
V_RNDNE_F64 = 409
V_FLOOR_F64 = 410
V_FRACT_F32 = 411
V_TRUNC_F32 = 412
V_CEIL_F32 = 413
V_RNDNE_F32 = 414
V_FLOOR_F32 = 415
V_EXP_F32 = 416
V_LOG_F32 = 417
V_RCP_F32 = 418
V_RCP_IFLAG_F32 = 419
V_RSQ_F32 = 420
V_RCP_F64 = 421
V_RSQ_F64 = 422
V_SQRT_F32 = 423
V_SQRT_F64 = 424
V_SIN_F32 = 425
V_COS_F32 = 426
V_NOT_B32 = 427
V_BFREV_B32 = 428
V_FFBH_U32 = 429
V_FFBL_B32 = 430
V_FFBH_I32 = 431
V_FREXP_EXP_I32_F64 = 432
V_FREXP_MANT_F64 = 433
V_FRACT_F64 = 434
V_FREXP_EXP_I32_F32 = 435
V_FREXP_MANT_F32 = 436
V_CLREXCP = 437
V_MOV_B64 = 440
V_CVT_F16_U16 = 441
V_CVT_F16_I16 = 442
V_CVT_U16_F16 = 443
V_CVT_I16_F16 = 444
V_RCP_F16 = 445
V_SQRT_F16 = 446
V_RSQ_F16 = 447
V_LOG_F16 = 448
V_EXP_F16 = 449
V_MAD_I32_I24 = 450
V_MAD_U32_U24 = 451
V_CUBEID_F32 = 452
V_CUBESC_F32 = 453
V_CUBETC_F32 = 454
V_CUBEMA_F32 = 455
V_BFE_U32 = 456
V_BFE_I32 = 457
V_BFI_B32 = 458
V_FMA_F32 = 459
V_FMA_F64 = 460
V_LERP_U8 = 461
V_ALIGNBIT_B32 = 462
V_ALIGNBYTE_B32 = 463
V_MIN3_F32 = 464
V_MIN3_I32 = 465
V_MIN3_U32 = 466
V_MAX3_F32 = 467
V_MAX3_I32 = 468
V_MAX3_U32 = 469
V_MED3_F32 = 470
V_MED3_I32 = 471
V_MED3_U32 = 472
V_SAD_U8 = 473
V_SAD_HI_U8 = 474
V_SAD_U16 = 475
V_SAD_U32 = 476
V_CVT_PK_U8_F32 = 477
V_DIV_FIXUP_F32 = 478
V_DIV_FIXUP_F64 = 479
V_DIV_FMAS_F32 = 482
V_DIV_FMAS_F64 = 483
V_MSAD_U8 = 484
V_QSAD_PK_U16_U8 = 485
V_MQSAD_PK_U16_U8 = 486
V_MQSAD_U32_U8 = 487
V_MAD_LEGACY_F16 = 490
V_MAD_LEGACY_U16 = 491
V_MAD_LEGACY_I16 = 492
V_PERM_B32 = 493
V_FMA_LEGACY_F16 = 494
V_DIV_FIXUP_LEGACY_F16 = 495
V_CVT_PKACCUM_U8_F32 = 496
V_MAD_U32_U16 = 497
V_MAD_I32_I16 = 498
V_XAD_U32 = 499
V_MIN3_F16 = 500
V_MIN3_I16 = 501
V_MIN3_U16 = 502
V_MAX3_F16 = 503
V_MAX3_I16 = 504
V_MAX3_U16 = 505
V_MED3_F16 = 506
V_MED3_I16 = 507
V_MED3_U16 = 508
V_LSHL_ADD_U32 = 509
V_ADD_LSHL_U32 = 510
V_ADD3_U32 = 511
V_LSHL_OR_B32 = 512
V_AND_OR_B32 = 513
V_OR3_B32 = 514
V_MAD_F16 = 515
V_MAD_U16 = 516
V_MAD_I16 = 517
V_FMA_F16 = 518
V_DIV_FIXUP_F16 = 519
V_LSHL_ADD_U64 = 520
V_BITOP3_B16 = 563
V_BITOP3_B32 = 564
V_CVT_SCALEF32_PK_FP8_F32 = 565
V_CVT_SCALEF32_PK_BF8_F32 = 566
V_CVT_SCALEF32_SR_FP8_F32 = 567
V_CVT_SCALEF32_SR_BF8_F32 = 568
V_CVT_SCALEF32_PK_F32_FP8 = 569
V_CVT_SCALEF32_PK_F32_BF8 = 570
V_CVT_SCALEF32_F32_FP8 = 571
V_CVT_SCALEF32_F32_BF8 = 572
V_CVT_SCALEF32_PK_FP4_F32 = 573
V_CVT_SCALEF32_SR_PK_FP4_F32 = 574
V_CVT_SCALEF32_PK_F32_FP4 = 575
V_CVT_SCALEF32_PK_FP8_F16 = 576
V_CVT_SCALEF32_PK_BF8_F16 = 577
V_CVT_SCALEF32_SR_FP8_F16 = 578
V_CVT_SCALEF32_SR_BF8_F16 = 579
V_CVT_SCALEF32_PK_FP8_BF16 = 580
V_CVT_SCALEF32_PK_BF8_BF16 = 581
V_CVT_SCALEF32_SR_FP8_BF16 = 582
V_CVT_SCALEF32_SR_BF8_BF16 = 583
V_CVT_SCALEF32_PK_F16_FP8 = 584
V_CVT_SCALEF32_PK_F16_BF8 = 585
V_CVT_SCALEF32_F16_FP8 = 586
V_CVT_SCALEF32_F16_BF8 = 587
V_CVT_SCALEF32_PK_FP4_F16 = 588
V_CVT_SCALEF32_PK_FP4_BF16 = 589
V_CVT_SCALEF32_SR_PK_FP4_F16 = 590
V_CVT_SCALEF32_SR_PK_FP4_BF16 = 591
V_CVT_SCALEF32_PK_F16_FP4 = 592
V_CVT_SCALEF32_PK_BF16_FP4 = 593
V_CVT_SCALEF32_2XPK16_FP6_F32 = 594
V_CVT_SCALEF32_2XPK16_BF6_F32 = 595
V_CVT_SCALEF32_SR_PK32_FP6_F32 = 596
V_CVT_SCALEF32_SR_PK32_BF6_F32 = 597
V_CVT_SCALEF32_PK32_F32_FP6 = 598
V_CVT_SCALEF32_PK32_F32_BF6 = 599
CDNA4 = 600
V_CVT_SCALEF32_PK32_FP6_BF16 = 601
V_CVT_SCALEF32_PK32_BF6_F16 = 602
V_CVT_SCALEF32_PK32_BF6_BF16 = 603
V_CVT_SCALEF32_SR_PK32_FP6_F16 = 604
V_CVT_SCALEF32_SR_PK32_FP6_BF16 = 605
V_CVT_SCALEF32_SR_PK32_BF6_F16 = 606
V_CVT_SCALEF32_SR_PK32_BF6_BF16 = 607
V_CVT_SCALEF32_PK32_F16_FP6 = 608
V_CVT_SCALEF32_PK32_BF16_FP6 = 609
V_CVT_SCALEF32_PK32_F16_BF6 = 610
V_CVT_SCALEF32_PK32_BF16_BF6 = 611
V_ASHR_PK_I8_I32 = 613
V_ASHR_PK_U8_I32 = 614
V_CVT_PK_F16_F32 = 615
V_CVT_PK_BF16_F32 = 616
V_CVT_SCALEF32_PK_BF16_FP8 = 617
V_CVT_SCALEF32_PK_BF16_BF8 = 618
V_ADD_F64 = 640
V_MUL_F64 = 641
V_MIN_F64 = 642
V_MAX_F64 = 643
V_LDEXP_F64 = 644
V_MUL_LO_U32 = 645
V_MUL_HI_U32 = 646
V_MUL_HI_I32 = 647
V_LDEXP_F32 = 648
V_READLANE_B32 = 649
V_WRITELANE_B32 = 650
V_BCNT_U32_B32 = 651
V_MBCNT_LO_U32_B32 = 652
V_MBCNT_HI_U32_B32 = 653
V_LSHLREV_B64 = 655
V_LSHRREV_B64 = 656
V_ASHRREV_I64 = 657
V_TRIG_PREOP_F64 = 658
V_BFM_B32 = 659
V_CVT_PKNORM_I16_F32 = 660
V_CVT_PKNORM_U16_F32 = 661
V_CVT_PKRTZ_F16_F32 = 662
V_CVT_PK_U16_U32 = 663
V_CVT_PK_I16_I32 = 664
V_CVT_PKNORM_I16_F16 = 665
V_CVT_PKNORM_U16_F16 = 666
V_ADD_I32 = 668
V_SUB_I32 = 669
V_ADD_I16 = 670
V_SUB_I16 = 671
V_PACK_B32_F16 = 672
V_MUL_LEGACY_F32 = 673
V_CVT_PK_FP8_F32 = 674
V_CVT_PK_BF8_F32 = 675
V_CVT_SR_FP8_F32 = 676
V_CVT_SR_BF8_F32 = 677
V_CVT_SR_F16_F32 = 678
V_CVT_SR_BF16_F32 = 679
V_MINIMUM3_F32 = 680
V_MAXIMUM3_F32 = 681
class VOP3BOp(IntEnum):
V_ADD_CO_U32 = 281
V_SUB_CO_U32 = 282
V_SUBREV_CO_U32 = 283
V_ADDC_CO_U32 = 284
V_SUBB_CO_U32 = 285
V_SUBBREV_CO_U32 = 286
V_DIV_SCALE_F32 = 480
V_DIV_SCALE_F64 = 481
V_MAD_U64_U32 = 488
V_MAD_I64_I32 = 489
CDNA4 = 600
class VOP3POp(IntEnum):
V_PK_MAD_I16 = 0
V_PK_MUL_LO_U16 = 1
V_PK_ADD_I16 = 2
V_PK_SUB_I16 = 3
V_PK_LSHLREV_B16 = 4
V_PK_LSHRREV_B16 = 5
V_PK_ASHRREV_I16 = 6
V_PK_MAX_I16 = 7
V_PK_MIN_I16 = 8
V_PK_MAD_U16 = 9
V_PK_ADD_U16 = 10
V_PK_SUB_U16 = 11
V_PK_MAX_U16 = 12
V_PK_MIN_U16 = 13
V_PK_FMA_F16 = 14
V_PK_ADD_F16 = 15
V_PK_MUL_F16 = 16
V_PK_MIN_F16 = 17
V_PK_MAX_F16 = 18
V_DOT2_F32_BF16 = 26
V_PK_MINIMUM3_F16 = 27
V_PK_MAXIMUM3_F16 = 28
V_MAD_MIX_F32 = 32
V_MAD_MIXLO_F16 = 33
V_MAD_MIXHI_F16 = 34
V_DOT2_F32_F16 = 35
V_DOT2_I32_I16 = 38
V_DOT2_U32_U16 = 39
V_DOT4_I32_I8 = 40
V_DOT4_U32_U8 = 41
V_DOT8_I32_I4 = 42
V_DOT8_U32_U4 = 43
V_MFMA_F32_16X16X128_F8F6F4 = 45
V_MFMA_F32_32X32X64_F8F6F4 = 46
V_PK_FMA_F32 = 48
V_PK_MUL_F32 = 49
V_PK_ADD_F32 = 50
V_PK_MOV_B32 = 51
V_MFMA_F32_16X16X32_BF16 = 53
V_MFMA_I32_16X16X64_I8 = 54
V_MFMA_F32_32X32X16_BF16 = 55
V_MFMA_I32_32X32X32_I8 = 56
V_SMFMAC_F32_16X16X64_BF16 = 57
V_SMFMAC_I32_16X16X128_I8 = 58
V_SMFMAC_F32_16X16X128_BF8_BF8 = 59
V_SMFMAC_F32_16X16X128_BF8_FP8 = 60
V_SMFMAC_F32_16X16X128_FP8_BF8 = 61
V_MFMA_F32_16X16X8_XF32 = 62
V_MFMA_F32_32X32X4_XF32 = 63
V_MFMA_F32_32X32X1_2B_F32 = 64
V_MFMA_F32_16X16X1_4B_F32 = 65
V_MFMA_F32_4X4X1_16B_F32 = 66
V_SMFMAC_F32_16X16X128_FP8_FP8 = 67
V_MFMA_F32_32X32X2_F32 = 68
V_MFMA_F32_16X16X4_F32 = 69
V_SMFMAC_F32_32X32X32_BF16 = 70
V_SMFMAC_I32_32X32X64_I8 = 71
V_MFMA_F32_32X32X4_2B_F16 = 72
V_MFMA_F32_16X16X4_4B_F16 = 73
V_MFMA_F32_4X4X4_16B_F16 = 74
V_SMFMAC_F32_32X32X64_BF8_BF8 = 75
V_MFMA_F32_32X32X8_F16 = 76
V_MFMA_F32_16X16X16_F16 = 77
V_SMFMAC_F32_32X32X64_BF8_FP8 = 78
V_SMFMAC_F32_32X32X64_FP8_BF8 = 79
V_MFMA_I32_32X32X4_2B_I8 = 80
V_MFMA_I32_16X16X4_4B_I8 = 81
V_MFMA_I32_4X4X4_16B_I8 = 82
V_SMFMAC_F32_32X32X64_FP8_FP8 = 83
V_MFMA_F32_16X16X32_F16 = 84
V_MFMA_F32_32X32X16_F16 = 85
V_MFMA_I32_32X32X16_I8 = 86
V_MFMA_I32_16X16X32_I8 = 87
V_ACCVGPR_READ = 88
V_ACCVGPR_WRITE = 89
V_SMFMAC_F32_16X16X64_F16 = 90
V_SMFMAC_F32_32X32X32_F16 = 91
V_MFMA_F32_32X32X4_2B_BF16 = 93
V_MFMA_F32_16X16X4_4B_BF16 = 94
V_MFMA_F32_4X4X4_16B_BF16 = 95
V_MFMA_F32_32X32X8_BF16 = 96
V_MFMA_F32_16X16X16_BF16 = 97
V_SMFMAC_F32_16X16X32_F16 = 98
V_SMFMAC_F32_32X32X16_F16 = 100
V_SMFMAC_F32_16X16X32_BF16 = 102
V_SMFMAC_F32_32X32X16_BF16 = 104
V_SMFMAC_I32_16X16X64_I8 = 106
V_SMFMAC_I32_32X32X32_I8 = 108
V_MFMA_F64_16X16X4_F64 = 110
V_MFMA_F64_4X4X4_4B_F64 = 111
V_MFMA_F32_16X16X32_BF8_BF8 = 112
V_MFMA_F32_16X16X32_BF8_FP8 = 113
V_MFMA_F32_16X16X32_FP8_BF8 = 114
V_MFMA_F32_16X16X32_FP8_FP8 = 115
V_MFMA_F32_32X32X16_BF8_BF8 = 116
V_MFMA_F32_32X32X16_BF8_FP8 = 117
V_MFMA_F32_32X32X16_FP8_BF8 = 118
V_MFMA_F32_32X32X16_FP8_FP8 = 119
V_SMFMAC_F32_16X16X64_BF8_BF8 = 120
V_SMFMAC_F32_16X16X64_BF8_FP8 = 121
V_SMFMAC_F32_16X16X64_FP8_BF8 = 122
V_SMFMAC_F32_16X16X64_FP8_FP8 = 123
V_SMFMAC_F32_32X32X32_BF8_BF8 = 124
V_SMFMAC_F32_32X32X32_BF8_FP8 = 125
V_SMFMAC_F32_32X32X32_FP8_BF8 = 126
V_SMFMAC_F32_32X32X32_FP8_FP8 = 127
CDNA4 = 600
class VOPCOp(IntEnum):
V_CMP_CLASS_F32 = 16
V_CMPX_CLASS_F32 = 17
V_CMP_CLASS_F64 = 18
V_CMPX_CLASS_F64 = 19
V_CMP_CLASS_F16 = 20
V_CMPX_CLASS_F16 = 21
V_CMP_F_F16 = 32
V_CMP_LT_F16 = 33
V_CMP_EQ_F16 = 34
V_CMP_LE_F16 = 35
V_CMP_GT_F16 = 36
V_CMP_LG_F16 = 37
V_CMP_GE_F16 = 38
V_CMP_O_F16 = 39
V_CMP_U_F16 = 40
V_CMP_NGE_F16 = 41
V_CMP_NLG_F16 = 42
V_CMP_NGT_F16 = 43
V_CMP_NLE_F16 = 44
V_CMP_NEQ_F16 = 45
V_CMP_NLT_F16 = 46
V_CMP_TRU_F16 = 47
V_CMPX_F_F16 = 48
V_CMPX_LT_F16 = 49
V_CMPX_EQ_F16 = 50
V_CMPX_LE_F16 = 51
V_CMPX_GT_F16 = 52
V_CMPX_LG_F16 = 53
V_CMPX_GE_F16 = 54
V_CMPX_O_F16 = 55
V_CMPX_U_F16 = 56
V_CMPX_NGE_F16 = 57
V_CMPX_NLG_F16 = 58
V_CMPX_NGT_F16 = 59
V_CMPX_NLE_F16 = 60
V_CMPX_NEQ_F16 = 61
V_CMPX_NLT_F16 = 62
V_CMPX_TRU_F16 = 63
V_CMP_F_F32 = 64
V_CMP_LT_F32 = 65
V_CMP_EQ_F32 = 66
V_CMP_LE_F32 = 67
V_CMP_GT_F32 = 68
V_CMP_LG_F32 = 69
V_CMP_GE_F32 = 70
V_CMP_O_F32 = 71
V_CMP_U_F32 = 72
V_CMP_NGE_F32 = 73
V_CMP_NLG_F32 = 74
V_CMP_NGT_F32 = 75
V_CMP_NLE_F32 = 76
V_CMP_NEQ_F32 = 77
V_CMP_NLT_F32 = 78
V_CMP_TRU_F32 = 79
V_CMPX_F_F32 = 80
V_CMPX_LT_F32 = 81
V_CMPX_EQ_F32 = 82
V_CMPX_LE_F32 = 83
V_CMPX_GT_F32 = 84
V_CMPX_LG_F32 = 85
V_CMPX_GE_F32 = 86
V_CMPX_O_F32 = 87
V_CMPX_U_F32 = 88
V_CMPX_NGE_F32 = 89
V_CMPX_NLG_F32 = 90
V_CMPX_NGT_F32 = 91
V_CMPX_NLE_F32 = 92
V_CMPX_NEQ_F32 = 93
V_CMPX_NLT_F32 = 94
V_CMPX_TRU_F32 = 95
V_CMP_F_F64 = 96
V_CMP_LT_F64 = 97
V_CMP_EQ_F64 = 98
V_CMP_LE_F64 = 99
V_CMP_GT_F64 = 100
V_CMP_LG_F64 = 101
V_CMP_GE_F64 = 102
V_CMP_O_F64 = 103
V_CMP_U_F64 = 104
V_CMP_NGE_F64 = 105
V_CMP_NLG_F64 = 106
V_CMP_NGT_F64 = 107
V_CMP_NLE_F64 = 108
V_CMP_NEQ_F64 = 109
V_CMP_NLT_F64 = 110
V_CMP_TRU_F64 = 111
V_CMPX_F_F64 = 112
V_CMPX_LT_F64 = 113
V_CMPX_EQ_F64 = 114
V_CMPX_LE_F64 = 115
V_CMPX_GT_F64 = 116
V_CMPX_LG_F64 = 117
V_CMPX_GE_F64 = 118
V_CMPX_O_F64 = 119
V_CMPX_U_F64 = 120
V_CMPX_NGE_F64 = 121
V_CMPX_NLG_F64 = 122
V_CMPX_NGT_F64 = 123
V_CMPX_NLE_F64 = 124
V_CMPX_NEQ_F64 = 125
V_CMPX_NLT_F64 = 126
V_CMPX_TRU_F64 = 127
V_CMP_F_I16 = 160
V_CMP_LT_I16 = 161
V_CMP_EQ_I16 = 162
V_CMP_LE_I16 = 163
V_CMP_GT_I16 = 164
V_CMP_NE_I16 = 165
V_CMP_GE_I16 = 166
V_CMP_T_I16 = 167
V_CMP_F_U16 = 168
V_CMP_LT_U16 = 169
V_CMP_EQ_U16 = 170
V_CMP_LE_U16 = 171
V_CMP_GT_U16 = 172
V_CMP_NE_U16 = 173
V_CMP_GE_U16 = 174
V_CMP_T_U16 = 175
V_CMPX_F_I16 = 176
V_CMPX_LT_I16 = 177
V_CMPX_EQ_I16 = 178
V_CMPX_LE_I16 = 179
V_CMPX_GT_I16 = 180
V_CMPX_NE_I16 = 181
V_CMPX_GE_I16 = 182
V_CMPX_T_I16 = 183
V_CMPX_F_U16 = 184
V_CMPX_LT_U16 = 185
V_CMPX_EQ_U16 = 186
V_CMPX_LE_U16 = 187
V_CMPX_GT_U16 = 188
V_CMPX_NE_U16 = 189
V_CMPX_GE_U16 = 190
V_CMPX_T_U16 = 191
V_CMP_F_I32 = 192
V_CMP_LT_I32 = 193
V_CMP_EQ_I32 = 194
V_CMP_LE_I32 = 195
V_CMP_GT_I32 = 196
V_CMP_NE_I32 = 197
V_CMP_GE_I32 = 198
V_CMP_T_I32 = 199
V_CMP_F_U32 = 200
V_CMP_LT_U32 = 201
V_CMP_EQ_U32 = 202
V_CMP_LE_U32 = 203
V_CMP_GT_U32 = 204
V_CMP_NE_U32 = 205
V_CMP_GE_U32 = 206
V_CMP_T_U32 = 207
V_CMPX_F_I32 = 208
V_CMPX_LT_I32 = 209
V_CMPX_EQ_I32 = 210
V_CMPX_LE_I32 = 211
V_CMPX_GT_I32 = 212
V_CMPX_NE_I32 = 213
V_CMPX_GE_I32 = 214
V_CMPX_T_I32 = 215
V_CMPX_F_U32 = 216
V_CMPX_LT_U32 = 217
V_CMPX_EQ_U32 = 218
V_CMPX_LE_U32 = 219
V_CMPX_GT_U32 = 220
V_CMPX_NE_U32 = 221
V_CMPX_GE_U32 = 222
V_CMPX_T_U32 = 223
V_CMP_F_I64 = 224
V_CMP_LT_I64 = 225
V_CMP_EQ_I64 = 226
V_CMP_LE_I64 = 227
V_CMP_GT_I64 = 228
V_CMP_NE_I64 = 229
V_CMP_GE_I64 = 230
V_CMP_T_I64 = 231
V_CMP_F_U64 = 232
V_CMP_LT_U64 = 233
V_CMP_EQ_U64 = 234
V_CMP_LE_U64 = 235
V_CMP_GT_U64 = 236
V_CMP_NE_U64 = 237
V_CMP_GE_U64 = 238
V_CMP_T_U64 = 239
V_CMPX_F_I64 = 240
V_CMPX_LT_I64 = 241
V_CMPX_EQ_I64 = 242
V_CMPX_LE_I64 = 243
V_CMPX_GT_I64 = 244
V_CMPX_NE_I64 = 245
V_CMPX_GE_I64 = 246
V_CMPX_T_I64 = 247
V_CMPX_F_U64 = 248
V_CMPX_LT_U64 = 249
V_CMPX_EQ_U64 = 250
V_CMPX_LE_U64 = 251
V_CMPX_GT_U64 = 252
V_CMPX_NE_U64 = 253
V_CMPX_GE_U64 = 254
V_CMPX_T_U64 = 255
CDNA4 = 600
# instruction formats
class DPP(Inst64):
encoding = bits[31:26] == 0b110110
src1_sel = bits[58:56]
src1_sext = bits[59]
src1_neg = bits[60]
src1_abs = bits[61]
s1 = bits[63]
offset0 = bits[7:0]
offset1 = bits[15:8]
op = bits[24:17]
acc = bits[25]
addr:VGPRField = bits[39:32]
data0:VGPRField = bits[47:40]
data1:VGPRField = bits[55:48]
vdst:VGPRField = bits[63:56]
row_mask = bits[63:60]
class DS(Inst64):
encoding = bits[31:26] == 0b110110
op:Annotated[BitField, DSOp] = bits[24:17]
vdst:VGPRField = bits[63:56]
addr:VGPRField = bits[39:32]
data0:VGPRField = bits[47:40]
data1:VGPRField = bits[55:48]
offset0 = bits[7:0]
offset1 = bits[15:8]
gds = bits[16]
acc = bits[25]
class FLAT(Inst64):
encoding = bits[31:26] == 0b110111
op:Annotated[BitField, FLATOp] = bits[24:18]
vdst:VGPRField = bits[63:56]
addr:VGPRField = bits[39:32]
data:VGPRField = bits[47:40]
saddr:SSrc = bits[54:48]
offset:Imm = bits[12:0]
seg = bits[15:14]
lds = bits[13]
sc0 = bits[16]
nt = bits[17]
sc1 = bits[25]
acc = bits[55]
class MTBUF(Inst64):
encoding = bits[31:26] == 0b111010
op:Annotated[BitField, MTBUFOp] = bits[18:15]
vdata:VGPRField = bits[47:40]
vaddr:VGPRField = bits[39:32]
srsrc:SGPRField = bits[52:48]
soffset:SSrc = bits[63:56]
offset:Imm = bits[11:0]
offen = bits[12]
idxen = bits[13]
sc1 = bits[53]
nt = bits[54]
acc = bits[55]
sc0 = bits[14]
class MUBUF(Inst64):
encoding = bits[31:26] == 0b111000
op:Annotated[BitField, MUBUFOp] = bits[24:18]
vdata:VGPRField = bits[47:40]
vaddr:VGPRField = bits[39:32]
srsrc:SGPRField = bits[52:48]
soffset:SSrc = bits[63:56]
offset:Imm = bits[11:0]
offen = bits[12]
idxen = bits[13]
sc0 = bits[14]
sc1 = bits[15]
lds = bits[16]
nt = bits[17]
acc = bits[55]
class SDWA(Inst64):
src0:Src = bits[39:32]
dst_sel = bits[42:40]
dst_u = bits[44:43]
clmp = bits[45]
omod = bits[47:46]
src0_sel = bits[50:48]
src0_sext = bits[51]
src0_neg = bits[52]
src0_abs = bits[53]
s0 = bits[55]
src1_sel = bits[58:56]
src1_sext = bits[59]
src1_neg = bits[60]
src1_abs = bits[61]
s1 = bits[63]
sdst:SGPRField = bits[46:40]
sd = bits[47]
row_mask = bits[63:60]
class SDWAB(Inst64):
src0:Src = bits[39:32]
dst_sel = bits[42:40]
dst_u = bits[44:43]
clmp = bits[45]
omod = bits[47:46]
src0_sel = bits[50:48]
src0_sext = bits[51]
src0_neg = bits[52]
src0_abs = bits[53]
s0 = bits[55]
src1_sel = bits[58:56]
src1_sext = bits[59]
src1_neg = bits[60]
src1_abs = bits[61]
s1 = bits[63]
class SMEM(Inst64):
encoding = bits[31:26] == 0b110000
op:Annotated[BitField, SMEMOp] = bits[25:18]
sdata:SGPRField = bits[12:6]
sbase:SGPRField = bits[5:0]
soffset:SSrc = bits[63:57]
offset:Imm = bits[52:32]
glc = bits[14]
soe = bits[14]
nv = bits[15]
imm = bits[17]
class SOP1(Inst32):
encoding = bits[31:23] == 0b101111101
op:Annotated[BitField, SOP1Op] = bits[15:8]
sdst:SGPRField = bits[22:16]
ssrc0:SSrc = bits[7:0]
class SOP2(Inst32):
encoding = bits[31:30] == 0b10
op:Annotated[BitField, SOP2Op] = bits[29:23]
sdst:SGPRField = bits[22:16]
ssrc0:SSrc = bits[7:0]
ssrc1:SSrc = bits[15:8]
class SOPC(Inst32):
encoding = bits[31:23] == 0b101111110
op:Annotated[BitField, SOPCOp] = bits[22:16]
ssrc0:SSrc = bits[7:0]
ssrc1:SSrc = bits[15:8]
class SOPK(Inst32):
encoding = bits[31:28] == 0b1011
op:Annotated[BitField, SOPKOp] = bits[27:23]
sdst:SGPRField = bits[22:16]
simm16:SImm = bits[15:0]
class SOPP(Inst32):
encoding = bits[31:23] == 0b101111111
op:Annotated[BitField, SOPPOp] = bits[22:16]
simm16:SImm = bits[15:0]
class VOP1(Inst32):
encoding = bits[31:25] == 0b111111
op:Annotated[BitField, VOP1Op] = bits[16:9]
vdst:VGPRField = bits[24:17]
src0:Src = bits[8:0]
class VOP2(Inst32):
encoding = bits[31] == 0
op:Annotated[BitField, VOP2Op] = bits[30:25]
vdst:VGPRField = bits[24:17]
src0:Src = bits[8:0]
vsrc1:VGPRField = bits[16:9]
class VOP3A(Inst64):
encoding = bits[31:26] == 0b110100
vdst:VGPRField = bits[7:0]
abs = bits[10:8]
opsel = bits[14:11]
clmp = bits[15]
op:Annotated[BitField, VOP3AOp] = bits[25:16]
src0:Src = bits[40:32]
src1:Src = bits[49:41]
src2:Src = bits[58:50]
omod = bits[60:59]
neg = bits[63:61]
class VOP3B(Inst64):
encoding = bits[31:26] == 0b110100
vdst:VGPRField = bits[7:0]
sdst:SGPRField = bits[14:8]
clmp = bits[15]
op:Annotated[BitField, VOP3BOp] = bits[25:16]
src0:Src = bits[40:32]
src1:Src = bits[49:41]
src2:Src = bits[58:50]
omod = bits[60:59]
neg = bits[63:61]
class VOP3P(Inst64):
encoding = bits[31:23] == 0b110100111
_defaults = {'opsel_hi': 3, 'opsel_hi2': 1}
op:Annotated[BitField, VOP3POp] = bits[22:16]
vdst:VGPRField = bits[7:0]
src0:Src = bits[40:32]
src1:Src = bits[49:41]
src2:Src = bits[58:50]
neg = bits[63:61]
neg_hi = bits[10:8]
opsel = bits[13:11]
opsel_hi = bits[60:59]
clmp = bits[15]
opsel_hi2 = bits[14]
class VOPC(Inst32):
encoding = bits[31:25] == 0b111110
op:Annotated[BitField, VOPCOp] = bits[24:17]
src0:Src = bits[8:0]
vsrc1:VGPRField = bits[16:9]
# instruction helpers
ds_add_u32 = functools.partial(DS, DSOp.DS_ADD_U32)
ds_sub_u32 = functools.partial(DS, DSOp.DS_SUB_U32)
ds_rsub_u32 = functools.partial(DS, DSOp.DS_RSUB_U32)
ds_inc_u32 = functools.partial(DS, DSOp.DS_INC_U32)
ds_dec_u32 = functools.partial(DS, DSOp.DS_DEC_U32)
ds_min_i32 = functools.partial(DS, DSOp.DS_MIN_I32)
ds_max_i32 = functools.partial(DS, DSOp.DS_MAX_I32)
ds_min_u32 = functools.partial(DS, DSOp.DS_MIN_U32)
ds_max_u32 = functools.partial(DS, DSOp.DS_MAX_U32)
ds_and_b32 = functools.partial(DS, DSOp.DS_AND_B32)
ds_or_b32 = functools.partial(DS, DSOp.DS_OR_B32)
ds_xor_b32 = functools.partial(DS, DSOp.DS_XOR_B32)
ds_mskor_b32 = functools.partial(DS, DSOp.DS_MSKOR_B32)
ds_write_b32 = functools.partial(DS, DSOp.DS_WRITE_B32)
ds_write2_b32 = functools.partial(DS, DSOp.DS_WRITE2_B32)
ds_write2st64_b32 = functools.partial(DS, DSOp.DS_WRITE2ST64_B32)
ds_cmpst_b32 = functools.partial(DS, DSOp.DS_CMPST_B32)
ds_cmpst_f32 = functools.partial(DS, DSOp.DS_CMPST_F32)
ds_min_f32 = functools.partial(DS, DSOp.DS_MIN_F32)
ds_max_f32 = functools.partial(DS, DSOp.DS_MAX_F32)
ds_nop = functools.partial(DS, DSOp.DS_NOP)
ds_add_f32 = functools.partial(DS, DSOp.DS_ADD_F32)
ds_pk_add_f16 = functools.partial(DS, DSOp.DS_PK_ADD_F16)
ds_pk_add_bf16 = functools.partial(DS, DSOp.DS_PK_ADD_BF16)
ds_write_addtid_b32 = functools.partial(DS, DSOp.DS_WRITE_ADDTID_B32)
ds_write_b8 = functools.partial(DS, DSOp.DS_WRITE_B8)
ds_write_b16 = functools.partial(DS, DSOp.DS_WRITE_B16)
ds_add_rtn_u32 = functools.partial(DS, DSOp.DS_ADD_RTN_U32)
ds_sub_rtn_u32 = functools.partial(DS, DSOp.DS_SUB_RTN_U32)
ds_rsub_rtn_u32 = functools.partial(DS, DSOp.DS_RSUB_RTN_U32)
ds_inc_rtn_u32 = functools.partial(DS, DSOp.DS_INC_RTN_U32)
ds_dec_rtn_u32 = functools.partial(DS, DSOp.DS_DEC_RTN_U32)
ds_min_rtn_i32 = functools.partial(DS, DSOp.DS_MIN_RTN_I32)
ds_max_rtn_i32 = functools.partial(DS, DSOp.DS_MAX_RTN_I32)
ds_min_rtn_u32 = functools.partial(DS, DSOp.DS_MIN_RTN_U32)
ds_max_rtn_u32 = functools.partial(DS, DSOp.DS_MAX_RTN_U32)
ds_and_rtn_b32 = functools.partial(DS, DSOp.DS_AND_RTN_B32)
ds_or_rtn_b32 = functools.partial(DS, DSOp.DS_OR_RTN_B32)
ds_xor_rtn_b32 = functools.partial(DS, DSOp.DS_XOR_RTN_B32)
ds_mskor_rtn_b32 = functools.partial(DS, DSOp.DS_MSKOR_RTN_B32)
ds_wrxchg_rtn_b32 = functools.partial(DS, DSOp.DS_WRXCHG_RTN_B32)
ds_wrxchg2_rtn_b32 = functools.partial(DS, DSOp.DS_WRXCHG2_RTN_B32)
ds_wrxchg2st64_rtn_b32 = functools.partial(DS, DSOp.DS_WRXCHG2ST64_RTN_B32)
ds_cmpst_rtn_b32 = functools.partial(DS, DSOp.DS_CMPST_RTN_B32)
ds_cmpst_rtn_f32 = functools.partial(DS, DSOp.DS_CMPST_RTN_F32)
ds_min_rtn_f32 = functools.partial(DS, DSOp.DS_MIN_RTN_F32)
ds_max_rtn_f32 = functools.partial(DS, DSOp.DS_MAX_RTN_F32)
ds_wrap_rtn_b32 = functools.partial(DS, DSOp.DS_WRAP_RTN_B32)
ds_add_rtn_f32 = functools.partial(DS, DSOp.DS_ADD_RTN_F32)
ds_read_b32 = functools.partial(DS, DSOp.DS_READ_B32)
ds_read2_b32 = functools.partial(DS, DSOp.DS_READ2_B32)
ds_read2st64_b32 = functools.partial(DS, DSOp.DS_READ2ST64_B32)
ds_read_i8 = functools.partial(DS, DSOp.DS_READ_I8)
ds_read_u8 = functools.partial(DS, DSOp.DS_READ_U8)
ds_read_i16 = functools.partial(DS, DSOp.DS_READ_I16)
ds_read_u16 = functools.partial(DS, DSOp.DS_READ_U16)
ds_swizzle_b32 = functools.partial(DS, DSOp.DS_SWIZZLE_B32)
ds_permute_b32 = functools.partial(DS, DSOp.DS_PERMUTE_B32)
ds_bpermute_b32 = functools.partial(DS, DSOp.DS_BPERMUTE_B32)
ds_add_u64 = functools.partial(DS, DSOp.DS_ADD_U64)
ds_sub_u64 = functools.partial(DS, DSOp.DS_SUB_U64)
ds_rsub_u64 = functools.partial(DS, DSOp.DS_RSUB_U64)
ds_inc_u64 = functools.partial(DS, DSOp.DS_INC_U64)
ds_dec_u64 = functools.partial(DS, DSOp.DS_DEC_U64)
ds_min_i64 = functools.partial(DS, DSOp.DS_MIN_I64)
ds_max_i64 = functools.partial(DS, DSOp.DS_MAX_I64)
ds_min_u64 = functools.partial(DS, DSOp.DS_MIN_U64)
ds_max_u64 = functools.partial(DS, DSOp.DS_MAX_U64)
ds_and_b64 = functools.partial(DS, DSOp.DS_AND_B64)
ds_or_b64 = functools.partial(DS, DSOp.DS_OR_B64)
ds_xor_b64 = functools.partial(DS, DSOp.DS_XOR_B64)
ds_mskor_b64 = functools.partial(DS, DSOp.DS_MSKOR_B64)
ds_write_b64 = functools.partial(DS, DSOp.DS_WRITE_B64)
ds_write2_b64 = functools.partial(DS, DSOp.DS_WRITE2_B64)
ds_write2st64_b64 = functools.partial(DS, DSOp.DS_WRITE2ST64_B64)
ds_cmpst_b64 = functools.partial(DS, DSOp.DS_CMPST_B64)
ds_cmpst_f64 = functools.partial(DS, DSOp.DS_CMPST_F64)
ds_min_f64 = functools.partial(DS, DSOp.DS_MIN_F64)
ds_max_f64 = functools.partial(DS, DSOp.DS_MAX_F64)
ds_write_b8_d16_hi = functools.partial(DS, DSOp.DS_WRITE_B8_D16_HI)
ds_write_b16_d16_hi = functools.partial(DS, DSOp.DS_WRITE_B16_D16_HI)
ds_read_u8_d16 = functools.partial(DS, DSOp.DS_READ_U8_D16)
ds_read_u8_d16_hi = functools.partial(DS, DSOp.DS_READ_U8_D16_HI)
ds_read_i8_d16 = functools.partial(DS, DSOp.DS_READ_I8_D16)
ds_read_i8_d16_hi = functools.partial(DS, DSOp.DS_READ_I8_D16_HI)
ds_read_u16_d16 = functools.partial(DS, DSOp.DS_READ_U16_D16)
ds_read_u16_d16_hi = functools.partial(DS, DSOp.DS_READ_U16_D16_HI)
ds_add_f64 = functools.partial(DS, DSOp.DS_ADD_F64)
ds_add_rtn_u64 = functools.partial(DS, DSOp.DS_ADD_RTN_U64)
ds_sub_rtn_u64 = functools.partial(DS, DSOp.DS_SUB_RTN_U64)
ds_rsub_rtn_u64 = functools.partial(DS, DSOp.DS_RSUB_RTN_U64)
ds_inc_rtn_u64 = functools.partial(DS, DSOp.DS_INC_RTN_U64)
ds_dec_rtn_u64 = functools.partial(DS, DSOp.DS_DEC_RTN_U64)
ds_min_rtn_i64 = functools.partial(DS, DSOp.DS_MIN_RTN_I64)
ds_max_rtn_i64 = functools.partial(DS, DSOp.DS_MAX_RTN_I64)
ds_min_rtn_u64 = functools.partial(DS, DSOp.DS_MIN_RTN_U64)
ds_max_rtn_u64 = functools.partial(DS, DSOp.DS_MAX_RTN_U64)
ds_and_rtn_b64 = functools.partial(DS, DSOp.DS_AND_RTN_B64)
ds_or_rtn_b64 = functools.partial(DS, DSOp.DS_OR_RTN_B64)
ds_xor_rtn_b64 = functools.partial(DS, DSOp.DS_XOR_RTN_B64)
ds_mskor_rtn_b64 = functools.partial(DS, DSOp.DS_MSKOR_RTN_B64)
ds_wrxchg_rtn_b64 = functools.partial(DS, DSOp.DS_WRXCHG_RTN_B64)
ds_wrxchg2_rtn_b64 = functools.partial(DS, DSOp.DS_WRXCHG2_RTN_B64)
ds_wrxchg2st64_rtn_b64 = functools.partial(DS, DSOp.DS_WRXCHG2ST64_RTN_B64)
ds_cmpst_rtn_b64 = functools.partial(DS, DSOp.DS_CMPST_RTN_B64)
ds_cmpst_rtn_f64 = functools.partial(DS, DSOp.DS_CMPST_RTN_F64)
ds_min_rtn_f64 = functools.partial(DS, DSOp.DS_MIN_RTN_F64)
ds_max_rtn_f64 = functools.partial(DS, DSOp.DS_MAX_RTN_F64)
ds_read_b64 = functools.partial(DS, DSOp.DS_READ_B64)
ds_read2_b64 = functools.partial(DS, DSOp.DS_READ2_B64)
ds_read2st64_b64 = functools.partial(DS, DSOp.DS_READ2ST64_B64)
ds_add_rtn_f64 = functools.partial(DS, DSOp.DS_ADD_RTN_F64)
ds_condxchg32_rtn_b64 = functools.partial(DS, DSOp.DS_CONDXCHG32_RTN_B64)
ds_gws_sema_release_all = functools.partial(DS, DSOp.DS_GWS_SEMA_RELEASE_ALL)
ds_gws_init = functools.partial(DS, DSOp.DS_GWS_INIT)
ds_gws_sema_v = functools.partial(DS, DSOp.DS_GWS_SEMA_V)
ds_gws_sema_br = functools.partial(DS, DSOp.DS_GWS_SEMA_BR)
ds_gws_sema_p = functools.partial(DS, DSOp.DS_GWS_SEMA_P)
ds_gws_barrier = functools.partial(DS, DSOp.DS_GWS_BARRIER)
ds_read_addtid_b32 = functools.partial(DS, DSOp.DS_READ_ADDTID_B32)
ds_pk_add_rtn_f16 = functools.partial(DS, DSOp.DS_PK_ADD_RTN_F16)
ds_pk_add_rtn_bf16 = functools.partial(DS, DSOp.DS_PK_ADD_RTN_BF16)
ds_consume = functools.partial(DS, DSOp.DS_CONSUME)
ds_append = functools.partial(DS, DSOp.DS_APPEND)
ds_write_b96 = functools.partial(DS, DSOp.DS_WRITE_B96)
ds_write_b128 = functools.partial(DS, DSOp.DS_WRITE_B128)
ds_read_b64_tr_b4 = functools.partial(DS, DSOp.DS_READ_B64_TR_B4)
ds_read_b96_tr_b6 = functools.partial(DS, DSOp.DS_READ_B96_TR_B6)
ds_read_b64_tr_b8 = functools.partial(DS, DSOp.DS_READ_B64_TR_B8)
ds_read_b64_tr_b16 = functools.partial(DS, DSOp.DS_READ_B64_TR_B16)
ds_read_b96 = functools.partial(DS, DSOp.DS_READ_B96)
ds_read_b128 = functools.partial(DS, DSOp.DS_READ_B128)
cdna4 = functools.partial(DS, DSOp.CDNA4)
flat_load_ubyte = functools.partial(FLAT, FLATOp.FLAT_LOAD_UBYTE)
flat_load_sbyte = functools.partial(FLAT, FLATOp.FLAT_LOAD_SBYTE)
flat_load_ushort = functools.partial(FLAT, FLATOp.FLAT_LOAD_USHORT)
flat_load_sshort = functools.partial(FLAT, FLATOp.FLAT_LOAD_SSHORT)
flat_load_dword = functools.partial(FLAT, FLATOp.FLAT_LOAD_DWORD)
flat_load_dwordx2 = functools.partial(FLAT, FLATOp.FLAT_LOAD_DWORDX2)
flat_load_dwordx3 = functools.partial(FLAT, FLATOp.FLAT_LOAD_DWORDX3)
flat_load_dwordx4 = functools.partial(FLAT, FLATOp.FLAT_LOAD_DWORDX4)
flat_store_byte = functools.partial(FLAT, FLATOp.FLAT_STORE_BYTE)
flat_store_byte_d16_hi = functools.partial(FLAT, FLATOp.FLAT_STORE_BYTE_D16_HI)
flat_store_short = functools.partial(FLAT, FLATOp.FLAT_STORE_SHORT)
flat_store_short_d16_hi = functools.partial(FLAT, FLATOp.FLAT_STORE_SHORT_D16_HI)
flat_store_dword = functools.partial(FLAT, FLATOp.FLAT_STORE_DWORD)
flat_store_dwordx2 = functools.partial(FLAT, FLATOp.FLAT_STORE_DWORDX2)
flat_store_dwordx3 = functools.partial(FLAT, FLATOp.FLAT_STORE_DWORDX3)
flat_store_dwordx4 = functools.partial(FLAT, FLATOp.FLAT_STORE_DWORDX4)
flat_load_ubyte_d16 = functools.partial(FLAT, FLATOp.FLAT_LOAD_UBYTE_D16)
flat_load_ubyte_d16_hi = functools.partial(FLAT, FLATOp.FLAT_LOAD_UBYTE_D16_HI)
flat_load_sbyte_d16 = functools.partial(FLAT, FLATOp.FLAT_LOAD_SBYTE_D16)
flat_load_sbyte_d16_hi = functools.partial(FLAT, FLATOp.FLAT_LOAD_SBYTE_D16_HI)
flat_load_short_d16 = functools.partial(FLAT, FLATOp.FLAT_LOAD_SHORT_D16)
flat_load_short_d16_hi = functools.partial(FLAT, FLATOp.FLAT_LOAD_SHORT_D16_HI)
flat_atomic_swap = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_SWAP)
flat_atomic_cmpswap = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_CMPSWAP)
flat_atomic_add = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_ADD)
flat_atomic_sub = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_SUB)
flat_atomic_smin = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_SMIN)
flat_atomic_umin = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_UMIN)
flat_atomic_smax = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_SMAX)
flat_atomic_umax = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_UMAX)
flat_atomic_and = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_AND)
flat_atomic_or = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_OR)
flat_atomic_xor = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_XOR)
flat_atomic_inc = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_INC)
flat_atomic_dec = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_DEC)
flat_atomic_add_f32 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_ADD_F32)
flat_atomic_pk_add_f16 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_PK_ADD_F16)
flat_atomic_add_f64 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_ADD_F64)
flat_atomic_min_f64 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_MIN_F64)
flat_atomic_max_f64 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_MAX_F64)
flat_atomic_pk_add_bf16 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_PK_ADD_BF16)
flat_atomic_swap_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_SWAP_X2)
flat_atomic_cmpswap_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_CMPSWAP_X2)
flat_atomic_add_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_ADD_X2)
flat_atomic_sub_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_SUB_X2)
flat_atomic_smin_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_SMIN_X2)
flat_atomic_umin_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_UMIN_X2)
flat_atomic_smax_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_SMAX_X2)
flat_atomic_umax_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_UMAX_X2)
flat_atomic_and_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_AND_X2)
flat_atomic_or_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_OR_X2)
flat_atomic_xor_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_XOR_X2)
flat_atomic_inc_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_INC_X2)
flat_atomic_dec_x2 = functools.partial(FLAT, FLATOp.FLAT_ATOMIC_DEC_X2)
cdna4 = functools.partial(FLAT, FLATOp.CDNA4)
global_load_ubyte = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_UBYTE, seg=2)
global_load_sbyte = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_SBYTE, seg=2)
global_load_ushort = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_USHORT, seg=2)
global_load_sshort = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_SSHORT, seg=2)
global_load_dword = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_DWORD, seg=2)
global_load_dwordx2 = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_DWORDX2, seg=2)
global_load_dwordx3 = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_DWORDX3, seg=2)
global_load_dwordx4 = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_DWORDX4, seg=2)
global_store_byte = functools.partial(FLAT, GLOBALOp.GLOBAL_STORE_BYTE, seg=2)
global_store_byte_d16_hi = functools.partial(FLAT, GLOBALOp.GLOBAL_STORE_BYTE_D16_HI, seg=2)
global_store_short = functools.partial(FLAT, GLOBALOp.GLOBAL_STORE_SHORT, seg=2)
global_store_short_d16_hi = functools.partial(FLAT, GLOBALOp.GLOBAL_STORE_SHORT_D16_HI, seg=2)
global_store_dword = functools.partial(FLAT, GLOBALOp.GLOBAL_STORE_DWORD, seg=2)
global_store_dwordx2 = functools.partial(FLAT, GLOBALOp.GLOBAL_STORE_DWORDX2, seg=2)
global_store_dwordx3 = functools.partial(FLAT, GLOBALOp.GLOBAL_STORE_DWORDX3, seg=2)
global_store_dwordx4 = functools.partial(FLAT, GLOBALOp.GLOBAL_STORE_DWORDX4, seg=2)
global_load_ubyte_d16 = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_UBYTE_D16, seg=2)
global_load_ubyte_d16_hi = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_UBYTE_D16_HI, seg=2)
global_load_sbyte_d16 = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_SBYTE_D16, seg=2)
global_load_sbyte_d16_hi = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_SBYTE_D16_HI, seg=2)
global_load_short_d16 = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_SHORT_D16, seg=2)
global_load_short_d16_hi = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_SHORT_D16_HI, seg=2)
global_load_lds_ubyte = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_LDS_UBYTE, seg=2)
global_load_lds_sbyte = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_LDS_SBYTE, seg=2)
global_load_lds_ushort = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_LDS_USHORT, seg=2)
global_load_lds_sshort = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_LDS_SSHORT, seg=2)
global_load_lds_dword = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_LDS_DWORD, seg=2)
global_atomic_swap = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_SWAP, seg=2)
global_atomic_cmpswap = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_CMPSWAP, seg=2)
global_atomic_add = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_ADD, seg=2)
global_atomic_sub = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_SUB, seg=2)
global_atomic_smin = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_SMIN, seg=2)
global_atomic_umin = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_UMIN, seg=2)
global_atomic_smax = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_SMAX, seg=2)
global_atomic_umax = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_UMAX, seg=2)
global_atomic_and = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_AND, seg=2)
global_atomic_or = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_OR, seg=2)
global_atomic_xor = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_XOR, seg=2)
global_atomic_inc = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_INC, seg=2)
global_atomic_dec = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_DEC, seg=2)
global_atomic_add_f32 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_ADD_F32, seg=2)
global_atomic_pk_add_f16 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_PK_ADD_F16, seg=2)
global_atomic_add_f64 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_ADD_F64, seg=2)
global_atomic_min_f64 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_MIN_F64, seg=2)
global_atomic_max_f64 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_MAX_F64, seg=2)
global_atomic_pk_add_bf16 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_PK_ADD_BF16, seg=2)
global_atomic_swap_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_SWAP_X2, seg=2)
global_atomic_cmpswap_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_CMPSWAP_X2, seg=2)
global_atomic_add_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_ADD_X2, seg=2)
global_atomic_sub_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_SUB_X2, seg=2)
global_atomic_smin_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_SMIN_X2, seg=2)
global_atomic_umin_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_UMIN_X2, seg=2)
global_atomic_smax_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_SMAX_X2, seg=2)
global_atomic_umax_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_UMAX_X2, seg=2)
global_atomic_and_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_AND_X2, seg=2)
global_atomic_or_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_OR_X2, seg=2)
global_atomic_xor_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_XOR_X2, seg=2)
global_atomic_inc_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_INC_X2, seg=2)
global_atomic_dec_x2 = functools.partial(FLAT, GLOBALOp.GLOBAL_ATOMIC_DEC_X2, seg=2)
global_load_lds_dwordx4 = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_LDS_DWORDX4, seg=2)
global_load_lds_dwordx3 = functools.partial(FLAT, GLOBALOp.GLOBAL_LOAD_LDS_DWORDX3, seg=2)
cdna4 = functools.partial(FLAT, GLOBALOp.CDNA4, seg=2)
tbuffer_load_format_x = functools.partial(MTBUF, MTBUFOp.TBUFFER_LOAD_FORMAT_X)
tbuffer_load_format_xy = functools.partial(MTBUF, MTBUFOp.TBUFFER_LOAD_FORMAT_XY)
tbuffer_load_format_xyz = functools.partial(MTBUF, MTBUFOp.TBUFFER_LOAD_FORMAT_XYZ)
tbuffer_load_format_xyzw = functools.partial(MTBUF, MTBUFOp.TBUFFER_LOAD_FORMAT_XYZW)
tbuffer_store_format_x = functools.partial(MTBUF, MTBUFOp.TBUFFER_STORE_FORMAT_X)
tbuffer_store_format_xy = functools.partial(MTBUF, MTBUFOp.TBUFFER_STORE_FORMAT_XY)
tbuffer_store_format_xyz = functools.partial(MTBUF, MTBUFOp.TBUFFER_STORE_FORMAT_XYZ)
tbuffer_store_format_xyzw = functools.partial(MTBUF, MTBUFOp.TBUFFER_STORE_FORMAT_XYZW)
tbuffer_load_format_d16_x = functools.partial(MTBUF, MTBUFOp.TBUFFER_LOAD_FORMAT_D16_X)
tbuffer_load_format_d16_xy = functools.partial(MTBUF, MTBUFOp.TBUFFER_LOAD_FORMAT_D16_XY)
tbuffer_load_format_d16_xyz = functools.partial(MTBUF, MTBUFOp.TBUFFER_LOAD_FORMAT_D16_XYZ)
tbuffer_load_format_d16_xyzw = functools.partial(MTBUF, MTBUFOp.TBUFFER_LOAD_FORMAT_D16_XYZW)
tbuffer_store_format_d16_x = functools.partial(MTBUF, MTBUFOp.TBUFFER_STORE_FORMAT_D16_X)
tbuffer_store_format_d16_xy = functools.partial(MTBUF, MTBUFOp.TBUFFER_STORE_FORMAT_D16_XY)
tbuffer_store_format_d16_xyz = functools.partial(MTBUF, MTBUFOp.TBUFFER_STORE_FORMAT_D16_XYZ)
tbuffer_store_format_d16_xyzw = functools.partial(MTBUF, MTBUFOp.TBUFFER_STORE_FORMAT_D16_XYZW)
buffer_load_format_x = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_FORMAT_X)
buffer_load_format_xy = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_FORMAT_XY)
buffer_load_format_xyz = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_FORMAT_XYZ)
buffer_load_format_xyzw = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_FORMAT_XYZW)
buffer_store_format_x = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_FORMAT_X)
buffer_store_format_xy = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_FORMAT_XY)
buffer_store_format_xyz = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_FORMAT_XYZ)
buffer_store_format_xyzw = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_FORMAT_XYZW)
buffer_load_format_d16_x = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_FORMAT_D16_X)
buffer_load_format_d16_xy = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_FORMAT_D16_XY)
buffer_load_format_d16_xyz = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_FORMAT_D16_XYZ)
buffer_load_format_d16_xyzw = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_FORMAT_D16_XYZW)
buffer_store_format_d16_x = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_FORMAT_D16_X)
buffer_store_format_d16_xy = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_FORMAT_D16_XY)
buffer_store_format_d16_xyz = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_FORMAT_D16_XYZ)
buffer_store_format_d16_xyzw = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_FORMAT_D16_XYZW)
buffer_load_ubyte = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_UBYTE)
buffer_load_sbyte = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_SBYTE)
buffer_load_ushort = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_USHORT)
buffer_load_sshort = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_SSHORT)
buffer_load_dword = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_DWORD)
buffer_load_dwordx2 = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_DWORDX2)
buffer_load_dwordx3 = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_DWORDX3)
buffer_load_dwordx4 = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_DWORDX4)
buffer_store_byte = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_BYTE)
buffer_store_byte_d16_hi = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_BYTE_D16_HI)
buffer_store_short = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_SHORT)
buffer_store_short_d16_hi = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_SHORT_D16_HI)
buffer_store_dword = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_DWORD)
buffer_store_dwordx2 = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_DWORDX2)
buffer_store_dwordx3 = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_DWORDX3)
buffer_store_dwordx4 = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_DWORDX4)
buffer_load_ubyte_d16 = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_UBYTE_D16)
buffer_load_ubyte_d16_hi = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_UBYTE_D16_HI)
buffer_load_sbyte_d16 = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_SBYTE_D16)
buffer_load_sbyte_d16_hi = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_SBYTE_D16_HI)
buffer_load_short_d16 = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_SHORT_D16)
buffer_load_short_d16_hi = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_SHORT_D16_HI)
buffer_load_format_d16_hi_x = functools.partial(MUBUF, MUBUFOp.BUFFER_LOAD_FORMAT_D16_HI_X)
buffer_store_format_d16_hi_x = functools.partial(MUBUF, MUBUFOp.BUFFER_STORE_FORMAT_D16_HI_X)
buffer_wbl2 = functools.partial(MUBUF, MUBUFOp.BUFFER_WBL2)
buffer_inv = functools.partial(MUBUF, MUBUFOp.BUFFER_INV)
buffer_atomic_swap = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_SWAP)
buffer_atomic_cmpswap = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_CMPSWAP)
buffer_atomic_add = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_ADD)
buffer_atomic_sub = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_SUB)
buffer_atomic_smin = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_SMIN)
buffer_atomic_umin = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_UMIN)
buffer_atomic_smax = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_SMAX)
buffer_atomic_umax = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_UMAX)
buffer_atomic_and = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_AND)
buffer_atomic_or = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_OR)
buffer_atomic_xor = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_XOR)
buffer_atomic_inc = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_INC)
buffer_atomic_dec = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_DEC)
buffer_atomic_add_f32 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_ADD_F32)
buffer_atomic_pk_add_f16 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_PK_ADD_F16)
buffer_atomic_add_f64 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_ADD_F64)
buffer_atomic_min_f64 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_MIN_F64)
buffer_atomic_max_f64 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_MAX_F64)
buffer_atomic_pk_add_bf16 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_PK_ADD_BF16)
buffer_atomic_swap_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_SWAP_X2)
buffer_atomic_cmpswap_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_CMPSWAP_X2)
buffer_atomic_add_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_ADD_X2)
buffer_atomic_sub_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_SUB_X2)
buffer_atomic_smin_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_SMIN_X2)
buffer_atomic_umin_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_UMIN_X2)
buffer_atomic_smax_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_SMAX_X2)
buffer_atomic_umax_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_UMAX_X2)
buffer_atomic_and_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_AND_X2)
buffer_atomic_or_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_OR_X2)
buffer_atomic_xor_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_XOR_X2)
buffer_atomic_inc_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_INC_X2)
buffer_atomic_dec_x2 = functools.partial(MUBUF, MUBUFOp.BUFFER_ATOMIC_DEC_X2)
cdna4 = functools.partial(MUBUF, MUBUFOp.CDNA4)
scratch_load_ubyte = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_UBYTE, seg=1)
scratch_load_sbyte = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_SBYTE, seg=1)
scratch_load_ushort = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_USHORT, seg=1)
scratch_load_sshort = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_SSHORT, seg=1)
scratch_load_dword = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_DWORD, seg=1)
scratch_load_dwordx2 = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_DWORDX2, seg=1)
scratch_load_dwordx3 = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_DWORDX3, seg=1)
scratch_load_dwordx4 = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_DWORDX4, seg=1)
scratch_store_byte = functools.partial(FLAT, SCRATCHOp.SCRATCH_STORE_BYTE, seg=1)
scratch_store_byte_d16_hi = functools.partial(FLAT, SCRATCHOp.SCRATCH_STORE_BYTE_D16_HI, seg=1)
scratch_store_short = functools.partial(FLAT, SCRATCHOp.SCRATCH_STORE_SHORT, seg=1)
scratch_store_short_d16_hi = functools.partial(FLAT, SCRATCHOp.SCRATCH_STORE_SHORT_D16_HI, seg=1)
scratch_store_dword = functools.partial(FLAT, SCRATCHOp.SCRATCH_STORE_DWORD, seg=1)
scratch_store_dwordx2 = functools.partial(FLAT, SCRATCHOp.SCRATCH_STORE_DWORDX2, seg=1)
scratch_store_dwordx3 = functools.partial(FLAT, SCRATCHOp.SCRATCH_STORE_DWORDX3, seg=1)
scratch_store_dwordx4 = functools.partial(FLAT, SCRATCHOp.SCRATCH_STORE_DWORDX4, seg=1)
scratch_load_ubyte_d16 = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_UBYTE_D16, seg=1)
scratch_load_ubyte_d16_hi = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_UBYTE_D16_HI, seg=1)
scratch_load_sbyte_d16 = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_SBYTE_D16, seg=1)
scratch_load_sbyte_d16_hi = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_SBYTE_D16_HI, seg=1)
scratch_load_short_d16 = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_SHORT_D16, seg=1)
scratch_load_short_d16_hi = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_SHORT_D16_HI, seg=1)
scratch_load_lds_ubyte = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_LDS_UBYTE, seg=1)
scratch_load_lds_sbyte = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_LDS_SBYTE, seg=1)
scratch_load_lds_ushort = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_LDS_USHORT, seg=1)
scratch_load_lds_sshort = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_LDS_SSHORT, seg=1)
scratch_load_lds_dword = functools.partial(FLAT, SCRATCHOp.SCRATCH_LOAD_LDS_DWORD, seg=1)
s_load_dword = functools.partial(SMEM, SMEMOp.S_LOAD_DWORD)
s_load_dwordx2 = functools.partial(SMEM, SMEMOp.S_LOAD_DWORDX2)
s_load_dwordx4 = functools.partial(SMEM, SMEMOp.S_LOAD_DWORDX4)
s_load_dwordx8 = functools.partial(SMEM, SMEMOp.S_LOAD_DWORDX8)
s_load_dwordx16 = functools.partial(SMEM, SMEMOp.S_LOAD_DWORDX16)
s_scratch_load_dword = functools.partial(SMEM, SMEMOp.S_SCRATCH_LOAD_DWORD)
s_scratch_load_dwordx2 = functools.partial(SMEM, SMEMOp.S_SCRATCH_LOAD_DWORDX2)
s_scratch_load_dwordx4 = functools.partial(SMEM, SMEMOp.S_SCRATCH_LOAD_DWORDX4)
s_buffer_load_dword = functools.partial(SMEM, SMEMOp.S_BUFFER_LOAD_DWORD)
s_buffer_load_dwordx2 = functools.partial(SMEM, SMEMOp.S_BUFFER_LOAD_DWORDX2)
s_buffer_load_dwordx4 = functools.partial(SMEM, SMEMOp.S_BUFFER_LOAD_DWORDX4)
s_buffer_load_dwordx8 = functools.partial(SMEM, SMEMOp.S_BUFFER_LOAD_DWORDX8)
s_buffer_load_dwordx16 = functools.partial(SMEM, SMEMOp.S_BUFFER_LOAD_DWORDX16)
s_store_dword = functools.partial(SMEM, SMEMOp.S_STORE_DWORD)
s_store_dwordx2 = functools.partial(SMEM, SMEMOp.S_STORE_DWORDX2)
s_store_dwordx4 = functools.partial(SMEM, SMEMOp.S_STORE_DWORDX4)
s_scratch_store_dword = functools.partial(SMEM, SMEMOp.S_SCRATCH_STORE_DWORD)
s_scratch_store_dwordx2 = functools.partial(SMEM, SMEMOp.S_SCRATCH_STORE_DWORDX2)
s_scratch_store_dwordx4 = functools.partial(SMEM, SMEMOp.S_SCRATCH_STORE_DWORDX4)
s_buffer_store_dword = functools.partial(SMEM, SMEMOp.S_BUFFER_STORE_DWORD)
s_buffer_store_dwordx2 = functools.partial(SMEM, SMEMOp.S_BUFFER_STORE_DWORDX2)
s_buffer_store_dwordx4 = functools.partial(SMEM, SMEMOp.S_BUFFER_STORE_DWORDX4)
s_dcache_inv = functools.partial(SMEM, SMEMOp.S_DCACHE_INV)
s_dcache_wb = functools.partial(SMEM, SMEMOp.S_DCACHE_WB)
s_dcache_inv_vol = functools.partial(SMEM, SMEMOp.S_DCACHE_INV_VOL)
s_dcache_wb_vol = functools.partial(SMEM, SMEMOp.S_DCACHE_WB_VOL)
s_memtime = functools.partial(SMEM, SMEMOp.S_MEMTIME)
s_memrealtime = functools.partial(SMEM, SMEMOp.S_MEMREALTIME)
s_dcache_discard = functools.partial(SMEM, SMEMOp.S_DCACHE_DISCARD)
s_dcache_discard_x2 = functools.partial(SMEM, SMEMOp.S_DCACHE_DISCARD_X2)
s_buffer_atomic_swap = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_SWAP)
s_buffer_atomic_cmpswap = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_CMPSWAP)
s_buffer_atomic_add = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_ADD)
s_buffer_atomic_sub = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_SUB)
s_buffer_atomic_smin = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_SMIN)
s_buffer_atomic_umin = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_UMIN)
s_buffer_atomic_smax = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_SMAX)
s_buffer_atomic_umax = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_UMAX)
s_buffer_atomic_and = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_AND)
s_buffer_atomic_or = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_OR)
s_buffer_atomic_xor = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_XOR)
s_buffer_atomic_inc = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_INC)
s_buffer_atomic_dec = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_DEC)
s_buffer_atomic_swap_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_SWAP_X2)
s_buffer_atomic_cmpswap_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_CMPSWAP_X2)
s_buffer_atomic_add_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_ADD_X2)
s_buffer_atomic_sub_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_SUB_X2)
s_buffer_atomic_smin_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_SMIN_X2)
s_buffer_atomic_umin_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_UMIN_X2)
s_buffer_atomic_smax_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_SMAX_X2)
s_buffer_atomic_umax_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_UMAX_X2)
s_buffer_atomic_and_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_AND_X2)
s_buffer_atomic_or_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_OR_X2)
s_buffer_atomic_xor_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_XOR_X2)
s_buffer_atomic_inc_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_INC_X2)
s_buffer_atomic_dec_x2 = functools.partial(SMEM, SMEMOp.S_BUFFER_ATOMIC_DEC_X2)
s_atomic_swap = functools.partial(SMEM, SMEMOp.S_ATOMIC_SWAP)
s_atomic_cmpswap = functools.partial(SMEM, SMEMOp.S_ATOMIC_CMPSWAP)
s_atomic_add = functools.partial(SMEM, SMEMOp.S_ATOMIC_ADD)
s_atomic_sub = functools.partial(SMEM, SMEMOp.S_ATOMIC_SUB)
s_atomic_smin = functools.partial(SMEM, SMEMOp.S_ATOMIC_SMIN)
s_atomic_umin = functools.partial(SMEM, SMEMOp.S_ATOMIC_UMIN)
s_atomic_smax = functools.partial(SMEM, SMEMOp.S_ATOMIC_SMAX)
s_atomic_umax = functools.partial(SMEM, SMEMOp.S_ATOMIC_UMAX)
s_atomic_and = functools.partial(SMEM, SMEMOp.S_ATOMIC_AND)
s_atomic_or = functools.partial(SMEM, SMEMOp.S_ATOMIC_OR)
s_atomic_xor = functools.partial(SMEM, SMEMOp.S_ATOMIC_XOR)
s_atomic_inc = functools.partial(SMEM, SMEMOp.S_ATOMIC_INC)
s_atomic_dec = functools.partial(SMEM, SMEMOp.S_ATOMIC_DEC)
s_atomic_swap_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_SWAP_X2)
s_atomic_cmpswap_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_CMPSWAP_X2)
s_atomic_add_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_ADD_X2)
s_atomic_sub_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_SUB_X2)
s_atomic_smin_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_SMIN_X2)
s_atomic_umin_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_UMIN_X2)
s_atomic_smax_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_SMAX_X2)
s_atomic_umax_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_UMAX_X2)
s_atomic_and_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_AND_X2)
s_atomic_or_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_OR_X2)
s_atomic_xor_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_XOR_X2)
s_atomic_inc_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_INC_X2)
s_atomic_dec_x2 = functools.partial(SMEM, SMEMOp.S_ATOMIC_DEC_X2)
cdna4 = functools.partial(SMEM, SMEMOp.CDNA4)
s_mov_b32 = functools.partial(SOP1, SOP1Op.S_MOV_B32)
s_mov_b64 = functools.partial(SOP1, SOP1Op.S_MOV_B64)
s_cmov_b32 = functools.partial(SOP1, SOP1Op.S_CMOV_B32)
s_cmov_b64 = functools.partial(SOP1, SOP1Op.S_CMOV_B64)
s_not_b32 = functools.partial(SOP1, SOP1Op.S_NOT_B32)
s_not_b64 = functools.partial(SOP1, SOP1Op.S_NOT_B64)
s_wqm_b32 = functools.partial(SOP1, SOP1Op.S_WQM_B32)
s_wqm_b64 = functools.partial(SOP1, SOP1Op.S_WQM_B64)
s_brev_b32 = functools.partial(SOP1, SOP1Op.S_BREV_B32)
s_brev_b64 = functools.partial(SOP1, SOP1Op.S_BREV_B64)
s_bcnt0_i32_b32 = functools.partial(SOP1, SOP1Op.S_BCNT0_I32_B32)
s_bcnt0_i32_b64 = functools.partial(SOP1, SOP1Op.S_BCNT0_I32_B64)
s_bcnt1_i32_b32 = functools.partial(SOP1, SOP1Op.S_BCNT1_I32_B32)
s_bcnt1_i32_b64 = functools.partial(SOP1, SOP1Op.S_BCNT1_I32_B64)
s_ff0_i32_b32 = functools.partial(SOP1, SOP1Op.S_FF0_I32_B32)
s_ff0_i32_b64 = functools.partial(SOP1, SOP1Op.S_FF0_I32_B64)
s_ff1_i32_b32 = functools.partial(SOP1, SOP1Op.S_FF1_I32_B32)
s_ff1_i32_b64 = functools.partial(SOP1, SOP1Op.S_FF1_I32_B64)
s_flbit_i32_b32 = functools.partial(SOP1, SOP1Op.S_FLBIT_I32_B32)
s_flbit_i32_b64 = functools.partial(SOP1, SOP1Op.S_FLBIT_I32_B64)
s_flbit_i32 = functools.partial(SOP1, SOP1Op.S_FLBIT_I32)
s_flbit_i32_i64 = functools.partial(SOP1, SOP1Op.S_FLBIT_I32_I64)
s_sext_i32_i8 = functools.partial(SOP1, SOP1Op.S_SEXT_I32_I8)
s_sext_i32_i16 = functools.partial(SOP1, SOP1Op.S_SEXT_I32_I16)
s_bitset0_b32 = functools.partial(SOP1, SOP1Op.S_BITSET0_B32)
s_bitset0_b64 = functools.partial(SOP1, SOP1Op.S_BITSET0_B64)
s_bitset1_b32 = functools.partial(SOP1, SOP1Op.S_BITSET1_B32)
s_bitset1_b64 = functools.partial(SOP1, SOP1Op.S_BITSET1_B64)
s_getpc_b64 = functools.partial(SOP1, SOP1Op.S_GETPC_B64)
s_setpc_b64 = functools.partial(SOP1, SOP1Op.S_SETPC_B64)
s_swappc_b64 = functools.partial(SOP1, SOP1Op.S_SWAPPC_B64)
s_rfe_b64 = functools.partial(SOP1, SOP1Op.S_RFE_B64)
s_and_saveexec_b64 = functools.partial(SOP1, SOP1Op.S_AND_SAVEEXEC_B64)
s_or_saveexec_b64 = functools.partial(SOP1, SOP1Op.S_OR_SAVEEXEC_B64)
s_xor_saveexec_b64 = functools.partial(SOP1, SOP1Op.S_XOR_SAVEEXEC_B64)
s_andn2_saveexec_b64 = functools.partial(SOP1, SOP1Op.S_ANDN2_SAVEEXEC_B64)
s_orn2_saveexec_b64 = functools.partial(SOP1, SOP1Op.S_ORN2_SAVEEXEC_B64)
s_nand_saveexec_b64 = functools.partial(SOP1, SOP1Op.S_NAND_SAVEEXEC_B64)
s_nor_saveexec_b64 = functools.partial(SOP1, SOP1Op.S_NOR_SAVEEXEC_B64)
s_xnor_saveexec_b64 = functools.partial(SOP1, SOP1Op.S_XNOR_SAVEEXEC_B64)
s_quadmask_b32 = functools.partial(SOP1, SOP1Op.S_QUADMASK_B32)
s_quadmask_b64 = functools.partial(SOP1, SOP1Op.S_QUADMASK_B64)
s_movrels_b32 = functools.partial(SOP1, SOP1Op.S_MOVRELS_B32)
s_movrels_b64 = functools.partial(SOP1, SOP1Op.S_MOVRELS_B64)
s_movreld_b32 = functools.partial(SOP1, SOP1Op.S_MOVRELD_B32)
s_movreld_b64 = functools.partial(SOP1, SOP1Op.S_MOVRELD_B64)
s_cbranch_join = functools.partial(SOP1, SOP1Op.S_CBRANCH_JOIN)
s_abs_i32 = functools.partial(SOP1, SOP1Op.S_ABS_I32)
s_set_gpr_idx_idx = functools.partial(SOP1, SOP1Op.S_SET_GPR_IDX_IDX)
s_andn1_saveexec_b64 = functools.partial(SOP1, SOP1Op.S_ANDN1_SAVEEXEC_B64)
s_orn1_saveexec_b64 = functools.partial(SOP1, SOP1Op.S_ORN1_SAVEEXEC_B64)
s_andn1_wrexec_b64 = functools.partial(SOP1, SOP1Op.S_ANDN1_WREXEC_B64)
s_andn2_wrexec_b64 = functools.partial(SOP1, SOP1Op.S_ANDN2_WREXEC_B64)
s_bitreplicate_b64_b32 = functools.partial(SOP1, SOP1Op.S_BITREPLICATE_B64_B32)
cdna4 = functools.partial(SOP1, SOP1Op.CDNA4)
s_add_u32 = functools.partial(SOP2, SOP2Op.S_ADD_U32)
s_sub_u32 = functools.partial(SOP2, SOP2Op.S_SUB_U32)
s_add_i32 = functools.partial(SOP2, SOP2Op.S_ADD_I32)
s_sub_i32 = functools.partial(SOP2, SOP2Op.S_SUB_I32)
s_addc_u32 = functools.partial(SOP2, SOP2Op.S_ADDC_U32)
s_subb_u32 = functools.partial(SOP2, SOP2Op.S_SUBB_U32)
s_min_i32 = functools.partial(SOP2, SOP2Op.S_MIN_I32)
s_min_u32 = functools.partial(SOP2, SOP2Op.S_MIN_U32)
s_max_i32 = functools.partial(SOP2, SOP2Op.S_MAX_I32)
s_max_u32 = functools.partial(SOP2, SOP2Op.S_MAX_U32)
s_cselect_b32 = functools.partial(SOP2, SOP2Op.S_CSELECT_B32)
s_cselect_b64 = functools.partial(SOP2, SOP2Op.S_CSELECT_B64)
s_and_b32 = functools.partial(SOP2, SOP2Op.S_AND_B32)
s_and_b64 = functools.partial(SOP2, SOP2Op.S_AND_B64)
s_or_b32 = functools.partial(SOP2, SOP2Op.S_OR_B32)
s_or_b64 = functools.partial(SOP2, SOP2Op.S_OR_B64)
s_xor_b32 = functools.partial(SOP2, SOP2Op.S_XOR_B32)
s_xor_b64 = functools.partial(SOP2, SOP2Op.S_XOR_B64)
s_andn2_b32 = functools.partial(SOP2, SOP2Op.S_ANDN2_B32)
s_andn2_b64 = functools.partial(SOP2, SOP2Op.S_ANDN2_B64)
s_orn2_b32 = functools.partial(SOP2, SOP2Op.S_ORN2_B32)
s_orn2_b64 = functools.partial(SOP2, SOP2Op.S_ORN2_B64)
s_nand_b32 = functools.partial(SOP2, SOP2Op.S_NAND_B32)
s_nand_b64 = functools.partial(SOP2, SOP2Op.S_NAND_B64)
s_nor_b32 = functools.partial(SOP2, SOP2Op.S_NOR_B32)
s_nor_b64 = functools.partial(SOP2, SOP2Op.S_NOR_B64)
s_xnor_b32 = functools.partial(SOP2, SOP2Op.S_XNOR_B32)
s_xnor_b64 = functools.partial(SOP2, SOP2Op.S_XNOR_B64)
s_lshl_b32 = functools.partial(SOP2, SOP2Op.S_LSHL_B32)
s_lshl_b64 = functools.partial(SOP2, SOP2Op.S_LSHL_B64)
s_lshr_b32 = functools.partial(SOP2, SOP2Op.S_LSHR_B32)
s_lshr_b64 = functools.partial(SOP2, SOP2Op.S_LSHR_B64)
s_ashr_i32 = functools.partial(SOP2, SOP2Op.S_ASHR_I32)
s_ashr_i64 = functools.partial(SOP2, SOP2Op.S_ASHR_I64)
s_bfm_b32 = functools.partial(SOP2, SOP2Op.S_BFM_B32)
s_bfm_b64 = functools.partial(SOP2, SOP2Op.S_BFM_B64)
s_mul_i32 = functools.partial(SOP2, SOP2Op.S_MUL_I32)
s_bfe_u32 = functools.partial(SOP2, SOP2Op.S_BFE_U32)
s_bfe_i32 = functools.partial(SOP2, SOP2Op.S_BFE_I32)
s_bfe_u64 = functools.partial(SOP2, SOP2Op.S_BFE_U64)
s_bfe_i64 = functools.partial(SOP2, SOP2Op.S_BFE_I64)
s_cbranch_g_fork = functools.partial(SOP2, SOP2Op.S_CBRANCH_G_FORK)
s_absdiff_i32 = functools.partial(SOP2, SOP2Op.S_ABSDIFF_I32)
s_mul_hi_u32 = functools.partial(SOP2, SOP2Op.S_MUL_HI_U32)
s_mul_hi_i32 = functools.partial(SOP2, SOP2Op.S_MUL_HI_I32)
s_lshl1_add_u32 = functools.partial(SOP2, SOP2Op.S_LSHL1_ADD_U32)
s_lshl2_add_u32 = functools.partial(SOP2, SOP2Op.S_LSHL2_ADD_U32)
s_lshl3_add_u32 = functools.partial(SOP2, SOP2Op.S_LSHL3_ADD_U32)
s_lshl4_add_u32 = functools.partial(SOP2, SOP2Op.S_LSHL4_ADD_U32)
s_pack_ll_b32_b16 = functools.partial(SOP2, SOP2Op.S_PACK_LL_B32_B16)
s_pack_lh_b32_b16 = functools.partial(SOP2, SOP2Op.S_PACK_LH_B32_B16)
s_pack_hh_b32_b16 = functools.partial(SOP2, SOP2Op.S_PACK_HH_B32_B16)
cdna4 = functools.partial(SOP2, SOP2Op.CDNA4)
s_cmp_eq_i32 = functools.partial(SOPC, SOPCOp.S_CMP_EQ_I32)
s_cmp_lg_i32 = functools.partial(SOPC, SOPCOp.S_CMP_LG_I32)
s_cmp_gt_i32 = functools.partial(SOPC, SOPCOp.S_CMP_GT_I32)
s_cmp_ge_i32 = functools.partial(SOPC, SOPCOp.S_CMP_GE_I32)
s_cmp_lt_i32 = functools.partial(SOPC, SOPCOp.S_CMP_LT_I32)
s_cmp_le_i32 = functools.partial(SOPC, SOPCOp.S_CMP_LE_I32)
s_cmp_eq_u32 = functools.partial(SOPC, SOPCOp.S_CMP_EQ_U32)
s_cmp_lg_u32 = functools.partial(SOPC, SOPCOp.S_CMP_LG_U32)
s_cmp_gt_u32 = functools.partial(SOPC, SOPCOp.S_CMP_GT_U32)
s_cmp_ge_u32 = functools.partial(SOPC, SOPCOp.S_CMP_GE_U32)
s_cmp_lt_u32 = functools.partial(SOPC, SOPCOp.S_CMP_LT_U32)
s_cmp_le_u32 = functools.partial(SOPC, SOPCOp.S_CMP_LE_U32)
s_bitcmp0_b32 = functools.partial(SOPC, SOPCOp.S_BITCMP0_B32)
s_bitcmp1_b32 = functools.partial(SOPC, SOPCOp.S_BITCMP1_B32)
s_bitcmp0_b64 = functools.partial(SOPC, SOPCOp.S_BITCMP0_B64)
s_bitcmp1_b64 = functools.partial(SOPC, SOPCOp.S_BITCMP1_B64)
s_setvskip = functools.partial(SOPC, SOPCOp.S_SETVSKIP)
s_set_gpr_idx_on = functools.partial(SOPC, SOPCOp.S_SET_GPR_IDX_ON)
s_cmp_eq_u64 = functools.partial(SOPC, SOPCOp.S_CMP_EQ_U64)
s_cmp_lg_u64 = functools.partial(SOPC, SOPCOp.S_CMP_LG_U64)
cdna4 = functools.partial(SOPC, SOPCOp.CDNA4)
s_movk_i32 = functools.partial(SOPK, SOPKOp.S_MOVK_I32)
s_cmovk_i32 = functools.partial(SOPK, SOPKOp.S_CMOVK_I32)
s_cmpk_eq_i32 = functools.partial(SOPK, SOPKOp.S_CMPK_EQ_I32)
s_cmpk_lg_i32 = functools.partial(SOPK, SOPKOp.S_CMPK_LG_I32)
s_cmpk_gt_i32 = functools.partial(SOPK, SOPKOp.S_CMPK_GT_I32)
s_cmpk_ge_i32 = functools.partial(SOPK, SOPKOp.S_CMPK_GE_I32)
s_cmpk_lt_i32 = functools.partial(SOPK, SOPKOp.S_CMPK_LT_I32)
s_cmpk_le_i32 = functools.partial(SOPK, SOPKOp.S_CMPK_LE_I32)
s_cmpk_eq_u32 = functools.partial(SOPK, SOPKOp.S_CMPK_EQ_U32)
s_cmpk_lg_u32 = functools.partial(SOPK, SOPKOp.S_CMPK_LG_U32)
s_cmpk_gt_u32 = functools.partial(SOPK, SOPKOp.S_CMPK_GT_U32)
s_cmpk_ge_u32 = functools.partial(SOPK, SOPKOp.S_CMPK_GE_U32)
s_cmpk_lt_u32 = functools.partial(SOPK, SOPKOp.S_CMPK_LT_U32)
s_cmpk_le_u32 = functools.partial(SOPK, SOPKOp.S_CMPK_LE_U32)
s_addk_i32 = functools.partial(SOPK, SOPKOp.S_ADDK_I32)
s_mulk_i32 = functools.partial(SOPK, SOPKOp.S_MULK_I32)
s_cbranch_i_fork = functools.partial(SOPK, SOPKOp.S_CBRANCH_I_FORK)
s_getreg_b32 = functools.partial(SOPK, SOPKOp.S_GETREG_B32)
s_setreg_b32 = functools.partial(SOPK, SOPKOp.S_SETREG_B32)
s_setreg_imm32_b32 = functools.partial(SOPK, SOPKOp.S_SETREG_IMM32_B32)
s_call_b64 = functools.partial(SOPK, SOPKOp.S_CALL_B64)
s_nop = functools.partial(SOPP, SOPPOp.S_NOP)
s_endpgm = functools.partial(SOPP, SOPPOp.S_ENDPGM)
s_branch = functools.partial(SOPP, SOPPOp.S_BRANCH)
s_wakeup = functools.partial(SOPP, SOPPOp.S_WAKEUP)
s_cbranch_scc0 = functools.partial(SOPP, SOPPOp.S_CBRANCH_SCC0)
s_cbranch_scc1 = functools.partial(SOPP, SOPPOp.S_CBRANCH_SCC1)
s_cbranch_vccz = functools.partial(SOPP, SOPPOp.S_CBRANCH_VCCZ)
s_cbranch_vccnz = functools.partial(SOPP, SOPPOp.S_CBRANCH_VCCNZ)
s_cbranch_execz = functools.partial(SOPP, SOPPOp.S_CBRANCH_EXECZ)
s_cbranch_execnz = functools.partial(SOPP, SOPPOp.S_CBRANCH_EXECNZ)
s_barrier = functools.partial(SOPP, SOPPOp.S_BARRIER)
s_setkill = functools.partial(SOPP, SOPPOp.S_SETKILL)
s_waitcnt = functools.partial(SOPP, SOPPOp.S_WAITCNT)
s_sethalt = functools.partial(SOPP, SOPPOp.S_SETHALT)
s_sleep = functools.partial(SOPP, SOPPOp.S_SLEEP)
s_setprio = functools.partial(SOPP, SOPPOp.S_SETPRIO)
s_sendmsg = functools.partial(SOPP, SOPPOp.S_SENDMSG)
s_sendmsghalt = functools.partial(SOPP, SOPPOp.S_SENDMSGHALT)
s_trap = functools.partial(SOPP, SOPPOp.S_TRAP)
s_icache_inv = functools.partial(SOPP, SOPPOp.S_ICACHE_INV)
s_incperflevel = functools.partial(SOPP, SOPPOp.S_INCPERFLEVEL)
s_decperflevel = functools.partial(SOPP, SOPPOp.S_DECPERFLEVEL)
s_ttracedata = functools.partial(SOPP, SOPPOp.S_TTRACEDATA)
s_cbranch_cdbgsys = functools.partial(SOPP, SOPPOp.S_CBRANCH_CDBGSYS)
s_cbranch_cdbguser = functools.partial(SOPP, SOPPOp.S_CBRANCH_CDBGUSER)
s_cbranch_cdbgsys_or_user = functools.partial(SOPP, SOPPOp.S_CBRANCH_CDBGSYS_OR_USER)
s_cbranch_cdbgsys_and_user = functools.partial(SOPP, SOPPOp.S_CBRANCH_CDBGSYS_AND_USER)
s_endpgm_saved = functools.partial(SOPP, SOPPOp.S_ENDPGM_SAVED)
s_set_gpr_idx_off = functools.partial(SOPP, SOPPOp.S_SET_GPR_IDX_OFF)
s_set_gpr_idx_mode = functools.partial(SOPP, SOPPOp.S_SET_GPR_IDX_MODE)
cdna4 = functools.partial(SOPP, SOPPOp.CDNA4)
v_nop_e32 = functools.partial(VOP1, VOP1Op.V_NOP)
v_mov_b32_e32 = functools.partial(VOP1, VOP1Op.V_MOV_B32)
v_readfirstlane_b32_e32 = functools.partial(VOP1, VOP1Op.V_READFIRSTLANE_B32)
v_cvt_i32_f64_e32 = functools.partial(VOP1, VOP1Op.V_CVT_I32_F64)
v_cvt_f64_i32_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F64_I32)
v_cvt_f32_i32_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F32_I32)
v_cvt_f32_u32_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F32_U32)
v_cvt_u32_f32_e32 = functools.partial(VOP1, VOP1Op.V_CVT_U32_F32)
v_cvt_i32_f32_e32 = functools.partial(VOP1, VOP1Op.V_CVT_I32_F32)
v_cvt_f16_f32_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F16_F32)
v_cvt_f32_f16_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F32_F16)
v_cvt_rpi_i32_f32_e32 = functools.partial(VOP1, VOP1Op.V_CVT_RPI_I32_F32)
v_cvt_flr_i32_f32_e32 = functools.partial(VOP1, VOP1Op.V_CVT_FLR_I32_F32)
v_cvt_off_f32_i4_e32 = functools.partial(VOP1, VOP1Op.V_CVT_OFF_F32_I4)
v_cvt_f32_f64_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F32_F64)
v_cvt_f64_f32_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F64_F32)
v_cvt_f32_ubyte0_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F32_UBYTE0)
v_cvt_f32_ubyte1_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F32_UBYTE1)
v_cvt_f32_ubyte2_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F32_UBYTE2)
v_cvt_f32_ubyte3_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F32_UBYTE3)
v_cvt_u32_f64_e32 = functools.partial(VOP1, VOP1Op.V_CVT_U32_F64)
v_cvt_f64_u32_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F64_U32)
v_trunc_f64_e32 = functools.partial(VOP1, VOP1Op.V_TRUNC_F64)
v_ceil_f64_e32 = functools.partial(VOP1, VOP1Op.V_CEIL_F64)
v_rndne_f64_e32 = functools.partial(VOP1, VOP1Op.V_RNDNE_F64)
v_floor_f64_e32 = functools.partial(VOP1, VOP1Op.V_FLOOR_F64)
v_fract_f32_e32 = functools.partial(VOP1, VOP1Op.V_FRACT_F32)
v_trunc_f32_e32 = functools.partial(VOP1, VOP1Op.V_TRUNC_F32)
v_ceil_f32_e32 = functools.partial(VOP1, VOP1Op.V_CEIL_F32)
v_rndne_f32_e32 = functools.partial(VOP1, VOP1Op.V_RNDNE_F32)
v_floor_f32_e32 = functools.partial(VOP1, VOP1Op.V_FLOOR_F32)
v_exp_f32_e32 = functools.partial(VOP1, VOP1Op.V_EXP_F32)
v_log_f32_e32 = functools.partial(VOP1, VOP1Op.V_LOG_F32)
v_rcp_f32_e32 = functools.partial(VOP1, VOP1Op.V_RCP_F32)
v_rcp_iflag_f32_e32 = functools.partial(VOP1, VOP1Op.V_RCP_IFLAG_F32)
v_rsq_f32_e32 = functools.partial(VOP1, VOP1Op.V_RSQ_F32)
v_rcp_f64_e32 = functools.partial(VOP1, VOP1Op.V_RCP_F64)
v_rsq_f64_e32 = functools.partial(VOP1, VOP1Op.V_RSQ_F64)
v_sqrt_f32_e32 = functools.partial(VOP1, VOP1Op.V_SQRT_F32)
v_sqrt_f64_e32 = functools.partial(VOP1, VOP1Op.V_SQRT_F64)
v_sin_f32_e32 = functools.partial(VOP1, VOP1Op.V_SIN_F32)
v_cos_f32_e32 = functools.partial(VOP1, VOP1Op.V_COS_F32)
v_not_b32_e32 = functools.partial(VOP1, VOP1Op.V_NOT_B32)
v_bfrev_b32_e32 = functools.partial(VOP1, VOP1Op.V_BFREV_B32)
v_ffbh_u32_e32 = functools.partial(VOP1, VOP1Op.V_FFBH_U32)
v_ffbl_b32_e32 = functools.partial(VOP1, VOP1Op.V_FFBL_B32)
v_ffbh_i32_e32 = functools.partial(VOP1, VOP1Op.V_FFBH_I32)
v_frexp_exp_i32_f64_e32 = functools.partial(VOP1, VOP1Op.V_FREXP_EXP_I32_F64)
v_frexp_mant_f64_e32 = functools.partial(VOP1, VOP1Op.V_FREXP_MANT_F64)
v_fract_f64_e32 = functools.partial(VOP1, VOP1Op.V_FRACT_F64)
v_frexp_exp_i32_f32_e32 = functools.partial(VOP1, VOP1Op.V_FREXP_EXP_I32_F32)
v_frexp_mant_f32_e32 = functools.partial(VOP1, VOP1Op.V_FREXP_MANT_F32)
v_clrexcp_e32 = functools.partial(VOP1, VOP1Op.V_CLREXCP)
v_mov_b64_e32 = functools.partial(VOP1, VOP1Op.V_MOV_B64)
v_cvt_f16_u16_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F16_U16)
v_cvt_f16_i16_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F16_I16)
v_cvt_u16_f16_e32 = functools.partial(VOP1, VOP1Op.V_CVT_U16_F16)
v_cvt_i16_f16_e32 = functools.partial(VOP1, VOP1Op.V_CVT_I16_F16)
v_rcp_f16_e32 = functools.partial(VOP1, VOP1Op.V_RCP_F16)
v_sqrt_f16_e32 = functools.partial(VOP1, VOP1Op.V_SQRT_F16)
v_rsq_f16_e32 = functools.partial(VOP1, VOP1Op.V_RSQ_F16)
v_log_f16_e32 = functools.partial(VOP1, VOP1Op.V_LOG_F16)
v_exp_f16_e32 = functools.partial(VOP1, VOP1Op.V_EXP_F16)
v_frexp_mant_f16_e32 = functools.partial(VOP1, VOP1Op.V_FREXP_MANT_F16)
v_frexp_exp_i16_f16_e32 = functools.partial(VOP1, VOP1Op.V_FREXP_EXP_I16_F16)
v_floor_f16_e32 = functools.partial(VOP1, VOP1Op.V_FLOOR_F16)
v_ceil_f16_e32 = functools.partial(VOP1, VOP1Op.V_CEIL_F16)
v_trunc_f16_e32 = functools.partial(VOP1, VOP1Op.V_TRUNC_F16)
v_rndne_f16_e32 = functools.partial(VOP1, VOP1Op.V_RNDNE_F16)
v_fract_f16_e32 = functools.partial(VOP1, VOP1Op.V_FRACT_F16)
v_sin_f16_e32 = functools.partial(VOP1, VOP1Op.V_SIN_F16)
v_cos_f16_e32 = functools.partial(VOP1, VOP1Op.V_COS_F16)
v_cvt_norm_i16_f16_e32 = functools.partial(VOP1, VOP1Op.V_CVT_NORM_I16_F16)
v_cvt_norm_u16_f16_e32 = functools.partial(VOP1, VOP1Op.V_CVT_NORM_U16_F16)
v_sat_pk_u8_i16_e32 = functools.partial(VOP1, VOP1Op.V_SAT_PK_U8_I16)
v_swap_b32_e32 = functools.partial(VOP1, VOP1Op.V_SWAP_B32)
v_accvgpr_mov_b32_e32 = functools.partial(VOP1, VOP1Op.V_ACCVGPR_MOV_B32)
v_cvt_f32_fp8_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F32_FP8)
v_cvt_f32_bf8_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F32_BF8)
v_cvt_pk_f32_fp8_e32 = functools.partial(VOP1, VOP1Op.V_CVT_PK_F32_FP8)
v_cvt_pk_f32_bf8_e32 = functools.partial(VOP1, VOP1Op.V_CVT_PK_F32_BF8)
v_prng_b32_e32 = functools.partial(VOP1, VOP1Op.V_PRNG_B32)
v_permlane16_swap_b32_e32 = functools.partial(VOP1, VOP1Op.V_PERMLANE16_SWAP_B32)
v_permlane32_swap_b32_e32 = functools.partial(VOP1, VOP1Op.V_PERMLANE32_SWAP_B32)
v_cvt_f32_bf16_e32 = functools.partial(VOP1, VOP1Op.V_CVT_F32_BF16)
cdna4_e32 = functools.partial(VOP1, VOP1Op.CDNA4)
v_cndmask_b32_e32 = functools.partial(VOP2, VOP2Op.V_CNDMASK_B32)
v_add_f32_e32 = functools.partial(VOP2, VOP2Op.V_ADD_F32)
v_sub_f32_e32 = functools.partial(VOP2, VOP2Op.V_SUB_F32)
v_subrev_f32_e32 = functools.partial(VOP2, VOP2Op.V_SUBREV_F32)
v_fmac_f64_e32 = functools.partial(VOP2, VOP2Op.V_FMAC_F64)
v_mul_f32_e32 = functools.partial(VOP2, VOP2Op.V_MUL_F32)
v_mul_i32_i24_e32 = functools.partial(VOP2, VOP2Op.V_MUL_I32_I24)
v_mul_hi_i32_i24_e32 = functools.partial(VOP2, VOP2Op.V_MUL_HI_I32_I24)
v_mul_u32_u24_e32 = functools.partial(VOP2, VOP2Op.V_MUL_U32_U24)
v_mul_hi_u32_u24_e32 = functools.partial(VOP2, VOP2Op.V_MUL_HI_U32_U24)
v_min_f32_e32 = functools.partial(VOP2, VOP2Op.V_MIN_F32)
v_max_f32_e32 = functools.partial(VOP2, VOP2Op.V_MAX_F32)
v_min_i32_e32 = functools.partial(VOP2, VOP2Op.V_MIN_I32)
v_max_i32_e32 = functools.partial(VOP2, VOP2Op.V_MAX_I32)
v_min_u32_e32 = functools.partial(VOP2, VOP2Op.V_MIN_U32)
v_max_u32_e32 = functools.partial(VOP2, VOP2Op.V_MAX_U32)
v_lshrrev_b32_e32 = functools.partial(VOP2, VOP2Op.V_LSHRREV_B32)
v_ashrrev_i32_e32 = functools.partial(VOP2, VOP2Op.V_ASHRREV_I32)
v_lshlrev_b32_e32 = functools.partial(VOP2, VOP2Op.V_LSHLREV_B32)
v_and_b32_e32 = functools.partial(VOP2, VOP2Op.V_AND_B32)
v_or_b32_e32 = functools.partial(VOP2, VOP2Op.V_OR_B32)
v_xor_b32_e32 = functools.partial(VOP2, VOP2Op.V_XOR_B32)
v_dot2c_f32_bf16_e32 = functools.partial(VOP2, VOP2Op.V_DOT2C_F32_BF16)
def v_fmamk_f32_e32(vdst, src0, K, vsrc1): return VOP2(VOP2Op.V_FMAMK_F32, vdst, src0, vsrc1, literal=K)
def v_fmaak_f32_e32(vdst, src0, vsrc1, K): return VOP2(VOP2Op.V_FMAAK_F32, vdst, src0, vsrc1, literal=K)
v_add_co_u32_e32 = functools.partial(VOP2, VOP2Op.V_ADD_CO_U32)
v_sub_co_u32_e32 = functools.partial(VOP2, VOP2Op.V_SUB_CO_U32)
v_subrev_co_u32_e32 = functools.partial(VOP2, VOP2Op.V_SUBREV_CO_U32)
v_addc_co_u32_e32 = functools.partial(VOP2, VOP2Op.V_ADDC_CO_U32)
v_subb_co_u32_e32 = functools.partial(VOP2, VOP2Op.V_SUBB_CO_U32)
v_subbrev_co_u32_e32 = functools.partial(VOP2, VOP2Op.V_SUBBREV_CO_U32)
v_add_f16_e32 = functools.partial(VOP2, VOP2Op.V_ADD_F16)
v_sub_f16_e32 = functools.partial(VOP2, VOP2Op.V_SUB_F16)
v_subrev_f16_e32 = functools.partial(VOP2, VOP2Op.V_SUBREV_F16)
v_mul_f16_e32 = functools.partial(VOP2, VOP2Op.V_MUL_F16)
v_mac_f16_e32 = functools.partial(VOP2, VOP2Op.V_MAC_F16)
v_madmk_f16_e32 = functools.partial(VOP2, VOP2Op.V_MADMK_F16)
v_madak_f16_e32 = functools.partial(VOP2, VOP2Op.V_MADAK_F16)
v_add_u16_e32 = functools.partial(VOP2, VOP2Op.V_ADD_U16)
v_sub_u16_e32 = functools.partial(VOP2, VOP2Op.V_SUB_U16)
v_subrev_u16_e32 = functools.partial(VOP2, VOP2Op.V_SUBREV_U16)
v_mul_lo_u16_e32 = functools.partial(VOP2, VOP2Op.V_MUL_LO_U16)
v_lshlrev_b16_e32 = functools.partial(VOP2, VOP2Op.V_LSHLREV_B16)
v_lshrrev_b16_e32 = functools.partial(VOP2, VOP2Op.V_LSHRREV_B16)
v_ashrrev_i16_e32 = functools.partial(VOP2, VOP2Op.V_ASHRREV_I16)
v_max_f16_e32 = functools.partial(VOP2, VOP2Op.V_MAX_F16)
v_min_f16_e32 = functools.partial(VOP2, VOP2Op.V_MIN_F16)
v_max_u16_e32 = functools.partial(VOP2, VOP2Op.V_MAX_U16)
v_max_i16_e32 = functools.partial(VOP2, VOP2Op.V_MAX_I16)
v_min_u16_e32 = functools.partial(VOP2, VOP2Op.V_MIN_U16)
v_min_i16_e32 = functools.partial(VOP2, VOP2Op.V_MIN_I16)
v_ldexp_f16_e32 = functools.partial(VOP2, VOP2Op.V_LDEXP_F16)
v_add_u32_e32 = functools.partial(VOP2, VOP2Op.V_ADD_U32)
v_sub_u32_e32 = functools.partial(VOP2, VOP2Op.V_SUB_U32)
v_subrev_u32_e32 = functools.partial(VOP2, VOP2Op.V_SUBREV_U32)
v_dot2c_f32_f16_e32 = functools.partial(VOP2, VOP2Op.V_DOT2C_F32_F16)
v_dot2c_i32_i16_e32 = functools.partial(VOP2, VOP2Op.V_DOT2C_I32_I16)
v_dot4c_i32_i8_e32 = functools.partial(VOP2, VOP2Op.V_DOT4C_I32_I8)
v_dot8c_i32_i4_e32 = functools.partial(VOP2, VOP2Op.V_DOT8C_I32_I4)
v_fmac_f32_e32 = functools.partial(VOP2, VOP2Op.V_FMAC_F32)
v_pk_fmac_f16_e32 = functools.partial(VOP2, VOP2Op.V_PK_FMAC_F16)
v_xnor_b32_e32 = functools.partial(VOP2, VOP2Op.V_XNOR_B32)
cdna4_e32 = functools.partial(VOP2, VOP2Op.CDNA4)
v_cmp_class_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_CLASS_F32)
v_cmpx_class_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_CLASS_F32)
v_cmp_class_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_CLASS_F64)
v_cmpx_class_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_CLASS_F64)
v_cmp_class_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_CLASS_F16)
v_cmpx_class_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_CLASS_F16)
v_cmp_f_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_F_F16)
v_cmp_lt_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_LT_F16)
v_cmp_eq_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_EQ_F16)
v_cmp_le_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_LE_F16)
v_cmp_gt_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_GT_F16)
v_cmp_lg_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_LG_F16)
v_cmp_ge_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_GE_F16)
v_cmp_o_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_O_F16)
v_cmp_u_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_U_F16)
v_cmp_nge_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_NGE_F16)
v_cmp_nlg_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_NLG_F16)
v_cmp_ngt_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_NGT_F16)
v_cmp_nle_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_NLE_F16)
v_cmp_neq_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_NEQ_F16)
v_cmp_nlt_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_NLT_F16)
v_cmp_tru_f16 = functools.partial(VOP3A, VOP3AOp.V_CMP_TRU_F16)
v_cmpx_f_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_F_F16)
v_cmpx_lt_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LT_F16)
v_cmpx_eq_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_EQ_F16)
v_cmpx_le_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LE_F16)
v_cmpx_gt_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GT_F16)
v_cmpx_lg_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LG_F16)
v_cmpx_ge_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GE_F16)
v_cmpx_o_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_O_F16)
v_cmpx_u_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_U_F16)
v_cmpx_nge_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NGE_F16)
v_cmpx_nlg_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NLG_F16)
v_cmpx_ngt_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NGT_F16)
v_cmpx_nle_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NLE_F16)
v_cmpx_neq_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NEQ_F16)
v_cmpx_nlt_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NLT_F16)
v_cmpx_tru_f16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_TRU_F16)
v_cmp_f_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_F_F32)
v_cmp_lt_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_LT_F32)
v_cmp_eq_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_EQ_F32)
v_cmp_le_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_LE_F32)
v_cmp_gt_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_GT_F32)
v_cmp_lg_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_LG_F32)
v_cmp_ge_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_GE_F32)
v_cmp_o_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_O_F32)
v_cmp_u_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_U_F32)
v_cmp_nge_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_NGE_F32)
v_cmp_nlg_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_NLG_F32)
v_cmp_ngt_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_NGT_F32)
v_cmp_nle_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_NLE_F32)
v_cmp_neq_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_NEQ_F32)
v_cmp_nlt_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_NLT_F32)
v_cmp_tru_f32 = functools.partial(VOP3A, VOP3AOp.V_CMP_TRU_F32)
v_cmpx_f_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_F_F32)
v_cmpx_lt_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LT_F32)
v_cmpx_eq_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_EQ_F32)
v_cmpx_le_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LE_F32)
v_cmpx_gt_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GT_F32)
v_cmpx_lg_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LG_F32)
v_cmpx_ge_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GE_F32)
v_cmpx_o_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_O_F32)
v_cmpx_u_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_U_F32)
v_cmpx_nge_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NGE_F32)
v_cmpx_nlg_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NLG_F32)
v_cmpx_ngt_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NGT_F32)
v_cmpx_nle_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NLE_F32)
v_cmpx_neq_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NEQ_F32)
v_cmpx_nlt_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NLT_F32)
v_cmpx_tru_f32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_TRU_F32)
v_cmp_f_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_F_F64)
v_cmp_lt_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_LT_F64)
v_cmp_eq_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_EQ_F64)
v_cmp_le_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_LE_F64)
v_cmp_gt_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_GT_F64)
v_cmp_lg_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_LG_F64)
v_cmp_ge_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_GE_F64)
v_cmp_o_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_O_F64)
v_cmp_u_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_U_F64)
v_cmp_nge_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_NGE_F64)
v_cmp_nlg_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_NLG_F64)
v_cmp_ngt_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_NGT_F64)
v_cmp_nle_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_NLE_F64)
v_cmp_neq_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_NEQ_F64)
v_cmp_nlt_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_NLT_F64)
v_cmp_tru_f64 = functools.partial(VOP3A, VOP3AOp.V_CMP_TRU_F64)
v_cmpx_f_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_F_F64)
v_cmpx_lt_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LT_F64)
v_cmpx_eq_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_EQ_F64)
v_cmpx_le_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LE_F64)
v_cmpx_gt_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GT_F64)
v_cmpx_lg_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LG_F64)
v_cmpx_ge_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GE_F64)
v_cmpx_o_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_O_F64)
v_cmpx_u_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_U_F64)
v_cmpx_nge_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NGE_F64)
v_cmpx_nlg_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NLG_F64)
v_cmpx_ngt_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NGT_F64)
v_cmpx_nle_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NLE_F64)
v_cmpx_neq_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NEQ_F64)
v_cmpx_nlt_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NLT_F64)
v_cmpx_tru_f64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_TRU_F64)
v_cmp_f_i16 = functools.partial(VOP3A, VOP3AOp.V_CMP_F_I16)
v_cmp_lt_i16 = functools.partial(VOP3A, VOP3AOp.V_CMP_LT_I16)
v_cmp_eq_i16 = functools.partial(VOP3A, VOP3AOp.V_CMP_EQ_I16)
v_cmp_le_i16 = functools.partial(VOP3A, VOP3AOp.V_CMP_LE_I16)
v_cmp_gt_i16 = functools.partial(VOP3A, VOP3AOp.V_CMP_GT_I16)
v_cmp_ne_i16 = functools.partial(VOP3A, VOP3AOp.V_CMP_NE_I16)
v_cmp_ge_i16 = functools.partial(VOP3A, VOP3AOp.V_CMP_GE_I16)
v_cmp_t_i16 = functools.partial(VOP3A, VOP3AOp.V_CMP_T_I16)
v_cmp_f_u16 = functools.partial(VOP3A, VOP3AOp.V_CMP_F_U16)
v_cmp_lt_u16 = functools.partial(VOP3A, VOP3AOp.V_CMP_LT_U16)
v_cmp_eq_u16 = functools.partial(VOP3A, VOP3AOp.V_CMP_EQ_U16)
v_cmp_le_u16 = functools.partial(VOP3A, VOP3AOp.V_CMP_LE_U16)
v_cmp_gt_u16 = functools.partial(VOP3A, VOP3AOp.V_CMP_GT_U16)
v_cmp_ne_u16 = functools.partial(VOP3A, VOP3AOp.V_CMP_NE_U16)
v_cmp_ge_u16 = functools.partial(VOP3A, VOP3AOp.V_CMP_GE_U16)
v_cmp_t_u16 = functools.partial(VOP3A, VOP3AOp.V_CMP_T_U16)
v_cmpx_f_i16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_F_I16)
v_cmpx_lt_i16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LT_I16)
v_cmpx_eq_i16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_EQ_I16)
v_cmpx_le_i16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LE_I16)
v_cmpx_gt_i16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GT_I16)
v_cmpx_ne_i16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NE_I16)
v_cmpx_ge_i16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GE_I16)
v_cmpx_t_i16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_T_I16)
v_cmpx_f_u16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_F_U16)
v_cmpx_lt_u16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LT_U16)
v_cmpx_eq_u16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_EQ_U16)
v_cmpx_le_u16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LE_U16)
v_cmpx_gt_u16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GT_U16)
v_cmpx_ne_u16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NE_U16)
v_cmpx_ge_u16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GE_U16)
v_cmpx_t_u16 = functools.partial(VOP3A, VOP3AOp.V_CMPX_T_U16)
v_cmp_f_i32 = functools.partial(VOP3A, VOP3AOp.V_CMP_F_I32)
v_cmp_lt_i32 = functools.partial(VOP3A, VOP3AOp.V_CMP_LT_I32)
v_cmp_eq_i32 = functools.partial(VOP3A, VOP3AOp.V_CMP_EQ_I32)
v_cmp_le_i32 = functools.partial(VOP3A, VOP3AOp.V_CMP_LE_I32)
v_cmp_gt_i32 = functools.partial(VOP3A, VOP3AOp.V_CMP_GT_I32)
v_cmp_ne_i32 = functools.partial(VOP3A, VOP3AOp.V_CMP_NE_I32)
v_cmp_ge_i32 = functools.partial(VOP3A, VOP3AOp.V_CMP_GE_I32)
v_cmp_t_i32 = functools.partial(VOP3A, VOP3AOp.V_CMP_T_I32)
v_cmp_f_u32 = functools.partial(VOP3A, VOP3AOp.V_CMP_F_U32)
v_cmp_lt_u32 = functools.partial(VOP3A, VOP3AOp.V_CMP_LT_U32)
v_cmp_eq_u32 = functools.partial(VOP3A, VOP3AOp.V_CMP_EQ_U32)
v_cmp_le_u32 = functools.partial(VOP3A, VOP3AOp.V_CMP_LE_U32)
v_cmp_gt_u32 = functools.partial(VOP3A, VOP3AOp.V_CMP_GT_U32)
v_cmp_ne_u32 = functools.partial(VOP3A, VOP3AOp.V_CMP_NE_U32)
v_cmp_ge_u32 = functools.partial(VOP3A, VOP3AOp.V_CMP_GE_U32)
v_cmp_t_u32 = functools.partial(VOP3A, VOP3AOp.V_CMP_T_U32)
v_cmpx_f_i32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_F_I32)
v_cmpx_lt_i32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LT_I32)
v_cmpx_eq_i32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_EQ_I32)
v_cmpx_le_i32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LE_I32)
v_cmpx_gt_i32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GT_I32)
v_cmpx_ne_i32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NE_I32)
v_cmpx_ge_i32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GE_I32)
v_cmpx_t_i32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_T_I32)
v_cmpx_f_u32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_F_U32)
v_cmpx_lt_u32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LT_U32)
v_cmpx_eq_u32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_EQ_U32)
v_cmpx_le_u32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LE_U32)
v_cmpx_gt_u32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GT_U32)
v_cmpx_ne_u32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NE_U32)
v_cmpx_ge_u32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GE_U32)
v_cmpx_t_u32 = functools.partial(VOP3A, VOP3AOp.V_CMPX_T_U32)
v_cmp_f_i64 = functools.partial(VOP3A, VOP3AOp.V_CMP_F_I64)
v_cmp_lt_i64 = functools.partial(VOP3A, VOP3AOp.V_CMP_LT_I64)
v_cmp_eq_i64 = functools.partial(VOP3A, VOP3AOp.V_CMP_EQ_I64)
v_cmp_le_i64 = functools.partial(VOP3A, VOP3AOp.V_CMP_LE_I64)
v_cmp_gt_i64 = functools.partial(VOP3A, VOP3AOp.V_CMP_GT_I64)
v_cmp_ne_i64 = functools.partial(VOP3A, VOP3AOp.V_CMP_NE_I64)
v_cmp_ge_i64 = functools.partial(VOP3A, VOP3AOp.V_CMP_GE_I64)
v_cmp_t_i64 = functools.partial(VOP3A, VOP3AOp.V_CMP_T_I64)
v_cmp_f_u64 = functools.partial(VOP3A, VOP3AOp.V_CMP_F_U64)
v_cmp_lt_u64 = functools.partial(VOP3A, VOP3AOp.V_CMP_LT_U64)
v_cmp_eq_u64 = functools.partial(VOP3A, VOP3AOp.V_CMP_EQ_U64)
v_cmp_le_u64 = functools.partial(VOP3A, VOP3AOp.V_CMP_LE_U64)
v_cmp_gt_u64 = functools.partial(VOP3A, VOP3AOp.V_CMP_GT_U64)
v_cmp_ne_u64 = functools.partial(VOP3A, VOP3AOp.V_CMP_NE_U64)
v_cmp_ge_u64 = functools.partial(VOP3A, VOP3AOp.V_CMP_GE_U64)
v_cmp_t_u64 = functools.partial(VOP3A, VOP3AOp.V_CMP_T_U64)
v_cmpx_f_i64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_F_I64)
v_cmpx_lt_i64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LT_I64)
v_cmpx_eq_i64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_EQ_I64)
v_cmpx_le_i64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LE_I64)
v_cmpx_gt_i64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GT_I64)
v_cmpx_ne_i64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NE_I64)
v_cmpx_ge_i64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GE_I64)
v_cmpx_t_i64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_T_I64)
v_cmpx_f_u64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_F_U64)
v_cmpx_lt_u64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LT_U64)
v_cmpx_eq_u64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_EQ_U64)
v_cmpx_le_u64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_LE_U64)
v_cmpx_gt_u64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GT_U64)
v_cmpx_ne_u64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_NE_U64)
v_cmpx_ge_u64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_GE_U64)
v_cmpx_t_u64 = functools.partial(VOP3A, VOP3AOp.V_CMPX_T_U64)
v_cndmask_b32 = functools.partial(VOP3A, VOP3AOp.V_CNDMASK_B32)
v_add_f32 = functools.partial(VOP3A, VOP3AOp.V_ADD_F32)
v_sub_f32 = functools.partial(VOP3A, VOP3AOp.V_SUB_F32)
v_subrev_f32 = functools.partial(VOP3A, VOP3AOp.V_SUBREV_F32)
v_fmac_f64 = functools.partial(VOP3A, VOP3AOp.V_FMAC_F64)
v_mul_f32 = functools.partial(VOP3A, VOP3AOp.V_MUL_F32)
v_mul_i32_i24 = functools.partial(VOP3A, VOP3AOp.V_MUL_I32_I24)
v_mul_hi_i32_i24 = functools.partial(VOP3A, VOP3AOp.V_MUL_HI_I32_I24)
v_mul_u32_u24 = functools.partial(VOP3A, VOP3AOp.V_MUL_U32_U24)
v_mul_hi_u32_u24 = functools.partial(VOP3A, VOP3AOp.V_MUL_HI_U32_U24)
v_min_f32 = functools.partial(VOP3A, VOP3AOp.V_MIN_F32)
v_max_f32 = functools.partial(VOP3A, VOP3AOp.V_MAX_F32)
v_min_i32 = functools.partial(VOP3A, VOP3AOp.V_MIN_I32)
v_max_i32 = functools.partial(VOP3A, VOP3AOp.V_MAX_I32)
v_min_u32 = functools.partial(VOP3A, VOP3AOp.V_MIN_U32)
v_max_u32 = functools.partial(VOP3A, VOP3AOp.V_MAX_U32)
v_lshrrev_b32 = functools.partial(VOP3A, VOP3AOp.V_LSHRREV_B32)
v_ashrrev_i32 = functools.partial(VOP3A, VOP3AOp.V_ASHRREV_I32)
v_lshlrev_b32 = functools.partial(VOP3A, VOP3AOp.V_LSHLREV_B32)
v_and_b32 = functools.partial(VOP3A, VOP3AOp.V_AND_B32)
v_or_b32 = functools.partial(VOP3A, VOP3AOp.V_OR_B32)
v_xor_b32 = functools.partial(VOP3A, VOP3AOp.V_XOR_B32)
v_dot2c_f32_bf16 = functools.partial(VOP3A, VOP3AOp.V_DOT2C_F32_BF16)
v_add_f16 = functools.partial(VOP3A, VOP3AOp.V_ADD_F16)
v_sub_f16 = functools.partial(VOP3A, VOP3AOp.V_SUB_F16)
v_subrev_f16 = functools.partial(VOP3A, VOP3AOp.V_SUBREV_F16)
v_mul_f16 = functools.partial(VOP3A, VOP3AOp.V_MUL_F16)
v_mac_f16 = functools.partial(VOP3A, VOP3AOp.V_MAC_F16)
v_add_u16 = functools.partial(VOP3A, VOP3AOp.V_ADD_U16)
v_sub_u16 = functools.partial(VOP3A, VOP3AOp.V_SUB_U16)
v_subrev_u16 = functools.partial(VOP3A, VOP3AOp.V_SUBREV_U16)
v_mul_lo_u16 = functools.partial(VOP3A, VOP3AOp.V_MUL_LO_U16)
v_lshlrev_b16 = functools.partial(VOP3A, VOP3AOp.V_LSHLREV_B16)
v_lshrrev_b16 = functools.partial(VOP3A, VOP3AOp.V_LSHRREV_B16)
v_ashrrev_i16 = functools.partial(VOP3A, VOP3AOp.V_ASHRREV_I16)
v_max_f16 = functools.partial(VOP3A, VOP3AOp.V_MAX_F16)
v_min_f16 = functools.partial(VOP3A, VOP3AOp.V_MIN_F16)
v_max_u16 = functools.partial(VOP3A, VOP3AOp.V_MAX_U16)
v_max_i16 = functools.partial(VOP3A, VOP3AOp.V_MAX_I16)
v_min_u16 = functools.partial(VOP3A, VOP3AOp.V_MIN_U16)
v_min_i16 = functools.partial(VOP3A, VOP3AOp.V_MIN_I16)
v_ldexp_f16 = functools.partial(VOP3A, VOP3AOp.V_LDEXP_F16)
v_add_u32 = functools.partial(VOP3A, VOP3AOp.V_ADD_U32)
v_sub_u32 = functools.partial(VOP3A, VOP3AOp.V_SUB_U32)
v_subrev_u32 = functools.partial(VOP3A, VOP3AOp.V_SUBREV_U32)
v_dot2c_f32_f16 = functools.partial(VOP3A, VOP3AOp.V_DOT2C_F32_F16)
v_dot2c_i32_i16 = functools.partial(VOP3A, VOP3AOp.V_DOT2C_I32_I16)
v_dot4c_i32_i8 = functools.partial(VOP3A, VOP3AOp.V_DOT4C_I32_I8)
v_dot8c_i32_i4 = functools.partial(VOP3A, VOP3AOp.V_DOT8C_I32_I4)
v_fmac_f32 = functools.partial(VOP3A, VOP3AOp.V_FMAC_F32)
v_pk_fmac_f16 = functools.partial(VOP3A, VOP3AOp.V_PK_FMAC_F16)
v_xnor_b32 = functools.partial(VOP3A, VOP3AOp.V_XNOR_B32)
v_nop = functools.partial(VOP3A, VOP3AOp.V_NOP)
v_mov_b32 = functools.partial(VOP3A, VOP3AOp.V_MOV_B32)
v_readfirstlane_b32 = functools.partial(VOP3A, VOP3AOp.V_READFIRSTLANE_B32)
v_cvt_i32_f64 = functools.partial(VOP3A, VOP3AOp.V_CVT_I32_F64)
v_cvt_f64_i32 = functools.partial(VOP3A, VOP3AOp.V_CVT_F64_I32)
v_cvt_f32_i32 = functools.partial(VOP3A, VOP3AOp.V_CVT_F32_I32)
v_cvt_f32_u32 = functools.partial(VOP3A, VOP3AOp.V_CVT_F32_U32)
v_cvt_u32_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_U32_F32)
v_cvt_i32_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_I32_F32)
v_cvt_f16_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_F16_F32)
v_cvt_f32_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_F32_F16)
v_cvt_rpi_i32_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_RPI_I32_F32)
v_cvt_flr_i32_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_FLR_I32_F32)
v_cvt_off_f32_i4 = functools.partial(VOP3A, VOP3AOp.V_CVT_OFF_F32_I4)
v_cvt_f32_f64 = functools.partial(VOP3A, VOP3AOp.V_CVT_F32_F64)
v_cvt_f64_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_F64_F32)
v_cvt_f32_ubyte0 = functools.partial(VOP3A, VOP3AOp.V_CVT_F32_UBYTE0)
v_cvt_f32_ubyte1 = functools.partial(VOP3A, VOP3AOp.V_CVT_F32_UBYTE1)
v_cvt_f32_ubyte2 = functools.partial(VOP3A, VOP3AOp.V_CVT_F32_UBYTE2)
v_cvt_f32_ubyte3 = functools.partial(VOP3A, VOP3AOp.V_CVT_F32_UBYTE3)
v_cvt_u32_f64 = functools.partial(VOP3A, VOP3AOp.V_CVT_U32_F64)
v_cvt_f64_u32 = functools.partial(VOP3A, VOP3AOp.V_CVT_F64_U32)
v_trunc_f64 = functools.partial(VOP3A, VOP3AOp.V_TRUNC_F64)
v_ceil_f64 = functools.partial(VOP3A, VOP3AOp.V_CEIL_F64)
v_rndne_f64 = functools.partial(VOP3A, VOP3AOp.V_RNDNE_F64)
v_floor_f64 = functools.partial(VOP3A, VOP3AOp.V_FLOOR_F64)
v_fract_f32 = functools.partial(VOP3A, VOP3AOp.V_FRACT_F32)
v_trunc_f32 = functools.partial(VOP3A, VOP3AOp.V_TRUNC_F32)
v_ceil_f32 = functools.partial(VOP3A, VOP3AOp.V_CEIL_F32)
v_rndne_f32 = functools.partial(VOP3A, VOP3AOp.V_RNDNE_F32)
v_floor_f32 = functools.partial(VOP3A, VOP3AOp.V_FLOOR_F32)
v_exp_f32 = functools.partial(VOP3A, VOP3AOp.V_EXP_F32)
v_log_f32 = functools.partial(VOP3A, VOP3AOp.V_LOG_F32)
v_rcp_f32 = functools.partial(VOP3A, VOP3AOp.V_RCP_F32)
v_rcp_iflag_f32 = functools.partial(VOP3A, VOP3AOp.V_RCP_IFLAG_F32)
v_rsq_f32 = functools.partial(VOP3A, VOP3AOp.V_RSQ_F32)
v_rcp_f64 = functools.partial(VOP3A, VOP3AOp.V_RCP_F64)
v_rsq_f64 = functools.partial(VOP3A, VOP3AOp.V_RSQ_F64)
v_sqrt_f32 = functools.partial(VOP3A, VOP3AOp.V_SQRT_F32)
v_sqrt_f64 = functools.partial(VOP3A, VOP3AOp.V_SQRT_F64)
v_sin_f32 = functools.partial(VOP3A, VOP3AOp.V_SIN_F32)
v_cos_f32 = functools.partial(VOP3A, VOP3AOp.V_COS_F32)
v_not_b32 = functools.partial(VOP3A, VOP3AOp.V_NOT_B32)
v_bfrev_b32 = functools.partial(VOP3A, VOP3AOp.V_BFREV_B32)
v_ffbh_u32 = functools.partial(VOP3A, VOP3AOp.V_FFBH_U32)
v_ffbl_b32 = functools.partial(VOP3A, VOP3AOp.V_FFBL_B32)
v_ffbh_i32 = functools.partial(VOP3A, VOP3AOp.V_FFBH_I32)
v_frexp_exp_i32_f64 = functools.partial(VOP3A, VOP3AOp.V_FREXP_EXP_I32_F64)
v_frexp_mant_f64 = functools.partial(VOP3A, VOP3AOp.V_FREXP_MANT_F64)
v_fract_f64 = functools.partial(VOP3A, VOP3AOp.V_FRACT_F64)
v_frexp_exp_i32_f32 = functools.partial(VOP3A, VOP3AOp.V_FREXP_EXP_I32_F32)
v_frexp_mant_f32 = functools.partial(VOP3A, VOP3AOp.V_FREXP_MANT_F32)
v_clrexcp = functools.partial(VOP3A, VOP3AOp.V_CLREXCP)
v_mov_b64 = functools.partial(VOP3A, VOP3AOp.V_MOV_B64)
v_cvt_f16_u16 = functools.partial(VOP3A, VOP3AOp.V_CVT_F16_U16)
v_cvt_f16_i16 = functools.partial(VOP3A, VOP3AOp.V_CVT_F16_I16)
v_cvt_u16_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_U16_F16)
v_cvt_i16_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_I16_F16)
v_rcp_f16 = functools.partial(VOP3A, VOP3AOp.V_RCP_F16)
v_sqrt_f16 = functools.partial(VOP3A, VOP3AOp.V_SQRT_F16)
v_rsq_f16 = functools.partial(VOP3A, VOP3AOp.V_RSQ_F16)
v_log_f16 = functools.partial(VOP3A, VOP3AOp.V_LOG_F16)
v_exp_f16 = functools.partial(VOP3A, VOP3AOp.V_EXP_F16)
v_mad_i32_i24 = functools.partial(VOP3A, VOP3AOp.V_MAD_I32_I24)
v_mad_u32_u24 = functools.partial(VOP3A, VOP3AOp.V_MAD_U32_U24)
v_cubeid_f32 = functools.partial(VOP3A, VOP3AOp.V_CUBEID_F32)
v_cubesc_f32 = functools.partial(VOP3A, VOP3AOp.V_CUBESC_F32)
v_cubetc_f32 = functools.partial(VOP3A, VOP3AOp.V_CUBETC_F32)
v_cubema_f32 = functools.partial(VOP3A, VOP3AOp.V_CUBEMA_F32)
v_bfe_u32 = functools.partial(VOP3A, VOP3AOp.V_BFE_U32)
v_bfe_i32 = functools.partial(VOP3A, VOP3AOp.V_BFE_I32)
v_bfi_b32 = functools.partial(VOP3A, VOP3AOp.V_BFI_B32)
v_fma_f32 = functools.partial(VOP3A, VOP3AOp.V_FMA_F32)
v_fma_f64 = functools.partial(VOP3A, VOP3AOp.V_FMA_F64)
v_lerp_u8 = functools.partial(VOP3A, VOP3AOp.V_LERP_U8)
v_alignbit_b32 = functools.partial(VOP3A, VOP3AOp.V_ALIGNBIT_B32)
v_alignbyte_b32 = functools.partial(VOP3A, VOP3AOp.V_ALIGNBYTE_B32)
v_min3_f32 = functools.partial(VOP3A, VOP3AOp.V_MIN3_F32)
v_min3_i32 = functools.partial(VOP3A, VOP3AOp.V_MIN3_I32)
v_min3_u32 = functools.partial(VOP3A, VOP3AOp.V_MIN3_U32)
v_max3_f32 = functools.partial(VOP3A, VOP3AOp.V_MAX3_F32)
v_max3_i32 = functools.partial(VOP3A, VOP3AOp.V_MAX3_I32)
v_max3_u32 = functools.partial(VOP3A, VOP3AOp.V_MAX3_U32)
v_med3_f32 = functools.partial(VOP3A, VOP3AOp.V_MED3_F32)
v_med3_i32 = functools.partial(VOP3A, VOP3AOp.V_MED3_I32)
v_med3_u32 = functools.partial(VOP3A, VOP3AOp.V_MED3_U32)
v_sad_u8 = functools.partial(VOP3A, VOP3AOp.V_SAD_U8)
v_sad_hi_u8 = functools.partial(VOP3A, VOP3AOp.V_SAD_HI_U8)
v_sad_u16 = functools.partial(VOP3A, VOP3AOp.V_SAD_U16)
v_sad_u32 = functools.partial(VOP3A, VOP3AOp.V_SAD_U32)
v_cvt_pk_u8_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_PK_U8_F32)
v_div_fixup_f32 = functools.partial(VOP3A, VOP3AOp.V_DIV_FIXUP_F32)
v_div_fixup_f64 = functools.partial(VOP3A, VOP3AOp.V_DIV_FIXUP_F64)
v_div_fmas_f32 = functools.partial(VOP3A, VOP3AOp.V_DIV_FMAS_F32)
v_div_fmas_f64 = functools.partial(VOP3A, VOP3AOp.V_DIV_FMAS_F64)
v_msad_u8 = functools.partial(VOP3A, VOP3AOp.V_MSAD_U8)
v_qsad_pk_u16_u8 = functools.partial(VOP3A, VOP3AOp.V_QSAD_PK_U16_U8)
v_mqsad_pk_u16_u8 = functools.partial(VOP3A, VOP3AOp.V_MQSAD_PK_U16_U8)
v_mqsad_u32_u8 = functools.partial(VOP3A, VOP3AOp.V_MQSAD_U32_U8)
v_mad_legacy_f16 = functools.partial(VOP3A, VOP3AOp.V_MAD_LEGACY_F16)
v_mad_legacy_u16 = functools.partial(VOP3A, VOP3AOp.V_MAD_LEGACY_U16)
v_mad_legacy_i16 = functools.partial(VOP3A, VOP3AOp.V_MAD_LEGACY_I16)
v_perm_b32 = functools.partial(VOP3A, VOP3AOp.V_PERM_B32)
v_fma_legacy_f16 = functools.partial(VOP3A, VOP3AOp.V_FMA_LEGACY_F16)
v_div_fixup_legacy_f16 = functools.partial(VOP3A, VOP3AOp.V_DIV_FIXUP_LEGACY_F16)
v_cvt_pkaccum_u8_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_PKACCUM_U8_F32)
v_mad_u32_u16 = functools.partial(VOP3A, VOP3AOp.V_MAD_U32_U16)
v_mad_i32_i16 = functools.partial(VOP3A, VOP3AOp.V_MAD_I32_I16)
v_xad_u32 = functools.partial(VOP3A, VOP3AOp.V_XAD_U32)
v_min3_f16 = functools.partial(VOP3A, VOP3AOp.V_MIN3_F16)
v_min3_i16 = functools.partial(VOP3A, VOP3AOp.V_MIN3_I16)
v_min3_u16 = functools.partial(VOP3A, VOP3AOp.V_MIN3_U16)
v_max3_f16 = functools.partial(VOP3A, VOP3AOp.V_MAX3_F16)
v_max3_i16 = functools.partial(VOP3A, VOP3AOp.V_MAX3_I16)
v_max3_u16 = functools.partial(VOP3A, VOP3AOp.V_MAX3_U16)
v_med3_f16 = functools.partial(VOP3A, VOP3AOp.V_MED3_F16)
v_med3_i16 = functools.partial(VOP3A, VOP3AOp.V_MED3_I16)
v_med3_u16 = functools.partial(VOP3A, VOP3AOp.V_MED3_U16)
v_lshl_add_u32 = functools.partial(VOP3A, VOP3AOp.V_LSHL_ADD_U32)
v_add_lshl_u32 = functools.partial(VOP3A, VOP3AOp.V_ADD_LSHL_U32)
v_add3_u32 = functools.partial(VOP3A, VOP3AOp.V_ADD3_U32)
v_lshl_or_b32 = functools.partial(VOP3A, VOP3AOp.V_LSHL_OR_B32)
v_and_or_b32 = functools.partial(VOP3A, VOP3AOp.V_AND_OR_B32)
v_or3_b32 = functools.partial(VOP3A, VOP3AOp.V_OR3_B32)
v_mad_f16 = functools.partial(VOP3A, VOP3AOp.V_MAD_F16)
v_mad_u16 = functools.partial(VOP3A, VOP3AOp.V_MAD_U16)
v_mad_i16 = functools.partial(VOP3A, VOP3AOp.V_MAD_I16)
v_fma_f16 = functools.partial(VOP3A, VOP3AOp.V_FMA_F16)
v_div_fixup_f16 = functools.partial(VOP3A, VOP3AOp.V_DIV_FIXUP_F16)
v_lshl_add_u64 = functools.partial(VOP3A, VOP3AOp.V_LSHL_ADD_U64)
v_bitop3_b16 = functools.partial(VOP3A, VOP3AOp.V_BITOP3_B16)
v_bitop3_b32 = functools.partial(VOP3A, VOP3AOp.V_BITOP3_B32)
v_cvt_scalef32_pk_fp8_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_FP8_F32)
v_cvt_scalef32_pk_bf8_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_BF8_F32)
v_cvt_scalef32_sr_fp8_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_FP8_F32)
v_cvt_scalef32_sr_bf8_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_BF8_F32)
v_cvt_scalef32_pk_f32_fp8 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_F32_FP8)
v_cvt_scalef32_pk_f32_bf8 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_F32_BF8)
v_cvt_scalef32_f32_fp8 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_F32_FP8)
v_cvt_scalef32_f32_bf8 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_F32_BF8)
v_cvt_scalef32_pk_fp4_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_FP4_F32)
v_cvt_scalef32_sr_pk_fp4_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_PK_FP4_F32)
v_cvt_scalef32_pk_f32_fp4 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_F32_FP4)
v_cvt_scalef32_pk_fp8_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_FP8_F16)
v_cvt_scalef32_pk_bf8_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_BF8_F16)
v_cvt_scalef32_sr_fp8_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_FP8_F16)
v_cvt_scalef32_sr_bf8_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_BF8_F16)
v_cvt_scalef32_pk_fp8_bf16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_FP8_BF16)
v_cvt_scalef32_pk_bf8_bf16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_BF8_BF16)
v_cvt_scalef32_sr_fp8_bf16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_FP8_BF16)
v_cvt_scalef32_sr_bf8_bf16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_BF8_BF16)
v_cvt_scalef32_pk_f16_fp8 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_F16_FP8)
v_cvt_scalef32_pk_f16_bf8 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_F16_BF8)
v_cvt_scalef32_f16_fp8 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_F16_FP8)
v_cvt_scalef32_f16_bf8 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_F16_BF8)
v_cvt_scalef32_pk_fp4_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_FP4_F16)
v_cvt_scalef32_pk_fp4_bf16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_FP4_BF16)
v_cvt_scalef32_sr_pk_fp4_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_PK_FP4_F16)
v_cvt_scalef32_sr_pk_fp4_bf16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_PK_FP4_BF16)
v_cvt_scalef32_pk_f16_fp4 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_F16_FP4)
v_cvt_scalef32_pk_bf16_fp4 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_BF16_FP4)
v_cvt_scalef32_2xpk16_fp6_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_2XPK16_FP6_F32)
v_cvt_scalef32_2xpk16_bf6_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_2XPK16_BF6_F32)
v_cvt_scalef32_sr_pk32_fp6_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_PK32_FP6_F32)
v_cvt_scalef32_sr_pk32_bf6_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_PK32_BF6_F32)
v_cvt_scalef32_pk32_f32_fp6 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK32_F32_FP6)
v_cvt_scalef32_pk32_f32_bf6 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK32_F32_BF6)
cdna4 = functools.partial(VOP3A, VOP3AOp.CDNA4)
v_cvt_scalef32_pk32_fp6_bf16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK32_FP6_BF16)
v_cvt_scalef32_pk32_bf6_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK32_BF6_F16)
v_cvt_scalef32_pk32_bf6_bf16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK32_BF6_BF16)
v_cvt_scalef32_sr_pk32_fp6_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_PK32_FP6_F16)
v_cvt_scalef32_sr_pk32_fp6_bf16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_PK32_FP6_BF16)
v_cvt_scalef32_sr_pk32_bf6_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_PK32_BF6_F16)
v_cvt_scalef32_sr_pk32_bf6_bf16 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_SR_PK32_BF6_BF16)
v_cvt_scalef32_pk32_f16_fp6 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK32_F16_FP6)
v_cvt_scalef32_pk32_bf16_fp6 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK32_BF16_FP6)
v_cvt_scalef32_pk32_f16_bf6 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK32_F16_BF6)
v_cvt_scalef32_pk32_bf16_bf6 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK32_BF16_BF6)
v_ashr_pk_i8_i32 = functools.partial(VOP3A, VOP3AOp.V_ASHR_PK_I8_I32)
v_ashr_pk_u8_i32 = functools.partial(VOP3A, VOP3AOp.V_ASHR_PK_U8_I32)
v_cvt_pk_f16_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_PK_F16_F32)
v_cvt_pk_bf16_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_PK_BF16_F32)
v_cvt_scalef32_pk_bf16_fp8 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_BF16_FP8)
v_cvt_scalef32_pk_bf16_bf8 = functools.partial(VOP3A, VOP3AOp.V_CVT_SCALEF32_PK_BF16_BF8)
v_add_f64 = functools.partial(VOP3A, VOP3AOp.V_ADD_F64)
v_mul_f64 = functools.partial(VOP3A, VOP3AOp.V_MUL_F64)
v_min_f64 = functools.partial(VOP3A, VOP3AOp.V_MIN_F64)
v_max_f64 = functools.partial(VOP3A, VOP3AOp.V_MAX_F64)
v_ldexp_f64 = functools.partial(VOP3A, VOP3AOp.V_LDEXP_F64)
v_mul_lo_u32 = functools.partial(VOP3A, VOP3AOp.V_MUL_LO_U32)
v_mul_hi_u32 = functools.partial(VOP3A, VOP3AOp.V_MUL_HI_U32)
v_mul_hi_i32 = functools.partial(VOP3A, VOP3AOp.V_MUL_HI_I32)
v_ldexp_f32 = functools.partial(VOP3A, VOP3AOp.V_LDEXP_F32)
v_readlane_b32 = functools.partial(VOP3A, VOP3AOp.V_READLANE_B32)
v_writelane_b32 = functools.partial(VOP3A, VOP3AOp.V_WRITELANE_B32)
v_bcnt_u32_b32 = functools.partial(VOP3A, VOP3AOp.V_BCNT_U32_B32)
v_mbcnt_lo_u32_b32 = functools.partial(VOP3A, VOP3AOp.V_MBCNT_LO_U32_B32)
v_mbcnt_hi_u32_b32 = functools.partial(VOP3A, VOP3AOp.V_MBCNT_HI_U32_B32)
v_lshlrev_b64 = functools.partial(VOP3A, VOP3AOp.V_LSHLREV_B64)
v_lshrrev_b64 = functools.partial(VOP3A, VOP3AOp.V_LSHRREV_B64)
v_ashrrev_i64 = functools.partial(VOP3A, VOP3AOp.V_ASHRREV_I64)
v_trig_preop_f64 = functools.partial(VOP3A, VOP3AOp.V_TRIG_PREOP_F64)
v_bfm_b32 = functools.partial(VOP3A, VOP3AOp.V_BFM_B32)
v_cvt_pknorm_i16_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_PKNORM_I16_F32)
v_cvt_pknorm_u16_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_PKNORM_U16_F32)
v_cvt_pkrtz_f16_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_PKRTZ_F16_F32)
v_cvt_pk_u16_u32 = functools.partial(VOP3A, VOP3AOp.V_CVT_PK_U16_U32)
v_cvt_pk_i16_i32 = functools.partial(VOP3A, VOP3AOp.V_CVT_PK_I16_I32)
v_cvt_pknorm_i16_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_PKNORM_I16_F16)
v_cvt_pknorm_u16_f16 = functools.partial(VOP3A, VOP3AOp.V_CVT_PKNORM_U16_F16)
v_add_i32 = functools.partial(VOP3A, VOP3AOp.V_ADD_I32)
v_sub_i32 = functools.partial(VOP3A, VOP3AOp.V_SUB_I32)
v_add_i16 = functools.partial(VOP3A, VOP3AOp.V_ADD_I16)
v_sub_i16 = functools.partial(VOP3A, VOP3AOp.V_SUB_I16)
v_pack_b32_f16 = functools.partial(VOP3A, VOP3AOp.V_PACK_B32_F16)
v_mul_legacy_f32 = functools.partial(VOP3A, VOP3AOp.V_MUL_LEGACY_F32)
v_cvt_pk_fp8_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_PK_FP8_F32)
v_cvt_pk_bf8_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_PK_BF8_F32)
v_cvt_sr_fp8_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SR_FP8_F32)
v_cvt_sr_bf8_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SR_BF8_F32)
v_cvt_sr_f16_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SR_F16_F32)
v_cvt_sr_bf16_f32 = functools.partial(VOP3A, VOP3AOp.V_CVT_SR_BF16_F32)
v_minimum3_f32 = functools.partial(VOP3A, VOP3AOp.V_MINIMUM3_F32)
v_maximum3_f32 = functools.partial(VOP3A, VOP3AOp.V_MAXIMUM3_F32)
v_add_co_u32 = functools.partial(VOP3B, VOP3BOp.V_ADD_CO_U32)
v_sub_co_u32 = functools.partial(VOP3B, VOP3BOp.V_SUB_CO_U32)
v_subrev_co_u32 = functools.partial(VOP3B, VOP3BOp.V_SUBREV_CO_U32)
v_addc_co_u32 = functools.partial(VOP3B, VOP3BOp.V_ADDC_CO_U32)
v_subb_co_u32 = functools.partial(VOP3B, VOP3BOp.V_SUBB_CO_U32)
v_subbrev_co_u32 = functools.partial(VOP3B, VOP3BOp.V_SUBBREV_CO_U32)
v_div_scale_f32 = functools.partial(VOP3B, VOP3BOp.V_DIV_SCALE_F32)
v_div_scale_f64 = functools.partial(VOP3B, VOP3BOp.V_DIV_SCALE_F64)
v_mad_u64_u32 = functools.partial(VOP3B, VOP3BOp.V_MAD_U64_U32)
v_mad_i64_i32 = functools.partial(VOP3B, VOP3BOp.V_MAD_I64_I32)
cdna4 = functools.partial(VOP3B, VOP3BOp.CDNA4)
v_pk_mad_i16 = functools.partial(VOP3P, VOP3POp.V_PK_MAD_I16)
v_pk_mul_lo_u16 = functools.partial(VOP3P, VOP3POp.V_PK_MUL_LO_U16)
v_pk_add_i16 = functools.partial(VOP3P, VOP3POp.V_PK_ADD_I16)
v_pk_sub_i16 = functools.partial(VOP3P, VOP3POp.V_PK_SUB_I16)
v_pk_lshlrev_b16 = functools.partial(VOP3P, VOP3POp.V_PK_LSHLREV_B16)
v_pk_lshrrev_b16 = functools.partial(VOP3P, VOP3POp.V_PK_LSHRREV_B16)
v_pk_ashrrev_i16 = functools.partial(VOP3P, VOP3POp.V_PK_ASHRREV_I16)
v_pk_max_i16 = functools.partial(VOP3P, VOP3POp.V_PK_MAX_I16)
v_pk_min_i16 = functools.partial(VOP3P, VOP3POp.V_PK_MIN_I16)
v_pk_mad_u16 = functools.partial(VOP3P, VOP3POp.V_PK_MAD_U16)
v_pk_add_u16 = functools.partial(VOP3P, VOP3POp.V_PK_ADD_U16)
v_pk_sub_u16 = functools.partial(VOP3P, VOP3POp.V_PK_SUB_U16)
v_pk_max_u16 = functools.partial(VOP3P, VOP3POp.V_PK_MAX_U16)
v_pk_min_u16 = functools.partial(VOP3P, VOP3POp.V_PK_MIN_U16)
v_pk_fma_f16 = functools.partial(VOP3P, VOP3POp.V_PK_FMA_F16)
v_pk_add_f16 = functools.partial(VOP3P, VOP3POp.V_PK_ADD_F16)
v_pk_mul_f16 = functools.partial(VOP3P, VOP3POp.V_PK_MUL_F16)
v_pk_min_f16 = functools.partial(VOP3P, VOP3POp.V_PK_MIN_F16)
v_pk_max_f16 = functools.partial(VOP3P, VOP3POp.V_PK_MAX_F16)
v_dot2_f32_bf16 = functools.partial(VOP3P, VOP3POp.V_DOT2_F32_BF16)
v_pk_minimum3_f16 = functools.partial(VOP3P, VOP3POp.V_PK_MINIMUM3_F16)
v_pk_maximum3_f16 = functools.partial(VOP3P, VOP3POp.V_PK_MAXIMUM3_F16)
v_mad_mix_f32 = functools.partial(VOP3P, VOP3POp.V_MAD_MIX_F32)
v_mad_mixlo_f16 = functools.partial(VOP3P, VOP3POp.V_MAD_MIXLO_F16)
v_mad_mixhi_f16 = functools.partial(VOP3P, VOP3POp.V_MAD_MIXHI_F16)
v_dot2_f32_f16 = functools.partial(VOP3P, VOP3POp.V_DOT2_F32_F16)
v_dot2_i32_i16 = functools.partial(VOP3P, VOP3POp.V_DOT2_I32_I16)
v_dot2_u32_u16 = functools.partial(VOP3P, VOP3POp.V_DOT2_U32_U16)
v_dot4_i32_i8 = functools.partial(VOP3P, VOP3POp.V_DOT4_I32_I8)
v_dot4_u32_u8 = functools.partial(VOP3P, VOP3POp.V_DOT4_U32_U8)
v_dot8_i32_i4 = functools.partial(VOP3P, VOP3POp.V_DOT8_I32_I4)
v_dot8_u32_u4 = functools.partial(VOP3P, VOP3POp.V_DOT8_U32_U4)
v_mfma_f32_16x16x128_f8f6f4 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X128_F8F6F4)
v_mfma_f32_32x32x64_f8f6f4 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X64_F8F6F4)
v_pk_fma_f32 = functools.partial(VOP3P, VOP3POp.V_PK_FMA_F32)
v_pk_mul_f32 = functools.partial(VOP3P, VOP3POp.V_PK_MUL_F32)
v_pk_add_f32 = functools.partial(VOP3P, VOP3POp.V_PK_ADD_F32)
v_pk_mov_b32 = functools.partial(VOP3P, VOP3POp.V_PK_MOV_B32)
v_mfma_f32_16x16x32_bf16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X32_BF16)
v_mfma_i32_16x16x64_i8 = functools.partial(VOP3P, VOP3POp.V_MFMA_I32_16X16X64_I8)
v_mfma_f32_32x32x16_bf16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X16_BF16)
v_mfma_i32_32x32x32_i8 = functools.partial(VOP3P, VOP3POp.V_MFMA_I32_32X32X32_I8)
v_smfmac_f32_16x16x64_bf16 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_16X16X64_BF16)
v_smfmac_i32_16x16x128_i8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_I32_16X16X128_I8)
v_smfmac_f32_16x16x128_bf8_bf8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_16X16X128_BF8_BF8)
v_smfmac_f32_16x16x128_bf8_fp8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_16X16X128_BF8_FP8)
v_smfmac_f32_16x16x128_fp8_bf8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_16X16X128_FP8_BF8)
v_mfma_f32_16x16x8_xf32 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X8_XF32)
v_mfma_f32_32x32x4_xf32 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X4_XF32)
v_mfma_f32_32x32x1_2b_f32 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X1_2B_F32)
v_mfma_f32_16x16x1_4b_f32 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X1_4B_F32)
v_mfma_f32_4x4x1_16b_f32 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_4X4X1_16B_F32)
v_smfmac_f32_16x16x128_fp8_fp8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_16X16X128_FP8_FP8)
v_mfma_f32_32x32x2_f32 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X2_F32)
v_mfma_f32_16x16x4_f32 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X4_F32)
v_smfmac_f32_32x32x32_bf16 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_32X32X32_BF16)
v_smfmac_i32_32x32x64_i8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_I32_32X32X64_I8)
v_mfma_f32_32x32x4_2b_f16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X4_2B_F16)
v_mfma_f32_16x16x4_4b_f16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X4_4B_F16)
v_mfma_f32_4x4x4_16b_f16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_4X4X4_16B_F16)
v_smfmac_f32_32x32x64_bf8_bf8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_32X32X64_BF8_BF8)
v_mfma_f32_32x32x8_f16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X8_F16)
v_mfma_f32_16x16x16_f16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X16_F16)
v_smfmac_f32_32x32x64_bf8_fp8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_32X32X64_BF8_FP8)
v_smfmac_f32_32x32x64_fp8_bf8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_32X32X64_FP8_BF8)
v_mfma_i32_32x32x4_2b_i8 = functools.partial(VOP3P, VOP3POp.V_MFMA_I32_32X32X4_2B_I8)
v_mfma_i32_16x16x4_4b_i8 = functools.partial(VOP3P, VOP3POp.V_MFMA_I32_16X16X4_4B_I8)
v_mfma_i32_4x4x4_16b_i8 = functools.partial(VOP3P, VOP3POp.V_MFMA_I32_4X4X4_16B_I8)
v_smfmac_f32_32x32x64_fp8_fp8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_32X32X64_FP8_FP8)
v_mfma_f32_16x16x32_f16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X32_F16)
v_mfma_f32_32x32x16_f16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X16_F16)
v_mfma_i32_32x32x16_i8 = functools.partial(VOP3P, VOP3POp.V_MFMA_I32_32X32X16_I8)
v_mfma_i32_16x16x32_i8 = functools.partial(VOP3P, VOP3POp.V_MFMA_I32_16X16X32_I8)
v_accvgpr_read = functools.partial(VOP3P, VOP3POp.V_ACCVGPR_READ)
v_accvgpr_write = functools.partial(VOP3P, VOP3POp.V_ACCVGPR_WRITE)
v_smfmac_f32_16x16x64_f16 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_16X16X64_F16)
v_smfmac_f32_32x32x32_f16 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_32X32X32_F16)
v_mfma_f32_32x32x4_2b_bf16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X4_2B_BF16)
v_mfma_f32_16x16x4_4b_bf16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X4_4B_BF16)
v_mfma_f32_4x4x4_16b_bf16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_4X4X4_16B_BF16)
v_mfma_f32_32x32x8_bf16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X8_BF16)
v_mfma_f32_16x16x16_bf16 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X16_BF16)
v_smfmac_f32_16x16x32_f16 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_16X16X32_F16)
v_smfmac_f32_32x32x16_f16 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_32X32X16_F16)
v_smfmac_f32_16x16x32_bf16 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_16X16X32_BF16)
v_smfmac_f32_32x32x16_bf16 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_32X32X16_BF16)
v_smfmac_i32_16x16x64_i8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_I32_16X16X64_I8)
v_smfmac_i32_32x32x32_i8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_I32_32X32X32_I8)
v_mfma_f64_16x16x4_f64 = functools.partial(VOP3P, VOP3POp.V_MFMA_F64_16X16X4_F64)
v_mfma_f64_4x4x4_4b_f64 = functools.partial(VOP3P, VOP3POp.V_MFMA_F64_4X4X4_4B_F64)
v_mfma_f32_16x16x32_bf8_bf8 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X32_BF8_BF8)
v_mfma_f32_16x16x32_bf8_fp8 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X32_BF8_FP8)
v_mfma_f32_16x16x32_fp8_bf8 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X32_FP8_BF8)
v_mfma_f32_16x16x32_fp8_fp8 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_16X16X32_FP8_FP8)
v_mfma_f32_32x32x16_bf8_bf8 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X16_BF8_BF8)
v_mfma_f32_32x32x16_bf8_fp8 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X16_BF8_FP8)
v_mfma_f32_32x32x16_fp8_bf8 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X16_FP8_BF8)
v_mfma_f32_32x32x16_fp8_fp8 = functools.partial(VOP3P, VOP3POp.V_MFMA_F32_32X32X16_FP8_FP8)
v_smfmac_f32_16x16x64_bf8_bf8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_16X16X64_BF8_BF8)
v_smfmac_f32_16x16x64_bf8_fp8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_16X16X64_BF8_FP8)
v_smfmac_f32_16x16x64_fp8_bf8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_16X16X64_FP8_BF8)
v_smfmac_f32_16x16x64_fp8_fp8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_16X16X64_FP8_FP8)
v_smfmac_f32_32x32x32_bf8_bf8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_32X32X32_BF8_BF8)
v_smfmac_f32_32x32x32_bf8_fp8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_32X32X32_BF8_FP8)
v_smfmac_f32_32x32x32_fp8_bf8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_32X32X32_FP8_BF8)
v_smfmac_f32_32x32x32_fp8_fp8 = functools.partial(VOP3P, VOP3POp.V_SMFMAC_F32_32X32X32_FP8_FP8)
cdna4 = functools.partial(VOP3P, VOP3POp.CDNA4)
v_cmp_class_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_CLASS_F32)
v_cmpx_class_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_CLASS_F32)
v_cmp_class_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_CLASS_F64)
v_cmpx_class_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_CLASS_F64)
v_cmp_class_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_CLASS_F16)
v_cmpx_class_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_CLASS_F16)
v_cmp_f_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_F_F16)
v_cmp_lt_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LT_F16)
v_cmp_eq_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_EQ_F16)
v_cmp_le_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LE_F16)
v_cmp_gt_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GT_F16)
v_cmp_lg_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LG_F16)
v_cmp_ge_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GE_F16)
v_cmp_o_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_O_F16)
v_cmp_u_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_U_F16)
v_cmp_nge_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NGE_F16)
v_cmp_nlg_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NLG_F16)
v_cmp_ngt_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NGT_F16)
v_cmp_nle_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NLE_F16)
v_cmp_neq_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NEQ_F16)
v_cmp_nlt_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NLT_F16)
v_cmp_tru_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_TRU_F16)
v_cmpx_f_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_F_F16)
v_cmpx_lt_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LT_F16)
v_cmpx_eq_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_EQ_F16)
v_cmpx_le_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LE_F16)
v_cmpx_gt_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GT_F16)
v_cmpx_lg_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LG_F16)
v_cmpx_ge_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GE_F16)
v_cmpx_o_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_O_F16)
v_cmpx_u_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_U_F16)
v_cmpx_nge_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NGE_F16)
v_cmpx_nlg_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NLG_F16)
v_cmpx_ngt_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NGT_F16)
v_cmpx_nle_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NLE_F16)
v_cmpx_neq_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NEQ_F16)
v_cmpx_nlt_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NLT_F16)
v_cmpx_tru_f16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_TRU_F16)
v_cmp_f_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_F_F32)
v_cmp_lt_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LT_F32)
v_cmp_eq_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_EQ_F32)
v_cmp_le_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LE_F32)
v_cmp_gt_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GT_F32)
v_cmp_lg_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LG_F32)
v_cmp_ge_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GE_F32)
v_cmp_o_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_O_F32)
v_cmp_u_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_U_F32)
v_cmp_nge_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NGE_F32)
v_cmp_nlg_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NLG_F32)
v_cmp_ngt_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NGT_F32)
v_cmp_nle_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NLE_F32)
v_cmp_neq_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NEQ_F32)
v_cmp_nlt_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NLT_F32)
v_cmp_tru_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_TRU_F32)
v_cmpx_f_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_F_F32)
v_cmpx_lt_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LT_F32)
v_cmpx_eq_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_EQ_F32)
v_cmpx_le_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LE_F32)
v_cmpx_gt_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GT_F32)
v_cmpx_lg_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LG_F32)
v_cmpx_ge_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GE_F32)
v_cmpx_o_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_O_F32)
v_cmpx_u_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_U_F32)
v_cmpx_nge_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NGE_F32)
v_cmpx_nlg_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NLG_F32)
v_cmpx_ngt_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NGT_F32)
v_cmpx_nle_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NLE_F32)
v_cmpx_neq_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NEQ_F32)
v_cmpx_nlt_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NLT_F32)
v_cmpx_tru_f32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_TRU_F32)
v_cmp_f_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_F_F64)
v_cmp_lt_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LT_F64)
v_cmp_eq_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_EQ_F64)
v_cmp_le_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LE_F64)
v_cmp_gt_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GT_F64)
v_cmp_lg_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LG_F64)
v_cmp_ge_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GE_F64)
v_cmp_o_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_O_F64)
v_cmp_u_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_U_F64)
v_cmp_nge_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NGE_F64)
v_cmp_nlg_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NLG_F64)
v_cmp_ngt_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NGT_F64)
v_cmp_nle_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NLE_F64)
v_cmp_neq_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NEQ_F64)
v_cmp_nlt_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NLT_F64)
v_cmp_tru_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_TRU_F64)
v_cmpx_f_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_F_F64)
v_cmpx_lt_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LT_F64)
v_cmpx_eq_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_EQ_F64)
v_cmpx_le_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LE_F64)
v_cmpx_gt_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GT_F64)
v_cmpx_lg_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LG_F64)
v_cmpx_ge_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GE_F64)
v_cmpx_o_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_O_F64)
v_cmpx_u_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_U_F64)
v_cmpx_nge_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NGE_F64)
v_cmpx_nlg_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NLG_F64)
v_cmpx_ngt_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NGT_F64)
v_cmpx_nle_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NLE_F64)
v_cmpx_neq_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NEQ_F64)
v_cmpx_nlt_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NLT_F64)
v_cmpx_tru_f64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_TRU_F64)
v_cmp_f_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_F_I16)
v_cmp_lt_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LT_I16)
v_cmp_eq_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_EQ_I16)
v_cmp_le_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LE_I16)
v_cmp_gt_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GT_I16)
v_cmp_ne_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NE_I16)
v_cmp_ge_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GE_I16)
v_cmp_t_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_T_I16)
v_cmp_f_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_F_U16)
v_cmp_lt_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LT_U16)
v_cmp_eq_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_EQ_U16)
v_cmp_le_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LE_U16)
v_cmp_gt_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GT_U16)
v_cmp_ne_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NE_U16)
v_cmp_ge_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GE_U16)
v_cmp_t_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMP_T_U16)
v_cmpx_f_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_F_I16)
v_cmpx_lt_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LT_I16)
v_cmpx_eq_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_EQ_I16)
v_cmpx_le_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LE_I16)
v_cmpx_gt_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GT_I16)
v_cmpx_ne_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NE_I16)
v_cmpx_ge_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GE_I16)
v_cmpx_t_i16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_T_I16)
v_cmpx_f_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_F_U16)
v_cmpx_lt_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LT_U16)
v_cmpx_eq_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_EQ_U16)
v_cmpx_le_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LE_U16)
v_cmpx_gt_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GT_U16)
v_cmpx_ne_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NE_U16)
v_cmpx_ge_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GE_U16)
v_cmpx_t_u16_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_T_U16)
v_cmp_f_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_F_I32)
v_cmp_lt_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LT_I32)
v_cmp_eq_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_EQ_I32)
v_cmp_le_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LE_I32)
v_cmp_gt_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GT_I32)
v_cmp_ne_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NE_I32)
v_cmp_ge_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GE_I32)
v_cmp_t_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_T_I32)
v_cmp_f_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_F_U32)
v_cmp_lt_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LT_U32)
v_cmp_eq_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_EQ_U32)
v_cmp_le_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LE_U32)
v_cmp_gt_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GT_U32)
v_cmp_ne_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NE_U32)
v_cmp_ge_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GE_U32)
v_cmp_t_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMP_T_U32)
v_cmpx_f_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_F_I32)
v_cmpx_lt_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LT_I32)
v_cmpx_eq_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_EQ_I32)
v_cmpx_le_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LE_I32)
v_cmpx_gt_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GT_I32)
v_cmpx_ne_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NE_I32)
v_cmpx_ge_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GE_I32)
v_cmpx_t_i32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_T_I32)
v_cmpx_f_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_F_U32)
v_cmpx_lt_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LT_U32)
v_cmpx_eq_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_EQ_U32)
v_cmpx_le_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LE_U32)
v_cmpx_gt_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GT_U32)
v_cmpx_ne_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NE_U32)
v_cmpx_ge_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GE_U32)
v_cmpx_t_u32_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_T_U32)
v_cmp_f_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_F_I64)
v_cmp_lt_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LT_I64)
v_cmp_eq_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_EQ_I64)
v_cmp_le_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LE_I64)
v_cmp_gt_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GT_I64)
v_cmp_ne_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NE_I64)
v_cmp_ge_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GE_I64)
v_cmp_t_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_T_I64)
v_cmp_f_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_F_U64)
v_cmp_lt_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LT_U64)
v_cmp_eq_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_EQ_U64)
v_cmp_le_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_LE_U64)
v_cmp_gt_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GT_U64)
v_cmp_ne_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_NE_U64)
v_cmp_ge_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_GE_U64)
v_cmp_t_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMP_T_U64)
v_cmpx_f_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_F_I64)
v_cmpx_lt_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LT_I64)
v_cmpx_eq_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_EQ_I64)
v_cmpx_le_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LE_I64)
v_cmpx_gt_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GT_I64)
v_cmpx_ne_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NE_I64)
v_cmpx_ge_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GE_I64)
v_cmpx_t_i64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_T_I64)
v_cmpx_f_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_F_U64)
v_cmpx_lt_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LT_U64)
v_cmpx_eq_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_EQ_U64)
v_cmpx_le_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_LE_U64)
v_cmpx_gt_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GT_U64)
v_cmpx_ne_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_NE_U64)
v_cmpx_ge_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_GE_U64)
v_cmpx_t_u64_e32 = functools.partial(VOPC, VOPCOp.V_CMPX_T_U64)
cdna4_e32 = functools.partial(VOPC, VOPCOp.CDNA4)
S_ADD_U32 = SrcEnum.S_ADD_U32
S_SUB_U32 = SrcEnum.S_SUB_U32
S_ADD_I32 = SrcEnum.S_ADD_I32
S_SUB_I32 = SrcEnum.S_SUB_I32
S_ADDC_U32 = SrcEnum.S_ADDC_U32
S_SUBB_U32 = SrcEnum.S_SUBB_U32
S_MIN_I32 = SrcEnum.S_MIN_I32
FLAT_SCRATCH_LO = SrcEnum.FLAT_SCRATCH_LO
FLAT_SCRATCH_HI = SrcEnum.FLAT_SCRATCH_HI
XNACK_MASK_LO = SrcEnum.XNACK_MASK_LO
XNACK_MASK_HI = SrcEnum.XNACK_MASK_HI
VCC_LO = SrcEnum.VCC_LO
VCC_HI = SrcEnum.VCC_HI
M0 = SrcEnum.M0
EXEC_LO = SrcEnum.EXEC_LO
EXEC_HI = SrcEnum.EXEC_HI
ZERO = SrcEnum.ZERO
DPP8FI = SrcEnum.DPP8FI
SHARED_BASE = SrcEnum.SHARED_BASE
SHARED_LIMIT = SrcEnum.SHARED_LIMIT
PRIVATE_BASE = SrcEnum.PRIVATE_BASE
PRIVATE_LIMIT = SrcEnum.PRIVATE_LIMIT
RESERVED = SrcEnum.RESERVED
POS_HALF = SrcEnum.POS_HALF
NEG_HALF = SrcEnum.NEG_HALF
POS_ONE = SrcEnum.POS_ONE
NEG_ONE = SrcEnum.NEG_ONE
POS_TWO = SrcEnum.POS_TWO
NEG_TWO = SrcEnum.NEG_TWO
POS_FOUR = SrcEnum.POS_FOUR
NEG_FOUR = SrcEnum.NEG_FOUR
INV_2PI = SrcEnum.INV_2PI
VCCZ = SrcEnum.VCCZ
EXECZ = SrcEnum.EXECZ
SCC = SrcEnum.SCC
LDS_DIRECT = SrcEnum.LDS_DIRECT
Reference in New Issue View Git Blame Copy Permalink