llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Mariusz Sikora (mariusz-sikora-at-amd) <details> <summary>Changes</summary> --- Patch is 75.67 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/75475.diff 19 Files Affected: - (modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+2) - (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl (+48) - (added) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx12-param.cl (+14) - (added) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx12.cl (+16) - (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+18) - (modified) llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp (+16-3) - (modified) llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp (+9) - (modified) llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td (+2) - (modified) llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp (+3-1) - (modified) llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp (+3-1) - (modified) llvm/lib/Target/AMDGPU/VOP3Instructions.td (+30) - (modified) llvm/test/Analysis/UniformityAnalysis/AMDGPU/intrinsics.ll (+16) - (added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane16.var.ll (+896) - (added) llvm/test/CodeGen/AMDGPU/permlane16_var-op-sel.ll (+15) - (added) llvm/test/CodeGen/AMDGPU/vcmpx-permlane16var-hazard.mir (+168) - (modified) llvm/test/MC/AMDGPU/gfx11_unsupported.s (+6) - (modified) llvm/test/MC/AMDGPU/gfx12_asm_vop3.s (+51) - (modified) llvm/test/MC/AMDGPU/gfx12_asm_vop3_err.s (+95) - (modified) llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3.txt (+51) ``````````diff diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 7465f13d552d6e..e562ef04a30194 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -410,6 +410,8 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32, "ifiiIi", "nc", "fp8-insts") // GFX12+ only builtins. //===----------------------------------------------------------------------===// +TARGET_BUILTIN(__builtin_amdgcn_permlane16_var, "UiUiUiUiIbIb", "nc", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_permlanex16_var, "UiUiUiUiIbIb", "nc", "gfx12-insts") TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal, "vIi", "n", "gfx12-insts") TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_var, "vi", "n", "gfx12-insts") TARGET_BUILTIN(__builtin_amdgcn_s_barrier_wait, "vIs", "n", "gfx12-insts") diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl index b8d281531e218e..2899d9e5c28898 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl @@ -1,6 +1,54 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py // RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -S -emit-llvm -o - %s | FileCheck %s +// REQUIRES: amdgpu-registered-target + +typedef unsigned int uint; + +// CHECK-LABEL: @test_permlane16_var( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5) +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: [[C_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8 +// CHECK-NEXT: store i32 [[A:%.*]], ptr addrspace(5) [[A_ADDR]], align 4 +// CHECK-NEXT: store i32 [[B:%.*]], ptr addrspace(5) [[B_ADDR]], align 4 +// CHECK-NEXT: store i32 [[C:%.*]], ptr addrspace(5) [[C_ADDR]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[A_ADDR]], align 4 +// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[B_ADDR]], align 4 +// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[C_ADDR]], align 4 +// CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.permlane16.var(i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]], i1 false, i1 false) +// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8 +// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(1) [[TMP4]], align 4 +// CHECK-NEXT: ret void +// +void test_permlane16_var(global uint* out, uint a, uint b, uint c) { + *out = __builtin_amdgcn_permlane16_var(a, b, c, 0, 0); +} + +// CHECK-LABEL: @test_permlanex16_var( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5) +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: [[C_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8 +// CHECK-NEXT: store i32 [[A:%.*]], ptr addrspace(5) [[A_ADDR]], align 4 +// CHECK-NEXT: store i32 [[B:%.*]], ptr addrspace(5) [[B_ADDR]], align 4 +// CHECK-NEXT: store i32 [[C:%.*]], ptr addrspace(5) [[C_ADDR]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[A_ADDR]], align 4 +// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[B_ADDR]], align 4 +// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[C_ADDR]], align 4 +// CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.permlanex16.var(i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]], i1 false, i1 false) +// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8 +// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(1) [[TMP4]], align 4 +// CHECK-NEXT: ret void +// +void test_permlanex16_var(global uint* out, uint a, uint b, uint c) { + *out = __builtin_amdgcn_permlanex16_var(a, b, c, 0, 0); +} + // CHECK-LABEL: @test_s_barrier_signal( // CHECK-NEXT: entry: // CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal(i32 -1) diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx12-param.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx12-param.cl new file mode 100644 index 00000000000000..0e0ea3646a2f1a --- /dev/null +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx12-param.cl @@ -0,0 +1,14 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1200 -verify -S -o - %s + +typedef unsigned int uint; + +void test_permlane16_var(global uint* out, uint a, uint b, uint c, uint d) { + *out = __builtin_amdgcn_permlane16_var(a, b, c, d, 1); // expected-error{{argument to '__builtin_amdgcn_permlane16_var' must be a constant integer}} + *out = __builtin_amdgcn_permlane16_var(a, b, c, 1, d); // expected-error{{argument to '__builtin_amdgcn_permlane16_var' must be a constant integer}} +} + +void test_permlanex16_var(global uint* out, uint a, uint b, uint c, uint d) { + *out = __builtin_amdgcn_permlanex16_var(a, b, c, d, 1); // expected-error{{argument to '__builtin_amdgcn_permlanex16_var' must be a constant integer}} + *out = __builtin_amdgcn_permlanex16_var(a, b, c, 1, d); // expected-error{{argument to '__builtin_amdgcn_permlanex16_var' must be a constant integer}} +} diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx12.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx12.cl new file mode 100644 index 00000000000000..34887a65021c33 --- /dev/null +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx12.cl @@ -0,0 +1,16 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu hawaii -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu fiji -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx900 -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx908 -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1010 -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1030 -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1100 -verify -S -o - %s + +typedef unsigned int uint; + +void test(global uint* out, uint a, uint b, uint c) { + *out = __builtin_amdgcn_permlane16_var(a, b, c, 1, 1); // expected-error {{'__builtin_amdgcn_permlane16_var' needs target feature gfx12-insts}} + *out = __builtin_amdgcn_permlanex16_var(a, b, c, 1, 1); // expected-error {{'__builtin_amdgcn_permlanex16_var' needs target feature gfx12-insts}} +} diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 09e88152e65d2a..cf054e89069d77 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -2460,6 +2460,24 @@ def int_amdgcn_s_wait_event_export_ready : Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn] >; +//===----------------------------------------------------------------------===// +// GFX12 Intrinsics +//===----------------------------------------------------------------------===// + +// llvm.amdgcn.permlane16.var <old> <src0> <src1> <fi> <bound_control> +def int_amdgcn_permlane16_var : ClangBuiltin<"__builtin_amdgcn_permlane16_var">, + Intrinsic<[llvm_i32_ty], + [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty], + [IntrNoMem, IntrConvergent, IntrWillReturn, + ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, IntrNoCallback, IntrNoFree]>; + +// llvm.amdgcn.permlanex16.var <old> <src0> <src1> <fi> <bound_control> +def int_amdgcn_permlanex16_var : ClangBuiltin<"__builtin_amdgcn_permlanex16_var">, + Intrinsic<[llvm_i32_ty], + [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty], + [IntrNoMem, IntrConvergent, IntrWillReturn, + ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, IntrNoCallback, IntrNoFree]>; + //===----------------------------------------------------------------------===// // Deep learning intrinsics. //===----------------------------------------------------------------------===// diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp index 5296415ab4c36d..ee93d9eb4c0a05 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp @@ -992,14 +992,27 @@ GCNTTIImpl::instCombineIntrinsic(InstCombiner &IC, IntrinsicInst &II) const { return IC.replaceOperand(II, 0, UndefValue::get(Old->getType())); } case Intrinsic::amdgcn_permlane16: - case Intrinsic::amdgcn_permlanex16: { + case Intrinsic::amdgcn_permlane16_var: + case Intrinsic::amdgcn_permlanex16: + case Intrinsic::amdgcn_permlanex16_var: { // Discard vdst_in if it's not going to be read. Value *VDstIn = II.getArgOperand(0); if (isa<UndefValue>(VDstIn)) break; - ConstantInt *FetchInvalid = cast<ConstantInt>(II.getArgOperand(4)); - ConstantInt *BoundCtrl = cast<ConstantInt>(II.getArgOperand(5)); + // FetchInvalid operand idx. + unsigned int FiIdx = (IID == Intrinsic::amdgcn_permlane16 || + IID == Intrinsic::amdgcn_permlanex16) + ? 4 /* for permlane16 and permlanex16 */ + : 3; /* for permlane16_var and permlanex16_var */ + + // BoundCtrl operand idx. + // For permlane16 and permlanex16 it should be 5 + // For Permlane16_var and permlanex16_var it should be 4 + unsigned int BcIdx = FiIdx + 1; + + ConstantInt *FetchInvalid = cast<ConstantInt>(II.getArgOperand(FiIdx)); + ConstantInt *BoundCtrl = cast<ConstantInt>(II.getArgOperand(BcIdx)); if (!FetchInvalid->getZExtValue() && !BoundCtrl->getZExtValue()) break; diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp index 03b6d19b2b3c06..269c065fdf7ab0 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp @@ -4435,6 +4435,15 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const { OpdsMapping[5] = getSGPROpMapping(MI.getOperand(4).getReg(), MRI, *TRI); break; } + case Intrinsic::amdgcn_permlane16_var: + case Intrinsic::amdgcn_permlanex16_var: { + unsigned Size = getSizeInBits(MI.getOperand(0).getReg(), MRI, *TRI); + OpdsMapping[0] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, Size); + OpdsMapping[2] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, Size); + OpdsMapping[3] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, Size); + OpdsMapping[4] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, Size); + break; + } case Intrinsic::amdgcn_mfma_f32_4x4x1f32: case Intrinsic::amdgcn_mfma_f32_4x4x4f16: case Intrinsic::amdgcn_mfma_i32_4x4x4i8: diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td b/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td index 317f3f21d24002..fc05f14744c0f9 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td +++ b/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td @@ -333,6 +333,8 @@ def : SourceOfDivergence<int_amdgcn_ds_ordered_add>; def : SourceOfDivergence<int_amdgcn_ds_ordered_swap>; def : SourceOfDivergence<int_amdgcn_permlane16>; def : SourceOfDivergence<int_amdgcn_permlanex16>; +def : SourceOfDivergence<int_amdgcn_permlane16_var>; +def : SourceOfDivergence<int_amdgcn_permlanex16_var>; def : SourceOfDivergence<int_amdgcn_mov_dpp>; def : SourceOfDivergence<int_amdgcn_mov_dpp8>; def : SourceOfDivergence<int_amdgcn_update_dpp>; diff --git a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp index 3af71727c5b741..a7d8ff0242b801 100644 --- a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp +++ b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp @@ -163,7 +163,9 @@ static bool isSendMsgTraceDataOrGDS(const SIInstrInfo &TII, static bool isPermlane(const MachineInstr &MI) { unsigned Opcode = MI.getOpcode(); return Opcode == AMDGPU::V_PERMLANE16_B32_e64 || - Opcode == AMDGPU::V_PERMLANEX16_B32_e64; + Opcode == AMDGPU::V_PERMLANEX16_B32_e64 || + Opcode == AMDGPU::V_PERMLANE16_VAR_B32_e64 || + Opcode == AMDGPU::V_PERMLANEX16_VAR_B32_e64; } static bool isLdsDma(const MachineInstr &MI) { diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp index 4edd7960bd8c40..0f92a56237acb3 100644 --- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp +++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp @@ -498,7 +498,9 @@ bool isPermlane16(unsigned Opc) { Opc == AMDGPU::V_PERMLANE16_B32_e64_gfx11 || Opc == AMDGPU::V_PERMLANEX16_B32_e64_gfx11 || Opc == AMDGPU::V_PERMLANE16_B32_e64_gfx12 || - Opc == AMDGPU::V_PERMLANEX16_B32_e64_gfx12; + Opc == AMDGPU::V_PERMLANEX16_B32_e64_gfx12 || + Opc == AMDGPU::V_PERMLANE16_VAR_B32_e64_gfx12 || + Opc == AMDGPU::V_PERMLANEX16_VAR_B32_e64_gfx12; } bool isGenericAtomic(unsigned Opc) { diff --git a/llvm/lib/Target/AMDGPU/VOP3Instructions.td b/llvm/lib/Target/AMDGPU/VOP3Instructions.td index 685c9ac6a2be40..2733f1d5634d80 100644 --- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td +++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td @@ -735,6 +735,15 @@ def VOP3_PERMLANE_Profile : VOP3_Profile<VOPProfile <[i32, i32, i32, i32]>, VOP3 let HasExtDPP = 0; } +def VOP3_PERMLANE_VAR_Profile : VOP3_Profile<VOPProfile <[i32, i32, i32, untyped]>, VOP3_OPSEL> { + let InsVOP3OpSel = (ins IntOpSelMods:$src0_modifiers, VRegSrc_32:$src0, + IntOpSelMods:$src1_modifiers, VRegSrc_32:$src1, + VGPR_32:$vdst_in, op_sel0:$op_sel); + let HasClamp = 0; + let HasExtVOP3DPP = 0; + let HasExtDPP = 0; +} + def opsel_i1timm : SDNodeXForm<timm, [{ return CurDAG->getTargetConstant( N->getZExtValue() ? SISrcMods::OP_SEL_0 : SISrcMods::NONE, @@ -751,6 +760,13 @@ class PermlanePat<SDPatternOperator permlane, SCSrc_b32:$src1, 0, SCSrc_b32:$src2, VGPR_32:$vdst_in) >; +class PermlaneVarPat<SDPatternOperator permlane, + Instruction inst> : GCNPat< + (permlane i32:$vdst_in, i32:$src0, i32:$src1, + timm:$fi, timm:$bc), + (inst (opsel_i1timm $fi), VGPR_32:$src0, (opsel_i1timm $bc), + VGPR_32:$src1, VGPR_32:$vdst_in) +>; let SubtargetPredicate = isGFX10Plus in { let isCommutable = 1, isReMaterializable = 1 in { @@ -781,6 +797,17 @@ let SubtargetPredicate = isGFX10Plus in { } // End SubtargetPredicate = isGFX10Plus +let SubtargetPredicate = isGFX12Plus in { + let Constraints = "$vdst = $vdst_in", DisableEncoding="$vdst_in" in { + defm V_PERMLANE16_VAR_B32 : VOP3Inst<"v_permlane16_var_b32", VOP3_PERMLANE_VAR_Profile>; + defm V_PERMLANEX16_VAR_B32 : VOP3Inst<"v_permlanex16_var_b32", VOP3_PERMLANE_VAR_Profile>; + } // End $vdst = $vdst_in, DisableEncoding $vdst_in + + def : PermlaneVarPat<int_amdgcn_permlane16_var, V_PERMLANE16_VAR_B32_e64>; + def : PermlaneVarPat<int_amdgcn_permlanex16_var, V_PERMLANEX16_VAR_B32_e64>; + +} // End SubtargetPredicate = isGFX12Plus + class DivFmasPat<ValueType vt, Instruction inst, Register CondReg> : GCNPat< (AMDGPUdiv_fmas (vt (VOP3Mods vt:$src0, i32:$src0_modifiers)), (vt (VOP3Mods vt:$src1, i32:$src1_modifiers)), @@ -915,6 +942,9 @@ defm V_MAXIMUM_F32 : VOP3Only_Realtriple_gfx12<0x366>; defm V_MINIMUM_F16 : VOP3Only_Realtriple_t16_gfx12<0x367>; defm V_MAXIMUM_F16 : VOP3Only_Realtriple_t16_gfx12<0x368>; +defm V_PERMLANE16_VAR_B32 : VOP3Only_Real_Base_gfx12<0x30f>; +defm V_PERMLANEX16_VAR_B32 : VOP3Only_Real_Base_gfx12<0x310>; + //===----------------------------------------------------------------------===// // GFX11, GFX12 //===----------------------------------------------------------------------===// diff --git a/llvm/test/Analysis/UniformityAnalysis/AMDGPU/intrinsics.ll b/llvm/test/Analysis/UniformityAnalysis/AMDGPU/intrinsics.ll index 1b3f7973c0d9e6..8826263eabb693 100644 --- a/llvm/test/Analysis/UniformityAnalysis/AMDGPU/intrinsics.ll +++ b/llvm/test/Analysis/UniformityAnalysis/AMDGPU/intrinsics.ll @@ -21,6 +21,20 @@ define amdgpu_kernel void @v_permlanex16_b32(ptr addrspace(1) %out, i32 %src0, i ret void } +; CHECK: DIVERGENT: %v = call i32 @llvm.amdgcn.permlane16.var(i32 %src0, i32 %src1, i32 %src2, i1 false, i1 false) #0 +define amdgpu_kernel void @v_permlane16_var_b32(ptr addrspace(1) %out, i32 %src0, i32 %src1, i32 %src2) #0 { + %v = call i32 @llvm.amdgcn.permlane16.var(i32 %src0, i32 %src1, i32 %src2, i1 false, i1 false) #0 + store i32 %v, ptr addrspace(1) %out + ret void +} + +; CHECK: DIVERGENT: %v = call i32 @llvm.amdgcn.permlanex16.var(i32 %src0, i32 %src1, i32 %src2, i1 false, i1 false) #0 +define amdgpu_kernel void @v_permlanex16_var_b32(ptr addrspace(1) %out, i32 %src0, i32 %src1, i32 %src2) #0 { + %v = call i32 @llvm.amdgcn.permlanex16.var(i32 %src0, i32 %src1, i32 %src2, i1 false, i1 false) #0 + store i32 %v, ptr addrspace(1) %out + ret void +} + ; CHECK: DIVERGENT: %tmp0 = call i32 @llvm.amdgcn.update.dpp.i32(i32 %in1, i32 %in2, i32 1, i32 1, i32 1, i1 false) #0 define amdgpu_kernel void @update_dpp(ptr addrspace(1) %out, i32 %in1, i32 %in2) #0 { %tmp0 = call i32 @llvm.amdgcn.update.dpp.i32(i32 %in1, i32 %in2, i32 1, i32 1, i32 1, i1 false) #0 @@ -98,6 +112,8 @@ bb: declare i32 @llvm.amdgcn.ds.swizzle(i32, i32) #1 declare i32 @llvm.amdgcn.permlane16(i32, i32, i32, i32, i1, i1) #1 declare i32 @llvm.amdgcn.permlanex16(i32, i32, i32, i32, i1, i1) #1 +declare i32 @llvm.amdgcn.permlane16.var(i32, i32, i32, i1, i1) #1 +declare i32 @llvm.amdgcn.permlanex16.var(i32, i32, i32, i1, i1) #1 declare i32 @llvm.amdgcn.mov.dpp.i32(i32, i32, i32, i32, i1) #1 declare i32 @llvm.amdgcn.mov.dpp8.i32(i32, i32) #1 declare i32 @llvm.amdgcn.update.dpp.i32(i32, i32, i32, i32, i32, i1) #1 diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane16.var.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane16.var.ll new file mode 100644 index 00000000000000..131a3951b2bf27 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane16.var.ll @@ -0,0 +1,896 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py +; RUN: llc -global-isel=0 -amdgpu-load-store-vectorizer=0 -march=amdgcn -mcpu=gfx1200 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX12-SDAG %s +; RUN: llc -global-isel=1 -amdgpu-load-store-vectorizer=0 -march=amdgcn -mcpu=gfx1200 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX12-GISEL %s + +declare i32 @llvm.amdgcn.permlane16.var(i32, i32, i32, i1, i1) +declare i32 @llvm.amdgcn.permlanex16.var(i32, i32, i32, i1, i1) +declare i32 @llvm.amdgcn.workitem.id.x() +declare i32 @llvm.amdgcn.workitem.id.y() + +define amdgpu_kernel void @v_permlane16var_b32_vv(ptr addrspace(1) %out, i32 %src0, i32 %src1) { +; GFX12-SDAG-LABEL: v_permlane16var_b32_vv: +; GFX12-SDAG: ; %bb.0: +; GFX12-SDAG-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-SDAG-NEXT: v_mov_b32_e32 v2, 0 +; GFX12-SDAG-NEXT: s_waitcnt lgkmcnt(0) +; GFX12-SDAG-NEXT: v_dual_mov_b32 v0, s2 :: v_dual_mov_b32 v1, s3 +; GFX12-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) +; GFX12-SDAG-NEXT: v_permlane16_var_b32 v0, v0, v1 +; GFX12-SDAG-NEXT: global_store_b32 v2, v0, s[0:1] +; GFX12-SDAG-NEXT: s_nop 0 +; GFX12-SDAG-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-SDAG-NEXT: s_endpgm +; +; GFX12-GISEL-LABEL: v_permlane16var_b32_vv: +; GFX12-GISEL: ; %bb.0: +; GFX12-GISEL-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-GISEL-NEXT: s_waitcnt lgkmcnt(0) +; GFX12-GISEL-NEXT: v_dual_mov_b32 v0, s2 :: v_dual_mov_b32 v1, s3 +; GFX12-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) +; GFX12-GISEL-NEXT: v_permlane16_var_b32 v0, v0, v1 +; GFX12-GISEL-NEXT: v_mov_b32_e32 v1, 0 +; GFX12-GISEL-NEXT: global_store_b32 v1, v0, s[0:1] +; GFX12-GISEL-NEXT: s_nop 0 +; GFX12-GISEL-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; ... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/75475 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
