gem5::VegaISA Namespace Reference

classes that represnt vector/scalar operands in VEGA ISA. More...

Classes
struct	BufferRsrcDescriptor
class	Decoder
class	GPUISA
class	GpuTLB
struct	InFmt_DS
struct	InFmt_DS_1
struct	InFmt_EXP
struct	InFmt_EXP_1
struct	InFmt_FLAT
struct	InFmt_FLAT_1
struct	InFmt_INST
struct	InFmt_MIMG
struct	InFmt_MIMG_1
struct	InFmt_MTBUF
struct	InFmt_MTBUF_1
struct	InFmt_MUBUF
struct	InFmt_MUBUF_1
struct	InFmt_SMEM
struct	InFmt_SMEM_1
struct	InFmt_SOP1
struct	InFmt_SOP2
struct	InFmt_SOPC
struct	InFmt_SOPK
struct	InFmt_SOPP
struct	InFmt_VINTRP
struct	InFmt_VOP1
struct	InFmt_VOP2
struct	InFmt_VOP3_1
struct	InFmt_VOP3A
struct	InFmt_VOP3B
struct	InFmt_VOP3P
struct	InFmt_VOP3P_1
struct	InFmt_VOP3P_MAI
struct	InFmt_VOP3P_MAI_1
struct	InFmt_VOP_DPP
struct	InFmt_VOP_SDWA
struct	InFmt_VOP_SDWAB
struct	InFmt_VOPC
class	Inst_DS
class	Inst_DS__DS_ADD_F32
class	Inst_DS__DS_ADD_RTN_F32
class	Inst_DS__DS_ADD_RTN_U32
class	Inst_DS__DS_ADD_RTN_U64
class	Inst_DS__DS_ADD_SRC2_F32
class	Inst_DS__DS_ADD_SRC2_U32
class	Inst_DS__DS_ADD_SRC2_U64
class	Inst_DS__DS_ADD_U32
class	Inst_DS__DS_ADD_U64
class	Inst_DS__DS_AND_B32
class	Inst_DS__DS_AND_B64
class	Inst_DS__DS_AND_RTN_B32
class	Inst_DS__DS_AND_RTN_B64
class	Inst_DS__DS_AND_SRC2_B32
class	Inst_DS__DS_AND_SRC2_B64
class	Inst_DS__DS_APPEND
class	Inst_DS__DS_BPERMUTE_B32
class	Inst_DS__DS_CMPST_B32
class	Inst_DS__DS_CMPST_B64
class	Inst_DS__DS_CMPST_F32
class	Inst_DS__DS_CMPST_F64
class	Inst_DS__DS_CMPST_RTN_B32
class	Inst_DS__DS_CMPST_RTN_B64
class	Inst_DS__DS_CMPST_RTN_F32
class	Inst_DS__DS_CMPST_RTN_F64
class	Inst_DS__DS_CONDXCHG32_RTN_B64
class	Inst_DS__DS_CONSUME
class	Inst_DS__DS_DEC_RTN_U32
class	Inst_DS__DS_DEC_RTN_U64
class	Inst_DS__DS_DEC_SRC2_U32
class	Inst_DS__DS_DEC_SRC2_U64
class	Inst_DS__DS_DEC_U32
class	Inst_DS__DS_DEC_U64
class	Inst_DS__DS_GWS_BARRIER
class	Inst_DS__DS_GWS_INIT
class	Inst_DS__DS_GWS_SEMA_BR
class	Inst_DS__DS_GWS_SEMA_P
class	Inst_DS__DS_GWS_SEMA_RELEASE_ALL
class	Inst_DS__DS_GWS_SEMA_V
class	Inst_DS__DS_INC_RTN_U32
class	Inst_DS__DS_INC_RTN_U64
class	Inst_DS__DS_INC_SRC2_U32
class	Inst_DS__DS_INC_SRC2_U64
class	Inst_DS__DS_INC_U32
class	Inst_DS__DS_INC_U64
class	Inst_DS__DS_MAX_F32
class	Inst_DS__DS_MAX_F64
class	Inst_DS__DS_MAX_I32
class	Inst_DS__DS_MAX_I64
class	Inst_DS__DS_MAX_RTN_F32
class	Inst_DS__DS_MAX_RTN_F64
class	Inst_DS__DS_MAX_RTN_I32
class	Inst_DS__DS_MAX_RTN_I64
class	Inst_DS__DS_MAX_RTN_U32
class	Inst_DS__DS_MAX_RTN_U64
class	Inst_DS__DS_MAX_SRC2_F32
class	Inst_DS__DS_MAX_SRC2_F64
class	Inst_DS__DS_MAX_SRC2_I32
class	Inst_DS__DS_MAX_SRC2_I64
class	Inst_DS__DS_MAX_SRC2_U32
class	Inst_DS__DS_MAX_SRC2_U64
class	Inst_DS__DS_MAX_U32
class	Inst_DS__DS_MAX_U64
class	Inst_DS__DS_MIN_F32
class	Inst_DS__DS_MIN_F64
class	Inst_DS__DS_MIN_I32
class	Inst_DS__DS_MIN_I64
class	Inst_DS__DS_MIN_RTN_F32
class	Inst_DS__DS_MIN_RTN_F64
class	Inst_DS__DS_MIN_RTN_I32
class	Inst_DS__DS_MIN_RTN_I64
class	Inst_DS__DS_MIN_RTN_U32
class	Inst_DS__DS_MIN_RTN_U64
class	Inst_DS__DS_MIN_SRC2_F32
class	Inst_DS__DS_MIN_SRC2_F64
class	Inst_DS__DS_MIN_SRC2_I32
class	Inst_DS__DS_MIN_SRC2_I64
class	Inst_DS__DS_MIN_SRC2_U32
class	Inst_DS__DS_MIN_SRC2_U64
class	Inst_DS__DS_MIN_U32
class	Inst_DS__DS_MIN_U64
class	Inst_DS__DS_MSKOR_B32
class	Inst_DS__DS_MSKOR_B64
class	Inst_DS__DS_MSKOR_RTN_B32
class	Inst_DS__DS_MSKOR_RTN_B64
class	Inst_DS__DS_NOP
class	Inst_DS__DS_OR_B32
class	Inst_DS__DS_OR_B64
class	Inst_DS__DS_OR_RTN_B32
class	Inst_DS__DS_OR_RTN_B64
class	Inst_DS__DS_OR_SRC2_B32
class	Inst_DS__DS_OR_SRC2_B64
class	Inst_DS__DS_ORDERED_COUNT
class	Inst_DS__DS_PERMUTE_B32
class	Inst_DS__DS_READ2_B32
class	Inst_DS__DS_READ2_B64
class	Inst_DS__DS_READ2ST64_B32
class	Inst_DS__DS_READ2ST64_B64
class	Inst_DS__DS_READ_B128
class	Inst_DS__DS_READ_B32
class	Inst_DS__DS_READ_B64
class	Inst_DS__DS_READ_B96
class	Inst_DS__DS_READ_I16
class	Inst_DS__DS_READ_I8
class	Inst_DS__DS_READ_U16
class	Inst_DS__DS_READ_U16_D16
class	Inst_DS__DS_READ_U16_D16_HI
class	Inst_DS__DS_READ_U8
class	Inst_DS__DS_RSUB_RTN_U32
class	Inst_DS__DS_RSUB_RTN_U64
class	Inst_DS__DS_RSUB_SRC2_U32
class	Inst_DS__DS_RSUB_SRC2_U64
class	Inst_DS__DS_RSUB_U32
class	Inst_DS__DS_RSUB_U64
class	Inst_DS__DS_SUB_RTN_U32
class	Inst_DS__DS_SUB_RTN_U64
class	Inst_DS__DS_SUB_SRC2_U32
class	Inst_DS__DS_SUB_SRC2_U64
class	Inst_DS__DS_SUB_U32
class	Inst_DS__DS_SUB_U64
class	Inst_DS__DS_SWIZZLE_B32
class	Inst_DS__DS_WRAP_RTN_B32
class	Inst_DS__DS_WRITE2_B32
class	Inst_DS__DS_WRITE2_B64
class	Inst_DS__DS_WRITE2ST64_B32
class	Inst_DS__DS_WRITE2ST64_B64
class	Inst_DS__DS_WRITE_B128
class	Inst_DS__DS_WRITE_B16
class	Inst_DS__DS_WRITE_B32
class	Inst_DS__DS_WRITE_B64
class	Inst_DS__DS_WRITE_B8
class	Inst_DS__DS_WRITE_B8_D16_HI
class	Inst_DS__DS_WRITE_B96
class	Inst_DS__DS_WRITE_SRC2_B32
class	Inst_DS__DS_WRITE_SRC2_B64
class	Inst_DS__DS_WRXCHG2_RTN_B32
class	Inst_DS__DS_WRXCHG2_RTN_B64
class	Inst_DS__DS_WRXCHG2ST64_RTN_B32
class	Inst_DS__DS_WRXCHG2ST64_RTN_B64
class	Inst_DS__DS_WRXCHG_RTN_B32
class	Inst_DS__DS_WRXCHG_RTN_B64
class	Inst_DS__DS_XOR_B32
class	Inst_DS__DS_XOR_B64
class	Inst_DS__DS_XOR_RTN_B32
class	Inst_DS__DS_XOR_RTN_B64
class	Inst_DS__DS_XOR_SRC2_B32
class	Inst_DS__DS_XOR_SRC2_B64
class	Inst_EXP
class	Inst_EXP__EXP
class	Inst_FLAT
class	Inst_FLAT__FLAT_ATOMIC_ADD
class	Inst_FLAT__FLAT_ATOMIC_ADD_F32
class	Inst_FLAT__FLAT_ATOMIC_ADD_F64
class	Inst_FLAT__FLAT_ATOMIC_ADD_X2
class	Inst_FLAT__FLAT_ATOMIC_AND
class	Inst_FLAT__FLAT_ATOMIC_AND_X2
class	Inst_FLAT__FLAT_ATOMIC_CMPSWAP
class	Inst_FLAT__FLAT_ATOMIC_CMPSWAP_X2
class	Inst_FLAT__FLAT_ATOMIC_DEC
class	Inst_FLAT__FLAT_ATOMIC_DEC_X2
class	Inst_FLAT__FLAT_ATOMIC_INC
class	Inst_FLAT__FLAT_ATOMIC_INC_X2
class	Inst_FLAT__FLAT_ATOMIC_MAX_F64
class	Inst_FLAT__FLAT_ATOMIC_MIN_F64
class	Inst_FLAT__FLAT_ATOMIC_OR
class	Inst_FLAT__FLAT_ATOMIC_OR_X2
class	Inst_FLAT__FLAT_ATOMIC_PK_ADD_F16
class	Inst_FLAT__FLAT_ATOMIC_SMAX
class	Inst_FLAT__FLAT_ATOMIC_SMAX_X2
class	Inst_FLAT__FLAT_ATOMIC_SMIN
class	Inst_FLAT__FLAT_ATOMIC_SMIN_X2
class	Inst_FLAT__FLAT_ATOMIC_SUB
class	Inst_FLAT__FLAT_ATOMIC_SUB_X2
class	Inst_FLAT__FLAT_ATOMIC_SWAP
class	Inst_FLAT__FLAT_ATOMIC_SWAP_X2
class	Inst_FLAT__FLAT_ATOMIC_UMAX
class	Inst_FLAT__FLAT_ATOMIC_UMAX_X2
class	Inst_FLAT__FLAT_ATOMIC_UMIN
class	Inst_FLAT__FLAT_ATOMIC_UMIN_X2
class	Inst_FLAT__FLAT_ATOMIC_XOR
class	Inst_FLAT__FLAT_ATOMIC_XOR_X2
class	Inst_FLAT__FLAT_LOAD_DWORD
class	Inst_FLAT__FLAT_LOAD_DWORDX2
class	Inst_FLAT__FLAT_LOAD_DWORDX3
class	Inst_FLAT__FLAT_LOAD_DWORDX4
class	Inst_FLAT__FLAT_LOAD_SBYTE
class	Inst_FLAT__FLAT_LOAD_SSHORT
class	Inst_FLAT__FLAT_LOAD_UBYTE
class	Inst_FLAT__FLAT_LOAD_USHORT
class	Inst_FLAT__FLAT_STORE_BYTE
class	Inst_FLAT__FLAT_STORE_DWORD
class	Inst_FLAT__FLAT_STORE_DWORDX2
class	Inst_FLAT__FLAT_STORE_DWORDX3
class	Inst_FLAT__FLAT_STORE_DWORDX4
class	Inst_FLAT__FLAT_STORE_SHORT
class	Inst_FLAT__FLAT_STORE_SHORT_D16_HI
class	Inst_MIMG
class	Inst_MIMG__IMAGE_ATOMIC_ADD
class	Inst_MIMG__IMAGE_ATOMIC_AND
class	Inst_MIMG__IMAGE_ATOMIC_CMPSWAP
class	Inst_MIMG__IMAGE_ATOMIC_DEC
class	Inst_MIMG__IMAGE_ATOMIC_INC
class	Inst_MIMG__IMAGE_ATOMIC_OR
class	Inst_MIMG__IMAGE_ATOMIC_SMAX
class	Inst_MIMG__IMAGE_ATOMIC_SMIN
class	Inst_MIMG__IMAGE_ATOMIC_SUB
class	Inst_MIMG__IMAGE_ATOMIC_SWAP
class	Inst_MIMG__IMAGE_ATOMIC_UMAX
class	Inst_MIMG__IMAGE_ATOMIC_UMIN
class	Inst_MIMG__IMAGE_ATOMIC_XOR
class	Inst_MIMG__IMAGE_GATHER4
class	Inst_MIMG__IMAGE_GATHER4_B
class	Inst_MIMG__IMAGE_GATHER4_B_CL
class	Inst_MIMG__IMAGE_GATHER4_B_CL_O
class	Inst_MIMG__IMAGE_GATHER4_B_O
class	Inst_MIMG__IMAGE_GATHER4_C
class	Inst_MIMG__IMAGE_GATHER4_C_B
class	Inst_MIMG__IMAGE_GATHER4_C_B_CL
class	Inst_MIMG__IMAGE_GATHER4_C_B_CL_O
class	Inst_MIMG__IMAGE_GATHER4_C_B_O
class	Inst_MIMG__IMAGE_GATHER4_C_CL
class	Inst_MIMG__IMAGE_GATHER4_C_CL_O
class	Inst_MIMG__IMAGE_GATHER4_C_L
class	Inst_MIMG__IMAGE_GATHER4_C_L_O
class	Inst_MIMG__IMAGE_GATHER4_C_LZ
class	Inst_MIMG__IMAGE_GATHER4_C_LZ_O
class	Inst_MIMG__IMAGE_GATHER4_C_O
class	Inst_MIMG__IMAGE_GATHER4_CL
class	Inst_MIMG__IMAGE_GATHER4_CL_O
class	Inst_MIMG__IMAGE_GATHER4_L
class	Inst_MIMG__IMAGE_GATHER4_L_O
class	Inst_MIMG__IMAGE_GATHER4_LZ
class	Inst_MIMG__IMAGE_GATHER4_LZ_O
class	Inst_MIMG__IMAGE_GATHER4_O
class	Inst_MIMG__IMAGE_GET_LOD
class	Inst_MIMG__IMAGE_GET_RESINFO
class	Inst_MIMG__IMAGE_LOAD
class	Inst_MIMG__IMAGE_LOAD_MIP
class	Inst_MIMG__IMAGE_LOAD_MIP_PCK
class	Inst_MIMG__IMAGE_LOAD_MIP_PCK_SGN
class	Inst_MIMG__IMAGE_LOAD_PCK
class	Inst_MIMG__IMAGE_LOAD_PCK_SGN
class	Inst_MIMG__IMAGE_SAMPLE
class	Inst_MIMG__IMAGE_SAMPLE_B
class	Inst_MIMG__IMAGE_SAMPLE_B_CL
class	Inst_MIMG__IMAGE_SAMPLE_B_CL_O
class	Inst_MIMG__IMAGE_SAMPLE_B_O
class	Inst_MIMG__IMAGE_SAMPLE_C
class	Inst_MIMG__IMAGE_SAMPLE_C_B
class	Inst_MIMG__IMAGE_SAMPLE_C_B_CL
class	Inst_MIMG__IMAGE_SAMPLE_C_B_CL_O
class	Inst_MIMG__IMAGE_SAMPLE_C_B_O
class	Inst_MIMG__IMAGE_SAMPLE_C_CD
class	Inst_MIMG__IMAGE_SAMPLE_C_CD_CL
class	Inst_MIMG__IMAGE_SAMPLE_C_CD_CL_O
class	Inst_MIMG__IMAGE_SAMPLE_C_CD_O
class	Inst_MIMG__IMAGE_SAMPLE_C_CL
class	Inst_MIMG__IMAGE_SAMPLE_C_CL_O
class	Inst_MIMG__IMAGE_SAMPLE_C_D
class	Inst_MIMG__IMAGE_SAMPLE_C_D_CL
class	Inst_MIMG__IMAGE_SAMPLE_C_D_CL_O
class	Inst_MIMG__IMAGE_SAMPLE_C_D_O
class	Inst_MIMG__IMAGE_SAMPLE_C_L
class	Inst_MIMG__IMAGE_SAMPLE_C_L_O
class	Inst_MIMG__IMAGE_SAMPLE_C_LZ
class	Inst_MIMG__IMAGE_SAMPLE_C_LZ_O
class	Inst_MIMG__IMAGE_SAMPLE_C_O
class	Inst_MIMG__IMAGE_SAMPLE_CD
class	Inst_MIMG__IMAGE_SAMPLE_CD_CL
class	Inst_MIMG__IMAGE_SAMPLE_CD_CL_O
class	Inst_MIMG__IMAGE_SAMPLE_CD_O
class	Inst_MIMG__IMAGE_SAMPLE_CL
class	Inst_MIMG__IMAGE_SAMPLE_CL_O
class	Inst_MIMG__IMAGE_SAMPLE_D
class	Inst_MIMG__IMAGE_SAMPLE_D_CL
class	Inst_MIMG__IMAGE_SAMPLE_D_CL_O
class	Inst_MIMG__IMAGE_SAMPLE_D_O
class	Inst_MIMG__IMAGE_SAMPLE_L
class	Inst_MIMG__IMAGE_SAMPLE_L_O
class	Inst_MIMG__IMAGE_SAMPLE_LZ
class	Inst_MIMG__IMAGE_SAMPLE_LZ_O
class	Inst_MIMG__IMAGE_SAMPLE_O
class	Inst_MIMG__IMAGE_STORE
class	Inst_MIMG__IMAGE_STORE_MIP
class	Inst_MIMG__IMAGE_STORE_MIP_PCK
class	Inst_MIMG__IMAGE_STORE_PCK
class	Inst_MTBUF
class	Inst_MTBUF__TBUFFER_LOAD_FORMAT_D16_X
class	Inst_MTBUF__TBUFFER_LOAD_FORMAT_D16_XY
class	Inst_MTBUF__TBUFFER_LOAD_FORMAT_D16_XYZ
class	Inst_MTBUF__TBUFFER_LOAD_FORMAT_D16_XYZW
class	Inst_MTBUF__TBUFFER_LOAD_FORMAT_X
class	Inst_MTBUF__TBUFFER_LOAD_FORMAT_XY
class	Inst_MTBUF__TBUFFER_LOAD_FORMAT_XYZ
class	Inst_MTBUF__TBUFFER_LOAD_FORMAT_XYZW
class	Inst_MTBUF__TBUFFER_STORE_FORMAT_D16_X
class	Inst_MTBUF__TBUFFER_STORE_FORMAT_D16_XY
class	Inst_MTBUF__TBUFFER_STORE_FORMAT_D16_XYZ
class	Inst_MTBUF__TBUFFER_STORE_FORMAT_D16_XYZW
class	Inst_MTBUF__TBUFFER_STORE_FORMAT_X
class	Inst_MTBUF__TBUFFER_STORE_FORMAT_XY
class	Inst_MTBUF__TBUFFER_STORE_FORMAT_XYZ
class	Inst_MTBUF__TBUFFER_STORE_FORMAT_XYZW
class	Inst_MUBUF
class	Inst_MUBUF__BUFFER_ATOMIC_ADD
class	Inst_MUBUF__BUFFER_ATOMIC_ADD_X2
class	Inst_MUBUF__BUFFER_ATOMIC_AND
class	Inst_MUBUF__BUFFER_ATOMIC_AND_X2
class	Inst_MUBUF__BUFFER_ATOMIC_CMPSWAP
class	Inst_MUBUF__BUFFER_ATOMIC_CMPSWAP_X2
class	Inst_MUBUF__BUFFER_ATOMIC_DEC
class	Inst_MUBUF__BUFFER_ATOMIC_DEC_X2
class	Inst_MUBUF__BUFFER_ATOMIC_INC
class	Inst_MUBUF__BUFFER_ATOMIC_INC_X2
class	Inst_MUBUF__BUFFER_ATOMIC_OR
class	Inst_MUBUF__BUFFER_ATOMIC_OR_X2
class	Inst_MUBUF__BUFFER_ATOMIC_SMAX
class	Inst_MUBUF__BUFFER_ATOMIC_SMAX_X2
class	Inst_MUBUF__BUFFER_ATOMIC_SMIN
class	Inst_MUBUF__BUFFER_ATOMIC_SMIN_X2
class	Inst_MUBUF__BUFFER_ATOMIC_SUB
class	Inst_MUBUF__BUFFER_ATOMIC_SUB_X2
class	Inst_MUBUF__BUFFER_ATOMIC_SWAP
class	Inst_MUBUF__BUFFER_ATOMIC_SWAP_X2
class	Inst_MUBUF__BUFFER_ATOMIC_UMAX
class	Inst_MUBUF__BUFFER_ATOMIC_UMAX_X2
class	Inst_MUBUF__BUFFER_ATOMIC_UMIN
class	Inst_MUBUF__BUFFER_ATOMIC_UMIN_X2
class	Inst_MUBUF__BUFFER_ATOMIC_XOR
class	Inst_MUBUF__BUFFER_ATOMIC_XOR_X2
class	Inst_MUBUF__BUFFER_LOAD_DWORD
class	Inst_MUBUF__BUFFER_LOAD_DWORDX2
class	Inst_MUBUF__BUFFER_LOAD_DWORDX3
class	Inst_MUBUF__BUFFER_LOAD_DWORDX4
class	Inst_MUBUF__BUFFER_LOAD_FORMAT_D16_X
class	Inst_MUBUF__BUFFER_LOAD_FORMAT_D16_XY
class	Inst_MUBUF__BUFFER_LOAD_FORMAT_D16_XYZ
class	Inst_MUBUF__BUFFER_LOAD_FORMAT_D16_XYZW
class	Inst_MUBUF__BUFFER_LOAD_FORMAT_X
class	Inst_MUBUF__BUFFER_LOAD_FORMAT_XY
class	Inst_MUBUF__BUFFER_LOAD_FORMAT_XYZ
class	Inst_MUBUF__BUFFER_LOAD_FORMAT_XYZW
class	Inst_MUBUF__BUFFER_LOAD_SBYTE
class	Inst_MUBUF__BUFFER_LOAD_SHORT_D16
class	Inst_MUBUF__BUFFER_LOAD_SHORT_D16_HI
class	Inst_MUBUF__BUFFER_LOAD_SSHORT
class	Inst_MUBUF__BUFFER_LOAD_UBYTE
class	Inst_MUBUF__BUFFER_LOAD_USHORT
class	Inst_MUBUF__BUFFER_STORE_BYTE
class	Inst_MUBUF__BUFFER_STORE_DWORD
class	Inst_MUBUF__BUFFER_STORE_DWORDX2
class	Inst_MUBUF__BUFFER_STORE_DWORDX3
class	Inst_MUBUF__BUFFER_STORE_DWORDX4
class	Inst_MUBUF__BUFFER_STORE_FORMAT_D16_X
class	Inst_MUBUF__BUFFER_STORE_FORMAT_D16_XY
class	Inst_MUBUF__BUFFER_STORE_FORMAT_D16_XYZ
class	Inst_MUBUF__BUFFER_STORE_FORMAT_D16_XYZW
class	Inst_MUBUF__BUFFER_STORE_FORMAT_X
class	Inst_MUBUF__BUFFER_STORE_FORMAT_XY
class	Inst_MUBUF__BUFFER_STORE_FORMAT_XYZ
class	Inst_MUBUF__BUFFER_STORE_FORMAT_XYZW
class	Inst_MUBUF__BUFFER_STORE_LDS_DWORD
class	Inst_MUBUF__BUFFER_STORE_SHORT
class	Inst_MUBUF__BUFFER_WBINVL1
class	Inst_MUBUF__BUFFER_WBINVL1_VOL
class	Inst_SMEM
class	Inst_SMEM__S_ATC_PROBE
class	Inst_SMEM__S_ATC_PROBE_BUFFER
class	Inst_SMEM__S_BUFFER_LOAD_DWORD
class	Inst_SMEM__S_BUFFER_LOAD_DWORDX16
class	Inst_SMEM__S_BUFFER_LOAD_DWORDX2
class	Inst_SMEM__S_BUFFER_LOAD_DWORDX4
class	Inst_SMEM__S_BUFFER_LOAD_DWORDX8
class	Inst_SMEM__S_BUFFER_STORE_DWORD
class	Inst_SMEM__S_BUFFER_STORE_DWORDX2
class	Inst_SMEM__S_BUFFER_STORE_DWORDX4
class	Inst_SMEM__S_DCACHE_INV
class	Inst_SMEM__S_DCACHE_INV_VOL
class	Inst_SMEM__S_DCACHE_WB
class	Inst_SMEM__S_DCACHE_WB_VOL
class	Inst_SMEM__S_LOAD_DWORD
class	Inst_SMEM__S_LOAD_DWORDX16
class	Inst_SMEM__S_LOAD_DWORDX2
class	Inst_SMEM__S_LOAD_DWORDX4
class	Inst_SMEM__S_LOAD_DWORDX8
class	Inst_SMEM__S_MEMREALTIME
class	Inst_SMEM__S_MEMTIME
class	Inst_SMEM__S_STORE_DWORD
class	Inst_SMEM__S_STORE_DWORDX2
class	Inst_SMEM__S_STORE_DWORDX4
class	Inst_SOP1
class	Inst_SOP1__S_ABS_I32
class	Inst_SOP1__S_AND_SAVEEXEC_B64
class	Inst_SOP1__S_ANDN2_SAVEEXEC_B64
class	Inst_SOP1__S_BCNT0_I32_B32
class	Inst_SOP1__S_BCNT0_I32_B64
class	Inst_SOP1__S_BCNT1_I32_B32
class	Inst_SOP1__S_BCNT1_I32_B64
class	Inst_SOP1__S_BITSET0_B32
class	Inst_SOP1__S_BITSET0_B64
class	Inst_SOP1__S_BITSET1_B32
class	Inst_SOP1__S_BITSET1_B64
class	Inst_SOP1__S_BREV_B32
class	Inst_SOP1__S_BREV_B64
class	Inst_SOP1__S_CBRANCH_JOIN
class	Inst_SOP1__S_CMOV_B32
class	Inst_SOP1__S_CMOV_B64
class	Inst_SOP1__S_FF0_I32_B32
class	Inst_SOP1__S_FF0_I32_B64
class	Inst_SOP1__S_FF1_I32_B32
class	Inst_SOP1__S_FF1_I32_B64
class	Inst_SOP1__S_FLBIT_I32
class	Inst_SOP1__S_FLBIT_I32_B32
class	Inst_SOP1__S_FLBIT_I32_B64
class	Inst_SOP1__S_FLBIT_I32_I64
class	Inst_SOP1__S_GETPC_B64
class	Inst_SOP1__S_MOV_B32
class	Inst_SOP1__S_MOV_B64
class	Inst_SOP1__S_MOV_FED_B32
class	Inst_SOP1__S_MOVRELD_B32
class	Inst_SOP1__S_MOVRELD_B64
class	Inst_SOP1__S_MOVRELS_B32
class	Inst_SOP1__S_MOVRELS_B64
class	Inst_SOP1__S_NAND_SAVEEXEC_B64
class	Inst_SOP1__S_NOR_SAVEEXEC_B64
class	Inst_SOP1__S_NOT_B32
class	Inst_SOP1__S_NOT_B64
class	Inst_SOP1__S_OR_SAVEEXEC_B64
class	Inst_SOP1__S_ORN2_SAVEEXEC_B64
class	Inst_SOP1__S_QUADMASK_B32
class	Inst_SOP1__S_QUADMASK_B64
class	Inst_SOP1__S_RFE_B64
class	Inst_SOP1__S_SET_GPR_IDX_IDX
class	Inst_SOP1__S_SETPC_B64
class	Inst_SOP1__S_SEXT_I32_I16
class	Inst_SOP1__S_SEXT_I32_I8
class	Inst_SOP1__S_SWAPPC_B64
class	Inst_SOP1__S_WQM_B32
class	Inst_SOP1__S_WQM_B64
class	Inst_SOP1__S_XNOR_SAVEEXEC_B64
class	Inst_SOP1__S_XOR_SAVEEXEC_B64
class	Inst_SOP2
class	Inst_SOP2__S_ABSDIFF_I32
class	Inst_SOP2__S_ADD_I32
class	Inst_SOP2__S_ADD_U32
class	Inst_SOP2__S_ADDC_U32
class	Inst_SOP2__S_AND_B32
class	Inst_SOP2__S_AND_B64
class	Inst_SOP2__S_ANDN2_B32
class	Inst_SOP2__S_ANDN2_B64
class	Inst_SOP2__S_ASHR_I32
class	Inst_SOP2__S_ASHR_I64
class	Inst_SOP2__S_BFE_I32
class	Inst_SOP2__S_BFE_I64
class	Inst_SOP2__S_BFE_U32
class	Inst_SOP2__S_BFE_U64
class	Inst_SOP2__S_BFM_B32
class	Inst_SOP2__S_BFM_B64
class	Inst_SOP2__S_CBRANCH_G_FORK
class	Inst_SOP2__S_CSELECT_B32
class	Inst_SOP2__S_CSELECT_B64
class	Inst_SOP2__S_LSHL_B32
class	Inst_SOP2__S_LSHL_B64
class	Inst_SOP2__S_LSHR_B32
class	Inst_SOP2__S_LSHR_B64
class	Inst_SOP2__S_MAX_I32
class	Inst_SOP2__S_MAX_U32
class	Inst_SOP2__S_MIN_I32
class	Inst_SOP2__S_MIN_U32
class	Inst_SOP2__S_MUL_HI_I32
class	Inst_SOP2__S_MUL_HI_U32
class	Inst_SOP2__S_MUL_I32
class	Inst_SOP2__S_NAND_B32
class	Inst_SOP2__S_NAND_B64
class	Inst_SOP2__S_NOR_B32
class	Inst_SOP2__S_NOR_B64
class	Inst_SOP2__S_OR_B32
class	Inst_SOP2__S_OR_B64
class	Inst_SOP2__S_ORN2_B32
class	Inst_SOP2__S_ORN2_B64
class	Inst_SOP2__S_RFE_RESTORE_B64
class	Inst_SOP2__S_SUB_I32
class	Inst_SOP2__S_SUB_U32
class	Inst_SOP2__S_SUBB_U32
class	Inst_SOP2__S_XNOR_B32
class	Inst_SOP2__S_XNOR_B64
class	Inst_SOP2__S_XOR_B32
class	Inst_SOP2__S_XOR_B64
class	Inst_SOPC
class	Inst_SOPC__S_BITCMP0_B32
class	Inst_SOPC__S_BITCMP0_B64
class	Inst_SOPC__S_BITCMP1_B32
class	Inst_SOPC__S_BITCMP1_B64
class	Inst_SOPC__S_CMP_EQ_I32
class	Inst_SOPC__S_CMP_EQ_U32
class	Inst_SOPC__S_CMP_EQ_U64
class	Inst_SOPC__S_CMP_GE_I32
class	Inst_SOPC__S_CMP_GE_U32
class	Inst_SOPC__S_CMP_GT_I32
class	Inst_SOPC__S_CMP_GT_U32
class	Inst_SOPC__S_CMP_LE_I32
class	Inst_SOPC__S_CMP_LE_U32
class	Inst_SOPC__S_CMP_LG_I32
class	Inst_SOPC__S_CMP_LG_U32
class	Inst_SOPC__S_CMP_LG_U64
class	Inst_SOPC__S_CMP_LT_I32
class	Inst_SOPC__S_CMP_LT_U32
class	Inst_SOPC__S_SET_GPR_IDX_ON
class	Inst_SOPC__S_SETVSKIP
class	Inst_SOPK
class	Inst_SOPK__S_ADDK_I32
class	Inst_SOPK__S_CBRANCH_I_FORK
class	Inst_SOPK__S_CMOVK_I32
class	Inst_SOPK__S_CMPK_EQ_I32
class	Inst_SOPK__S_CMPK_EQ_U32
class	Inst_SOPK__S_CMPK_GE_I32
class	Inst_SOPK__S_CMPK_GE_U32
class	Inst_SOPK__S_CMPK_GT_I32
class	Inst_SOPK__S_CMPK_GT_U32
class	Inst_SOPK__S_CMPK_LE_I32
class	Inst_SOPK__S_CMPK_LE_U32
class	Inst_SOPK__S_CMPK_LG_I32
class	Inst_SOPK__S_CMPK_LG_U32
class	Inst_SOPK__S_CMPK_LT_I32
class	Inst_SOPK__S_CMPK_LT_U32
class	Inst_SOPK__S_GETREG_B32
class	Inst_SOPK__S_MOVK_I32
class	Inst_SOPK__S_MULK_I32
class	Inst_SOPK__S_SETREG_B32
class	Inst_SOPK__S_SETREG_IMM32_B32
class	Inst_SOPP
class	Inst_SOPP__S_BARRIER
class	Inst_SOPP__S_BRANCH
class	Inst_SOPP__S_CBRANCH_CDBGSYS
class	Inst_SOPP__S_CBRANCH_CDBGSYS_AND_USER
class	Inst_SOPP__S_CBRANCH_CDBGSYS_OR_USER
class	Inst_SOPP__S_CBRANCH_CDBGUSER
class	Inst_SOPP__S_CBRANCH_EXECNZ
class	Inst_SOPP__S_CBRANCH_EXECZ
class	Inst_SOPP__S_CBRANCH_SCC0
class	Inst_SOPP__S_CBRANCH_SCC1
class	Inst_SOPP__S_CBRANCH_VCCNZ
class	Inst_SOPP__S_CBRANCH_VCCZ
class	Inst_SOPP__S_DECPERFLEVEL
class	Inst_SOPP__S_ENDPGM
class	Inst_SOPP__S_ENDPGM_SAVED
class	Inst_SOPP__S_ICACHE_INV
class	Inst_SOPP__S_INCPERFLEVEL
class	Inst_SOPP__S_NOP
class	Inst_SOPP__S_SENDMSG
class	Inst_SOPP__S_SENDMSGHALT
class	Inst_SOPP__S_SET_GPR_IDX_MODE
class	Inst_SOPP__S_SET_GPR_IDX_OFF
class	Inst_SOPP__S_SETHALT
class	Inst_SOPP__S_SETKILL
class	Inst_SOPP__S_SETPRIO
class	Inst_SOPP__S_SLEEP
class	Inst_SOPP__S_TRAP
class	Inst_SOPP__S_TTRACEDATA
class	Inst_SOPP__S_WAITCNT
class	Inst_SOPP__S_WAKEUP
class	Inst_VINTRP
class	Inst_VINTRP__V_INTERP_MOV_F32
class	Inst_VINTRP__V_INTERP_P1_F32
class	Inst_VINTRP__V_INTERP_P2_F32
class	Inst_VOP1
class	Inst_VOP1__V_ACCVGPR_MOV_B32
class	Inst_VOP1__V_BFREV_B32
class	Inst_VOP1__V_CEIL_F16
class	Inst_VOP1__V_CEIL_F32
class	Inst_VOP1__V_CEIL_F64
class	Inst_VOP1__V_CLREXCP
class	Inst_VOP1__V_COS_F16
class	Inst_VOP1__V_COS_F32
class	Inst_VOP1__V_CVT_F16_F32
class	Inst_VOP1__V_CVT_F16_I16
class	Inst_VOP1__V_CVT_F16_U16
class	Inst_VOP1__V_CVT_F32_F16
class	Inst_VOP1__V_CVT_F32_F64
class	Inst_VOP1__V_CVT_F32_I32
class	Inst_VOP1__V_CVT_F32_U32
class	Inst_VOP1__V_CVT_F32_UBYTE0
class	Inst_VOP1__V_CVT_F32_UBYTE1
class	Inst_VOP1__V_CVT_F32_UBYTE2
class	Inst_VOP1__V_CVT_F32_UBYTE3
class	Inst_VOP1__V_CVT_F64_F32
class	Inst_VOP1__V_CVT_F64_I32
class	Inst_VOP1__V_CVT_F64_U32
class	Inst_VOP1__V_CVT_FLR_I32_F32
class	Inst_VOP1__V_CVT_I16_F16
class	Inst_VOP1__V_CVT_I32_F32
class	Inst_VOP1__V_CVT_I32_F64
class	Inst_VOP1__V_CVT_OFF_F32_I4
class	Inst_VOP1__V_CVT_RPI_I32_F32
class	Inst_VOP1__V_CVT_U16_F16
class	Inst_VOP1__V_CVT_U32_F32
class	Inst_VOP1__V_CVT_U32_F64
class	Inst_VOP1__V_EXP_F16
class	Inst_VOP1__V_EXP_F32
class	Inst_VOP1__V_EXP_LEGACY_F32
class	Inst_VOP1__V_FFBH_I32
class	Inst_VOP1__V_FFBH_U32
class	Inst_VOP1__V_FFBL_B32
class	Inst_VOP1__V_FLOOR_F16
class	Inst_VOP1__V_FLOOR_F32
class	Inst_VOP1__V_FLOOR_F64
class	Inst_VOP1__V_FRACT_F16
class	Inst_VOP1__V_FRACT_F32
class	Inst_VOP1__V_FRACT_F64
class	Inst_VOP1__V_FREXP_EXP_I16_F16
class	Inst_VOP1__V_FREXP_EXP_I32_F32
class	Inst_VOP1__V_FREXP_EXP_I32_F64
class	Inst_VOP1__V_FREXP_MANT_F16
class	Inst_VOP1__V_FREXP_MANT_F32
class	Inst_VOP1__V_FREXP_MANT_F64
class	Inst_VOP1__V_LOG_F16
class	Inst_VOP1__V_LOG_F32
class	Inst_VOP1__V_LOG_LEGACY_F32
class	Inst_VOP1__V_MOV_B32
class	Inst_VOP1__V_MOV_B64
class	Inst_VOP1__V_MOV_FED_B32
class	Inst_VOP1__V_NOP
class	Inst_VOP1__V_NOT_B32
class	Inst_VOP1__V_RCP_F16
class	Inst_VOP1__V_RCP_F32
class	Inst_VOP1__V_RCP_F64
class	Inst_VOP1__V_RCP_IFLAG_F32
class	Inst_VOP1__V_READFIRSTLANE_B32
class	Inst_VOP1__V_RNDNE_F16
class	Inst_VOP1__V_RNDNE_F32
class	Inst_VOP1__V_RNDNE_F64
class	Inst_VOP1__V_RSQ_F16
class	Inst_VOP1__V_RSQ_F32
class	Inst_VOP1__V_RSQ_F64
class	Inst_VOP1__V_SIN_F16
class	Inst_VOP1__V_SIN_F32
class	Inst_VOP1__V_SQRT_F16
class	Inst_VOP1__V_SQRT_F32
class	Inst_VOP1__V_SQRT_F64
class	Inst_VOP1__V_TRUNC_F16
class	Inst_VOP1__V_TRUNC_F32
class	Inst_VOP1__V_TRUNC_F64
class	Inst_VOP2
class	Inst_VOP2__V_ADD_CO_U32
class	Inst_VOP2__V_ADD_F16
class	Inst_VOP2__V_ADD_F32
class	Inst_VOP2__V_ADD_U16
class	Inst_VOP2__V_ADD_U32
class	Inst_VOP2__V_ADDC_CO_U32
class	Inst_VOP2__V_AND_B32
class	Inst_VOP2__V_ASHRREV_I16
class	Inst_VOP2__V_ASHRREV_I32
class	Inst_VOP2__V_CNDMASK_B32
class	Inst_VOP2__V_FMAC_F32
class	Inst_VOP2__V_LDEXP_F16
class	Inst_VOP2__V_LSHLREV_B16
class	Inst_VOP2__V_LSHLREV_B32
class	Inst_VOP2__V_LSHRREV_B16
class	Inst_VOP2__V_LSHRREV_B32
class	Inst_VOP2__V_MAC_F16
class	Inst_VOP2__V_MAC_F32
class	Inst_VOP2__V_MADAK_F16
class	Inst_VOP2__V_MADAK_F32
class	Inst_VOP2__V_MADMK_F16
class	Inst_VOP2__V_MADMK_F32
class	Inst_VOP2__V_MAX_F16
class	Inst_VOP2__V_MAX_F32
class	Inst_VOP2__V_MAX_I16
class	Inst_VOP2__V_MAX_I32
class	Inst_VOP2__V_MAX_U16
class	Inst_VOP2__V_MAX_U32
class	Inst_VOP2__V_MIN_F16
class	Inst_VOP2__V_MIN_F32
class	Inst_VOP2__V_MIN_I16
class	Inst_VOP2__V_MIN_I32
class	Inst_VOP2__V_MIN_U16
class	Inst_VOP2__V_MIN_U32
class	Inst_VOP2__V_MUL_F16
class	Inst_VOP2__V_MUL_F32
class	Inst_VOP2__V_MUL_HI_I32_I24
class	Inst_VOP2__V_MUL_HI_U32_U24
class	Inst_VOP2__V_MUL_I32_I24
class	Inst_VOP2__V_MUL_LEGACY_F32
class	Inst_VOP2__V_MUL_LO_U16
class	Inst_VOP2__V_MUL_U32_U24
class	Inst_VOP2__V_OR_B32
class	Inst_VOP2__V_SUB_CO_U32
class	Inst_VOP2__V_SUB_F16
class	Inst_VOP2__V_SUB_F32
class	Inst_VOP2__V_SUB_U16
class	Inst_VOP2__V_SUB_U32
class	Inst_VOP2__V_SUBB_CO_U32
class	Inst_VOP2__V_SUBBREV_CO_U32
class	Inst_VOP2__V_SUBREV_CO_U32
class	Inst_VOP2__V_SUBREV_F16
class	Inst_VOP2__V_SUBREV_F32
class	Inst_VOP2__V_SUBREV_U16
class	Inst_VOP2__V_SUBREV_U32
class	Inst_VOP2__V_XNOR_B32
class	Inst_VOP2__V_XOR_B32
class	Inst_VOP3__V_ADD3_U32
class	Inst_VOP3__V_ADD_CO_U32
class	Inst_VOP3__V_ADD_F16
class	Inst_VOP3__V_ADD_F32
class	Inst_VOP3__V_ADD_F64
class	Inst_VOP3__V_ADD_LSHL_U32
class	Inst_VOP3__V_ADD_U16
class	Inst_VOP3__V_ADD_U32
class	Inst_VOP3__V_ADDC_CO_U32
class	Inst_VOP3__V_ALIGNBIT_B32
class	Inst_VOP3__V_ALIGNBYTE_B32
class	Inst_VOP3__V_AND_B32
class	Inst_VOP3__V_AND_OR_B32
class	Inst_VOP3__V_ASHRREV_I16
class	Inst_VOP3__V_ASHRREV_I32
class	Inst_VOP3__V_ASHRREV_I64
class	Inst_VOP3__V_BCNT_U32_B32
class	Inst_VOP3__V_BFE_I32
class	Inst_VOP3__V_BFE_U32
class	Inst_VOP3__V_BFI_B32
class	Inst_VOP3__V_BFM_B32
class	Inst_VOP3__V_BFREV_B32
class	Inst_VOP3__V_CEIL_F16
class	Inst_VOP3__V_CEIL_F32
class	Inst_VOP3__V_CEIL_F64
class	Inst_VOP3__V_CLREXCP
class	Inst_VOP3__V_CMP_CLASS_F16
class	Inst_VOP3__V_CMP_CLASS_F32
class	Inst_VOP3__V_CMP_CLASS_F64
class	Inst_VOP3__V_CMP_EQ_F16
class	Inst_VOP3__V_CMP_EQ_F32
class	Inst_VOP3__V_CMP_EQ_F64
class	Inst_VOP3__V_CMP_EQ_I16
class	Inst_VOP3__V_CMP_EQ_I32
class	Inst_VOP3__V_CMP_EQ_I64
class	Inst_VOP3__V_CMP_EQ_U16
class	Inst_VOP3__V_CMP_EQ_U32
class	Inst_VOP3__V_CMP_EQ_U64
class	Inst_VOP3__V_CMP_F_F16
class	Inst_VOP3__V_CMP_F_F32
class	Inst_VOP3__V_CMP_F_F64
class	Inst_VOP3__V_CMP_F_I16
class	Inst_VOP3__V_CMP_F_I32
class	Inst_VOP3__V_CMP_F_I64
class	Inst_VOP3__V_CMP_F_U16
class	Inst_VOP3__V_CMP_F_U32
class	Inst_VOP3__V_CMP_F_U64
class	Inst_VOP3__V_CMP_GE_F16
class	Inst_VOP3__V_CMP_GE_F32
class	Inst_VOP3__V_CMP_GE_F64
class	Inst_VOP3__V_CMP_GE_I16
class	Inst_VOP3__V_CMP_GE_I32
class	Inst_VOP3__V_CMP_GE_I64
class	Inst_VOP3__V_CMP_GE_U16
class	Inst_VOP3__V_CMP_GE_U32
class	Inst_VOP3__V_CMP_GE_U64
class	Inst_VOP3__V_CMP_GT_F16
class	Inst_VOP3__V_CMP_GT_F32
class	Inst_VOP3__V_CMP_GT_F64
class	Inst_VOP3__V_CMP_GT_I16
class	Inst_VOP3__V_CMP_GT_I32
class	Inst_VOP3__V_CMP_GT_I64
class	Inst_VOP3__V_CMP_GT_U16
class	Inst_VOP3__V_CMP_GT_U32
class	Inst_VOP3__V_CMP_GT_U64
class	Inst_VOP3__V_CMP_LE_F16
class	Inst_VOP3__V_CMP_LE_F32
class	Inst_VOP3__V_CMP_LE_F64
class	Inst_VOP3__V_CMP_LE_I16
class	Inst_VOP3__V_CMP_LE_I32
class	Inst_VOP3__V_CMP_LE_I64
class	Inst_VOP3__V_CMP_LE_U16
class	Inst_VOP3__V_CMP_LE_U32
class	Inst_VOP3__V_CMP_LE_U64
class	Inst_VOP3__V_CMP_LG_F16
class	Inst_VOP3__V_CMP_LG_F32
class	Inst_VOP3__V_CMP_LG_F64
class	Inst_VOP3__V_CMP_LT_F16
class	Inst_VOP3__V_CMP_LT_F32
class	Inst_VOP3__V_CMP_LT_F64
class	Inst_VOP3__V_CMP_LT_I16
class	Inst_VOP3__V_CMP_LT_I32
class	Inst_VOP3__V_CMP_LT_I64
class	Inst_VOP3__V_CMP_LT_U16
class	Inst_VOP3__V_CMP_LT_U32
class	Inst_VOP3__V_CMP_LT_U64
class	Inst_VOP3__V_CMP_NE_I16
class	Inst_VOP3__V_CMP_NE_I32
class	Inst_VOP3__V_CMP_NE_I64
class	Inst_VOP3__V_CMP_NE_U16
class	Inst_VOP3__V_CMP_NE_U32
class	Inst_VOP3__V_CMP_NE_U64
class	Inst_VOP3__V_CMP_NEQ_F16
class	Inst_VOP3__V_CMP_NEQ_F32
class	Inst_VOP3__V_CMP_NEQ_F64
class	Inst_VOP3__V_CMP_NGE_F16
class	Inst_VOP3__V_CMP_NGE_F32
class	Inst_VOP3__V_CMP_NGE_F64
class	Inst_VOP3__V_CMP_NGT_F16
class	Inst_VOP3__V_CMP_NGT_F32
class	Inst_VOP3__V_CMP_NGT_F64
class	Inst_VOP3__V_CMP_NLE_F16
class	Inst_VOP3__V_CMP_NLE_F32
class	Inst_VOP3__V_CMP_NLE_F64
class	Inst_VOP3__V_CMP_NLG_F16
class	Inst_VOP3__V_CMP_NLG_F32
class	Inst_VOP3__V_CMP_NLG_F64
class	Inst_VOP3__V_CMP_NLT_F16
class	Inst_VOP3__V_CMP_NLT_F32
class	Inst_VOP3__V_CMP_NLT_F64
class	Inst_VOP3__V_CMP_O_F16
class	Inst_VOP3__V_CMP_O_F32
class	Inst_VOP3__V_CMP_O_F64
class	Inst_VOP3__V_CMP_T_I16
class	Inst_VOP3__V_CMP_T_I32
class	Inst_VOP3__V_CMP_T_I64
class	Inst_VOP3__V_CMP_T_U16
class	Inst_VOP3__V_CMP_T_U32
class	Inst_VOP3__V_CMP_T_U64
class	Inst_VOP3__V_CMP_TRU_F16
class	Inst_VOP3__V_CMP_TRU_F32
class	Inst_VOP3__V_CMP_TRU_F64
class	Inst_VOP3__V_CMP_U_F16
class	Inst_VOP3__V_CMP_U_F32
class	Inst_VOP3__V_CMP_U_F64
class	Inst_VOP3__V_CMPX_CLASS_F16
class	Inst_VOP3__V_CMPX_CLASS_F32
class	Inst_VOP3__V_CMPX_CLASS_F64
class	Inst_VOP3__V_CMPX_EQ_F16
class	Inst_VOP3__V_CMPX_EQ_F32
class	Inst_VOP3__V_CMPX_EQ_F64
class	Inst_VOP3__V_CMPX_EQ_I16
class	Inst_VOP3__V_CMPX_EQ_I32
class	Inst_VOP3__V_CMPX_EQ_I64
class	Inst_VOP3__V_CMPX_EQ_U16
class	Inst_VOP3__V_CMPX_EQ_U32
class	Inst_VOP3__V_CMPX_EQ_U64
class	Inst_VOP3__V_CMPX_F_F16
class	Inst_VOP3__V_CMPX_F_F32
class	Inst_VOP3__V_CMPX_F_F64
class	Inst_VOP3__V_CMPX_F_I16
class	Inst_VOP3__V_CMPX_F_I32
class	Inst_VOP3__V_CMPX_F_I64
class	Inst_VOP3__V_CMPX_F_U16
class	Inst_VOP3__V_CMPX_F_U32
class	Inst_VOP3__V_CMPX_F_U64
class	Inst_VOP3__V_CMPX_GE_F16
class	Inst_VOP3__V_CMPX_GE_F32
class	Inst_VOP3__V_CMPX_GE_F64
class	Inst_VOP3__V_CMPX_GE_I16
class	Inst_VOP3__V_CMPX_GE_I32
class	Inst_VOP3__V_CMPX_GE_I64
class	Inst_VOP3__V_CMPX_GE_U16
class	Inst_VOP3__V_CMPX_GE_U32
class	Inst_VOP3__V_CMPX_GE_U64
class	Inst_VOP3__V_CMPX_GT_F16
class	Inst_VOP3__V_CMPX_GT_F32
class	Inst_VOP3__V_CMPX_GT_F64
class	Inst_VOP3__V_CMPX_GT_I16
class	Inst_VOP3__V_CMPX_GT_I32
class	Inst_VOP3__V_CMPX_GT_I64
class	Inst_VOP3__V_CMPX_GT_U16
class	Inst_VOP3__V_CMPX_GT_U32
class	Inst_VOP3__V_CMPX_GT_U64
class	Inst_VOP3__V_CMPX_LE_F16
class	Inst_VOP3__V_CMPX_LE_F32
class	Inst_VOP3__V_CMPX_LE_F64
class	Inst_VOP3__V_CMPX_LE_I16
class	Inst_VOP3__V_CMPX_LE_I32
class	Inst_VOP3__V_CMPX_LE_I64
class	Inst_VOP3__V_CMPX_LE_U16
class	Inst_VOP3__V_CMPX_LE_U32
class	Inst_VOP3__V_CMPX_LE_U64
class	Inst_VOP3__V_CMPX_LG_F16
class	Inst_VOP3__V_CMPX_LG_F32
class	Inst_VOP3__V_CMPX_LG_F64
class	Inst_VOP3__V_CMPX_LT_F16
class	Inst_VOP3__V_CMPX_LT_F32
class	Inst_VOP3__V_CMPX_LT_F64
class	Inst_VOP3__V_CMPX_LT_I16
class	Inst_VOP3__V_CMPX_LT_I32
class	Inst_VOP3__V_CMPX_LT_I64
class	Inst_VOP3__V_CMPX_LT_U16
class	Inst_VOP3__V_CMPX_LT_U32
class	Inst_VOP3__V_CMPX_LT_U64
class	Inst_VOP3__V_CMPX_NE_I16
class	Inst_VOP3__V_CMPX_NE_I32
class	Inst_VOP3__V_CMPX_NE_I64
class	Inst_VOP3__V_CMPX_NE_U16
class	Inst_VOP3__V_CMPX_NE_U32
class	Inst_VOP3__V_CMPX_NE_U64
class	Inst_VOP3__V_CMPX_NEQ_F16
class	Inst_VOP3__V_CMPX_NEQ_F32
class	Inst_VOP3__V_CMPX_NEQ_F64
class	Inst_VOP3__V_CMPX_NGE_F16
class	Inst_VOP3__V_CMPX_NGE_F32
class	Inst_VOP3__V_CMPX_NGE_F64
class	Inst_VOP3__V_CMPX_NGT_F16
class	Inst_VOP3__V_CMPX_NGT_F32
class	Inst_VOP3__V_CMPX_NGT_F64
class	Inst_VOP3__V_CMPX_NLE_F16
class	Inst_VOP3__V_CMPX_NLE_F32
class	Inst_VOP3__V_CMPX_NLE_F64
class	Inst_VOP3__V_CMPX_NLG_F16
class	Inst_VOP3__V_CMPX_NLG_F32
class	Inst_VOP3__V_CMPX_NLG_F64
class	Inst_VOP3__V_CMPX_NLT_F16
class	Inst_VOP3__V_CMPX_NLT_F32
class	Inst_VOP3__V_CMPX_NLT_F64
class	Inst_VOP3__V_CMPX_O_F16
class	Inst_VOP3__V_CMPX_O_F32
class	Inst_VOP3__V_CMPX_O_F64
class	Inst_VOP3__V_CMPX_T_I16
class	Inst_VOP3__V_CMPX_T_I32
class	Inst_VOP3__V_CMPX_T_I64
class	Inst_VOP3__V_CMPX_T_U16
class	Inst_VOP3__V_CMPX_T_U32
class	Inst_VOP3__V_CMPX_T_U64
class	Inst_VOP3__V_CMPX_TRU_F16
class	Inst_VOP3__V_CMPX_TRU_F32
class	Inst_VOP3__V_CMPX_TRU_F64
class	Inst_VOP3__V_CMPX_U_F16
class	Inst_VOP3__V_CMPX_U_F32
class	Inst_VOP3__V_CMPX_U_F64
class	Inst_VOP3__V_CNDMASK_B32
class	Inst_VOP3__V_COS_F16
class	Inst_VOP3__V_COS_F32
class	Inst_VOP3__V_CUBEID_F32
class	Inst_VOP3__V_CUBEMA_F32
class	Inst_VOP3__V_CUBESC_F32
class	Inst_VOP3__V_CUBETC_F32
class	Inst_VOP3__V_CVT_F16_F32
class	Inst_VOP3__V_CVT_F16_I16
class	Inst_VOP3__V_CVT_F16_U16
class	Inst_VOP3__V_CVT_F32_F16
class	Inst_VOP3__V_CVT_F32_F64
class	Inst_VOP3__V_CVT_F32_I32
class	Inst_VOP3__V_CVT_F32_U32
class	Inst_VOP3__V_CVT_F32_UBYTE0
class	Inst_VOP3__V_CVT_F32_UBYTE1
class	Inst_VOP3__V_CVT_F32_UBYTE2
class	Inst_VOP3__V_CVT_F32_UBYTE3
class	Inst_VOP3__V_CVT_F64_F32
class	Inst_VOP3__V_CVT_F64_I32
class	Inst_VOP3__V_CVT_F64_U32
class	Inst_VOP3__V_CVT_FLR_I32_F32
class	Inst_VOP3__V_CVT_I16_F16
class	Inst_VOP3__V_CVT_I32_F32
class	Inst_VOP3__V_CVT_I32_F64
class	Inst_VOP3__V_CVT_OFF_F32_I4
class	Inst_VOP3__V_CVT_PK_FP8_F32
class	Inst_VOP3__V_CVT_PK_I16_I32
class	Inst_VOP3__V_CVT_PK_U16_U32
class	Inst_VOP3__V_CVT_PK_U8_F32
class	Inst_VOP3__V_CVT_PKACCUM_U8_F32
class	Inst_VOP3__V_CVT_PKNORM_I16_F32
class	Inst_VOP3__V_CVT_PKNORM_U16_F32
class	Inst_VOP3__V_CVT_PKRTZ_F16_F32
class	Inst_VOP3__V_CVT_RPI_I32_F32
class	Inst_VOP3__V_CVT_U16_F16
class	Inst_VOP3__V_CVT_U32_F32
class	Inst_VOP3__V_CVT_U32_F64
class	Inst_VOP3__V_DIV_FIXUP_F16
class	Inst_VOP3__V_DIV_FIXUP_F32
class	Inst_VOP3__V_DIV_FIXUP_F64
class	Inst_VOP3__V_DIV_FMAS_F32
class	Inst_VOP3__V_DIV_FMAS_F64
class	Inst_VOP3__V_DIV_SCALE_F32
class	Inst_VOP3__V_DIV_SCALE_F64
class	Inst_VOP3__V_EXP_F16
class	Inst_VOP3__V_EXP_F32
class	Inst_VOP3__V_EXP_LEGACY_F32
class	Inst_VOP3__V_FFBH_I32
class	Inst_VOP3__V_FFBH_U32
class	Inst_VOP3__V_FFBL_B32
class	Inst_VOP3__V_FLOOR_F16
class	Inst_VOP3__V_FLOOR_F32
class	Inst_VOP3__V_FLOOR_F64
class	Inst_VOP3__V_FMA_F16
class	Inst_VOP3__V_FMA_F32
class	Inst_VOP3__V_FMA_F64
class	Inst_VOP3__V_FMAC_F32
class	Inst_VOP3__V_FRACT_F16
class	Inst_VOP3__V_FRACT_F32
class	Inst_VOP3__V_FRACT_F64
class	Inst_VOP3__V_FREXP_EXP_I16_F16
class	Inst_VOP3__V_FREXP_EXP_I32_F32
class	Inst_VOP3__V_FREXP_EXP_I32_F64
class	Inst_VOP3__V_FREXP_MANT_F16
class	Inst_VOP3__V_FREXP_MANT_F32
class	Inst_VOP3__V_FREXP_MANT_F64
class	Inst_VOP3__V_INTERP_MOV_F32
class	Inst_VOP3__V_INTERP_P1_F32
class	Inst_VOP3__V_INTERP_P1LL_F16
class	Inst_VOP3__V_INTERP_P1LV_F16
class	Inst_VOP3__V_INTERP_P2_F16
class	Inst_VOP3__V_INTERP_P2_F32
class	Inst_VOP3__V_LDEXP_F16
class	Inst_VOP3__V_LDEXP_F32
class	Inst_VOP3__V_LDEXP_F64
class	Inst_VOP3__V_LERP_U8
class	Inst_VOP3__V_LOG_F16
class	Inst_VOP3__V_LOG_F32
class	Inst_VOP3__V_LOG_LEGACY_F32
class	Inst_VOP3__V_LSHL_ADD_U32
class	Inst_VOP3__V_LSHL_ADD_U64
class	Inst_VOP3__V_LSHL_OR_B32
class	Inst_VOP3__V_LSHLREV_B16
class	Inst_VOP3__V_LSHLREV_B32
class	Inst_VOP3__V_LSHLREV_B64
class	Inst_VOP3__V_LSHRREV_B16
class	Inst_VOP3__V_LSHRREV_B32
class	Inst_VOP3__V_LSHRREV_B64
class	Inst_VOP3__V_MAC_F16
class	Inst_VOP3__V_MAC_F32
class	Inst_VOP3__V_MAD_F16
class	Inst_VOP3__V_MAD_F32
class	Inst_VOP3__V_MAD_I16
class	Inst_VOP3__V_MAD_I32_I24
class	Inst_VOP3__V_MAD_I64_I32
class	Inst_VOP3__V_MAD_LEGACY_F32
class	Inst_VOP3__V_MAD_U16
class	Inst_VOP3__V_MAD_U32_U24
class	Inst_VOP3__V_MAD_U64_U32
class	Inst_VOP3__V_MAX3_F32
class	Inst_VOP3__V_MAX3_I32
class	Inst_VOP3__V_MAX3_U32
class	Inst_VOP3__V_MAX_F16
class	Inst_VOP3__V_MAX_F32
class	Inst_VOP3__V_MAX_F64
class	Inst_VOP3__V_MAX_I16
class	Inst_VOP3__V_MAX_I32
class	Inst_VOP3__V_MAX_U16
class	Inst_VOP3__V_MAX_U32
class	Inst_VOP3__V_MBCNT_HI_U32_B32
class	Inst_VOP3__V_MBCNT_LO_U32_B32
class	Inst_VOP3__V_MED3_F32
class	Inst_VOP3__V_MED3_I32
class	Inst_VOP3__V_MED3_U32
class	Inst_VOP3__V_MIN3_F32
class	Inst_VOP3__V_MIN3_I32
class	Inst_VOP3__V_MIN3_U32
class	Inst_VOP3__V_MIN_F16
class	Inst_VOP3__V_MIN_F32
class	Inst_VOP3__V_MIN_F64
class	Inst_VOP3__V_MIN_I16
class	Inst_VOP3__V_MIN_I32
class	Inst_VOP3__V_MIN_U16
class	Inst_VOP3__V_MIN_U32
class	Inst_VOP3__V_MOV_B32
class	Inst_VOP3__V_MOV_FED_B32
class	Inst_VOP3__V_MQSAD_PK_U16_U8
class	Inst_VOP3__V_MQSAD_U32_U8
class	Inst_VOP3__V_MSAD_U8
class	Inst_VOP3__V_MUL_F16
class	Inst_VOP3__V_MUL_F32
class	Inst_VOP3__V_MUL_F64
class	Inst_VOP3__V_MUL_HI_I32
class	Inst_VOP3__V_MUL_HI_I32_I24
class	Inst_VOP3__V_MUL_HI_U32
class	Inst_VOP3__V_MUL_HI_U32_U24
class	Inst_VOP3__V_MUL_I32_I24
class	Inst_VOP3__V_MUL_LEGACY_F32
class	Inst_VOP3__V_MUL_LO_U16
class	Inst_VOP3__V_MUL_LO_U32
class	Inst_VOP3__V_MUL_U32_U24
class	Inst_VOP3__V_NOP
class	Inst_VOP3__V_NOT_B32
class	Inst_VOP3__V_OR3_B32
class	Inst_VOP3__V_OR_B32
class	Inst_VOP3__V_PERM_B32
class	Inst_VOP3__V_QSAD_PK_U16_U8
class	Inst_VOP3__V_RCP_F16
class	Inst_VOP3__V_RCP_F32
class	Inst_VOP3__V_RCP_F64
class	Inst_VOP3__V_RCP_IFLAG_F32
class	Inst_VOP3__V_READLANE_B32
class	Inst_VOP3__V_RNDNE_F16
class	Inst_VOP3__V_RNDNE_F32
class	Inst_VOP3__V_RNDNE_F64
class	Inst_VOP3__V_RSQ_F16
class	Inst_VOP3__V_RSQ_F32
class	Inst_VOP3__V_RSQ_F64
class	Inst_VOP3__V_SAD_HI_U8
class	Inst_VOP3__V_SAD_U16
class	Inst_VOP3__V_SAD_U32
class	Inst_VOP3__V_SAD_U8
class	Inst_VOP3__V_SIN_F16
class	Inst_VOP3__V_SIN_F32
class	Inst_VOP3__V_SQRT_F16
class	Inst_VOP3__V_SQRT_F32
class	Inst_VOP3__V_SQRT_F64
class	Inst_VOP3__V_SUB_CO_U32
class	Inst_VOP3__V_SUB_F16
class	Inst_VOP3__V_SUB_F32
class	Inst_VOP3__V_SUB_U16
class	Inst_VOP3__V_SUB_U32
class	Inst_VOP3__V_SUBB_CO_U32
class	Inst_VOP3__V_SUBBREV_CO_U32
class	Inst_VOP3__V_SUBREV_CO_U32
class	Inst_VOP3__V_SUBREV_F16
class	Inst_VOP3__V_SUBREV_F32
class	Inst_VOP3__V_SUBREV_U16
class	Inst_VOP3__V_SUBREV_U32
class	Inst_VOP3__V_TRIG_PREOP_F64
class	Inst_VOP3__V_TRUNC_F16
class	Inst_VOP3__V_TRUNC_F32
class	Inst_VOP3__V_TRUNC_F64
class	Inst_VOP3__V_WRITELANE_B32
class	Inst_VOP3__V_XAD_U32
class	Inst_VOP3__V_XOR_B32
class	Inst_VOP3A
class	Inst_VOP3B
class	Inst_VOP3P
class	Inst_VOP3P__1OP
class	Inst_VOP3P__2OP_X16
class	Inst_VOP3P__3OP_X16
class	Inst_VOP3P__V_ACCVGPR_READ
class	Inst_VOP3P__V_ACCVGPR_WRITE
class	Inst_VOP3P__V_DOT2_F32_F16
class	Inst_VOP3P__V_DOT2_I32_I16
class	Inst_VOP3P__V_DOT2_U32_U16
class	Inst_VOP3P__V_DOT4_I32_I8
class	Inst_VOP3P__V_DOT4_U32_U8
class	Inst_VOP3P__V_DOT8_I32_I4
class	Inst_VOP3P__V_DOT8_U32_U4
class	Inst_VOP3P__V_PK_ADD_F16
class	Inst_VOP3P__V_PK_ADD_F32
class	Inst_VOP3P__V_PK_ADD_I16
class	Inst_VOP3P__V_PK_ADD_U16
class	Inst_VOP3P__V_PK_ASHRREV_B16
class	Inst_VOP3P__V_PK_FMA_F16
class	Inst_VOP3P__V_PK_FMA_F32
class	Inst_VOP3P__V_PK_LSHLREV_B16
class	Inst_VOP3P__V_PK_LSHRREV_B16
class	Inst_VOP3P__V_PK_MAD_I16
class	Inst_VOP3P__V_PK_MAD_U16
class	Inst_VOP3P__V_PK_MAX_F16
class	Inst_VOP3P__V_PK_MAX_I16
class	Inst_VOP3P__V_PK_MAX_U16
class	Inst_VOP3P__V_PK_MIN_F16
class	Inst_VOP3P__V_PK_MIN_I16
class	Inst_VOP3P__V_PK_MIN_U16
class	Inst_VOP3P__V_PK_MOV_B32
class	Inst_VOP3P__V_PK_MUL_F16
class	Inst_VOP3P__V_PK_MUL_F32
class	Inst_VOP3P__V_PK_MUL_LO_U16
class	Inst_VOP3P__V_PK_SUB_I16
class	Inst_VOP3P__V_PK_SUB_U16
class	Inst_VOP3P_MAI
class	Inst_VOP3P_MAI__V_MFMA
class	Inst_VOP3P_MAI__V_MFMA_I8
class	Inst_VOP3P_MAI__V_MFMA_MXFP
class	Inst_VOPC
class	Inst_VOPC__V_CMP_CLASS_F16
class	Inst_VOPC__V_CMP_CLASS_F32
class	Inst_VOPC__V_CMP_CLASS_F64
class	Inst_VOPC__V_CMP_EQ_F16
class	Inst_VOPC__V_CMP_EQ_F32
class	Inst_VOPC__V_CMP_EQ_F64
class	Inst_VOPC__V_CMP_EQ_I16
class	Inst_VOPC__V_CMP_EQ_I32
class	Inst_VOPC__V_CMP_EQ_I64
class	Inst_VOPC__V_CMP_EQ_U16
class	Inst_VOPC__V_CMP_EQ_U32
class	Inst_VOPC__V_CMP_EQ_U64
class	Inst_VOPC__V_CMP_F_F16
class	Inst_VOPC__V_CMP_F_F32
class	Inst_VOPC__V_CMP_F_F64
class	Inst_VOPC__V_CMP_F_I16
class	Inst_VOPC__V_CMP_F_I32
class	Inst_VOPC__V_CMP_F_I64
class	Inst_VOPC__V_CMP_F_U16
class	Inst_VOPC__V_CMP_F_U32
class	Inst_VOPC__V_CMP_F_U64
class	Inst_VOPC__V_CMP_GE_F16
class	Inst_VOPC__V_CMP_GE_F32
class	Inst_VOPC__V_CMP_GE_F64
class	Inst_VOPC__V_CMP_GE_I16
class	Inst_VOPC__V_CMP_GE_I32
class	Inst_VOPC__V_CMP_GE_I64
class	Inst_VOPC__V_CMP_GE_U16
class	Inst_VOPC__V_CMP_GE_U32
class	Inst_VOPC__V_CMP_GE_U64
class	Inst_VOPC__V_CMP_GT_F16
class	Inst_VOPC__V_CMP_GT_F32
class	Inst_VOPC__V_CMP_GT_F64
class	Inst_VOPC__V_CMP_GT_I16
class	Inst_VOPC__V_CMP_GT_I32
class	Inst_VOPC__V_CMP_GT_I64
class	Inst_VOPC__V_CMP_GT_U16
class	Inst_VOPC__V_CMP_GT_U32
class	Inst_VOPC__V_CMP_GT_U64
class	Inst_VOPC__V_CMP_LE_F16
class	Inst_VOPC__V_CMP_LE_F32
class	Inst_VOPC__V_CMP_LE_F64
class	Inst_VOPC__V_CMP_LE_I16
class	Inst_VOPC__V_CMP_LE_I32
class	Inst_VOPC__V_CMP_LE_I64
class	Inst_VOPC__V_CMP_LE_U16
class	Inst_VOPC__V_CMP_LE_U32
class	Inst_VOPC__V_CMP_LE_U64
class	Inst_VOPC__V_CMP_LG_F16
class	Inst_VOPC__V_CMP_LG_F32
class	Inst_VOPC__V_CMP_LG_F64
class	Inst_VOPC__V_CMP_LT_F16
class	Inst_VOPC__V_CMP_LT_F32
class	Inst_VOPC__V_CMP_LT_F64
class	Inst_VOPC__V_CMP_LT_I16
class	Inst_VOPC__V_CMP_LT_I32
class	Inst_VOPC__V_CMP_LT_I64
class	Inst_VOPC__V_CMP_LT_U16
class	Inst_VOPC__V_CMP_LT_U32
class	Inst_VOPC__V_CMP_LT_U64
class	Inst_VOPC__V_CMP_NE_I16
class	Inst_VOPC__V_CMP_NE_I32
class	Inst_VOPC__V_CMP_NE_I64
class	Inst_VOPC__V_CMP_NE_U16
class	Inst_VOPC__V_CMP_NE_U32
class	Inst_VOPC__V_CMP_NE_U64
class	Inst_VOPC__V_CMP_NEQ_F16
class	Inst_VOPC__V_CMP_NEQ_F32
class	Inst_VOPC__V_CMP_NEQ_F64
class	Inst_VOPC__V_CMP_NGE_F16
class	Inst_VOPC__V_CMP_NGE_F32
class	Inst_VOPC__V_CMP_NGE_F64
class	Inst_VOPC__V_CMP_NGT_F16
class	Inst_VOPC__V_CMP_NGT_F32
class	Inst_VOPC__V_CMP_NGT_F64
class	Inst_VOPC__V_CMP_NLE_F16
class	Inst_VOPC__V_CMP_NLE_F32
class	Inst_VOPC__V_CMP_NLE_F64
class	Inst_VOPC__V_CMP_NLG_F16
class	Inst_VOPC__V_CMP_NLG_F32
class	Inst_VOPC__V_CMP_NLG_F64
class	Inst_VOPC__V_CMP_NLT_F16
class	Inst_VOPC__V_CMP_NLT_F32
class	Inst_VOPC__V_CMP_NLT_F64
class	Inst_VOPC__V_CMP_O_F16
class	Inst_VOPC__V_CMP_O_F32
class	Inst_VOPC__V_CMP_O_F64
class	Inst_VOPC__V_CMP_T_I16
class	Inst_VOPC__V_CMP_T_I32
class	Inst_VOPC__V_CMP_T_I64
class	Inst_VOPC__V_CMP_T_U16
class	Inst_VOPC__V_CMP_T_U32
class	Inst_VOPC__V_CMP_T_U64
class	Inst_VOPC__V_CMP_TRU_F16
class	Inst_VOPC__V_CMP_TRU_F32
class	Inst_VOPC__V_CMP_TRU_F64
class	Inst_VOPC__V_CMP_U_F16
class	Inst_VOPC__V_CMP_U_F32
class	Inst_VOPC__V_CMP_U_F64
class	Inst_VOPC__V_CMPX_CLASS_F16
class	Inst_VOPC__V_CMPX_CLASS_F32
class	Inst_VOPC__V_CMPX_CLASS_F64
class	Inst_VOPC__V_CMPX_EQ_F16
class	Inst_VOPC__V_CMPX_EQ_F32
class	Inst_VOPC__V_CMPX_EQ_F64
class	Inst_VOPC__V_CMPX_EQ_I16
class	Inst_VOPC__V_CMPX_EQ_I32
class	Inst_VOPC__V_CMPX_EQ_I64
class	Inst_VOPC__V_CMPX_EQ_U16
class	Inst_VOPC__V_CMPX_EQ_U32
class	Inst_VOPC__V_CMPX_EQ_U64
class	Inst_VOPC__V_CMPX_F_F16
class	Inst_VOPC__V_CMPX_F_F32
class	Inst_VOPC__V_CMPX_F_F64
class	Inst_VOPC__V_CMPX_F_I16
class	Inst_VOPC__V_CMPX_F_I32
class	Inst_VOPC__V_CMPX_F_I64
class	Inst_VOPC__V_CMPX_F_U16
class	Inst_VOPC__V_CMPX_F_U32
class	Inst_VOPC__V_CMPX_F_U64
class	Inst_VOPC__V_CMPX_GE_F16
class	Inst_VOPC__V_CMPX_GE_F32
class	Inst_VOPC__V_CMPX_GE_F64
class	Inst_VOPC__V_CMPX_GE_I16
class	Inst_VOPC__V_CMPX_GE_I32
class	Inst_VOPC__V_CMPX_GE_I64
class	Inst_VOPC__V_CMPX_GE_U16
class	Inst_VOPC__V_CMPX_GE_U32
class	Inst_VOPC__V_CMPX_GE_U64
class	Inst_VOPC__V_CMPX_GT_F16
class	Inst_VOPC__V_CMPX_GT_F32
class	Inst_VOPC__V_CMPX_GT_F64
class	Inst_VOPC__V_CMPX_GT_I16
class	Inst_VOPC__V_CMPX_GT_I32
class	Inst_VOPC__V_CMPX_GT_I64
class	Inst_VOPC__V_CMPX_GT_U16
class	Inst_VOPC__V_CMPX_GT_U32
class	Inst_VOPC__V_CMPX_GT_U64
class	Inst_VOPC__V_CMPX_LE_F16
class	Inst_VOPC__V_CMPX_LE_F32
class	Inst_VOPC__V_CMPX_LE_F64
class	Inst_VOPC__V_CMPX_LE_I16
class	Inst_VOPC__V_CMPX_LE_I32
class	Inst_VOPC__V_CMPX_LE_I64
class	Inst_VOPC__V_CMPX_LE_U16
class	Inst_VOPC__V_CMPX_LE_U32
class	Inst_VOPC__V_CMPX_LE_U64
class	Inst_VOPC__V_CMPX_LG_F16
class	Inst_VOPC__V_CMPX_LG_F32
class	Inst_VOPC__V_CMPX_LG_F64
class	Inst_VOPC__V_CMPX_LT_F16
class	Inst_VOPC__V_CMPX_LT_F32
class	Inst_VOPC__V_CMPX_LT_F64
class	Inst_VOPC__V_CMPX_LT_I16
class	Inst_VOPC__V_CMPX_LT_I32
class	Inst_VOPC__V_CMPX_LT_I64
class	Inst_VOPC__V_CMPX_LT_U16
class	Inst_VOPC__V_CMPX_LT_U32
class	Inst_VOPC__V_CMPX_LT_U64
class	Inst_VOPC__V_CMPX_NE_I16
class	Inst_VOPC__V_CMPX_NE_I32
class	Inst_VOPC__V_CMPX_NE_I64
class	Inst_VOPC__V_CMPX_NE_U16
class	Inst_VOPC__V_CMPX_NE_U32
class	Inst_VOPC__V_CMPX_NE_U64
class	Inst_VOPC__V_CMPX_NEQ_F16
class	Inst_VOPC__V_CMPX_NEQ_F32
class	Inst_VOPC__V_CMPX_NEQ_F64
class	Inst_VOPC__V_CMPX_NGE_F16
class	Inst_VOPC__V_CMPX_NGE_F32
class	Inst_VOPC__V_CMPX_NGE_F64
class	Inst_VOPC__V_CMPX_NGT_F16
class	Inst_VOPC__V_CMPX_NGT_F32
class	Inst_VOPC__V_CMPX_NGT_F64
class	Inst_VOPC__V_CMPX_NLE_F16
class	Inst_VOPC__V_CMPX_NLE_F32
class	Inst_VOPC__V_CMPX_NLE_F64
class	Inst_VOPC__V_CMPX_NLG_F16
class	Inst_VOPC__V_CMPX_NLG_F32
class	Inst_VOPC__V_CMPX_NLG_F64
class	Inst_VOPC__V_CMPX_NLT_F16
class	Inst_VOPC__V_CMPX_NLT_F32
class	Inst_VOPC__V_CMPX_NLT_F64
class	Inst_VOPC__V_CMPX_O_F16
class	Inst_VOPC__V_CMPX_O_F32
class	Inst_VOPC__V_CMPX_O_F64
class	Inst_VOPC__V_CMPX_T_I16
class	Inst_VOPC__V_CMPX_T_I32
class	Inst_VOPC__V_CMPX_T_I64
class	Inst_VOPC__V_CMPX_T_U16
class	Inst_VOPC__V_CMPX_T_U32
class	Inst_VOPC__V_CMPX_T_U64
class	Inst_VOPC__V_CMPX_TRU_F16
class	Inst_VOPC__V_CMPX_TRU_F32
class	Inst_VOPC__V_CMPX_TRU_F64
class	Inst_VOPC__V_CMPX_U_F16
class	Inst_VOPC__V_CMPX_U_F32
class	Inst_VOPC__V_CMPX_U_F64
union	InstFormat
class	Operand
struct	OpTraits
	convenience traits so we can automatically infer the correct FP type without looking at the number of dwords (i.e., to determine if we need a float or a double when creating FP constants). More...
struct	OpTraits< ScalarRegF64 >
struct	OpTraits< ScalarRegU64 >
class	PackedReg
class	PageFault
class	ScalarOperand
struct	StatusReg
class	VecOperand
class	VegaFault
class	VEGAGPUStaticInst
class	Walker

Typedefs
using	IsaDecodeMethod = GPUStaticInst *(Decoder::*)(MachInst)
typedef uint8_t	ScalarRegU8
typedef int8_t	ScalarRegI8
typedef uint16_t	ScalarRegU16
typedef int16_t	ScalarRegI16
typedef uint32_t	ScalarRegU32
typedef int32_t	ScalarRegI32
typedef float	ScalarRegF32
typedef uint64_t	ScalarRegU64
typedef int64_t	ScalarRegI64
typedef double	ScalarRegF64
typedef uint8_t	VecElemU8
typedef int8_t	VecElemI8
typedef uint16_t	VecElemU16
typedef int16_t	VecElemI16
typedef uint32_t	VecElemU32
typedef int32_t	VecElemI32
typedef float	VecElemF32
typedef uint64_t	VecElemU64
typedef int64_t	VecElemI64
typedef double	VecElemF64
using	VecRegContainerU32 = VecRegContainer< sizeof(VecElemU32) *NumVecElemPerVecReg >
using	VecRegContainerU64 = VecRegContainer< sizeof(VecElemU64) *NumVecElemPerVecReg >
typedef uint64_t	RawMachInst
	used to represnt a GPU inst in its raw format.
typedef InstFormat *	MachInst
	used to represent the encoding of a VEGA inst.
using	Inst_VOP3P_MAI__V_MFMA_F32_4X4X1_16B_F32 = Inst_VOP3P_MAI__V_MFMA< 1, 4, 4, 1, 16, ConstVecOperandF32, VecOperandF32, &MNEM__V_MFMA_F32_4X4X1_16B_F32 >
using	Inst_VOP3P_MAI__V_MFMA_F32_32X32X1_2B_F32 = Inst_VOP3P_MAI__V_MFMA< 1, 32, 32, 1, 2, ConstVecOperandF32, VecOperandF32, &MNEM__V_MFMA_F32_32X32X1_2B_F32 >
using	Inst_VOP3P_MAI__V_MFMA_F32_32X32X2_F32 = Inst_VOP3P_MAI__V_MFMA< 1, 32, 32, 2, 1, ConstVecOperandF32, VecOperandF32, &MNEM__V_MFMA_F32_32X32X2_F32 >
using	Inst_VOP3P_MAI__V_MFMA_F32_16X16X4_F32 = Inst_VOP3P_MAI__V_MFMA< 1, 16, 16, 4, 1, ConstVecOperandF32, VecOperandF32, &MNEM__V_MFMA_F32_16X16X4_F32 >
using	Inst_VOP3P_MAI__V_MFMA_F32_16X16X1_4B_F32 = Inst_VOP3P_MAI__V_MFMA< 1, 16, 16, 1, 4, ConstVecOperandF32, VecOperandF32, &MNEM__V_MFMA_F32_16X16X1_4B_F32 >
using	Inst_VOP3P_MAI__V_MFMA_F64_4X4X4_4B_F64 = Inst_VOP3P_MAI__V_MFMA< 2, 4, 4, 4, 4, ConstVecOperandF64, VecOperandF64, &MNEM__V_MFMA_F64_4X4X4_4B_F64 >
using	Inst_VOP3P_MAI__V_MFMA_F64_16X16X4_F64 = Inst_VOP3P_MAI__V_MFMA< 2, 16, 16, 4, 1, ConstVecOperandF64, VecOperandF64, &MNEM__V_MFMA_F64_16X16X4_F64 >
using	Inst_VOP3P_MAI__V_MFMA_F32_16X16X16_F16 = Inst_VOP3P_MAI__V_MFMA_MXFP< 16, 16, 16, 1, AMDGPU::mxfloat16, &MNEM__V_MFMA_F32_16X16X16_F16 >
using	Inst_VOP3P_MAI__V_MFMA_F32_16X16X4_4B_F16 = Inst_VOP3P_MAI__V_MFMA_MXFP< 16, 16, 4, 4, AMDGPU::mxfloat16, &MNEM__V_MFMA_F32_16X16X4_4B_F16 >
using	Inst_VOP3P_MAI__V_MFMA_F32_32X32X4_2B_F16 = Inst_VOP3P_MAI__V_MFMA_MXFP< 32, 32, 4, 2, AMDGPU::mxfloat16, &MNEM__V_MFMA_F32_32X32X4_2B_F16 >
using	Inst_VOP3P_MAI__V_MFMA_F32_32X32X8_F16 = Inst_VOP3P_MAI__V_MFMA_MXFP< 32, 32, 8, 1, AMDGPU::mxfloat16, &NMEM__V_MFMA_F32_32X32X8_F16 >
using	Inst_VOP3P_MAI__V_MFMA_F32_4X4X4_16B_F16 = Inst_VOP3P_MAI__V_MFMA_MXFP< 4, 4, 4, 16, AMDGPU::mxfloat16, &MNEM__V_MFMA_F32_4X4X4_16B_F16 >
using	Inst_VOP3P_MAI__V_MFMA_F32_32X32X8_BF16 = Inst_VOP3P_MAI__V_MFMA_MXFP< 32, 32, 8, 1, AMDGPU::mxbfloat16, &MNEM__V_MFMA_F32_32X32X8_BF16 >
using	Inst_VOP3P_MAI__V_MFMA_I32_16X16X16_I8 = Inst_VOP3P_MAI__V_MFMA_I8< 16, 16, 16, 1, &MNEM__V_MFMA_I32_16X16X16_I8 >
using	half = uint16_t
using	ScalarOperandU8 = ScalarOperand< ScalarRegU8, false, 1 >
using	ScalarOperandI8 = ScalarOperand< ScalarRegI8, false, 1 >
using	ScalarOperandU16 = ScalarOperand< ScalarRegU16, false, 1 >
using	ScalarOperandI16 = ScalarOperand< ScalarRegI16, false, 1 >
using	ScalarOperandU32 = ScalarOperand< ScalarRegU32, false >
using	ScalarOperandI32 = ScalarOperand< ScalarRegI32, false >
using	ScalarOperandF32 = ScalarOperand< ScalarRegF32, false >
using	ScalarOperandU64 = ScalarOperand< ScalarRegU64, false >
using	ScalarOperandI64 = ScalarOperand< ScalarRegI64, false >
using	ScalarOperandF64 = ScalarOperand< ScalarRegF64, false >
using	ScalarOperandU128 = ScalarOperand< ScalarRegU32, false, 4 >
using	ScalarOperandU256 = ScalarOperand< ScalarRegU32, false, 8 >
using	ScalarOperandU512 = ScalarOperand< ScalarRegU32, false, 16 >
using	ConstScalarOperandU8 = ScalarOperand< ScalarRegU8, true, 1 >
using	ConstScalarOperandI8 = ScalarOperand< ScalarRegI8, true, 1 >
using	ConstScalarOperandU16 = ScalarOperand< ScalarRegU16, true, 1 >
using	ConstScalarOperandI16 = ScalarOperand< ScalarRegI16, true, 1 >
using	ConstScalarOperandU32 = ScalarOperand< ScalarRegU32, true >
using	ConstScalarOperandI32 = ScalarOperand< ScalarRegI32, true >
using	ConstScalarOperandF32 = ScalarOperand< ScalarRegF32, true >
using	ConstScalarOperandU64 = ScalarOperand< ScalarRegU64, true >
using	ConstScalarOperandI64 = ScalarOperand< ScalarRegI64, true >
using	ConstScalarOperandF64 = ScalarOperand< ScalarRegF64, true >
using	ConstScalarOperandU128 = ScalarOperand< ScalarRegU32, true, 4 >
using	ConstScalarOperandU256 = ScalarOperand< ScalarRegU32, true, 8 >
using	ConstScalarOperandU512 = ScalarOperand< ScalarRegU32, true, 16 >
using	VecOperandU8 = VecOperand< VecElemU8, false, 1 >
using	VecOperandI8 = VecOperand< VecElemI8, false, 1 >
using	VecOperandU16 = VecOperand< VecElemU16, false, 1 >
using	VecOperandI16 = VecOperand< VecElemI16, false, 1 >
using	VecOperandU32 = VecOperand< VecElemU32, false >
using	VecOperandI32 = VecOperand< VecElemI32, false >
using	VecOperandF32 = VecOperand< VecElemF32, false >
using	VecOperandU64 = VecOperand< VecElemU64, false >
using	VecOperandF64 = VecOperand< VecElemF64, false >
using	VecOperandI64 = VecOperand< VecElemI64, false >
using	VecOperandU96 = VecOperand< VecElemU32, false, 3 >
using	VecOperandU128 = VecOperand< VecElemU32, false, 4 >
using	VecOperandU256 = VecOperand< VecElemU32, false, 8 >
using	VecOperandU512 = VecOperand< VecElemU32, false, 16 >
using	ConstVecOperandU8 = VecOperand< VecElemU8, true, 1 >
using	ConstVecOperandI8 = VecOperand< VecElemI8, true, 1 >
using	ConstVecOperandU16 = VecOperand< VecElemU16, true, 1 >
using	ConstVecOperandI16 = VecOperand< VecElemI16, true, 1 >
using	ConstVecOperandU32 = VecOperand< VecElemU32, true >
using	ConstVecOperandI32 = VecOperand< VecElemI32, true >
using	ConstVecOperandF32 = VecOperand< VecElemF32, true >
using	ConstVecOperandU64 = VecOperand< VecElemU64, true >
using	ConstVecOperandI64 = VecOperand< VecElemI64, true >
using	ConstVecOperandF64 = VecOperand< VecElemF64, true >
using	ConstVecOperandU96 = VecOperand< VecElemU32, true, 3 >
using	ConstVecOperandU128 = VecOperand< VecElemU32, true, 4 >
using	ConstVecOperandU256 = VecOperand< VecElemU32, true, 8 >
using	ConstVecOperandU512 = VecOperand< VecElemU32, true, 16 >

Enumerations
enum	ExceptionCode : uint64_t { INST_PAGE = 0 , LOAD_PAGE = 1 , STORE_PAGE = 2 }
enum	OpSelector : int {   REG_SGPR_MIN = 0 , REG_SGPR_MAX = 101 , REG_FLAT_SCRATCH_LO = 102 , REG_FLAT_SCRATCH_HI = 103 ,   REG_XNACK_MASK_LO = 104 , REG_XNACK_MASK_HI = 105 , REG_VCC_LO = 106 , REG_VCC_HI = 107 ,   REG_TBA_LO = 108 , REG_TBA_HI = 109 , REG_TMA_LO = 110 , REG_TMA_HI = 111 ,   REG_TTMP_0 = 112 , REG_TTMP_1 = 113 , REG_TTMP_2 = 114 , REG_TTMP_3 = 115 ,   REG_TTMP_4 = 116 , REG_TTMP_5 = 117 , REG_TTMP_6 = 118 , REG_TTMP_7 = 119 ,   REG_TTMP_8 = 120 , REG_TTMP_9 = 121 , REG_TTMP_10 = 122 , REG_TTMP_11 = 123 ,   REG_M0 = 124 , REG_RESERVED_1 = 125 , REG_EXEC_LO = 126 , REG_EXEC_HI = 127 ,   REG_ZERO = 128 , REG_INT_CONST_POS_MIN = 129 , REG_INT_CONST_POS_MAX = 192 , REG_INT_CONST_NEG_MIN = 193 ,   REG_INT_CONST_NEG_MAX = 208 , REG_RESERVED_2 = 209 , REG_RESERVED_3 = 210 , REG_RESERVED_4 = 211 ,   REG_RESERVED_5 = 212 , REG_RESERVED_6 = 213 , REG_RESERVED_7 = 214 , REG_RESERVED_8 = 215 ,   REG_RESERVED_9 = 216 , REG_RESERVED_10 = 217 , REG_RESERVED_11 = 218 , REG_RESERVED_12 = 219 ,   REG_RESERVED_13 = 220 , REG_RESERVED_14 = 221 , REG_RESERVED_15 = 222 , REG_RESERVED_16 = 223 ,   REG_RESERVED_17 = 224 , REG_RESERVED_18 = 225 , REG_RESERVED_19 = 226 , REG_RESERVED_20 = 227 ,   REG_RESERVED_21 = 228 , REG_RESERVED_22 = 229 , REG_RESERVED_23 = 230 , REG_RESERVED_24 = 231 ,   REG_RESERVED_25 = 232 , REG_RESERVED_26 = 233 , REG_RESERVED_27 = 234 , REG_SHARED_BASE = 235 ,   REG_SHARED_LIMIT = 236 , REG_PRIVATE_BASE = 237 , REG_PRIVATE_LIMIT = 238 , REG_RESERVED_32 = 239 ,   REG_POS_HALF = 240 , REG_NEG_HALF = 241 , REG_POS_ONE = 242 , REG_NEG_ONE = 243 ,   REG_POS_TWO = 244 , REG_NEG_TWO = 245 , REG_POS_FOUR = 246 , REG_NEG_FOUR = 247 ,   REG_PI = 248 , REG_SRC_SWDA = 249 , REG_SRC_DPP = 250 , REG_VCCZ = 251 ,   REG_EXECZ = 252 , REG_SCC = 253 , REG_LDS_DIRECT = 254 , REG_SRC_LITERAL = 255 ,   REG_VGPR_MIN = 256 , REG_VGPR_MAX = 767 }

Functions
std::string	opSelectorToRegSym (int idx, int numRegs)
int	opSelectorToRegIdx (int idx, int numScalarRegs)
bool	isPosConstVal (int opIdx)
bool	isNegConstVal (int opIdx)
bool	isConstVal (int opIdx)
bool	isLiteral (int opIdx)
bool	isExecMask (int opIdx)
bool	isVccReg (int opIdx)
bool	isFlatScratchReg (int opIdx)
bool	isScalarReg (int opIdx)
bool	isVectorReg (int opIdx)
constexpr size_t	MaxOperandDwords (16)
const int	NumVecElemPerVecReg (64)
template<typename T >
T	wholeQuadMode (T val)
template<typename T >
T	quadMask (T val)
template<typename T >
ScalarRegI32	countZeroBits (T val)
template<typename T >
ScalarRegI32	findFirstZero (T val)
template<typename T >
ScalarRegI32	findFirstOne (T val)
template<typename T >
ScalarRegI32	findFirstOneMsb (T val)
template<typename T >
ScalarRegI32	countZeroBitsMsb (T val)
ScalarRegI32	firstOppositeSignBit (ScalarRegI32 val)
ScalarRegI32	firstOppositeSignBit (ScalarRegI64 val)
template<typename T >
T	median (T val_0, T val_1, T val_2)
template<typename T >
T	roundNearestEven (T val)
VecElemU32	muladd (VecElemU64 &dst, VecElemU32 val_0, VecElemU32 val_1, VecElemU64 val_2)
VecElemU32	muladd (VecElemI64 &dst, VecElemI32 val_0, VecElemI32 val_1, VecElemI64 val_2)
int	dppInstImpl (SqDPPVals dppCtrl, int currLane, int rowNum, int rowOffset, bool &outOfBounds)
	dppInstImpl is a helper function that performs the inputted operation on the inputted vector register lane.
template<typename T >
void	processDPP (GPUDynInstPtr gpuDynInst, InFmt_VOP_DPP dppInst, T &src0)
	processDPP is a helper function for implementing Data Parallel Primitive instructions.
template<typename T >
void	processDPP (GPUDynInstPtr gpuDynInst, InFmt_VOP_DPP dppInst, T &src0, T &src1)
	processDPP is a helper function for implementing Data Parallel Primitive instructions.
template<typename T >
T	sdwaInstSrcImpl_helper (T currOperVal, const T origOperVal, const SDWASelVals sel, const bool signExt)
	sdwaInstSrcImpl_helper contains the per-lane code for selecting the appropriate bytes/words of the lane and doing the appropriate masking/padding/sign extending.
template<typename T >
void	sdwaInstSrcImpl (T &currOper, T &origCurrOper, const SDWASelVals sel, const bool signExt)
	sdwaInstSrcImpl is a helper function that selects the appropriate bits/bytes for each lane of the inputted source operand of an SDWA instruction, does the appropriate masking/padding/sign extending for the non-selected bits/bytes, and updates the operands values with the resultant value.
template<typename T >
T	sdwaInstDstImpl_helper (T currDstVal, const T origDstVal, const bool clamp, const SDWASelVals sel, const SDWADstVals unusedBits_format)
	sdwaInstDstImpl_helper contains the per-lane code for selecting the appropriate bytes/words of the lane and doing the appropriate masking/padding/sign extending.
template<typename T >
void	sdwaInstDstImpl (T &dstOper, T &origDstOper, const bool clamp, const SDWASelVals sel, const SDWADstVals unusedBits_format)
	sdwaInstDestImpl is a helper function that selects the appropriate bits/bytes for the inputted dest operand of an SDWA instruction, does the appropriate masking/padding/sign extending for the non-selected bits/bytes, and updates the operands values with the resultant value.
template<typename T >
void	processSDWA_src_helper (T &currSrc, T &origCurrSrc, const SDWASelVals src_sel, const bool src_signExt, const bool src_abs, const bool src_neg)
	processSDWA_srcHelper is a helper function for implementing sub d-word addressing instructions for the src operands.
template<typename T >
void	processSDWA_src (InFmt_VOP_SDWA sdwaInst, T &src0, T &origSrc0)
	processSDWA_src is a helper function for implementing sub d-word addressing instructions for the src operands.
template<typename T >
void	processSDWA_src (InFmt_VOP_SDWA sdwaInst, T &src0, T &origSrc0, T &src1, T &origSrc1)
	processSDWA_src is a helper function for implementing sub d-word addressing instructions.
template<typename T >
void	processSDWA_dst (InFmt_VOP_SDWA sdwaInst, T &dst, T &origDst)
	processSDWA_dst is a helper function for implementing sub d-word addressing instructions for the dst operand.
template<int N>
int32_t	dotClampI (int32_t value, bool clamp)
template<int N>
uint32_t	dotClampU (uint32_t value, bool clamp)
int16_t	clampI16 (int32_t value, bool clamp)
uint16_t	clampU16 (uint32_t value, bool clamp)
uint16_t	clampF16 (uint16_t value, bool clamp)
float	clampF32 (float value, bool clamp)
	BitUnion64 (PageTableEntry) Bitfield< 58
	The page table entry is reverse engineered from the macros here:
	EndBitUnion (PageTableEntry) BitUnion64(PageDirectoryEntry) Bitfield< 63

Variables
const int	NumPosConstRegs
const int	NumNegConstRegs
const int	BITS_PER_BYTE = 8
const int	BITS_PER_WORD = 16
const int	MSB_PER_BYTE = (BITS_PER_BYTE - 1)
const int	MSB_PER_WORD = (BITS_PER_WORD - 1)
const int	DWordSize = sizeof(VecElemU32)
const int	RegSizeDWords = sizeof(VecElemU32) / DWordSize
	Size of a single-precision register in DWords.
static const char *	MNEM__V_MFMA_F32_4X4X1_16B_F32
static const char *	MNEM__V_MFMA_F32_32X32X1_2B_F32
static const char *	MNEM__V_MFMA_F32_32X32X2_F32
static const char *	MNEM__V_MFMA_F32_16X16X4_F32
static const char *	MNEM__V_MFMA_F32_16X16X1_4B_F32
static const char *	MNEM__V_MFMA_F64_4X4X4_4B_F64
static const char *	MNEM__V_MFMA_F64_16X16X4_F64
static const char *	MNEM__V_MFMA_F32_16X16X16_F16
static const char *	MNEM__V_MFMA_F32_16X16X4_4B_F16
static const char *	MNEM__V_MFMA_F32_32X32X4_2B_F16
static const char *	NMEM__V_MFMA_F32_32X32X8_F16
static const char *	MNEM__V_MFMA_F32_4X4X4_16B_F16
static const char *	MNEM__V_MFMA_F32_32X32X8_BF16
static const char *	MNEM__V_MFMA_I32_16X16X16_I8
const Addr	PageShift = 12
const Addr	PageBytes = 1ULL << PageShift
	m
Bitfield< 56 >	f
Bitfield< 55 >	l
Bitfield< 53, 52 >	sw
Bitfield< 51 >	t
Bitfield< 47, 12 >	ppn
Bitfield< 11, 7 >	fragment
Bitfield< 6 >	w
Bitfield< 5 >	r
Bitfield< 4 >	x
Bitfield< 3 >	z
Bitfield< 2 >	c
Bitfield< 1 >	s
Bitfield< 0 >	v
	blockFragmentSize
Bitfield< 54 >	p
Bitfield< 47, 6 >	baseAddr

Detailed Description

classes that represnt vector/scalar operands in VEGA ISA.

these classes wrap the generic vector register type (i.e., src/arch/generic/vec_reg.hh) and allow them to be manipulated in ways that are unique to VEGA insts.

Typedef Documentation

ConstScalarOperandF32

using gem5::VegaISA::ConstScalarOperandF32 = typedef ScalarOperand<ScalarRegF32, true>

Definition at line 801 of file operand.hh.

ConstScalarOperandF64

using gem5::VegaISA::ConstScalarOperandF64 = typedef ScalarOperand<ScalarRegF64, true>

Definition at line 804 of file operand.hh.

ConstScalarOperandI16

using gem5::VegaISA::ConstScalarOperandI16 = typedef ScalarOperand<ScalarRegI16, true, 1>

Definition at line 798 of file operand.hh.

ConstScalarOperandI32

using gem5::VegaISA::ConstScalarOperandI32 = typedef ScalarOperand<ScalarRegI32, true>

Definition at line 800 of file operand.hh.

ConstScalarOperandI64

using gem5::VegaISA::ConstScalarOperandI64 = typedef ScalarOperand<ScalarRegI64, true>

Definition at line 803 of file operand.hh.

ConstScalarOperandI8

using gem5::VegaISA::ConstScalarOperandI8 = typedef ScalarOperand<ScalarRegI8, true, 1>

Definition at line 796 of file operand.hh.

ConstScalarOperandU128

using gem5::VegaISA::ConstScalarOperandU128 = typedef ScalarOperand<ScalarRegU32, true, 4>

Definition at line 805 of file operand.hh.

ConstScalarOperandU16

using gem5::VegaISA::ConstScalarOperandU16 = typedef ScalarOperand<ScalarRegU16, true, 1>

Definition at line 797 of file operand.hh.

ConstScalarOperandU256

using gem5::VegaISA::ConstScalarOperandU256 = typedef ScalarOperand<ScalarRegU32, true, 8>

Definition at line 806 of file operand.hh.

ConstScalarOperandU32

using gem5::VegaISA::ConstScalarOperandU32 = typedef ScalarOperand<ScalarRegU32, true>

Definition at line 799 of file operand.hh.

ConstScalarOperandU512

using gem5::VegaISA::ConstScalarOperandU512 = typedef ScalarOperand<ScalarRegU32, true, 16>

Definition at line 807 of file operand.hh.

ConstScalarOperandU64

using gem5::VegaISA::ConstScalarOperandU64 = typedef ScalarOperand<ScalarRegU64, true>

Definition at line 802 of file operand.hh.

ConstScalarOperandU8

using gem5::VegaISA::ConstScalarOperandU8 = typedef ScalarOperand<ScalarRegU8, true, 1>

Definition at line 795 of file operand.hh.

ConstVecOperandF32

using gem5::VegaISA::ConstVecOperandF32 = typedef VecOperand<VecElemF32, true>

Definition at line 830 of file operand.hh.

ConstVecOperandF64

using gem5::VegaISA::ConstVecOperandF64 = typedef VecOperand<VecElemF64, true>

Definition at line 833 of file operand.hh.

ConstVecOperandI16

using gem5::VegaISA::ConstVecOperandI16 = typedef VecOperand<VecElemI16, true, 1>

Definition at line 827 of file operand.hh.

ConstVecOperandI32

using gem5::VegaISA::ConstVecOperandI32 = typedef VecOperand<VecElemI32, true>

Definition at line 829 of file operand.hh.

ConstVecOperandI64

using gem5::VegaISA::ConstVecOperandI64 = typedef VecOperand<VecElemI64, true>

Definition at line 832 of file operand.hh.

ConstVecOperandI8

using gem5::VegaISA::ConstVecOperandI8 = typedef VecOperand<VecElemI8, true, 1>

Definition at line 825 of file operand.hh.

ConstVecOperandU128

using gem5::VegaISA::ConstVecOperandU128 = typedef VecOperand<VecElemU32, true, 4>

Definition at line 835 of file operand.hh.

ConstVecOperandU16

using gem5::VegaISA::ConstVecOperandU16 = typedef VecOperand<VecElemU16, true, 1>

Definition at line 826 of file operand.hh.

ConstVecOperandU256

using gem5::VegaISA::ConstVecOperandU256 = typedef VecOperand<VecElemU32, true, 8>

Definition at line 836 of file operand.hh.

ConstVecOperandU32

using gem5::VegaISA::ConstVecOperandU32 = typedef VecOperand<VecElemU32, true>

Definition at line 828 of file operand.hh.

ConstVecOperandU512

using gem5::VegaISA::ConstVecOperandU512 = typedef VecOperand<VecElemU32, true, 16>

Definition at line 837 of file operand.hh.

ConstVecOperandU64

using gem5::VegaISA::ConstVecOperandU64 = typedef VecOperand<VecElemU64, true>

Definition at line 831 of file operand.hh.

ConstVecOperandU8

using gem5::VegaISA::ConstVecOperandU8 = typedef VecOperand<VecElemU8, true, 1>

Definition at line 824 of file operand.hh.

ConstVecOperandU96

using gem5::VegaISA::ConstVecOperandU96 = typedef VecOperand<VecElemU32, true, 3>

Definition at line 834 of file operand.hh.

half

using gem5::VegaISA::half = typedef uint16_t

Definition at line 43 of file vop3p.cc.

Inst_VOP3P_MAI__V_MFMA_F32_16X16X16_F16

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F32_16X16X16_F16 = typedef Inst_VOP3P_MAI__V_MFMA_MXFP<16, 16, 16, 1, AMDGPU::mxfloat16, &MNEM__V_MFMA_F32_16X16X16_F16>

Definition at line 44606 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_F32_16X16X1_4B_F32

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F32_16X16X1_4B_F32 = typedef Inst_VOP3P_MAI__V_MFMA<1, 16, 16, 1, 4, ConstVecOperandF32, VecOperandF32, &MNEM__V_MFMA_F32_16X16X1_4B_F32>

Definition at line 44377 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_F32_16X16X4_4B_F16

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F32_16X16X4_4B_F16 = typedef Inst_VOP3P_MAI__V_MFMA_MXFP<16, 16, 4, 4, AMDGPU::mxfloat16, &MNEM__V_MFMA_F32_16X16X4_4B_F16>

Definition at line 44612 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_F32_16X16X4_F32

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F32_16X16X4_F32 = typedef Inst_VOP3P_MAI__V_MFMA<1, 16, 16, 4, 1, ConstVecOperandF32, VecOperandF32, &MNEM__V_MFMA_F32_16X16X4_F32>

Definition at line 44371 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_F32_32X32X1_2B_F32

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F32_32X32X1_2B_F32 = typedef Inst_VOP3P_MAI__V_MFMA<1, 32, 32, 1, 2, ConstVecOperandF32, VecOperandF32, &MNEM__V_MFMA_F32_32X32X1_2B_F32>

Definition at line 44358 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_F32_32X32X2_F32

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F32_32X32X2_F32 = typedef Inst_VOP3P_MAI__V_MFMA<1, 32, 32, 2, 1, ConstVecOperandF32, VecOperandF32, &MNEM__V_MFMA_F32_32X32X2_F32>

Definition at line 44365 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_F32_32X32X4_2B_F16

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F32_32X32X4_2B_F16 = typedef Inst_VOP3P_MAI__V_MFMA_MXFP<32, 32, 4, 2, AMDGPU::mxfloat16, &MNEM__V_MFMA_F32_32X32X4_2B_F16>

Definition at line 44618 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_F32_32X32X8_BF16

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F32_32X32X8_BF16 = typedef Inst_VOP3P_MAI__V_MFMA_MXFP<32, 32, 8, 1, AMDGPU::mxbfloat16, &MNEM__V_MFMA_F32_32X32X8_BF16>

Definition at line 44636 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_F32_32X32X8_F16

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F32_32X32X8_F16 = typedef Inst_VOP3P_MAI__V_MFMA_MXFP<32, 32, 8, 1, AMDGPU::mxfloat16, &NMEM__V_MFMA_F32_32X32X8_F16>

Definition at line 44624 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_F32_4X4X1_16B_F32

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F32_4X4X1_16B_F32 = typedef Inst_VOP3P_MAI__V_MFMA<1, 4, 4, 1, 16, ConstVecOperandF32, VecOperandF32, &MNEM__V_MFMA_F32_4X4X1_16B_F32>

Definition at line 44352 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_F32_4X4X4_16B_F16

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F32_4X4X4_16B_F16 = typedef Inst_VOP3P_MAI__V_MFMA_MXFP<4, 4, 4, 16, AMDGPU::mxfloat16, &MNEM__V_MFMA_F32_4X4X4_16B_F16>

Definition at line 44630 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_F64_16X16X4_F64

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F64_16X16X4_F64 = typedef Inst_VOP3P_MAI__V_MFMA<2, 16, 16, 4, 1, ConstVecOperandF64, VecOperandF64, &MNEM__V_MFMA_F64_16X16X4_F64>

Definition at line 44390 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_F64_4X4X4_4B_F64

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_F64_4X4X4_4B_F64 = typedef Inst_VOP3P_MAI__V_MFMA<2, 4, 4, 4, 4, ConstVecOperandF64, VecOperandF64, &MNEM__V_MFMA_F64_4X4X4_4B_F64>

Definition at line 44384 of file instructions.hh.

Inst_VOP3P_MAI__V_MFMA_I32_16X16X16_I8

using gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_I32_16X16X16_I8 = typedef Inst_VOP3P_MAI__V_MFMA_I8<16, 16, 16, 1, &MNEM__V_MFMA_I32_16X16X16_I8>

Definition at line 44854 of file instructions.hh.

IsaDecodeMethod

using gem5::VegaISA::IsaDecodeMethod = typedef GPUStaticInst*(Decoder::*)(MachInst)

Definition at line 50 of file gpu_decoder.hh.

MachInst

typedef InstFormat* gem5::VegaISA::MachInst

used to represent the encoding of a VEGA inst.

each portion of a VEGA inst must be 1 DWORD (32b), so we use a pointer to InstFormat type (which is 32b). for the case in which we need multiple DWORDS to represnt a single inst, this pointer essentialy acts as an array of the DWORDs needed to represent the entire inst encoding.

Definition at line 61 of file gpu_types.hh.

RawMachInst

typedef uint64_t gem5::VegaISA::RawMachInst

used to represnt a GPU inst in its raw format.

VEGA instructions may be 32b or 64b, therefore we represent a raw inst with 64b to ensure that all of its inst data, including potential immediate values, may be represented in the worst case.

Definition at line 51 of file gpu_types.hh.

ScalarOperandF32

using gem5::VegaISA::ScalarOperandF32 = typedef ScalarOperand<ScalarRegF32, false>

Definition at line 787 of file operand.hh.

ScalarOperandF64

using gem5::VegaISA::ScalarOperandF64 = typedef ScalarOperand<ScalarRegF64, false>

Definition at line 790 of file operand.hh.

ScalarOperandI16

using gem5::VegaISA::ScalarOperandI16 = typedef ScalarOperand<ScalarRegI16, false, 1>

Definition at line 784 of file operand.hh.

ScalarOperandI32

using gem5::VegaISA::ScalarOperandI32 = typedef ScalarOperand<ScalarRegI32, false>

Definition at line 786 of file operand.hh.

ScalarOperandI64

using gem5::VegaISA::ScalarOperandI64 = typedef ScalarOperand<ScalarRegI64, false>

Definition at line 789 of file operand.hh.

ScalarOperandI8

using gem5::VegaISA::ScalarOperandI8 = typedef ScalarOperand<ScalarRegI8, false, 1>

Definition at line 782 of file operand.hh.

ScalarOperandU128

using gem5::VegaISA::ScalarOperandU128 = typedef ScalarOperand<ScalarRegU32, false, 4>

Definition at line 791 of file operand.hh.

ScalarOperandU16

using gem5::VegaISA::ScalarOperandU16 = typedef ScalarOperand<ScalarRegU16, false, 1>

Definition at line 783 of file operand.hh.

ScalarOperandU256

using gem5::VegaISA::ScalarOperandU256 = typedef ScalarOperand<ScalarRegU32, false, 8>

Definition at line 792 of file operand.hh.

ScalarOperandU32

using gem5::VegaISA::ScalarOperandU32 = typedef ScalarOperand<ScalarRegU32, false>

Definition at line 785 of file operand.hh.

ScalarOperandU512

using gem5::VegaISA::ScalarOperandU512 = typedef ScalarOperand<ScalarRegU32, false, 16>

Definition at line 793 of file operand.hh.

ScalarOperandU64

using gem5::VegaISA::ScalarOperandU64 = typedef ScalarOperand<ScalarRegU64, false>

Definition at line 788 of file operand.hh.

ScalarOperandU8

using gem5::VegaISA::ScalarOperandU8 = typedef ScalarOperand<ScalarRegU8, false, 1>

Definition at line 781 of file operand.hh.

ScalarRegF32

typedef float gem5::VegaISA::ScalarRegF32

Definition at line 155 of file gpu_registers.hh.

ScalarRegF64

typedef double gem5::VegaISA::ScalarRegF64

Definition at line 158 of file gpu_registers.hh.

ScalarRegI16

typedef int16_t gem5::VegaISA::ScalarRegI16

Definition at line 152 of file gpu_registers.hh.

ScalarRegI32

typedef int32_t gem5::VegaISA::ScalarRegI32

Definition at line 154 of file gpu_registers.hh.

ScalarRegI64

typedef int64_t gem5::VegaISA::ScalarRegI64

Definition at line 157 of file gpu_registers.hh.

ScalarRegI8

typedef int8_t gem5::VegaISA::ScalarRegI8

Definition at line 150 of file gpu_registers.hh.

ScalarRegU16

typedef uint16_t gem5::VegaISA::ScalarRegU16

Definition at line 151 of file gpu_registers.hh.

ScalarRegU32

typedef uint32_t gem5::VegaISA::ScalarRegU32

Definition at line 153 of file gpu_registers.hh.

ScalarRegU64

typedef uint64_t gem5::VegaISA::ScalarRegU64

Definition at line 156 of file gpu_registers.hh.

ScalarRegU8

typedef uint8_t gem5::VegaISA::ScalarRegU8

Definition at line 149 of file gpu_registers.hh.

VecElemF32

typedef float gem5::VegaISA::VecElemF32

Definition at line 167 of file gpu_registers.hh.

VecElemF64

typedef double gem5::VegaISA::VecElemF64

Definition at line 170 of file gpu_registers.hh.

VecElemI16

typedef int16_t gem5::VegaISA::VecElemI16

Definition at line 164 of file gpu_registers.hh.

VecElemI32

typedef int32_t gem5::VegaISA::VecElemI32

Definition at line 166 of file gpu_registers.hh.

VecElemI64

typedef int64_t gem5::VegaISA::VecElemI64

Definition at line 169 of file gpu_registers.hh.

VecElemI8

typedef int8_t gem5::VegaISA::VecElemI8

Definition at line 162 of file gpu_registers.hh.

VecElemU16

typedef uint16_t gem5::VegaISA::VecElemU16

Definition at line 163 of file gpu_registers.hh.

VecElemU32

typedef uint32_t gem5::VegaISA::VecElemU32

Definition at line 165 of file gpu_registers.hh.

VecElemU64

typedef uint64_t gem5::VegaISA::VecElemU64

Definition at line 168 of file gpu_registers.hh.

VecElemU8

typedef uint8_t gem5::VegaISA::VecElemU8

Definition at line 161 of file gpu_registers.hh.

VecOperandF32

using gem5::VegaISA::VecOperandF32 = typedef VecOperand<VecElemF32, false>

Definition at line 815 of file operand.hh.

VecOperandF64

using gem5::VegaISA::VecOperandF64 = typedef VecOperand<VecElemF64, false>

Definition at line 817 of file operand.hh.

VecOperandI16

using gem5::VegaISA::VecOperandI16 = typedef VecOperand<VecElemI16, false, 1>

Definition at line 812 of file operand.hh.

VecOperandI32

using gem5::VegaISA::VecOperandI32 = typedef VecOperand<VecElemI32, false>

Definition at line 814 of file operand.hh.

VecOperandI64

using gem5::VegaISA::VecOperandI64 = typedef VecOperand<VecElemI64, false>

Definition at line 818 of file operand.hh.

VecOperandI8

using gem5::VegaISA::VecOperandI8 = typedef VecOperand<VecElemI8, false, 1>

Definition at line 810 of file operand.hh.

VecOperandU128

using gem5::VegaISA::VecOperandU128 = typedef VecOperand<VecElemU32, false, 4>

Definition at line 820 of file operand.hh.

VecOperandU16

using gem5::VegaISA::VecOperandU16 = typedef VecOperand<VecElemU16, false, 1>

Definition at line 811 of file operand.hh.

VecOperandU256

using gem5::VegaISA::VecOperandU256 = typedef VecOperand<VecElemU32, false, 8>

Definition at line 821 of file operand.hh.

VecOperandU32

using gem5::VegaISA::VecOperandU32 = typedef VecOperand<VecElemU32, false>

Definition at line 813 of file operand.hh.

VecOperandU512

using gem5::VegaISA::VecOperandU512 = typedef VecOperand<VecElemU32, false, 16>

Definition at line 822 of file operand.hh.

VecOperandU64

using gem5::VegaISA::VecOperandU64 = typedef VecOperand<VecElemU64, false>

Definition at line 816 of file operand.hh.

VecOperandU8

using gem5::VegaISA::VecOperandU8 = typedef VecOperand<VecElemU8, false, 1>

Definition at line 809 of file operand.hh.

VecOperandU96

using gem5::VegaISA::VecOperandU96 = typedef VecOperand<VecElemU32, false, 3>

Definition at line 819 of file operand.hh.

VecRegContainerU32

using gem5::VegaISA::VecRegContainerU32 = typedef VecRegContainer<sizeof(VecElemU32) * NumVecElemPerVecReg>

Definition at line 178 of file gpu_registers.hh.

VecRegContainerU64

using gem5::VegaISA::VecRegContainerU64 = typedef VecRegContainer<sizeof(VecElemU64) * NumVecElemPerVecReg>

Definition at line 180 of file gpu_registers.hh.

Enumeration Type Documentation

ExceptionCode

enum gem5::VegaISA::ExceptionCode : uint64_t

Enumerator
INST_PAGE
LOAD_PAGE
STORE_PAGE

Definition at line 45 of file faults.hh.

OpSelector

enum gem5::VegaISA::OpSelector : int

Enumerator
REG_SGPR_MIN
REG_SGPR_MAX
REG_FLAT_SCRATCH_LO
REG_FLAT_SCRATCH_HI
REG_XNACK_MASK_LO
REG_XNACK_MASK_HI
REG_VCC_LO
REG_VCC_HI
REG_TBA_LO
REG_TBA_HI
REG_TMA_LO
REG_TMA_HI
REG_TTMP_0
REG_TTMP_1
REG_TTMP_2
REG_TTMP_3
REG_TTMP_4
REG_TTMP_5
REG_TTMP_6
REG_TTMP_7
REG_TTMP_8
REG_TTMP_9
REG_TTMP_10
REG_TTMP_11
REG_M0
REG_RESERVED_1
REG_EXEC_LO
REG_EXEC_HI
REG_ZERO
REG_INT_CONST_POS_MIN
REG_INT_CONST_POS_MAX
REG_INT_CONST_NEG_MIN
REG_INT_CONST_NEG_MAX
REG_RESERVED_2
REG_RESERVED_3
REG_RESERVED_4
REG_RESERVED_5
REG_RESERVED_6
REG_RESERVED_7
REG_RESERVED_8
REG_RESERVED_9
REG_RESERVED_10
REG_RESERVED_11
REG_RESERVED_12
REG_RESERVED_13
REG_RESERVED_14
REG_RESERVED_15
REG_RESERVED_16
REG_RESERVED_17
REG_RESERVED_18
REG_RESERVED_19
REG_RESERVED_20
REG_RESERVED_21
REG_RESERVED_22
REG_RESERVED_23
REG_RESERVED_24
REG_RESERVED_25
REG_RESERVED_26
REG_RESERVED_27
REG_SHARED_BASE
REG_SHARED_LIMIT
REG_PRIVATE_BASE
REG_PRIVATE_LIMIT
REG_RESERVED_32
REG_POS_HALF
REG_NEG_HALF
REG_POS_ONE
REG_NEG_ONE
REG_POS_TWO
REG_NEG_TWO
REG_POS_FOUR
REG_NEG_FOUR
REG_PI
REG_SRC_SWDA
REG_SRC_DPP
REG_VCCZ
REG_EXECZ
REG_SCC
REG_LDS_DIRECT
REG_SRC_LITERAL
REG_VGPR_MIN
REG_VGPR_MAX

Definition at line 48 of file gpu_registers.hh.

Function Documentation

BitUnion64()

gem5::VegaISA::BitUnion64

(

PageTableEntry

)

The page table entry is reverse engineered from the macros here:

https://github.com/RadeonOpenCompute/ROCK-Kernel-Driver/blob/roc-4.3.x/ drivers/gpu/drm/amd/amdgpu/amdgpu_vm.h#L53

clampF16()

uint16_t gem5::VegaISA::clampF16	(	uint16_t	value,
		bool	clamp
	)

Definition at line 103 of file vop3p.cc.

References gem5::ArmISA::fplibMax(), gem5::ArmISA::fplibMin(), and gem5::ArmISA::imm.

Referenced by gem5::VegaISA::Inst_VOP3P__V_PK_FMA_F16::execute(), gem5::VegaISA::Inst_VOP3P__V_PK_ADD_F16::execute(), gem5::VegaISA::Inst_VOP3P__V_PK_MUL_F16::execute(), gem5::VegaISA::Inst_VOP3P__V_PK_MIN_F16::execute(), and gem5::VegaISA::Inst_VOP3P__V_PK_MAX_F16::execute().

clampF32()

float gem5::VegaISA::clampF32	(	float	value,
		bool	clamp
	)

Definition at line 120 of file vop3p.cc.

Referenced by gem5::VegaISA::Inst_VOP3P__V_DOT2_F32_F16::execute().

clampI16()

int16_t gem5::VegaISA::clampI16	(	int32_t	value,
		bool	clamp
	)

Definition at line 79 of file vop3p.cc.

Referenced by gem5::VegaISA::Inst_VOP3P__V_PK_MAD_I16::execute(), gem5::VegaISA::Inst_VOP3P__V_PK_ADD_I16::execute(), gem5::VegaISA::Inst_VOP3P__V_PK_SUB_I16::execute(), gem5::VegaISA::Inst_VOP3P__V_PK_MAX_I16::execute(), and gem5::VegaISA::Inst_VOP3P__V_PK_MIN_I16::execute().

clampU16()

uint16_t gem5::VegaISA::clampU16	(	uint32_t	value,
		bool	clamp
	)

Definition at line 91 of file vop3p.cc.

Referenced by gem5::VegaISA::Inst_VOP3P__V_PK_MAD_U16::execute(), gem5::VegaISA::Inst_VOP3P__V_PK_ADD_U16::execute(), gem5::VegaISA::Inst_VOP3P__V_PK_SUB_U16::execute(), gem5::VegaISA::Inst_VOP3P__V_PK_MAX_U16::execute(), and gem5::VegaISA::Inst_VOP3P__V_PK_MIN_U16::execute().

countZeroBits()

template<typename T >

ScalarRegI32 gem5::VegaISA::countZeroBits ( T  val )

inline

Definition at line 121 of file inst_util.hh.

References gem5::popCount(), and gem5::X86ISA::val.

Referenced by gem5::VegaISA::Inst_SOP1__S_BCNT0_I32_B32::execute(), and gem5::VegaISA::Inst_SOP1__S_BCNT0_I32_B64::execute().

countZeroBitsMsb()

template<typename T >

ScalarRegI32 gem5::VegaISA::countZeroBitsMsb ( T  val )

inline

Definition at line 164 of file inst_util.hh.

References gem5::findMsbSet(), and gem5::X86ISA::val.

Referenced by gem5::VegaISA::Inst_SOP1__S_FLBIT_I32_B32::execute(), and gem5::VegaISA::Inst_SOP1__S_FLBIT_I32_B64::execute().

dotClampI()

template<int N>

int32_t gem5::VegaISA::dotClampI	(	int32_t	value,
		bool	clamp
	)

Definition at line 48 of file vop3p.cc.

dotClampU()

template<int N>

uint32_t gem5::VegaISA::dotClampU	(	uint32_t	value,
		bool	clamp
	)

Definition at line 64 of file vop3p.cc.

dppInstImpl()

int gem5::VegaISA::dppInstImpl ( SqDPPVals  dppCtrl, int  currLane, int  rowNum, int  rowOffset, bool &  outOfBounds  )

inline

dppInstImpl is a helper function that performs the inputted operation on the inputted vector register lane.

The returned output lane represents the input lane given the destination lane and DPP_CTRL word.

Currently the values are: 0x0 - 0xFF: full permute of four threads 0x100: reserved 0x101 - 0x10F: row shift left by 1-15 threads 0x111 - 0x11F: row shift right by 1-15 threads 0x121 - 0x12F: row rotate right by 1-15 threads 0x130: wavefront left shift by 1 thread 0x134: wavefront left rotate by 1 thread 0x138: wavefront right shift by 1 thread 0x13C: wavefront right rotate by 1 thread 0x140: mirror threads within row 0x141: mirror threads within 1/2 row (8 threads) 0x142: broadcast 15th thread of each row to next row 0x143: broadcast thread 31 to rows 2 and 3

Definition at line 320 of file inst_util.hh.

References gem5::X86ISA::count, gem5::ArmISA::NumVecElemPerVecReg, panic, gem5::ROW_SIZE, gem5::SQ_DPP_QUAD_PERM_MAX, gem5::SQ_DPP_RESERVED, gem5::SQ_DPP_ROW_BCAST15, gem5::SQ_DPP_ROW_BCAST31, gem5::SQ_DPP_ROW_HALF_MIRROR, gem5::SQ_DPP_ROW_MIRROR, gem5::SQ_DPP_ROW_RR1, gem5::SQ_DPP_ROW_RR15, gem5::SQ_DPP_ROW_SL1, gem5::SQ_DPP_ROW_SL15, gem5::SQ_DPP_ROW_SR1, gem5::SQ_DPP_ROW_SR15, gem5::SQ_DPP_WF_RL1, gem5::SQ_DPP_WF_RR1, gem5::SQ_DPP_WF_SL1, and gem5::SQ_DPP_WF_SR1.

Referenced by processDPP().

EndBitUnion()

gem5::VegaISA::EndBitUnion

(

PageTableEntry

)

findFirstOne()

template<typename T >

ScalarRegI32 gem5::VegaISA::findFirstOne ( T  val )

inline

Definition at line 142 of file inst_util.hh.

References gem5::findLsbSet(), and gem5::X86ISA::val.

Referenced by gem5::VegaISA::Inst_SOP1__S_FF1_I32_B32::execute(), gem5::VegaISA::Inst_SOP1__S_FF1_I32_B64::execute(), gem5::VegaISA::Inst_VOP1__V_FFBL_B32::execute(), and gem5::VegaISA::Inst_VOP3__V_FFBL_B32::execute().

findFirstOneMsb()

template<typename T >

ScalarRegI32 gem5::VegaISA::findFirstOneMsb ( T  val )

inline

Definition at line 153 of file inst_util.hh.

References gem5::findMsbSet(), and gem5::X86ISA::val.

Referenced by gem5::VegaISA::Inst_VOP1__V_FFBH_U32::execute(), and gem5::VegaISA::Inst_VOP3__V_FFBH_U32::execute().

findFirstZero()

template<typename T >

ScalarRegI32 gem5::VegaISA::findFirstZero ( T  val )

inline

Definition at line 131 of file inst_util.hh.

References gem5::findLsbSet(), and gem5::X86ISA::val.

Referenced by gem5::VegaISA::Inst_SOP1__S_FF0_I32_B32::execute(), and gem5::VegaISA::Inst_SOP1__S_FF0_I32_B64::execute().

firstOppositeSignBit() [1/2]

ScalarRegI32 gem5::VegaISA::firstOppositeSignBit ( ScalarRegI32  val )

inline

Definition at line 174 of file inst_util.hh.

References gem5::X86ISA::count, gem5::ArmISA::i, and gem5::X86ISA::val.

Referenced by gem5::VegaISA::Inst_SOP1__S_FLBIT_I32::execute(), gem5::VegaISA::Inst_SOP1__S_FLBIT_I32_I64::execute(), gem5::VegaISA::Inst_VOP1__V_FFBH_I32::execute(), and gem5::VegaISA::Inst_VOP3__V_FFBH_I32::execute().

firstOppositeSignBit() [2/2]

ScalarRegI32 gem5::VegaISA::firstOppositeSignBit ( ScalarRegI64  val )

inline

Definition at line 210 of file inst_util.hh.

References gem5::X86ISA::count, gem5::ArmISA::i, and gem5::X86ISA::val.

isConstVal()

bool gem5::VegaISA::isConstVal

(

int

opIdx

)

Definition at line 209 of file gpu_registers.cc.

References isNegConstVal(), and isPosConstVal().

Referenced by gem5::VegaISA::ScalarOperand< DataType, Const, NumDwords >::readSpecialVal().

isExecMask()

bool gem5::VegaISA::isExecMask

(

int

opIdx

)

Definition at line 222 of file gpu_registers.cc.

References REG_EXEC_HI, and REG_EXEC_LO.

Referenced by gem5::VegaISA::VEGAGPUStaticInst::isExecMaskRegister().

isFlatScratchReg()

bool gem5::VegaISA::isFlatScratchReg

(

int

opIdx

)

Definition at line 234 of file gpu_registers.cc.

References REG_FLAT_SCRATCH_HI, and REG_FLAT_SCRATCH_LO.

Referenced by gem5::VegaISA::VEGAGPUStaticInst::isFlatScratchRegister().

isLiteral()

bool gem5::VegaISA::isLiteral

(

int

opIdx

)

Definition at line 216 of file gpu_registers.cc.

References REG_SRC_LITERAL.

isNegConstVal()

bool gem5::VegaISA::isNegConstVal

(

int

opIdx

)

Definition at line 200 of file gpu_registers.cc.

References REG_INT_CONST_NEG_MAX, and REG_INT_CONST_NEG_MIN.

Referenced by isConstVal(), and gem5::VegaISA::GPUISA::readConstVal().

isPosConstVal()

bool gem5::VegaISA::isPosConstVal

(

int

opIdx

)

Definition at line 191 of file gpu_registers.cc.

References REG_INT_CONST_POS_MAX, and REG_INT_CONST_POS_MIN.

Referenced by isConstVal(), and gem5::VegaISA::GPUISA::readConstVal().

isScalarReg()

bool gem5::VegaISA::isScalarReg

(

int

opIdx

)

Definition at line 240 of file gpu_registers.cc.

References REG_FLAT_SCRATCH_HI, REG_FLAT_SCRATCH_LO, REG_SGPR_MAX, REG_VCC_HI, and REG_VCC_LO.

Referenced by gem5::VegaISA::Inst_SOP2::initOperandInfo(), gem5::VegaISA::Inst_SOPK::initOperandInfo(), gem5::VegaISA::Inst_SOP1::initOperandInfo(), gem5::VegaISA::Inst_SOPC::initOperandInfo(), gem5::VegaISA::Inst_SMEM::initOperandInfo(), gem5::VegaISA::Inst_VOP2::initOperandInfo(), gem5::VegaISA::Inst_VOP1::initOperandInfo(), gem5::VegaISA::Inst_VOPC::initOperandInfo(), gem5::VegaISA::Inst_VOP3A::initOperandInfo(), gem5::VegaISA::Inst_VOP3B::initOperandInfo(), gem5::VegaISA::Inst_VOP3P::initOperandInfo(), gem5::VegaISA::Inst_VOP3P_MAI::initOperandInfo(), gem5::VegaISA::Inst_MUBUF::initOperandInfo(), gem5::VegaISA::Inst_MTBUF::initOperandInfo(), gem5::VegaISA::Inst_MIMG::initOperandInfo(), gem5::VegaISA::ScalarOperand< DataType, Const, NumDwords >::read(), and gem5::VegaISA::ScalarOperand< DataType, Const, NumDwords >::write().

isVccReg()

bool gem5::VegaISA::isVccReg

(

int

opIdx

)

Definition at line 228 of file gpu_registers.cc.

References REG_VCC_HI, and REG_VCC_LO.

isVectorReg()

bool gem5::VegaISA::isVectorReg

(

int

opIdx

)

Definition at line 253 of file gpu_registers.cc.

References REG_VGPR_MAX, and REG_VGPR_MIN.

Referenced by gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA< _delta, M, N, K, B, T1, T2, MNEMONIC >::execute(), gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_MXFP< M, N, K, B, MXFPT, MNEMONIC >::execute(), gem5::VegaISA::Inst_VOP3P_MAI__V_MFMA_I8< M, N, K, B, MNEMONIC >::execute(), gem5::VegaISA::Inst_VOP2::initOperandInfo(), gem5::VegaISA::Inst_VOP1::initOperandInfo(), gem5::VegaISA::Inst_VOPC::initOperandInfo(), gem5::VegaISA::Inst_VOP3A::initOperandInfo(), gem5::VegaISA::Inst_VOP3B::initOperandInfo(), gem5::VegaISA::Inst_VOP3P::initOperandInfo(), gem5::VegaISA::Inst_VOP3P_MAI::initOperandInfo(), and gem5::VegaISA::VecOperand< DataType, Const, NumDwords >::readSrc().

MaxOperandDwords()

constexpr size_t gem5::VegaISA::MaxOperandDwords ( 16  )

constexpr

median()

template<typename T >

T gem5::VegaISA::median ( T  val_0, T  val_1, T  val_2  )

inline

Definition at line 247 of file inst_util.hh.

Referenced by gem5::VegaISA::Inst_VOP3__V_MED3_F32::execute(), gem5::VegaISA::Inst_VOP3__V_MED3_I32::execute(), and gem5::VegaISA::Inst_VOP3__V_MED3_U32::execute().

muladd() [1/2]

VecElemU32 gem5::VegaISA::muladd ( VecElemI64 &  dst, VecElemI32  val_0, VecElemI32  val_1, VecElemI64  val_2  )

inline

Definition at line 286 of file inst_util.hh.

muladd() [2/2]

VecElemU32 gem5::VegaISA::muladd ( VecElemU64 &  dst, VecElemU32  val_0, VecElemU32  val_1, VecElemU64  val_2  )

inline

Definition at line 272 of file inst_util.hh.

Referenced by gem5::VegaISA::Inst_VOP3__V_MAD_U64_U32::execute(), and gem5::VegaISA::Inst_VOP3__V_MAD_I64_I32::execute().

NumVecElemPerVecReg()

const int gem5::VegaISA::NumVecElemPerVecReg

(

64

)

Referenced by gem5::initMemReqHelper(), gem5::initScratchReqHelper(), and gem5::GPUCommandProcessor::MQDDmaEvent().

opSelectorToRegIdx()

int gem5::VegaISA::opSelectorToRegIdx	(	int	idx,
		int	numScalarRegs
	)

the VCC register occupies the two highest numbered SRF entries. VCC is typically indexed by specifying VCC_LO (simply called VCC) in the instruction encoding and reading it as a 64b value so we only return the index to the lower half of the VCC register.

VCC_LO = s[NUM_SGPRS - 2] VCC_HI = s[NUM_SGPRS - 1]

the FLAT_SCRATCH register occupies the two SRF entries just below VCC. FLAT_SCRATCH is typically indexed by specifying FLAT_SCRATCH_LO (simply called FLAT_SCRATCH) in the instruction encoding and reading it as a 64b value so we only return the index to the lower half of the FLAT_SCRATCH register.

FLAT_SCRATCH_LO = s[NUM_SGPRS - 4] FLAT_SCRATCH_HI = s[NUM_SGPRS - 3]

If the operand is the EXEC mask we just return the op selector value indicating it is the EXEC mask, which is not part of any RF. Higher-level calls will understand that this resolves to a special system register, not an index into an RF.

Definition at line 137 of file gpu_registers.cc.

References REG_EXEC_HI, REG_EXEC_LO, REG_FLAT_SCRATCH_HI, REG_FLAT_SCRATCH_LO, REG_SGPR_MAX, REG_VCC_HI, REG_VCC_LO, REG_VGPR_MAX, and REG_VGPR_MIN.

Referenced by gem5::VegaISA::VecOperand< DataType, Const, NumDwords >::readSrc().

opSelectorToRegSym()

std::string gem5::VegaISA::opSelectorToRegSym	(	int	idx,
		int	numRegs
	)

Definition at line 40 of file gpu_registers.cc.

References fatal, REG_EXEC_LO, REG_FLAT_SCRATCH_HI, REG_FLAT_SCRATCH_LO, REG_INT_CONST_NEG_MAX, REG_INT_CONST_NEG_MIN, REG_INT_CONST_POS_MAX, REG_INT_CONST_POS_MIN, REG_M0, REG_NEG_FOUR, REG_NEG_HALF, REG_NEG_ONE, REG_NEG_TWO, REG_POS_FOUR, REG_POS_HALF, REG_POS_ONE, REG_POS_TWO, REG_PRIVATE_BASE, REG_PRIVATE_LIMIT, REG_SGPR_MAX, REG_SHARED_BASE, REG_SHARED_LIMIT, REG_VCC_HI, REG_VCC_LO, REG_VGPR_MAX, REG_VGPR_MIN, and REG_ZERO.

Referenced by gem5::VegaISA::Inst_SOP2::generateDisassembly(), gem5::VegaISA::Inst_SOPK::generateDisassembly(), gem5::VegaISA::Inst_SOP1::generateDisassembly(), gem5::VegaISA::Inst_SOPC::generateDisassembly(), gem5::VegaISA::Inst_VOP2::generateDisassembly(), gem5::VegaISA::Inst_VOP1::generateDisassembly(), gem5::VegaISA::Inst_VOPC::generateDisassembly(), gem5::VegaISA::Inst_VOP3A::generateDisassembly(), gem5::VegaISA::Inst_VOP3B::generateDisassembly(), gem5::VegaISA::Inst_VOP3P::generateDisassembly(), gem5::VegaISA::Inst_VOP3P_MAI::generateDisassembly(), gem5::VegaISA::Inst_FLAT::generateFlatDisassembly(), and gem5::VegaISA::Inst_FLAT::generateGlobalScratchDisassembly().

processDPP() [1/2]

template<typename T >

void gem5::VegaISA::processDPP	(	GPUDynInstPtr	gpuDynInst,
		InFmt_VOP_DPP	dppInst,
		T &	src0
	)

processDPP is a helper function for implementing Data Parallel Primitive instructions.

This function may be called by many different VOP1 instructions to do operations within a register.

STEP 1a: check if the absolute value (ABS) or negation (NEG) tags are set. If so, do the appropriate action(s) on src0 and/or src1.

NOTE: ABS takes priority over NEG.

STEP 2: check the row and bank mask values. These determine which threads are enabled for the subsequent DPP_CTRL operations.

STEP 4: Handle the potential values of DPP_CTRL: 0x0 - 0xFF: full permute of four threads 0x100: reserved 0x101 - 0x10F: row shift right by 1-15 threads 0x111 - 0x11F: row shift right by 1-15 threads 0x121 - 0x12F: row shift right by 1-15 threads 0x130: wavefront left shift by 1 thread 0x134: wavefront left rotate by 1 thread 0x138: wavefront right shift by 1 thread 0x13C: wavefront right rotate by 1 thread 0x140: mirror threads within row 0x141: mirror threads within 1/2 row (8 threads) 0x142: broadcast 15th thread of each row to next row 0x143: broadcast thread 31 to rows 2 and 3

STEP 4: Implement bound control for disabled threads. If thread is disabled but boundCtrl is set, then we need to set the source data to 0 (i.e., set this lane to 0).

Definition at line 424 of file inst_util.hh.

References gem5::VegaISA::InFmt_VOP_DPP::BANK_MASK, gem5::VegaISA::InFmt_VOP_DPP::BC, gem5::VegaISA::InFmt_VOP_DPP::DPP_CTRL, dppInstImpl(), gem5::NUM_BANKS, gem5::ArmISA::NumVecElemPerVecReg, gem5::VegaISA::InFmt_VOP_DPP::ROW_MASK, gem5::ROW_SIZE, gem5::VegaISA::InFmt_VOP_DPP::SRC0_ABS, and gem5::VegaISA::InFmt_VOP_DPP::SRC0_NEG.

Referenced by gem5::VegaISA::Inst_VOP2::dppHelper(), gem5::VegaISA::Inst_VOP2__V_ADD_F32::execute(), gem5::VegaISA::Inst_VOP2__V_AND_B32::execute(), gem5::VegaISA::Inst_VOP2__V_MAC_F32::execute(), gem5::VegaISA::Inst_VOP1__V_MOV_B32::execute(), and processDPP().

processDPP() [2/2]

template<typename T >

void gem5::VegaISA::processDPP	(	GPUDynInstPtr	gpuDynInst,
		InFmt_VOP_DPP	dppInst,
		T &	src0,
		T &	src1
	)

processDPP is a helper function for implementing Data Parallel Primitive instructions.

This function may be called by many different VOP2/VOPC instructions to do operations within a register.

STEP 1b: check if the absolute value (ABS) or negation (NEG) tags are set. If so, do the appropriate action(s) on src0 and/or src1.

NOTE: ABS takes priority over NEG.

Definition at line 537 of file inst_util.hh.

References processDPP(), gem5::VegaISA::InFmt_VOP_DPP::SRC1_ABS, and gem5::VegaISA::InFmt_VOP_DPP::SRC1_NEG.

processSDWA_dst()

template<typename T >

void gem5::VegaISA::processSDWA_dst	(	InFmt_VOP_SDWA	sdwaInst,
		T &	dst,
		T &	origDst
	)

processSDWA_dst is a helper function for implementing sub d-word addressing instructions for the dst operand.

This function may be called by many different VOP1/VOP2/VOPC instructions to do operations within a register. processSDWA_dst is called after the math, while processSDWA_src is called before the math.

STEP 1: select the appropriate bits for dst and pad/sign-extend as appropriate.

Definition at line 892 of file inst_util.hh.

References gem5::VegaISA::InFmt_VOP_SDWA::CLMP, gem5::VegaISA::InFmt_VOP_SDWA::DST_SEL, gem5::VegaISA::InFmt_VOP_SDWA::DST_U, and sdwaInstDstImpl().

Referenced by gem5::VegaISA::Inst_VOP2__V_LSHLREV_B32::execute(), gem5::VegaISA::Inst_VOP2__V_OR_B32::execute(), gem5::VegaISA::Inst_VOP2__V_ADD_CO_U32::execute(), gem5::VegaISA::Inst_VOP2__V_ADD_U32::execute(), and gem5::VegaISA::Inst_VOP2::sdwaDstHelper().

processSDWA_src() [1/2]

template<typename T >

void gem5::VegaISA::processSDWA_src	(	InFmt_VOP_SDWA	sdwaInst,
		T &	src0,
		T &	origSrc0
	)

processSDWA_src is a helper function for implementing sub d-word addressing instructions for the src operands.

This function may be called by many different VOP1 instructions to do operations within a register. processSDWA_dst is called after the math, while processSDWA_src is called before the math.

Definition at line 836 of file inst_util.hh.

References processSDWA_src_helper(), gem5::VegaISA::InFmt_VOP_SDWA::SRC0_ABS, gem5::VegaISA::InFmt_VOP_SDWA::SRC0_NEG, gem5::VegaISA::InFmt_VOP_SDWA::SRC0_SEL, gem5::VegaISA::InFmt_VOP_SDWA::SRC0_SEXT, gem5::VegaISA::InFmt_VOP_SDWA::SRC1_ABS, gem5::VegaISA::InFmt_VOP_SDWA::SRC1_NEG, and gem5::VegaISA::InFmt_VOP_SDWA::SRC1_SEXT.

Referenced by gem5::VegaISA::Inst_VOP2__V_LSHLREV_B32::execute(), gem5::VegaISA::Inst_VOP2__V_OR_B32::execute(), gem5::VegaISA::Inst_VOP2__V_ADD_CO_U32::execute(), gem5::VegaISA::Inst_VOP2__V_ADD_U32::execute(), and gem5::VegaISA::Inst_VOP2::sdwaSrcHelper().

processSDWA_src() [2/2]

template<typename T >

void gem5::VegaISA::processSDWA_src	(	InFmt_VOP_SDWA	sdwaInst,
		T &	src0,
		T &	origSrc0,
		T &	src1,
		T &	origSrc1
	)

processSDWA_src is a helper function for implementing sub d-word addressing instructions.

This function may be called by many different VOP2/VOPC instructions to do operations within a register. processSDWA_dst is called after the math, while processSDWA_src is called before the math.

Definition at line 864 of file inst_util.hh.

References processSDWA_src_helper(), gem5::VegaISA::InFmt_VOP_SDWA::SRC0_ABS, gem5::VegaISA::InFmt_VOP_SDWA::SRC0_NEG, gem5::VegaISA::InFmt_VOP_SDWA::SRC0_SEL, gem5::VegaISA::InFmt_VOP_SDWA::SRC0_SEXT, gem5::VegaISA::InFmt_VOP_SDWA::SRC1_ABS, gem5::VegaISA::InFmt_VOP_SDWA::SRC1_NEG, gem5::VegaISA::InFmt_VOP_SDWA::SRC1_SEL, and gem5::VegaISA::InFmt_VOP_SDWA::SRC1_SEXT.

processSDWA_src_helper()

template<typename T >

void gem5::VegaISA::processSDWA_src_helper	(	T &	currSrc,
		T &	origCurrSrc,
		const SDWASelVals	src_sel,
		const bool	src_signExt,
		const bool	src_abs,
		const bool	src_neg
	)

processSDWA_srcHelper is a helper function for implementing sub d-word addressing instructions for the src operands.

This function may be called by many different VOP1/VOP2/VOPC instructions to do operations within a register. This function is also agnostic of which operand it is operating on, so that it can be called for any src operand.

STEP 1: check if the absolute value (ABS) or negation (NEG) tags are set. If so, do the appropriate action(s) on the src operand.

NOTE: According to the CSim implementation, ABS takes priority over NEG.

STEP 2: select the appropriate bits for each lane of source operand.

Definition at line 801 of file inst_util.hh.

References sdwaInstSrcImpl().

Referenced by processSDWA_src(), and processSDWA_src().

quadMask()

template<typename T >

T gem5::VegaISA::quadMask ( T  val )

inline

Definition at line 104 of file inst_util.hh.

References gem5::bits(), gem5::ArmISA::mask, and gem5::X86ISA::val.

Referenced by gem5::VegaISA::Inst_SOP1__S_QUADMASK_B32::execute(), and gem5::VegaISA::Inst_SOP1__S_QUADMASK_B64::execute().

roundNearestEven()

template<typename T >

T gem5::VegaISA::roundNearestEven ( T  val )

inline

Definition at line 259 of file inst_util.hh.

References gem5::X86ISA::val.

Referenced by gem5::VegaISA::Inst_VOP1__V_RNDNE_F64::execute(), gem5::VegaISA::Inst_VOP1__V_RNDNE_F32::execute(), gem5::VegaISA::Inst_VOP3__V_RNDNE_F64::execute(), and gem5::VegaISA::Inst_VOP3__V_RNDNE_F32::execute().

sdwaInstDstImpl()

template<typename T >

void gem5::VegaISA::sdwaInstDstImpl	(	T &	dstOper,
		T &	origDstOper,
		const bool	clamp,
		const SDWASelVals	sel,
		const SDWADstVals	unusedBits_format
	)

sdwaInstDestImpl is a helper function that selects the appropriate bits/bytes for the inputted dest operand of an SDWA instruction, does the appropriate masking/padding/sign extending for the non-selected bits/bytes, and updates the operands values with the resultant value.

The desired behavior is:

1. Select the appropriate bits/bytes based on sel: 0 (SDWA_BYTE_0): select data[7:0] 1 (SDWA_BYTE_1): select data[15:8] 2 (SDWA_BYTE_2): select data[23:16] 3 (SDWA_BYTE_3): select data[31:24] 4 (SDWA_WORD_0): select data[15:0] 5 (SDWA_WORD_1): select data[31:16] 6 (SDWA_DWORD): select data[31:0]
2. either pad, sign extend, or select all bits based on the value of unusedBits_format: 0 (SDWA_UNUSED_PAD): pad all unused bits with 0 1 (SDWA_UNUSED_SEXT): sign-extend upper bits; pad lower bits w/ 0 2 (SDWA_UNUSED_PRESERVE): select data[31:0]

Definition at line 780 of file inst_util.hh.

References gem5::ArmISA::NumVecElemPerVecReg, sdwaInstDstImpl_helper(), and gem5::ArmISA::sel.

Referenced by processSDWA_dst().

sdwaInstDstImpl_helper()

template<typename T >

T gem5::VegaISA::sdwaInstDstImpl_helper	(	T	currDstVal,
		const T	origDstVal,
		const bool	clamp,
		const SDWASelVals	sel,
		const SDWADstVals	unusedBits_format
	)

sdwaInstDstImpl_helper contains the per-lane code for selecting the appropriate bytes/words of the lane and doing the appropriate masking/padding/sign extending.

It returns the value after these operations are done on it.

Definition at line 679 of file inst_util.hh.

References gem5::bits(), BITS_PER_BYTE, BITS_PER_WORD, gem5::insertBits(), MSB_PER_BYTE, MSB_PER_WORD, panic, gem5::SDWA_DWORD, gem5::SDWA_UNUSED_PRESERVE, gem5::SDWA_UNUSED_SEXT, gem5::SDWA_WORD_0, and gem5::ArmISA::sel.

Referenced by sdwaInstDstImpl().

sdwaInstSrcImpl()

template<typename T >

void gem5::VegaISA::sdwaInstSrcImpl	(	T &	currOper,
		T &	origCurrOper,
		const SDWASelVals	sel,
		const bool	signExt
	)

sdwaInstSrcImpl is a helper function that selects the appropriate bits/bytes for each lane of the inputted source operand of an SDWA instruction, does the appropriate masking/padding/sign extending for the non-selected bits/bytes, and updates the operands values with the resultant value.

The desired behavior is:

1. Select the appropriate bits/bytes based on sel: 0 (SDWA_BYTE_0): select data[7:0] 1 (SDWA_BYTE_1): select data[15:8] 2 (SDWA_BYTE_2): select data[23:16] 3 (SDWA_BYTE_3): select data[31:24] 4 (SDWA_WORD_0): select data[15:0] 5 (SDWA_WORD_1): select data[31:16] 6 (SDWA_DWORD): select data[31:0]
2. if sign extend is set, then sign extend the value

Definition at line 660 of file inst_util.hh.

References gem5::ArmISA::NumVecElemPerVecReg, sdwaInstSrcImpl_helper(), and gem5::ArmISA::sel.

Referenced by processSDWA_src_helper().

sdwaInstSrcImpl_helper()

template<typename T >

T gem5::VegaISA::sdwaInstSrcImpl_helper	(	T	currOperVal,
		const T	origOperVal,
		const SDWASelVals	sel,
		const bool	signExt
	)

sdwaInstSrcImpl_helper contains the per-lane code for selecting the appropriate bytes/words of the lane and doing the appropriate masking/padding/sign extending.

It returns the value after these operations are done on it.

Definition at line 567 of file inst_util.hh.

References gem5::bits(), BITS_PER_BYTE, BITS_PER_WORD, MSB_PER_BYTE, MSB_PER_WORD, panic, panic_if, gem5::SDWA_DWORD, gem5::SDWA_WORD_0, and gem5::ArmISA::sel.

Referenced by sdwaInstSrcImpl().

wholeQuadMode()

template<typename T >

T gem5::VegaISA::wholeQuadMode ( T  val )

inline

Definition at line 90 of file inst_util.hh.

References gem5::bits(), gem5::ArmISA::mask, and gem5::X86ISA::val.

Referenced by gem5::VegaISA::Inst_SOP1__S_WQM_B32::execute(), and gem5::VegaISA::Inst_SOP1__S_WQM_B64::execute().

Variable Documentation

baseAddr

Bitfield<47, 6> gem5::VegaISA::baseAddr

Definition at line 71 of file pagetable.hh.

BITS_PER_BYTE

const int gem5::VegaISA::BITS_PER_BYTE = 8

Definition at line 143 of file gpu_registers.hh.

Referenced by gem5::VegaISA::Inst_VOPC::sdwabSelect(), sdwaInstDstImpl_helper(), and sdwaInstSrcImpl_helper().

BITS_PER_WORD

const int gem5::VegaISA::BITS_PER_WORD = 16

Definition at line 144 of file gpu_registers.hh.

Referenced by gem5::VegaISA::Inst_VOPC::sdwabSelect(), sdwaInstDstImpl_helper(), and sdwaInstSrcImpl_helper().

blockFragmentSize

gem5::VegaISA::blockFragmentSize

Definition at line 69 of file pagetable.hh.

Referenced by gem5::VegaISA::Walker::WalkerState::walkStateMachine().

c

Bitfield< 2 > gem5::VegaISA::c

Definition at line 63 of file pagetable.hh.

DWordSize

const int gem5::VegaISA::DWordSize = sizeof(VecElemU32)

Definition at line 172 of file gpu_registers.hh.

f

Bitfield<56> gem5::VegaISA::f

Definition at line 53 of file pagetable.hh.

fragment

Bitfield<11, 7> gem5::VegaISA::fragment

Definition at line 58 of file pagetable.hh.

Referenced by gem5::VegaISA::Walker::WalkerState::walkStateMachine().

l

Bitfield<55> gem5::VegaISA::l

Definition at line 54 of file pagetable.hh.

m

gem5::VegaISA::m

Definition at line 52 of file pagetable.hh.

MNEM__V_MFMA_F32_16X16X16_F16

const char* gem5::VegaISA::MNEM__V_MFMA_F32_16X16X16_F16

static

Initial value:

=
        "v_mfma_f32_16x16x16_f16"

Definition at line 44604 of file instructions.hh.

MNEM__V_MFMA_F32_16X16X1_4B_F32

const char* gem5::VegaISA::MNEM__V_MFMA_F32_16X16X1_4B_F32

static

Initial value:

=
        "v_mfma_f32_16x16x1_4b_f32"

Definition at line 44375 of file instructions.hh.

MNEM__V_MFMA_F32_16X16X4_4B_F16

const char* gem5::VegaISA::MNEM__V_MFMA_F32_16X16X4_4B_F16

static

Initial value:

=
        "v_mfma_f32_16x16x4_4b_f16"

Definition at line 44610 of file instructions.hh.

MNEM__V_MFMA_F32_16X16X4_F32

const char* gem5::VegaISA::MNEM__V_MFMA_F32_16X16X4_F32

static

Initial value:

=
        "v_mfma_f32_16x16x4_f32"

Definition at line 44369 of file instructions.hh.

MNEM__V_MFMA_F32_32X32X1_2B_F32

const char* gem5::VegaISA::MNEM__V_MFMA_F32_32X32X1_2B_F32

static

Initial value:

=
        "v_mfma_f32_32x32x1_2b_f32"

Definition at line 44356 of file instructions.hh.

MNEM__V_MFMA_F32_32X32X2_F32

const char* gem5::VegaISA::MNEM__V_MFMA_F32_32X32X2_F32

static

Initial value:

=
        "v_mfma_f32_32x32x2_f32"

Definition at line 44363 of file instructions.hh.

MNEM__V_MFMA_F32_32X32X4_2B_F16

const char* gem5::VegaISA::MNEM__V_MFMA_F32_32X32X4_2B_F16

static

Initial value:

=
        "v_mfma_f32_32x32x4_2b_f16"

Definition at line 44616 of file instructions.hh.

MNEM__V_MFMA_F32_32X32X8_BF16

const char* gem5::VegaISA::MNEM__V_MFMA_F32_32X32X8_BF16

static

Initial value:

=
        "v_mfma_f32_32x32x8_bf16"

Definition at line 44634 of file instructions.hh.

MNEM__V_MFMA_F32_4X4X1_16B_F32

const char* gem5::VegaISA::MNEM__V_MFMA_F32_4X4X1_16B_F32

static

Initial value:

=
        "v_mfma_f32_4x4x1_16b_f32"

Definition at line 44350 of file instructions.hh.

MNEM__V_MFMA_F32_4X4X4_16B_F16

const char* gem5::VegaISA::MNEM__V_MFMA_F32_4X4X4_16B_F16

static

Initial value:

=
        "v_mfma_f32_4x4x4_16b_f16"

Definition at line 44628 of file instructions.hh.

MNEM__V_MFMA_F64_16X16X4_F64

const char* gem5::VegaISA::MNEM__V_MFMA_F64_16X16X4_F64

static

Initial value:

=
        "v_mfma_f64_16x16x4_f64"

Definition at line 44388 of file instructions.hh.

MNEM__V_MFMA_F64_4X4X4_4B_F64

const char* gem5::VegaISA::MNEM__V_MFMA_F64_4X4X4_4B_F64

static

Initial value:

=
        "v_mfma_f64_4x4x4_4b_f64"

Definition at line 44382 of file instructions.hh.

MNEM__V_MFMA_I32_16X16X16_I8

const char* gem5::VegaISA::MNEM__V_MFMA_I32_16X16X16_I8

static

Initial value:

=
        "v_mfma_i32_16x16x16_i8"

Definition at line 44852 of file instructions.hh.

MSB_PER_BYTE

const int gem5::VegaISA::MSB_PER_BYTE = (BITS_PER_BYTE - 1)

Definition at line 145 of file gpu_registers.hh.

Referenced by sdwaInstDstImpl_helper(), and sdwaInstSrcImpl_helper().

MSB_PER_WORD

const int gem5::VegaISA::MSB_PER_WORD = (BITS_PER_WORD - 1)

Definition at line 146 of file gpu_registers.hh.

Referenced by gem5::VegaISA::Inst_VOPC::sdwabSelect(), sdwaInstDstImpl_helper(), and sdwaInstSrcImpl_helper().

NMEM__V_MFMA_F32_32X32X8_F16

const char* gem5::VegaISA::NMEM__V_MFMA_F32_32X32X8_F16

static

Initial value:

=
        "v_mfma_f32_32x32x8_f16"

Definition at line 44622 of file instructions.hh.

NumNegConstRegs

const int gem5::VegaISA::NumNegConstRegs

Initial value:

= REG_INT_CONST_NEG_MAX
        - REG_INT_CONST_NEG_MIN + 1

Definition at line 141 of file gpu_registers.hh.

NumPosConstRegs

const int gem5::VegaISA::NumPosConstRegs

Initial value:

= REG_INT_CONST_POS_MAX
        - REG_INT_CONST_POS_MIN + 1

Definition at line 138 of file gpu_registers.hh.

p

Bitfield<54> gem5::VegaISA::p

Definition at line 70 of file pagetable.hh.

PageBytes

const Addr gem5::VegaISA::PageBytes = 1ULL << PageShift

Definition at line 42 of file page_size.hh.

Referenced by gem5::VegaTLBCoalescer::canCoalesce(), gem5::VegaISA::GpuTLB::handleFuncTranslationReturn(), gem5::VegaISA::GpuTLB::issueTLBLookup(), gem5::VegaTLBCoalescer::processProbeTLBEvent(), gem5::VegaISA::GpuTLB::CpuSidePort::recvFunctional(), gem5::VegaTLBCoalescer::CpuSidePort::recvFunctional(), gem5::VegaISA::GpuTLB::CpuSidePort::recvTimingReq(), gem5::VegaISA::GpuTLB::MemSidePort::recvTimingResp(), gem5::VegaTLBCoalescer::updatePhysAddresses(), and gem5::VegaISA::GpuTLB::walkerResponse().

PageShift

const Addr gem5::VegaISA::PageShift = 12

Definition at line 41 of file page_size.hh.

Referenced by gem5::VegaISA::GpuTLB::demapPage(), gem5::VegaISA::Walker::WalkerState::initState(), gem5::VegaISA::GpuTLB::insert(), gem5::VegaISA::GpuTLB::lookup(), gem5::VegaISA::GpuTLB::lookupIt(), gem5::VegaISA::GpuTLB::pageAlign(), gem5::VegaISA::GpuTLB::CpuSidePort::recvFunctional(), gem5::VegaISA::Walker::startFunctional(), gem5::VegaISA::Walker::WalkerState::startFunctional(), gem5::VegaISA::Walker::startTiming(), gem5::VegaISA::Walker::WalkerState::startWalk(), gem5::VegaISA::GpuTLB::walkerResponse(), and gem5::VegaISA::Walker::WalkerState::walkStateMachine().

ppn

Bitfield<47, 12> gem5::VegaISA::ppn

Definition at line 57 of file pagetable.hh.

r

Bitfield<5> gem5::VegaISA::r

Definition at line 60 of file pagetable.hh.

RegSizeDWords

const int gem5::VegaISA::RegSizeDWords = sizeof(VecElemU32) / DWordSize

Size of a single-precision register in DWords.

Definition at line 176 of file gpu_registers.hh.

s

Bitfield< 1 > gem5::VegaISA::s

Definition at line 64 of file pagetable.hh.

sw

Bitfield<53,52> gem5::VegaISA::sw

Definition at line 55 of file pagetable.hh.

t

Bitfield<51> gem5::VegaISA::t

Definition at line 56 of file pagetable.hh.

v

Bitfield< 0 > gem5::VegaISA::v

Definition at line 65 of file pagetable.hh.

w

Bitfield<6> gem5::VegaISA::w

Definition at line 59 of file pagetable.hh.

x

Bitfield<4> gem5::VegaISA::x

Definition at line 61 of file pagetable.hh.

z

Bitfield<3> gem5::VegaISA::z

Definition at line 62 of file pagetable.hh.