[AMDGPU] Add support for `v_sat_pk4_i4_[i8,u8]` on gfx1250 #149528

shiltian · 2025-07-18T15:07:50Z

Co-authored-by: Mekhanoshin, Stanislav Stanislav.Mekhanoshin@amd.com
Co-authored-by: Foad, Jay Jay.Foad@amd.com

shiltian · 2025-07-18T15:08:10Z

This stack of pull requests is managed by Graphite. Learn more about stacking.

llvmbot · 2025-07-18T15:08:21Z

@llvm/pr-subscribers-mc

@llvm/pr-subscribers-clang

Author: Shilei Tian (shiltian)

Changes

Co-authored-by: Mekhanoshin, Stanislav <Stanislav.Mekhanoshin@amd.com>
Co-authored-by: Foad, Jay <Jay.Foad@amd.com>

Patch is 60.85 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/149528.diff

25 Files Affected:

(modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+2)
(modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl (+25)
(modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+6)
(modified) llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp (+2)
(modified) llvm/lib/Target/AMDGPU/SIInstrInfo.td (+1)
(modified) llvm/lib/Target/AMDGPU/VOP1Instructions.td (+5)
(added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll (+305)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s (+51)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s (+30)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s (+24)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s (+24)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s (+24)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s (+30)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s (+24)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s (+24)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt (+42)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt (+22)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt (+22)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt (+56)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt (+24)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt (+24)

diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index a916af7e0c2df..d4fef5d46af73 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -684,6 +684,8 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_fp8, "hiIi", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_bf8, "hiIi", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_fp8, "V2hs", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_bf8, "V2hs", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_sat_pk4_i4_i8, "UsUi", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_sat_pk4_u4_u8, "UsUi", "nc", "gfx1250-insts")
 
 // GFX1250 WMMA builtins
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x4_f32, "V8fIbV2fIbV2fIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
index 4c3f308a6cf75..a21862c4a9395 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
@@ -5,6 +5,7 @@
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 typedef unsigned int uint;
+typedef unsigned short int ushort;
 typedef unsigned int __attribute__((ext_vector_type(2))) uint2;
 typedef half __attribute__((ext_vector_type(2))) half2;
 
@@ -369,6 +370,30 @@ void test_cvt_pk_f16_bf8(global half2* out, short a)
   out[0] = __builtin_amdgcn_cvt_pk_f16_bf8(a);
 }
 
+// CHECK-LABEL: @test_sat_pk4_i4_i8(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT:    [[SRC_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
+// CHECK-NEXT:    [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr
+// CHECK-NEXT:    store ptr [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i32 [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr [[SRC_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP1:%.*]] = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 [[TMP0]])
+// CHECK-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i16 [[TMP1]], ptr [[TMP2]], align 2
+// CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[SRC_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP4:%.*]] = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 [[TMP3]])
+// CHECK-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i16 [[TMP4]], ptr [[TMP5]], align 2
+// CHECK-NEXT:    ret void
+//
+void test_sat_pk4_i4_i8(ushort *out, uint src)
+{
+  *out = __builtin_amdgcn_sat_pk4_i4_i8(src);
+  *out = __builtin_amdgcn_sat_pk4_u4_u8(src);
+}
+
 // CHECK-LABEL: @test_permlane16_swap(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index d8fda0e2bcfa3..ecda6c4efefe3 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -3572,6 +3572,12 @@ def int_amdgcn_cvt_f16_bf8 : ClangBuiltin<"__builtin_amdgcn_cvt_f16_bf8">,
             [llvm_i32_ty, llvm_i32_ty],
             [IntrNoMem, ImmArg<ArgIndex<1>>]>;
 
+def int_amdgcn_sat_pk4_i4_i8 : ClangBuiltin<"__builtin_amdgcn_sat_pk4_i4_i8">,
+  DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i32_ty], [IntrNoMem, IntrSpeculatable]>;
+
+def int_amdgcn_sat_pk4_u4_u8 : ClangBuiltin<"__builtin_amdgcn_sat_pk4_u4_u8">,
+  DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i32_ty], [IntrNoMem, IntrSpeculatable]>;
+
 //===----------------------------------------------------------------------===//
 // Special Intrinsics for backend internal use only. No frontend
 // should emit calls to these.
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index cbbb57c6f8122..bf2f37bddb9ed 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4558,6 +4558,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
     case Intrinsic::amdgcn_cvt_pk_u16:
     case Intrinsic::amdgcn_cvt_pk_f16_fp8:
     case Intrinsic::amdgcn_cvt_pk_f16_bf8:
+    case Intrinsic::amdgcn_sat_pk4_i4_i8:
+    case Intrinsic::amdgcn_sat_pk4_u4_u8:
     case Intrinsic::amdgcn_fmed3:
     case Intrinsic::amdgcn_cubeid:
     case Intrinsic::amdgcn_cubema:
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.td b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
index ab7d34002e9f1..9e1951e2946c4 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.td
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
@@ -2850,6 +2850,7 @@ def VOP_F16_I16 : VOPProfile <[f16, i16, untyped, untyped]>;
 def VOP_I16_F16 : VOPProfile <[i16, f16, untyped, untyped]>;
 def VOP_I16_I16 : VOPProfile <[i16, i16, untyped, untyped]>;
 def VOP_BF16_BF16 : VOPProfile<[bf16, bf16, untyped, untyped]>;
+def VOP1_I16_I32 :  VOPProfile<[i16, i32, untyped, untyped]>;
 
 def VOP_F16_F16_F16 : VOPProfile <[f16, f16, f16, untyped]>;
 def VOP_F16_F16_I16 : VOPProfile <[f16, f16, i16, untyped]>;
diff --git a/llvm/lib/Target/AMDGPU/VOP1Instructions.td b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
index 80eb5d8b7d571..f621f8581f778 100644
--- a/llvm/lib/Target/AMDGPU/VOP1Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
@@ -803,6 +803,9 @@ let SubtargetPredicate = isGFX1250Plus in {
     def : Cvt_F_F8_Pat_ByteSel<int_amdgcn_cvt_f16_fp8, V_CVT_F16_FP8_fake16_e64, 1>;
     def : Cvt_F_F8_Pat_ByteSel<int_amdgcn_cvt_f16_bf8, V_CVT_F16_BF8_fake16_e64, 1>;
   }
+
+  defm V_SAT_PK4_I4_I8 : VOP1Inst_t16<"v_sat_pk4_i4_i8", VOP1_I16_I32, int_amdgcn_sat_pk4_i4_i8>;
+  defm V_SAT_PK4_U4_U8 : VOP1Inst_t16<"v_sat_pk4_u4_u8", VOP1_I16_I32, int_amdgcn_sat_pk4_u4_u8>;
 } // End SubtargetPredicate = isGFX1250Plus
 
 let SubtargetPredicate = isGFX10Plus in {
@@ -1158,6 +1161,8 @@ defm V_PERMLANE16_SWAP_B32   : VOP1_Real_OpSelIsDPP_gfx1250<0x049>;
 defm V_TANH_BF16             : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x04a>;
 defm V_PRNG_B32              : VOP1_Real_FULL<GFX1250Gen, 0x04b>;
 defm V_CVT_F32_BF16          : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x072, "v_cvt_f32_bf16", "V_CVT_F32_BF16_gfx1250">;
+defm V_SAT_PK4_I4_I8         : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x073>;
+defm V_SAT_PK4_U4_U8         : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x074>;
 defm V_CVT_PK_F16_FP8        : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x075>;
 defm V_CVT_PK_F16_BF8        : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x076>;
 defm V_CVT_F16_FP8           : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x077>;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll
new file mode 100644
index 0000000000000..3a5507063b834
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll
@@ -0,0 +1,305 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=+real-true16 %s -o - | FileCheck -check-prefix=SDAG-REAL16 %s
+; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=-real-true16 %s -o - | FileCheck -check-prefix=SDAG-FAKE16 %s
+; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=+real-true16 %s -o - | FileCheck -check-prefix=GISEL-REAL16 %s
+; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=-real-true16 %s -o - | FileCheck -check-prefix=GISEL-FAKE16 %s
+
+declare i16 @llvm.amdgcn.sat.pk4.i4.i8(i32) #0
+declare i16 @llvm.amdgcn.sat.pk4.u4.u8(i32) #0
+
+define amdgpu_kernel void @sat_pk4_i4_i8_f32_v(i32 %src, ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_i4_i8_f32_v:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_clause 0x1
+; SDAG-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, s2
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_i4_i8_f32_v:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_clause 0x1
+; SDAG-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v1, s2
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_i4_i8_f32_v:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_clause 0x1
+; GISEL-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, s2
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_i4_i8_f32_v:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_clause 0x1
+; GISEL-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v0, s2
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 %src) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_i4_i8_f32_s(i32 inreg %src, ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_i4_i8_f32_s:
+; SDAG-REAL16:       ; %bb.1:
+; SDAG-REAL16-NEXT:    s_load_b32 s8, s[4:5], 0x0
+; SDAG-REAL16-NEXT:    s_waitcnt lgkmcnt(0)
+; SDAG-REAL16-NEXT:    s_branch .LBB1_0
+; SDAG-REAL16-NEXT:    .p2align 8
+; SDAG-REAL16-NEXT:  ; %bb.2:
+; SDAG-REAL16-NEXT:  .LBB1_0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, s8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_i4_i8_f32_s:
+; SDAG-FAKE16:       ; %bb.1:
+; SDAG-FAKE16-NEXT:    s_load_b32 s8, s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    s_waitcnt lgkmcnt(0)
+; SDAG-FAKE16-NEXT:    s_branch .LBB1_0
+; SDAG-FAKE16-NEXT:    .p2align 8
+; SDAG-FAKE16-NEXT:  ; %bb.2:
+; SDAG-FAKE16-NEXT:  .LBB1_0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v1, s8
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_i4_i8_f32_s:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_clause 0x1
+; GISEL-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, s2
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_i4_i8_f32_s:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_clause 0x1
+; GISEL-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v0, s2
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 %src) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_i4_i8_f32_i(ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_i4_i8_f32_i:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, 0x64
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_i4_i8_f32_i:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v1, 0x64
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_i4_i8_f32_i:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, 0x64
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_i4_i8_f32_i:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v0, 0x64
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 100) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_u4_u8_f32_v(i32 %src, ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_u4_u8_f32_v:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_clause 0x1
+; SDAG-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, s2
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_u4_u8_f32_v:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_clause 0x1
+; SDAG-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v1, s2
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_u4_u8_f32_v:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_clause 0x1
+; GISEL-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, s2
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_u4_u8_f32_v:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_clause 0x1
+; GISEL-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v0, s2
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 %src) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_u4_u8_f32_s(i32 inreg %src, ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_u4_u8_f32_s:
+; SDAG-REAL16:       ; %bb.1:
+; SDAG-REAL16-NEXT:    s_load_b32 s8, s[4:5], 0x0
+; SDAG-REAL16-NEXT:    s_waitcnt lgkmcnt(0)
+; SDAG-REAL16-NEXT:    s_branch .LBB4_0
+; SDAG-REAL16-NEXT:    .p2align 8
+; SDAG-REAL16-NEXT:  ; %bb.2:
+; SDAG-REAL16-NEXT:  .LBB4_0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, s8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_u4_u8_f32_s:
+; SDAG-FAKE16:       ; %bb.1:
+; SDAG-FAKE16-NEXT:    s_load_b32 s8, s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    s_waitcnt lgkmcnt(0)
+; SDAG-FAKE16-NEXT:    s_branch .LBB4_0
+; SDAG-FAKE16-NEXT:    .p2align 8
+; SDAG-FAKE16-NEXT:  ; %bb.2:
+; SDAG-FAKE16-NEXT:  .LBB4_0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v1, s8
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_u4_u8_f32_s:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_clause 0x1
+; GISEL-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, s2
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_u4_u8_f32_s:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_clause 0x1
+; GISEL-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v0, s2
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 %src) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_u4_u8_f32_i(ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_u4_u8_f32_i:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, 0x64
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_u4_u8_f32_i:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v1, 0x64
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_u4_u8_f32_i:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, 0x64
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_u4_u8_f32_i:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v0, 0x64
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 100) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+attributes #0 = { nounwind memory(none) }
+attributes #1 = { nounwind }
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
index f2cf3d58fb0cf..811c6ebfe0161 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
@@ -628,6 +628,57 @@ v_cvt_f32_fp8_e32 v1, 3
 v_cvt_f32_fp8_e32 v1, v3
 // GFX1250: v_cvt_f32_fp8_e32 v1, v3                ; encoding: [0x03,0xd9,0x02,0x7e]
 
+v_cvt_pk_f32_bf8_e32 v[2:3], s3
+// GFX1250: v_cvt_pk_f32_bf8_e32 v[2:3], s3         ; encoding: [0x03,0xde,0x04,0x7e]
+
+v_cvt_pk_f32_bf8_e32 v[4:5], s5
+// GFX1250: v_cvt_pk_f32_bf8_e32 v[4:5], s5         ; encoding: [0x05,0xde,0x08,0x7e]
+
+v_cvt_pk_f32_bf8_e32 v[2:3], 3
+// GFX1250: v_cvt_pk_f32_bf8_e32 v[2:3], 3          ; encoding: [0x83,0xde,0x04,0x7e]
+
+v_cvt_pk_f32_bf8...
[truncated]

llvmbot · 2025-07-18T15:08:22Z

@llvm/pr-subscribers-backend-amdgpu

Author: Shilei Tian (shiltian)

Changes

Co-authored-by: Mekhanoshin, Stanislav <Stanislav.Mekhanoshin@amd.com>
Co-authored-by: Foad, Jay <Jay.Foad@amd.com>

Patch is 60.85 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/149528.diff

25 Files Affected:

(modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+2)
(modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl (+25)
(modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+6)
(modified) llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp (+2)
(modified) llvm/lib/Target/AMDGPU/SIInstrInfo.td (+1)
(modified) llvm/lib/Target/AMDGPU/VOP1Instructions.td (+5)
(added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll (+305)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s (+51)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s (+30)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s (+24)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s (+24)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s (+24)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s (+30)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s (+24)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s (+24)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt (+42)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt (+22)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt (+22)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt (+56)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt (+24)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt (+24)

diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index a916af7e0c2df..d4fef5d46af73 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -684,6 +684,8 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_fp8, "hiIi", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_bf8, "hiIi", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_fp8, "V2hs", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_bf8, "V2hs", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_sat_pk4_i4_i8, "UsUi", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_sat_pk4_u4_u8, "UsUi", "nc", "gfx1250-insts")
 
 // GFX1250 WMMA builtins
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x4_f32, "V8fIbV2fIbV2fIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
index 4c3f308a6cf75..a21862c4a9395 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
@@ -5,6 +5,7 @@
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 typedef unsigned int uint;
+typedef unsigned short int ushort;
 typedef unsigned int __attribute__((ext_vector_type(2))) uint2;
 typedef half __attribute__((ext_vector_type(2))) half2;
 
@@ -369,6 +370,30 @@ void test_cvt_pk_f16_bf8(global half2* out, short a)
   out[0] = __builtin_amdgcn_cvt_pk_f16_bf8(a);
 }
 
+// CHECK-LABEL: @test_sat_pk4_i4_i8(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT:    [[SRC_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
+// CHECK-NEXT:    [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr
+// CHECK-NEXT:    store ptr [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i32 [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr [[SRC_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP1:%.*]] = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 [[TMP0]])
+// CHECK-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i16 [[TMP1]], ptr [[TMP2]], align 2
+// CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[SRC_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP4:%.*]] = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 [[TMP3]])
+// CHECK-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i16 [[TMP4]], ptr [[TMP5]], align 2
+// CHECK-NEXT:    ret void
+//
+void test_sat_pk4_i4_i8(ushort *out, uint src)
+{
+  *out = __builtin_amdgcn_sat_pk4_i4_i8(src);
+  *out = __builtin_amdgcn_sat_pk4_u4_u8(src);
+}
+
 // CHECK-LABEL: @test_permlane16_swap(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index d8fda0e2bcfa3..ecda6c4efefe3 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -3572,6 +3572,12 @@ def int_amdgcn_cvt_f16_bf8 : ClangBuiltin<"__builtin_amdgcn_cvt_f16_bf8">,
             [llvm_i32_ty, llvm_i32_ty],
             [IntrNoMem, ImmArg<ArgIndex<1>>]>;
 
+def int_amdgcn_sat_pk4_i4_i8 : ClangBuiltin<"__builtin_amdgcn_sat_pk4_i4_i8">,
+  DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i32_ty], [IntrNoMem, IntrSpeculatable]>;
+
+def int_amdgcn_sat_pk4_u4_u8 : ClangBuiltin<"__builtin_amdgcn_sat_pk4_u4_u8">,
+  DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i32_ty], [IntrNoMem, IntrSpeculatable]>;
+
 //===----------------------------------------------------------------------===//
 // Special Intrinsics for backend internal use only. No frontend
 // should emit calls to these.
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index cbbb57c6f8122..bf2f37bddb9ed 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4558,6 +4558,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
     case Intrinsic::amdgcn_cvt_pk_u16:
     case Intrinsic::amdgcn_cvt_pk_f16_fp8:
     case Intrinsic::amdgcn_cvt_pk_f16_bf8:
+    case Intrinsic::amdgcn_sat_pk4_i4_i8:
+    case Intrinsic::amdgcn_sat_pk4_u4_u8:
     case Intrinsic::amdgcn_fmed3:
     case Intrinsic::amdgcn_cubeid:
     case Intrinsic::amdgcn_cubema:
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.td b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
index ab7d34002e9f1..9e1951e2946c4 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.td
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
@@ -2850,6 +2850,7 @@ def VOP_F16_I16 : VOPProfile <[f16, i16, untyped, untyped]>;
 def VOP_I16_F16 : VOPProfile <[i16, f16, untyped, untyped]>;
 def VOP_I16_I16 : VOPProfile <[i16, i16, untyped, untyped]>;
 def VOP_BF16_BF16 : VOPProfile<[bf16, bf16, untyped, untyped]>;
+def VOP1_I16_I32 :  VOPProfile<[i16, i32, untyped, untyped]>;
 
 def VOP_F16_F16_F16 : VOPProfile <[f16, f16, f16, untyped]>;
 def VOP_F16_F16_I16 : VOPProfile <[f16, f16, i16, untyped]>;
diff --git a/llvm/lib/Target/AMDGPU/VOP1Instructions.td b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
index 80eb5d8b7d571..f621f8581f778 100644
--- a/llvm/lib/Target/AMDGPU/VOP1Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
@@ -803,6 +803,9 @@ let SubtargetPredicate = isGFX1250Plus in {
     def : Cvt_F_F8_Pat_ByteSel<int_amdgcn_cvt_f16_fp8, V_CVT_F16_FP8_fake16_e64, 1>;
     def : Cvt_F_F8_Pat_ByteSel<int_amdgcn_cvt_f16_bf8, V_CVT_F16_BF8_fake16_e64, 1>;
   }
+
+  defm V_SAT_PK4_I4_I8 : VOP1Inst_t16<"v_sat_pk4_i4_i8", VOP1_I16_I32, int_amdgcn_sat_pk4_i4_i8>;
+  defm V_SAT_PK4_U4_U8 : VOP1Inst_t16<"v_sat_pk4_u4_u8", VOP1_I16_I32, int_amdgcn_sat_pk4_u4_u8>;
 } // End SubtargetPredicate = isGFX1250Plus
 
 let SubtargetPredicate = isGFX10Plus in {
@@ -1158,6 +1161,8 @@ defm V_PERMLANE16_SWAP_B32   : VOP1_Real_OpSelIsDPP_gfx1250<0x049>;
 defm V_TANH_BF16             : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x04a>;
 defm V_PRNG_B32              : VOP1_Real_FULL<GFX1250Gen, 0x04b>;
 defm V_CVT_F32_BF16          : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x072, "v_cvt_f32_bf16", "V_CVT_F32_BF16_gfx1250">;
+defm V_SAT_PK4_I4_I8         : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x073>;
+defm V_SAT_PK4_U4_U8         : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x074>;
 defm V_CVT_PK_F16_FP8        : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x075>;
 defm V_CVT_PK_F16_BF8        : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x076>;
 defm V_CVT_F16_FP8           : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x077>;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll
new file mode 100644
index 0000000000000..3a5507063b834
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll
@@ -0,0 +1,305 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=+real-true16 %s -o - | FileCheck -check-prefix=SDAG-REAL16 %s
+; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=-real-true16 %s -o - | FileCheck -check-prefix=SDAG-FAKE16 %s
+; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=+real-true16 %s -o - | FileCheck -check-prefix=GISEL-REAL16 %s
+; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=-real-true16 %s -o - | FileCheck -check-prefix=GISEL-FAKE16 %s
+
+declare i16 @llvm.amdgcn.sat.pk4.i4.i8(i32) #0
+declare i16 @llvm.amdgcn.sat.pk4.u4.u8(i32) #0
+
+define amdgpu_kernel void @sat_pk4_i4_i8_f32_v(i32 %src, ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_i4_i8_f32_v:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_clause 0x1
+; SDAG-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, s2
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_i4_i8_f32_v:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_clause 0x1
+; SDAG-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v1, s2
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_i4_i8_f32_v:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_clause 0x1
+; GISEL-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, s2
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_i4_i8_f32_v:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_clause 0x1
+; GISEL-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v0, s2
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 %src) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_i4_i8_f32_s(i32 inreg %src, ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_i4_i8_f32_s:
+; SDAG-REAL16:       ; %bb.1:
+; SDAG-REAL16-NEXT:    s_load_b32 s8, s[4:5], 0x0
+; SDAG-REAL16-NEXT:    s_waitcnt lgkmcnt(0)
+; SDAG-REAL16-NEXT:    s_branch .LBB1_0
+; SDAG-REAL16-NEXT:    .p2align 8
+; SDAG-REAL16-NEXT:  ; %bb.2:
+; SDAG-REAL16-NEXT:  .LBB1_0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, s8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_i4_i8_f32_s:
+; SDAG-FAKE16:       ; %bb.1:
+; SDAG-FAKE16-NEXT:    s_load_b32 s8, s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    s_waitcnt lgkmcnt(0)
+; SDAG-FAKE16-NEXT:    s_branch .LBB1_0
+; SDAG-FAKE16-NEXT:    .p2align 8
+; SDAG-FAKE16-NEXT:  ; %bb.2:
+; SDAG-FAKE16-NEXT:  .LBB1_0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v1, s8
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_i4_i8_f32_s:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_clause 0x1
+; GISEL-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, s2
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_i4_i8_f32_s:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_clause 0x1
+; GISEL-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v0, s2
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 %src) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_i4_i8_f32_i(ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_i4_i8_f32_i:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, 0x64
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_i4_i8_f32_i:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v1, 0x64
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_i4_i8_f32_i:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, 0x64
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_i4_i8_f32_i:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v0, 0x64
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 100) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_u4_u8_f32_v(i32 %src, ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_u4_u8_f32_v:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_clause 0x1
+; SDAG-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, s2
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_u4_u8_f32_v:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_clause 0x1
+; SDAG-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v1, s2
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_u4_u8_f32_v:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_clause 0x1
+; GISEL-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, s2
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_u4_u8_f32_v:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_clause 0x1
+; GISEL-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v0, s2
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 %src) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_u4_u8_f32_s(i32 inreg %src, ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_u4_u8_f32_s:
+; SDAG-REAL16:       ; %bb.1:
+; SDAG-REAL16-NEXT:    s_load_b32 s8, s[4:5], 0x0
+; SDAG-REAL16-NEXT:    s_waitcnt lgkmcnt(0)
+; SDAG-REAL16-NEXT:    s_branch .LBB4_0
+; SDAG-REAL16-NEXT:    .p2align 8
+; SDAG-REAL16-NEXT:  ; %bb.2:
+; SDAG-REAL16-NEXT:  .LBB4_0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, s8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_u4_u8_f32_s:
+; SDAG-FAKE16:       ; %bb.1:
+; SDAG-FAKE16-NEXT:    s_load_b32 s8, s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    s_waitcnt lgkmcnt(0)
+; SDAG-FAKE16-NEXT:    s_branch .LBB4_0
+; SDAG-FAKE16-NEXT:    .p2align 8
+; SDAG-FAKE16-NEXT:  ; %bb.2:
+; SDAG-FAKE16-NEXT:  .LBB4_0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v1, s8
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_u4_u8_f32_s:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_clause 0x1
+; GISEL-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, s2
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_u4_u8_f32_s:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_clause 0x1
+; GISEL-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v0, s2
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 %src) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_u4_u8_f32_i(ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_u4_u8_f32_i:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, 0x64
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_u4_u8_f32_i:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v1, 0x64
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_u4_u8_f32_i:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, 0x64
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_u4_u8_f32_i:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v0, 0x64
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 100) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+attributes #0 = { nounwind memory(none) }
+attributes #1 = { nounwind }
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
index f2cf3d58fb0cf..811c6ebfe0161 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
@@ -628,6 +628,57 @@ v_cvt_f32_fp8_e32 v1, 3
 v_cvt_f32_fp8_e32 v1, v3
 // GFX1250: v_cvt_f32_fp8_e32 v1, v3                ; encoding: [0x03,0xd9,0x02,0x7e]
 
+v_cvt_pk_f32_bf8_e32 v[2:3], s3
+// GFX1250: v_cvt_pk_f32_bf8_e32 v[2:3], s3         ; encoding: [0x03,0xde,0x04,0x7e]
+
+v_cvt_pk_f32_bf8_e32 v[4:5], s5
+// GFX1250: v_cvt_pk_f32_bf8_e32 v[4:5], s5         ; encoding: [0x05,0xde,0x08,0x7e]
+
+v_cvt_pk_f32_bf8_e32 v[2:3], 3
+// GFX1250: v_cvt_pk_f32_bf8_e32 v[2:3], 3          ; encoding: [0x83,0xde,0x04,0x7e]
+
+v_cvt_pk_f32_bf8...
[truncated]

llvmbot · 2025-07-18T15:08:22Z

@llvm/pr-subscribers-llvm-ir

Author: Shilei Tian (shiltian)

Changes

Co-authored-by: Mekhanoshin, Stanislav <Stanislav.Mekhanoshin@amd.com>
Co-authored-by: Foad, Jay <Jay.Foad@amd.com>

Patch is 60.85 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/149528.diff

25 Files Affected:

(modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+2)
(modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl (+25)
(modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+6)
(modified) llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp (+2)
(modified) llvm/lib/Target/AMDGPU/SIInstrInfo.td (+1)
(modified) llvm/lib/Target/AMDGPU/VOP1Instructions.td (+5)
(added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll (+305)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s (+51)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s (+30)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s (+24)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s (+24)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s (+24)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s (+30)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s (+24)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s (+24)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt (+42)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt (+22)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt (+22)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt (+56)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt (+24)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt (+24)

diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index a916af7e0c2df..d4fef5d46af73 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -684,6 +684,8 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_fp8, "hiIi", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_bf8, "hiIi", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_fp8, "V2hs", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_bf8, "V2hs", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_sat_pk4_i4_i8, "UsUi", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_sat_pk4_u4_u8, "UsUi", "nc", "gfx1250-insts")
 
 // GFX1250 WMMA builtins
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x4_f32, "V8fIbV2fIbV2fIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
index 4c3f308a6cf75..a21862c4a9395 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
@@ -5,6 +5,7 @@
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 typedef unsigned int uint;
+typedef unsigned short int ushort;
 typedef unsigned int __attribute__((ext_vector_type(2))) uint2;
 typedef half __attribute__((ext_vector_type(2))) half2;
 
@@ -369,6 +370,30 @@ void test_cvt_pk_f16_bf8(global half2* out, short a)
   out[0] = __builtin_amdgcn_cvt_pk_f16_bf8(a);
 }
 
+// CHECK-LABEL: @test_sat_pk4_i4_i8(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT:    [[SRC_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
+// CHECK-NEXT:    [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr
+// CHECK-NEXT:    store ptr [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i32 [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr [[SRC_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP1:%.*]] = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 [[TMP0]])
+// CHECK-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i16 [[TMP1]], ptr [[TMP2]], align 2
+// CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[SRC_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP4:%.*]] = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 [[TMP3]])
+// CHECK-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i16 [[TMP4]], ptr [[TMP5]], align 2
+// CHECK-NEXT:    ret void
+//
+void test_sat_pk4_i4_i8(ushort *out, uint src)
+{
+  *out = __builtin_amdgcn_sat_pk4_i4_i8(src);
+  *out = __builtin_amdgcn_sat_pk4_u4_u8(src);
+}
+
 // CHECK-LABEL: @test_permlane16_swap(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index d8fda0e2bcfa3..ecda6c4efefe3 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -3572,6 +3572,12 @@ def int_amdgcn_cvt_f16_bf8 : ClangBuiltin<"__builtin_amdgcn_cvt_f16_bf8">,
             [llvm_i32_ty, llvm_i32_ty],
             [IntrNoMem, ImmArg<ArgIndex<1>>]>;
 
+def int_amdgcn_sat_pk4_i4_i8 : ClangBuiltin<"__builtin_amdgcn_sat_pk4_i4_i8">,
+  DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i32_ty], [IntrNoMem, IntrSpeculatable]>;
+
+def int_amdgcn_sat_pk4_u4_u8 : ClangBuiltin<"__builtin_amdgcn_sat_pk4_u4_u8">,
+  DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i32_ty], [IntrNoMem, IntrSpeculatable]>;
+
 //===----------------------------------------------------------------------===//
 // Special Intrinsics for backend internal use only. No frontend
 // should emit calls to these.
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index cbbb57c6f8122..bf2f37bddb9ed 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4558,6 +4558,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
     case Intrinsic::amdgcn_cvt_pk_u16:
     case Intrinsic::amdgcn_cvt_pk_f16_fp8:
     case Intrinsic::amdgcn_cvt_pk_f16_bf8:
+    case Intrinsic::amdgcn_sat_pk4_i4_i8:
+    case Intrinsic::amdgcn_sat_pk4_u4_u8:
     case Intrinsic::amdgcn_fmed3:
     case Intrinsic::amdgcn_cubeid:
     case Intrinsic::amdgcn_cubema:
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.td b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
index ab7d34002e9f1..9e1951e2946c4 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.td
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
@@ -2850,6 +2850,7 @@ def VOP_F16_I16 : VOPProfile <[f16, i16, untyped, untyped]>;
 def VOP_I16_F16 : VOPProfile <[i16, f16, untyped, untyped]>;
 def VOP_I16_I16 : VOPProfile <[i16, i16, untyped, untyped]>;
 def VOP_BF16_BF16 : VOPProfile<[bf16, bf16, untyped, untyped]>;
+def VOP1_I16_I32 :  VOPProfile<[i16, i32, untyped, untyped]>;
 
 def VOP_F16_F16_F16 : VOPProfile <[f16, f16, f16, untyped]>;
 def VOP_F16_F16_I16 : VOPProfile <[f16, f16, i16, untyped]>;
diff --git a/llvm/lib/Target/AMDGPU/VOP1Instructions.td b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
index 80eb5d8b7d571..f621f8581f778 100644
--- a/llvm/lib/Target/AMDGPU/VOP1Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
@@ -803,6 +803,9 @@ let SubtargetPredicate = isGFX1250Plus in {
     def : Cvt_F_F8_Pat_ByteSel<int_amdgcn_cvt_f16_fp8, V_CVT_F16_FP8_fake16_e64, 1>;
     def : Cvt_F_F8_Pat_ByteSel<int_amdgcn_cvt_f16_bf8, V_CVT_F16_BF8_fake16_e64, 1>;
   }
+
+  defm V_SAT_PK4_I4_I8 : VOP1Inst_t16<"v_sat_pk4_i4_i8", VOP1_I16_I32, int_amdgcn_sat_pk4_i4_i8>;
+  defm V_SAT_PK4_U4_U8 : VOP1Inst_t16<"v_sat_pk4_u4_u8", VOP1_I16_I32, int_amdgcn_sat_pk4_u4_u8>;
 } // End SubtargetPredicate = isGFX1250Plus
 
 let SubtargetPredicate = isGFX10Plus in {
@@ -1158,6 +1161,8 @@ defm V_PERMLANE16_SWAP_B32   : VOP1_Real_OpSelIsDPP_gfx1250<0x049>;
 defm V_TANH_BF16             : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x04a>;
 defm V_PRNG_B32              : VOP1_Real_FULL<GFX1250Gen, 0x04b>;
 defm V_CVT_F32_BF16          : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x072, "v_cvt_f32_bf16", "V_CVT_F32_BF16_gfx1250">;
+defm V_SAT_PK4_I4_I8         : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x073>;
+defm V_SAT_PK4_U4_U8         : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x074>;
 defm V_CVT_PK_F16_FP8        : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x075>;
 defm V_CVT_PK_F16_BF8        : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x076>;
 defm V_CVT_F16_FP8           : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x077>;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll
new file mode 100644
index 0000000000000..3a5507063b834
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll
@@ -0,0 +1,305 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=+real-true16 %s -o - | FileCheck -check-prefix=SDAG-REAL16 %s
+; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=-real-true16 %s -o - | FileCheck -check-prefix=SDAG-FAKE16 %s
+; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=+real-true16 %s -o - | FileCheck -check-prefix=GISEL-REAL16 %s
+; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=-real-true16 %s -o - | FileCheck -check-prefix=GISEL-FAKE16 %s
+
+declare i16 @llvm.amdgcn.sat.pk4.i4.i8(i32) #0
+declare i16 @llvm.amdgcn.sat.pk4.u4.u8(i32) #0
+
+define amdgpu_kernel void @sat_pk4_i4_i8_f32_v(i32 %src, ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_i4_i8_f32_v:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_clause 0x1
+; SDAG-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, s2
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_i4_i8_f32_v:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_clause 0x1
+; SDAG-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v1, s2
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_i4_i8_f32_v:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_clause 0x1
+; GISEL-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, s2
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_i4_i8_f32_v:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_clause 0x1
+; GISEL-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v0, s2
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 %src) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_i4_i8_f32_s(i32 inreg %src, ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_i4_i8_f32_s:
+; SDAG-REAL16:       ; %bb.1:
+; SDAG-REAL16-NEXT:    s_load_b32 s8, s[4:5], 0x0
+; SDAG-REAL16-NEXT:    s_waitcnt lgkmcnt(0)
+; SDAG-REAL16-NEXT:    s_branch .LBB1_0
+; SDAG-REAL16-NEXT:    .p2align 8
+; SDAG-REAL16-NEXT:  ; %bb.2:
+; SDAG-REAL16-NEXT:  .LBB1_0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, s8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_i4_i8_f32_s:
+; SDAG-FAKE16:       ; %bb.1:
+; SDAG-FAKE16-NEXT:    s_load_b32 s8, s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    s_waitcnt lgkmcnt(0)
+; SDAG-FAKE16-NEXT:    s_branch .LBB1_0
+; SDAG-FAKE16-NEXT:    .p2align 8
+; SDAG-FAKE16-NEXT:  ; %bb.2:
+; SDAG-FAKE16-NEXT:  .LBB1_0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v1, s8
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_i4_i8_f32_s:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_clause 0x1
+; GISEL-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, s2
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_i4_i8_f32_s:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_clause 0x1
+; GISEL-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v0, s2
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 %src) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_i4_i8_f32_i(ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_i4_i8_f32_i:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, 0x64
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_i4_i8_f32_i:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v1, 0x64
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_i4_i8_f32_i:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_i4_i8_e32 v0.l, 0x64
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_i4_i8_f32_i:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_i4_i8_e32 v0, 0x64
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 100) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_u4_u8_f32_v(i32 %src, ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_u4_u8_f32_v:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_clause 0x1
+; SDAG-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, s2
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_u4_u8_f32_v:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_clause 0x1
+; SDAG-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v1, s2
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_u4_u8_f32_v:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_clause 0x1
+; GISEL-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, s2
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_u4_u8_f32_v:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_clause 0x1
+; GISEL-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v0, s2
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 %src) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_u4_u8_f32_s(i32 inreg %src, ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_u4_u8_f32_s:
+; SDAG-REAL16:       ; %bb.1:
+; SDAG-REAL16-NEXT:    s_load_b32 s8, s[4:5], 0x0
+; SDAG-REAL16-NEXT:    s_waitcnt lgkmcnt(0)
+; SDAG-REAL16-NEXT:    s_branch .LBB4_0
+; SDAG-REAL16-NEXT:    .p2align 8
+; SDAG-REAL16-NEXT:  ; %bb.2:
+; SDAG-REAL16-NEXT:  .LBB4_0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, s8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_u4_u8_f32_s:
+; SDAG-FAKE16:       ; %bb.1:
+; SDAG-FAKE16-NEXT:    s_load_b32 s8, s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    s_waitcnt lgkmcnt(0)
+; SDAG-FAKE16-NEXT:    s_branch .LBB4_0
+; SDAG-FAKE16-NEXT:    .p2align 8
+; SDAG-FAKE16-NEXT:  ; %bb.2:
+; SDAG-FAKE16-NEXT:  .LBB4_0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v1, s8
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_u4_u8_f32_s:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_clause 0x1
+; GISEL-REAL16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, s2
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_u4_u8_f32_s:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_clause 0x1
+; GISEL-FAKE16-NEXT:    s_load_b32 s2, s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x8
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v0, s2
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 %src) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @sat_pk4_u4_u8_f32_i(ptr %out) #1 {
+; SDAG-REAL16-LABEL: sat_pk4_u4_u8_f32_i:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, 0x64
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: sat_pk4_u4_u8_f32_i:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v0, 0
+; SDAG-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v1, 0x64
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    flat_store_b16 v0, v1, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+;
+; GISEL-REAL16-LABEL: sat_pk4_u4_u8_f32_i:
+; GISEL-REAL16:       ; %bb.0:
+; GISEL-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; GISEL-REAL16-NEXT:    v_sat_pk4_u4_u8_e32 v0.l, 0x64
+; GISEL-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-REAL16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-REAL16-NEXT:    s_endpgm
+;
+; GISEL-FAKE16-LABEL: sat_pk4_u4_u8_f32_i:
+; GISEL-FAKE16:       ; %bb.0:
+; GISEL-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; GISEL-FAKE16-NEXT:    v_sat_pk4_u4_u8_e32 v0, 0x64
+; GISEL-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; GISEL-FAKE16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; GISEL-FAKE16-NEXT:    s_endpgm
+  %cvt = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 100) #0
+  store i16 %cvt, ptr %out, align 2
+  ret void
+}
+
+attributes #0 = { nounwind memory(none) }
+attributes #1 = { nounwind }
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
index f2cf3d58fb0cf..811c6ebfe0161 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
@@ -628,6 +628,57 @@ v_cvt_f32_fp8_e32 v1, 3
 v_cvt_f32_fp8_e32 v1, v3
 // GFX1250: v_cvt_f32_fp8_e32 v1, v3                ; encoding: [0x03,0xd9,0x02,0x7e]
 
+v_cvt_pk_f32_bf8_e32 v[2:3], s3
+// GFX1250: v_cvt_pk_f32_bf8_e32 v[2:3], s3         ; encoding: [0x03,0xde,0x04,0x7e]
+
+v_cvt_pk_f32_bf8_e32 v[4:5], s5
+// GFX1250: v_cvt_pk_f32_bf8_e32 v[4:5], s5         ; encoding: [0x05,0xde,0x08,0x7e]
+
+v_cvt_pk_f32_bf8_e32 v[2:3], 3
+// GFX1250: v_cvt_pk_f32_bf8_e32 v[2:3], 3          ; encoding: [0x83,0xde,0x04,0x7e]
+
+v_cvt_pk_f32_bf8...
[truncated]

shiltian · 2025-07-18T16:46:40Z

Merge activity

Jul 18, 4:46 PM UTC: A user started a stack merge that includes this pull request via Graphite.
Jul 18, 5:05 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 18, 5:08 PM UTC: @shiltian merged this pull request with Graphite.

Co-authored-by: Mekhanoshin, Stanislav <Stanislav.Mekhanoshin@amd.com> Co-authored-by: Foad, Jay <Jay.Foad@amd.com>

llvm-ci · 2025-07-18T19:08:38Z

LLVM Buildbot has detected a new failure on builder amdgpu-offload-rhel-9-cmake-build-only running on rocm-docker-rhel-9 while building clang,llvm at step 2 "update-annotated-scripts".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/205/builds/16004

Here is the relevant piece of the build log for the reference

Step 2 (update-annotated-scripts) failure: update (failure)
git version 2.43.5
fatal: unable to access 'https://github.com/llvm/llvm-zorg.git/': Failed to connect to github.com port 443: Connection timed out
fatal: unable to access 'https://github.com/llvm/llvm-zorg.git/': Failed to connect to github.com port 443: Connection timed out

shiltian requested review from changpeng and rampitec July 18, 2025 15:07

shiltian mentioned this pull request Jul 18, 2025

[AMDGPU] Add support for v_permlane16_swap_b32 on gfx1250 #149518

Merged

llvmbot added clang Clang issues not falling into any other category backend:AMDGPU clang:frontend Language frontend issues, e.g. anything involving "Sema" llvm:mc Machine (object) code llvm:ir labels Jul 18, 2025

This was referenced Jul 18, 2025

[Clang][AMDGPU] Add the missing builtin __builtin_amdgcn_sqrt_bf16 #149447

Merged

[AMDGPU] Add support for v_prng_b32 on gfx1250 #149450

Merged

rampitec approved these changes Jul 18, 2025

View reviewed changes

shiltian force-pushed the users/shiltian/v_permlane16_swap_b32 branch 6 times, most recently from f7c11e6 to d58f960 Compare July 18, 2025 17:02

Base automatically changed from users/shiltian/v_permlane16_swap_b32 to main July 18, 2025 17:05

[AMDGPU] Add support for v_sat_pk4_i4_[i8,u8] on gfx1250

3fb65bc

Co-authored-by: Mekhanoshin, Stanislav <Stanislav.Mekhanoshin@amd.com> Co-authored-by: Foad, Jay <Jay.Foad@amd.com>

shiltian force-pushed the users/shiltian/v_sat_pk4_i4_i8 branch from 74b03a1 to 3fb65bc Compare July 18, 2025 17:05

shiltian merged commit 2c50e4c into main Jul 18, 2025
7 of 9 checks passed

shiltian deleted the users/shiltian/v_sat_pk4_i4_i8 branch July 18, 2025 17:08

This was referenced Jul 23, 2025

test abhinavgaba/llvm-project#2

Closed

Add dataFence plugin interface abhinavgaba/llvm-project#3

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[AMDGPU] Add support for `v_sat_pk4_i4_[i8,u8]` on gfx1250 #149528

[AMDGPU] Add support for `v_sat_pk4_i4_[i8,u8]` on gfx1250 #149528

Uh oh!

shiltian commented Jul 18, 2025

Uh oh!

shiltian commented Jul 18, 2025 •

edited

Loading

Uh oh!

llvmbot commented Jul 18, 2025 •

edited

Loading

Uh oh!

llvmbot commented Jul 18, 2025

Uh oh!

llvmbot commented Jul 18, 2025

Uh oh!

shiltian commented Jul 18, 2025 •

edited

Loading

Uh oh!

Uh oh!

llvm-ci commented Jul 18, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

[AMDGPU] Add support for v_sat_pk4_i4_[i8,u8] on gfx1250 #149528

[AMDGPU] Add support for v_sat_pk4_i4_[i8,u8] on gfx1250 #149528

Uh oh!

Conversation

shiltian commented Jul 18, 2025

Uh oh!

shiltian commented Jul 18, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Jul 18, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Jul 18, 2025

Uh oh!

llvmbot commented Jul 18, 2025

Uh oh!

shiltian commented Jul 18, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Merge activity

Uh oh!

Uh oh!

llvm-ci commented Jul 18, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

[AMDGPU] Add support for `v_sat_pk4_i4_[i8,u8]` on gfx1250 #149528

[AMDGPU] Add support for `v_sat_pk4_i4_[i8,u8]` on gfx1250 #149528

shiltian commented Jul 18, 2025 •

edited

Loading

llvmbot commented Jul 18, 2025 •

edited

Loading

shiltian commented Jul 18, 2025 •

edited

Loading