Vector Packed Arithmetic Instructions
176 AMDGPU instructions in this category - showing 100 per page, page 1 of 2 - click any row for encoding, pseudocode, and full documentation.
| Mnemonic | Syntax | Format | Summary |
|---|---|---|---|
| v_dot2_i32_i16 | v_dot2_i32_i16 | VOP3P | Compute the dot product of two packed 2-D signed 16-bit integer inputs in the signed 32-bit integer domain, add a signed 32-bit integer value from… |
| v_dot2_u32_u16 | v_dot2_u32_u16 | VOP3P | Compute the dot product of two packed 2-D unsigned 16-bit integer inputs in the unsigned 32-bit integer domain, add an unsigned 32-bit integer value… |
| v_dot4_f32_bf8_bf8 | v_dot4_f32_bf8_bf8 | VOP3P | Compute the dot product of two packed 4-D BF8 float inputs in the single-precision float domain, add a single-precision float value from the third… |
| v_dot4_f32_bf8_fp8 | v_dot4_f32_bf8_fp8 | VOP3P | Compute the dot product of a packed 4-D BF8 float input and a packed 4-D FP8 float input in the single-precision float domain, add a single-precision… |
| v_dot4_f32_fp8_bf8 | v_dot4_f32_fp8_bf8 | VOP3P | Compute the dot product of a packed 4-D FP8 float input and a packed 4-D BF8 float input in the single-precision float domain, add a single-precision… |
| v_dot4_f32_fp8_fp8 | v_dot4_f32_fp8_fp8 | VOP3P | Compute the dot product of two packed 4-D FP8 float inputs in the single-precision float domain, add a single-precision float value from the third… |
| v_dot4_i32_i8 | v_dot4_i32_i8 | VOP3P | Compute the dot product of two packed 4-D signed 8-bit integer inputs in the signed 32-bit integer domain, add a signed 32-bit integer value from the… |
| v_dot4_i32_iu8 | v_dot4_i32_iu8 | VOP3P | Compute the dot product of two packed 4-D signed or unsigned 8-bit integer inputs in the signed 32-bit integer domain, add a signed 32-bit integer… |
| v_dot4_u32_u8 | v_dot4_u32_u8 | VOP3P | Compute the dot product of two packed 4-D unsigned 8-bit integer inputs in the unsigned 32-bit integer domain, add an unsigned 32-bit integer value… |
| v_dot8_i32_i4 | v_dot8_i32_i4 | VOP3P | Compute the dot product of two packed 8-D signed 4-bit integer inputs in the signed 32-bit integer domain, add a signed 32-bit integer value from the… |
| v_dot8_i32_iu4 | v_dot8_i32_iu4 | VOP3P | Compute the dot product of two packed 8-D signed or unsigned 4-bit integer inputs in the signed 32-bit integer domain, add a signed 32-bit integer… |
| v_dot8_u32_u4 | v_dot8_u32_u4 | VOP3P | Compute the dot product of two packed 8-D unsigned 4-bit integer inputs in the unsigned 32-bit integer domain, add an unsigned 32-bit integer value… |
| v_fma_mix_bf16_t16 | v_fma_mix_bf16_t16 | VOP3P | AMDGPU VOP3P vector instruction. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_fma_mix_f16_t16 | v_fma_mix_f16_t16 | VOP3P | AMDGPU VOP3P vector instruction operating on f16 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_fma_mix_f32 | v_fma_mix_f32 | VOP3P | Multiply two inputs and add a third input using fused multiply add where the inputs are a mix of half-precision float and single-precision float… |
| v_fma_mix_f32_bf16 | v_fma_mix_f32_bf16 | VOP3P | AMDGPU VOP3P vector instruction operating on f32 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_fma_mixhi_bf16 | v_fma_mixhi_bf16 | VOP3P | AMDGPU VOP3P vector instruction. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_fma_mixhi_f16 | v_fma_mixhi_f16 | VOP3P | Multiply two inputs and add a third input using fused multiply add where the inputs are a mix of half-precision float and single-precision float… |
| v_fma_mixlo_bf16 | v_fma_mixlo_bf16 | VOP3P | AMDGPU VOP3P vector instruction. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_fma_mixlo_f16 | v_fma_mixlo_f16 | VOP3P | Multiply two inputs and add a third input using fused multiply add where the inputs are a mix of half-precision float and single-precision float… |
| v_mad_mix_f32 | v_mad_mix_f32 | VOP3P | Multiply two inputs and add a third input where the inputs are a mix of half-precision float and single- precision float values. |
| v_mad_mixhi_f16 | v_mad_mixhi_f16 | VOP3P | Multiply two inputs and add a third input where the inputs are a mix of half-precision float and single- precision float values. |
| v_mad_mixlo_f16 | v_mad_mixlo_f16 | VOP3P | Multiply two inputs and add a third input where the inputs are a mix of half-precision float and single- precision float values. |
| v_mfma_f32_16x16x128_f8f6f4 | v_mfma_f32_16x16x128_f8f6f4 | VOP3P | Multiply the 16x128 matrix in the first input by the 128x16 matrix in the second input and add the 16x16 matrix in the third input using fused… |
| v_mfma_f32_16x16x16_bf16 | v_mfma_f32_16x16x16_bf16 | VOP3P | Multiply the 16x16 matrix in the first input by the 16x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_f32_16x16x16_f16 | v_mfma_f32_16x16x16_f16 | VOP3P | Multiply the 16x16 matrix in the first input by the 16x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_f32_16x16x1_4b_f32 | v_mfma_f32_16x16x1_4b_f32 | VOP3P | Multiply the 16x1 matrix in the first input by the 1x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_f32_16x16x32_bf16 | v_mfma_f32_16x16x32_bf16 | VOP3P | Multiply the 16x32 matrix in the first input by the 32x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_f32_16x16x32_f16 | v_mfma_f32_16x16x32_f16 | VOP3P | Multiply the 16x32 matrix in the first input by the 32x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_f32_16x16x4_4b_bf16 | v_mfma_f32_16x16x4_4b_bf16 | VOP3P | Multiply the 16x4 matrix in the first input by the 4x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_f32_16x16x4_4b_f16 | v_mfma_f32_16x16x4_4b_f16 | VOP3P | Multiply the 16x4 matrix in the first input by the 4x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_f32_16x16x4_f32 | v_mfma_f32_16x16x4_f32 | VOP3P | Multiply the 16x4 matrix in the first input by the 4x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_f32_16x16x8_xf32 | v_mfma_f32_16x16x8_xf32 | VOP3P | Multiply the 16x8 matrix in the first input by the 8x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_f32_32x32x16_bf16 | v_mfma_f32_32x32x16_bf16 | VOP3P | Multiply the 32x16 matrix in the first input by the 16x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_f32_32x32x16_f16 | v_mfma_f32_32x32x16_f16 | VOP3P | Multiply the 32x16 matrix in the first input by the 16x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_f32_32x32x1_2b_f32 | v_mfma_f32_32x32x1_2b_f32 | VOP3P | Multiply the 32x1 matrix in the first input by the 1x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_f32_32x32x2_f32 | v_mfma_f32_32x32x2_f32 | VOP3P | Multiply the 32x2 matrix in the first input by the 2x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_f32_32x32x4_2b_bf16 | v_mfma_f32_32x32x4_2b_bf16 | VOP3P | Multiply the 32x4 matrix in the first input by the 4x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_f32_32x32x4_2b_f16 | v_mfma_f32_32x32x4_2b_f16 | VOP3P | Multiply the 32x4 matrix in the first input by the 4x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_f32_32x32x4_xf32 | v_mfma_f32_32x32x4_xf32 | VOP3P | Multiply the 32x4 matrix in the first input by the 4x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_f32_32x32x64_f8f6f4 | v_mfma_f32_32x32x64_f8f6f4 | VOP3P | Multiply the 32x64 matrix in the first input by the 64x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_f32_32x32x8_bf16 | v_mfma_f32_32x32x8_bf16 | VOP3P | Multiply the 32x8 matrix in the first input by the 8x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_f32_32x32x8_f16 | v_mfma_f32_32x32x8_f16 | VOP3P | Multiply the 32x8 matrix in the first input by the 8x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_f32_4x4x1_16b_f32 | v_mfma_f32_4x4x1_16b_f32 | VOP3P | Multiply the 4x1 matrix in the first input by the 1x4 matrix in the second input and add the 4x4 matrix in the third input using fused multiply add. |
| v_mfma_f32_4x4x4_16b_bf16 | v_mfma_f32_4x4x4_16b_bf16 | VOP3P | Multiply the 4x4 matrix in the first input by the 4x4 matrix in the second input and add the 4x4 matrix in the third input using fused multiply add. |
| v_mfma_f32_4x4x4_16b_f16 | v_mfma_f32_4x4x4_16b_f16 | VOP3P | Multiply the 4x4 matrix in the first input by the 4x4 matrix in the second input and add the 4x4 matrix in the third input using fused multiply add. |
| v_mfma_f64_16x16x4_f64 | v_mfma_f64_16x16x4_f64 | VOP3P | Multiply the 16x4 matrix in the first input by the 4x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_f64_4x4x4_4b_f64 | v_mfma_f64_4x4x4_4b_f64 | VOP3P | Multiply the 4x4 matrix in the first input by the 4x4 matrix in the second input and add the 4x4 matrix in the third input using fused multiply add. |
| v_mfma_i32_16x16x32_i8 | v_mfma_i32_16x16x32_i8 | VOP3P | Multiply the 16x32 matrix in the first input by the 32x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_i32_16x16x4_4b_i8 | v_mfma_i32_16x16x4_4b_i8 | VOP3P | Multiply the 16x4 matrix in the first input by the 4x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_i32_16x16x64_i8 | v_mfma_i32_16x16x64_i8 | VOP3P | Multiply the 16x64 matrix in the first input by the 64x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… |
| v_mfma_i32_32x32x16_i8 | v_mfma_i32_32x32x16_i8 | VOP3P | Multiply the 32x16 matrix in the first input by the 16x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_i32_32x32x32_i8 | v_mfma_i32_32x32x32_i8 | VOP3P | Multiply the 32x32 matrix in the first input by the 32x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_i32_32x32x4_2b_i8 | v_mfma_i32_32x32x4_2b_i8 | VOP3P | Multiply the 32x4 matrix in the first input by the 4x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… |
| v_mfma_i32_4x4x4_16b_i8 | v_mfma_i32_4x4x4_16b_i8 | VOP3P | Multiply the 4x4 matrix in the first input by the 4x4 matrix in the second input and add the 4x4 matrix in the third input using fused multiply add. |
| v_mfma_ld_scale_b32 | v_mfma_ld_scale_b32 | VOP3P | AMDGPU VOP3P matrix instruction operating on b32 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_add_bf16 | v_pk_add_bf16 | VOP3P | AMDGPU VOP3P vector instruction. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_add_f16 | v_pk_add_f16 | VOP3P | Add two packed half-precision float inputs component-wise and store the result into a vector register. No carry- in or carry-out support. |
| v_pk_add_f32 | v_pk_add_f32 | VOP3P | Add two packed single-precision float inputs component-wise and store the result into a vector register. No carry-in or carry-out support. |
| v_pk_add_f64 | v_pk_add_f64 | VOP3P | AMDGPU VOP3P vector instruction operating on f64 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_add_i16 | v_pk_add_i16 | VOP3P | Add two packed signed 16-bit integer inputs component-wise and store the result into a vector register. No carry-in or carry-out support. |
| v_pk_add_max_i16 | v_pk_add_max_i16 | VOP3P | AMDGPU VOP3P vector instruction operating on i16 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_add_max_u16 | v_pk_add_max_u16 | VOP3P | AMDGPU VOP3P vector instruction operating on u16 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_add_min_i16 | v_pk_add_min_i16 | VOP3P | AMDGPU VOP3P vector instruction operating on i16 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_add_min_u16 | v_pk_add_min_u16 | VOP3P | AMDGPU VOP3P vector instruction operating on u16 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_add_nc_u64 | v_pk_add_nc_u64 | VOP3P | AMDGPU VOP3P vector instruction operating on u64 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_add_u16 | v_pk_add_u16 | VOP3P | Add two packed unsigned 16-bit integer inputs component-wise and store the result into a vector register. No carry-in or carry-out support. |
| v_pk_ashrrev_i16 | v_pk_ashrrev_i16 | VOP3P | Given a packed shift count in the first vector input, calculate the component-wise arithmetic shift right (preserving sign bit) of the second packed… |
| v_pk_fma_bf16 | v_pk_fma_bf16 | VOP3P | AMDGPU VOP3P vector instruction. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_fma_f16 | v_pk_fma_f16 | VOP3P | Multiply two packed half-precision float inputs component-wise and add a third input component-wise using fused multiply add, and store the result… |
| v_pk_fma_f32 | v_pk_fma_f32 | VOP3P | Multiply two packed single-precision float inputs component-wise and add a third input component-wise using fused multiply add, and store the result… |
| v_pk_fma_f64 | v_pk_fma_f64 | VOP3P | AMDGPU VOP3P vector instruction operating on f64 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_lshl_add_u64 | v_pk_lshl_add_u64 | VOP3P | AMDGPU VOP3P vector instruction operating on u64 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_lshlrev_b16 | v_pk_lshlrev_b16 | VOP3P | Given a packed shift count in the first vector input, calculate the component-wise logical shift left of the second packed vector input and store the… |
| v_pk_lshrrev_b16 | v_pk_lshrrev_b16 | VOP3P | Given a packed shift count in the first vector input, calculate the component-wise logical shift right of the second packed vector input and store… |
| v_pk_mad_i16 | v_pk_mad_i16 | VOP3P | Multiply two packed signed 16-bit integer inputs component-wise, add a packed signed 16-bit integer value from a third input component-wise, and… |
| v_pk_mad_u16 | v_pk_mad_u16 | VOP3P | Multiply two packed unsigned 16-bit integer inputs component-wise, add a packed unsigned 16-bit integer value from a third input component-wise, and… |
| v_pk_max3_i16 | v_pk_max3_i16 | VOP3P | AMDGPU VOP3P vector instruction operating on i16 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_max3_num_f16 | v_pk_max3_num_f16 | VOP3P | AMDGPU VOP3P vector instruction operating on f16 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_max3_u16 | v_pk_max3_u16 | VOP3P | AMDGPU VOP3P vector instruction operating on u16 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_max_f16 | v_pk_max_f16 | VOP3P | Select the component-wise maximum of two packed half-precision float inputs and store the result into a vector register. |
| v_pk_max_i16 | v_pk_max_i16 | VOP3P | Select the component-wise maximum of two packed signed 16-bit integer inputs and store the selected values into a vector register. |
| v_pk_max_num_bf16 | v_pk_max_num_bf16 | VOP3P | AMDGPU VOP3P vector instruction. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_max_num_f16 | v_pk_max_num_f16 | VOP3P | Select the component-wise IEEE maximumNumber() of two packed half-precision float inputs and store the selected values into a vector register. |
| v_pk_max_num_f64 | v_pk_max_num_f64 | VOP3P | AMDGPU VOP3P vector instruction operating on f64 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_max_u16 | v_pk_max_u16 | VOP3P | Select the component-wise maximum of two packed unsigned 16-bit integer inputs and store the selected values into a vector register. |
| v_pk_maximum3_f16 | v_pk_maximum3_f16 | VOP3P | Select the component-wise IEEE maximum() of three half-precision float inputs and store the result into a vector register. |
| v_pk_maximum_f16 | v_pk_maximum_f16 | VOP3P | Select the component-wise IEEE maximum() of two packed half-precision float inputs and store the selected values into a vector register. |
| v_pk_min3_i16 | v_pk_min3_i16 | VOP3P | AMDGPU VOP3P vector instruction operating on i16 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_min3_num_f16 | v_pk_min3_num_f16 | VOP3P | AMDGPU VOP3P vector instruction operating on f16 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_min3_u16 | v_pk_min3_u16 | VOP3P | AMDGPU VOP3P vector instruction operating on u16 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_min_f16 | v_pk_min_f16 | VOP3P | Select the component-wise minimum of two packed half-precision float inputs and store the result into a vector register. |
| v_pk_min_i16 | v_pk_min_i16 | VOP3P | Select the component-wise minimum of two packed signed 16-bit integer inputs and store the selected values into a vector register. |
| v_pk_min_num_bf16 | v_pk_min_num_bf16 | VOP3P | AMDGPU VOP3P vector instruction. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_min_num_f16 | v_pk_min_num_f16 | VOP3P | Select the component-wise IEEE minimumNumber() of two packed half-precision float inputs and store the selected values into a vector register. |
| v_pk_min_num_f64 | v_pk_min_num_f64 | VOP3P | AMDGPU VOP3P vector instruction operating on f64 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) |
| v_pk_min_u16 | v_pk_min_u16 | VOP3P | Select the component-wise minimum of two packed unsigned 16-bit integer inputs and store the selected values into a vector register. |
| v_pk_minimum3_f16 | v_pk_minimum3_f16 | VOP3P | Select the component-wise IEEE minimum() of three half-precision float inputs and store the result into a vector register. |
| v_pk_minimum_f16 | v_pk_minimum_f16 | VOP3P | Select the component-wise IEEE minimum() of two packed half-precision float inputs and store the selected values into a vector register. |
| v_pk_mov_b32 | v_pk_mov_b32 | VOP3P | Move data from two vector inputs into two vector registers. |