diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index a0f365e04b217..d208342d9c516 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -428,6 +428,13 @@ TARGET_BUILTIN(__builtin_amdgcn_s_wakeup_barrier, "vi", "n", "gfx12-insts") TARGET_BUILTIN(__builtin_amdgcn_s_barrier_leave, "b", "n", "gfx12-insts") TARGET_BUILTIN(__builtin_amdgcn_s_get_barrier_state, "Uii", "n", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_v2i32, "V2iV2i*1", "nc", "gfx12-insts,wavefrontsize32") +TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_v8i16, "V8sV8s*1", "nc", "gfx12-insts,wavefrontsize32") +TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_v8f16, "V8hV8h*1", "nc", "gfx12-insts,wavefrontsize32") + +TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_i32, "ii*1", "nc", "gfx12-insts,wavefrontsize64") +TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_v4i16, "V4sV4s*1", "nc", "gfx12-insts,wavefrontsize64") +TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_v4f16, "V4hV4h*1", "nc", "gfx12-insts,wavefrontsize64") #undef BUILTIN #undef TARGET_BUILTIN diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index f4246c5e8f68e..cb65c7373a34c 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -18178,6 +18178,45 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, llvm::Function *F = CGM.getIntrinsic(IID, {ArgTy}); return Builder.CreateCall(F, {Addr, Val, ZeroI32, ZeroI32, ZeroI1}); } + case AMDGPU::BI__builtin_amdgcn_global_load_tr_i32: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v2i32: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v4f16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v4i16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v8f16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v8i16: { + + llvm::Type *ArgTy; + switch (BuiltinID) { + case AMDGPU::BI__builtin_amdgcn_global_load_tr_i32: + ArgTy = llvm::Type::getInt32Ty(getLLVMContext()); + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v2i32: + ArgTy = llvm::FixedVectorType::get( + llvm::Type::getInt32Ty(getLLVMContext()), 2); + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v4f16: + ArgTy = llvm::FixedVectorType::get( + llvm::Type::getHalfTy(getLLVMContext()), 4); + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v4i16: + ArgTy = llvm::FixedVectorType::get( + llvm::Type::getInt16Ty(getLLVMContext()), 4); + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v8f16: + ArgTy = llvm::FixedVectorType::get( + llvm::Type::getHalfTy(getLLVMContext()), 8); + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v8i16: + ArgTy = llvm::FixedVectorType::get( + llvm::Type::getInt16Ty(getLLVMContext()), 8); + break; + } + + llvm::Value *Addr = EmitScalarExpr(E->getArg(0)); + llvm::Function *F = + CGM.getIntrinsic(Intrinsic::amdgcn_global_load_tr, {ArgTy}); + return Builder.CreateCall(F, {Addr}); + } case AMDGPU::BI__builtin_amdgcn_read_exec: return EmitAMDGCNBallotForExec(*this, E, Int64Ty, Int64Ty, false); case AMDGPU::BI__builtin_amdgcn_read_exec_lo: diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx11-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx11-err.cl new file mode 100644 index 0000000000000..f7afb7cb97eda --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx11-err.cl @@ -0,0 +1,26 @@ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1100 -emit-llvm \ +// RUN: -verify -S -o - %s + +// REQUIRES: amdgpu-registered-target + +typedef int v2i __attribute__((ext_vector_type(2))); +typedef half v8h __attribute__((ext_vector_type(8))); +typedef short v8s __attribute__((ext_vector_type(8))); + +typedef half v4h __attribute__((ext_vector_type(4))); +typedef short v4s __attribute__((ext_vector_type(4))); + + + +void amdgcn_global_load_tr(global v2i* v2i_inptr, global v8s* v8s_inptr, global v8h* v8h_inptr, + global int* int_inptr, global v4s* v4s_inptr, global v4h* v4h_inptr) +{ + v2i out_1 = __builtin_amdgcn_global_load_tr_v2i32(v2i_inptr); // expected-error{{'__builtin_amdgcn_global_load_tr_v2i32' needs target feature gfx12-insts,wavefrontsize32}} + v8s out_2 = __builtin_amdgcn_global_load_tr_v8i16(v8s_inptr); // expected-error{{'__builtin_amdgcn_global_load_tr_v8i16' needs target feature gfx12-insts,wavefrontsize32}} + v8h out_3 = __builtin_amdgcn_global_load_tr_v8f16(v8h_inptr); // expected-error{{'__builtin_amdgcn_global_load_tr_v8f16' needs target feature gfx12-insts,wavefrontsize32}} + + int out_4 = __builtin_amdgcn_global_load_tr_i32(int_inptr); // expected-error{{'__builtin_amdgcn_global_load_tr_i32' needs target feature gfx12-insts,wavefrontsize64}} + v4s out_5 = __builtin_amdgcn_global_load_tr_v4i16(v4s_inptr); // expected-error{{'__builtin_amdgcn_global_load_tr_v4i16' needs target feature gfx12-insts,wavefrontsize64}} + v4h out_6 = __builtin_amdgcn_global_load_tr_v4f16(v4h_inptr); // expected-error{{'__builtin_amdgcn_global_load_tr_v4f16' needs target feature gfx12-insts,wavefrontsize64}} +} + diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w32-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w32-err.cl new file mode 100644 index 0000000000000..04ac0a66db7ce --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w32-err.cl @@ -0,0 +1,15 @@ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1200 -target-feature +wavefrontsize32 -emit-llvm \ +// RUN: -verify -S -o - %s + +// REQUIRES: amdgpu-registered-target + +typedef half v4h __attribute__((ext_vector_type(4))); +typedef short v4s __attribute__((ext_vector_type(4))); + +void amdgcn_global_load_tr(global int* int_inptr, global v4s* v4s_inptr, global v4h* v4h_inptr) +{ + int out_4 = __builtin_amdgcn_global_load_tr_i32(int_inptr); // expected-error{{'__builtin_amdgcn_global_load_tr_i32' needs target feature gfx12-insts,wavefrontsize64}} + v4s out_5 = __builtin_amdgcn_global_load_tr_v4i16(v4s_inptr); // expected-error{{'__builtin_amdgcn_global_load_tr_v4i16' needs target feature gfx12-insts,wavefrontsize64}} + v4h out_6 = __builtin_amdgcn_global_load_tr_v4f16(v4h_inptr); // expected-error{{'__builtin_amdgcn_global_load_tr_v4f16' needs target feature gfx12-insts,wavefrontsize64}} +} + diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w64-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w64-err.cl new file mode 100644 index 0000000000000..113b54b853a9f --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w64-err.cl @@ -0,0 +1,16 @@ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1200 -target-feature +wavefrontsize64 -emit-llvm \ +// RUN: -verify -S -o - %s + +// REQUIRES: amdgpu-registered-target + +typedef int v2i __attribute__((ext_vector_type(2))); +typedef half v8h __attribute__((ext_vector_type(8))); +typedef short v8s __attribute__((ext_vector_type(8))); + +void amdgcn_global_load_tr(global v2i* v2i_inptr, global v8s* v8s_inptr, global v8h* v8h_inptr) +{ + v2i out_1 = __builtin_amdgcn_global_load_tr_v2i32(v2i_inptr); // expected-error{{'__builtin_amdgcn_global_load_tr_v2i32' needs target feature gfx12-insts,wavefrontsize32}} + v8s out_2 = __builtin_amdgcn_global_load_tr_v8i16(v8s_inptr); // expected-error{{'__builtin_amdgcn_global_load_tr_v8i16' needs target feature gfx12-insts,wavefrontsize32}} + v8h out_3 = __builtin_amdgcn_global_load_tr_v8f16(v8h_inptr); // expected-error{{'__builtin_amdgcn_global_load_tr_v8f16' needs target feature gfx12-insts,wavefrontsize32}} +} + diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-w32.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-w32.cl new file mode 100644 index 0000000000000..b5fcad68a4702 --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-w32.cl @@ -0,0 +1,48 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -target-feature +wavefrontsize32 -S -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1200 + +typedef int v2i __attribute__((ext_vector_type(2))); +typedef half v8h __attribute__((ext_vector_type(8))); +typedef short v8s __attribute__((ext_vector_type(8))); + +// Wave32 + +// +// amdgcn_global_load_tr +// + +// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_v2i32( +// CHECK-GFX1200-NEXT: entry: +// CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.global.load.tr.v2i32(ptr addrspace(1) [[INPTR:%.*]]) +// CHECK-GFX1200-NEXT: ret <2 x i32> [[TMP0]] +// +v2i test_amdgcn_global_load_tr_v2i32(global v2i* inptr) +{ + return __builtin_amdgcn_global_load_tr_v2i32(inptr); +} + +// +// amdgcn_global_load_tr +// + +// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_v8i16( +// CHECK-GFX1200-NEXT: entry: +// CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call <8 x i16> @llvm.amdgcn.global.load.tr.v8i16(ptr addrspace(1) [[INPTR:%.*]]) +// CHECK-GFX1200-NEXT: ret <8 x i16> [[TMP0]] +// +v8s test_amdgcn_global_load_tr_v8i16(global v8s* inptr) +{ + return __builtin_amdgcn_global_load_tr_v8i16(inptr); +} + +// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_v8f16( +// CHECK-GFX1200-NEXT: entry: +// CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call <8 x half> @llvm.amdgcn.global.load.tr.v8f16(ptr addrspace(1) [[INPTR:%.*]]) +// CHECK-GFX1200-NEXT: ret <8 x half> [[TMP0]] +// +v8h test_amdgcn_global_load_tr_v8f16(global v8h* inptr) +{ + return __builtin_amdgcn_global_load_tr_v8f16(inptr); +} + diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-w64.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-w64.cl new file mode 100644 index 0000000000000..9c48ac071b4d3 --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-w64.cl @@ -0,0 +1,47 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -target-feature +wavefrontsize64 -S -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1200 + +typedef half v4h __attribute__((ext_vector_type(4))); +typedef short v4s __attribute__((ext_vector_type(4))); + +// Wave64 + +// +// amdgcn_global_load_tr +// + +// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_i32( +// CHECK-GFX1200-NEXT: entry: +// CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.tr.i32(ptr addrspace(1) [[INPTR:%.*]]) +// CHECK-GFX1200-NEXT: ret i32 [[TMP0]] +// +int test_amdgcn_global_load_tr_i32(global int* inptr) +{ + return __builtin_amdgcn_global_load_tr_i32(inptr); +} + +// +// amdgcn_global_load_tr +// + +// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_v4i16( +// CHECK-GFX1200-NEXT: entry: +// CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call <4 x i16> @llvm.amdgcn.global.load.tr.v4i16(ptr addrspace(1) [[INPTR:%.*]]) +// CHECK-GFX1200-NEXT: ret <4 x i16> [[TMP0]] +// +v4s test_amdgcn_global_load_tr_v4i16(global v4s* inptr) +{ + return __builtin_amdgcn_global_load_tr_v4i16(inptr); +} + +// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_v4f16( +// CHECK-GFX1200-NEXT: entry: +// CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call <4 x half> @llvm.amdgcn.global.load.tr.v4f16(ptr addrspace(1) [[INPTR:%.*]]) +// CHECK-GFX1200-NEXT: ret <4 x half> [[TMP0]] +// +v4h test_amdgcn_global_load_tr_v4f16(global v4h* inptr) +{ + return __builtin_amdgcn_global_load_tr_v4f16(inptr); +} + diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 61c4463f66cff..9302e590a6fc9 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -2577,6 +2577,28 @@ def int_amdgcn_global_atomic_fmax_num : AMDGPUAtomicRtn; def int_amdgcn_atomic_cond_sub_u32 : AMDGPUAtomicRtn; +class AMDGPULoadTr: + Intrinsic< + [llvm_any_ty], + [ptr_ty], + [IntrReadMem, IntrWillReturn, IntrConvergent, NoCapture>, IntrNoCallback, IntrNoFree], + "", + [SDNPMemOperand] + >; + +// Wave32 +// <2 x i32> @llvm.amdgcn.global.load.tr.v2i32(ptr addrspace(1)) -> global_load_tr_b64 +// <8 x i16> @llvm.amdgcn.global.load.tr.v8i16(ptr addrspace(1)) -> global_load_tr_b128 +// <8 x half> @llvm.amdgcn.global.load.tr.v8f16(ptr addrspace(1)) -> global_load_tr_b128 +// <8 x bfloat> @llvm.amdgcn.global.load.tr.v8bf16(ptr addrspace(1)) -> global_load_tr_b128 +// Wave64 +// i32 @llvm.amdgcn.global.load.tr.i32(ptr addrspace(1)) -> global_load_tr_b64 +// <4 x i16> @llvm.amdgcn.global.load.tr.v4i16(ptr addrspace(1)) -> global_load_tr_b128 +// <4 x half> @llvm.amdgcn.global.load.tr.v4f16(ptr addrspace(1)) -> global_load_tr_b128 +// <4 x bfloat> @llvm.amdgcn.global.load.tr.v4bf16(ptr addrspace(1)) -> global_load_tr_b128 + +def int_amdgcn_global_load_tr : AMDGPULoadTr; + //===----------------------------------------------------------------------===// // Deep learning intrinsics. //===----------------------------------------------------------------------===// diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp index 216d96fced568..bdd4e891f1589 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp @@ -4864,6 +4864,7 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const { case Intrinsic::amdgcn_flat_atomic_fadd_v2bf16: case Intrinsic::amdgcn_atomic_cond_sub_u32: case Intrinsic::amdgcn_global_atomic_ordered_add_b64: + case Intrinsic::amdgcn_global_load_tr: return getDefaultMappingAllVGPR(MI); case Intrinsic::amdgcn_ds_ordered_add: case Intrinsic::amdgcn_ds_ordered_swap: diff --git a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp index 9dff3f6c2efd0..441032a37dfd9 100644 --- a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp +++ b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp @@ -544,6 +544,10 @@ DecodeStatus AMDGPUDisassembler::getInstruction(MCInst &MI, uint64_t &Size, Res = tryDecodeInst(DecoderTableGFX1296, MI, DecW, Address, CS); if (Res) break; + + Res = tryDecodeInst(DecoderTableGFX12W6496, MI, DecW, Address, CS); + if (Res) + break; } // Reinitialize Bytes Bytes = Bytes_.slice(0, MaxInstBytesNum); diff --git a/llvm/lib/Target/AMDGPU/FLATInstructions.td b/llvm/lib/Target/AMDGPU/FLATInstructions.td index edd17a4cadb3b..688a440e2c396 100644 --- a/llvm/lib/Target/AMDGPU/FLATInstructions.td +++ b/llvm/lib/Target/AMDGPU/FLATInstructions.td @@ -992,6 +992,17 @@ defm SCRATCH_LOAD_LDS_DWORD : FLAT_Scratch_Load_LDS_Pseudo <"scratch_load_lds_d } // End SubtargetPredicate = HasFlatScratchInsts +let SubtargetPredicate = isGFX12Plus in { + let WaveSizePredicate = isWave32 in { + defm GLOBAL_LOAD_TR_B128_w32 : FLAT_Global_Load_Pseudo <"global_load_tr_b128_w32", VReg_128>; + defm GLOBAL_LOAD_TR_B64_w32 : FLAT_Global_Load_Pseudo <"global_load_tr_b64_w32", VReg_64>; + } + let WaveSizePredicate = isWave64 in { + defm GLOBAL_LOAD_TR_B128_w64 : FLAT_Global_Load_Pseudo <"global_load_tr_b128_w64", VReg_64>; + defm GLOBAL_LOAD_TR_B64_w64 : FLAT_Global_Load_Pseudo <"global_load_tr_b64_w64", VGPR_32>; + } +} // End SubtargetPredicate = isGFX12Plus + let SubtargetPredicate = isGFX10Plus, is_flat_global = 1 in { defm GLOBAL_ATOMIC_FCMPSWAP : FLAT_Global_Atomic_Pseudo<"global_atomic_fcmpswap", VGPR_32, f32, v2f32, VReg_64>; @@ -1590,6 +1601,19 @@ let SubtargetPredicate = isGFX12Plus in { let OtherPredicates = [isGFX12Plus] in { defm : GlobalFLATAtomicPatsRtn <"GLOBAL_ATOMIC_ORDERED_ADD_B64", "int_amdgcn_global_atomic_ordered_add_b64", i64, i64, /* isIntr */ 1>; + + let WaveSizePredicate = isWave32 in { + defm : GlobalFLATLoadPats ; + defm : GlobalFLATLoadPats ; + defm : GlobalFLATLoadPats ; + defm : GlobalFLATLoadPats ; + } + let WaveSizePredicate = isWave64 in { + defm : GlobalFLATLoadPats ; + defm : GlobalFLATLoadPats ; + defm : GlobalFLATLoadPats ; + defm : GlobalFLATLoadPats ; + } } let OtherPredicates = [isGFX10Plus] in { @@ -2726,6 +2750,17 @@ defm GLOBAL_ATOMIC_COND_SUB_U32 : VGLOBAL_Real_Atomics_gfx12<0x050, "GLOBAL_A defm GLOBAL_ATOMIC_MIN_NUM_F32 : VGLOBAL_Real_Atomics_gfx12<0x051, "GLOBAL_ATOMIC_FMIN", "global_atomic_min_num_f32", true, "global_atomic_min_f32">; defm GLOBAL_ATOMIC_MAX_NUM_F32 : VGLOBAL_Real_Atomics_gfx12<0x052, "GLOBAL_ATOMIC_FMAX", "global_atomic_max_num_f32", true, "global_atomic_max_f32">; defm GLOBAL_ATOMIC_ADD_F32 : VGLOBAL_Real_Atomics_gfx12<0x056>; + +let WaveSizePredicate = isWave32, DecoderNamespace = "GFX12" in { + defm GLOBAL_LOAD_TR_B128_w32 : VGLOBAL_Real_AllAddr_gfx12<0x057, "GLOBAL_LOAD_TR_B128_w32", "global_load_tr_b128">; + defm GLOBAL_LOAD_TR_B64_w32 : VGLOBAL_Real_AllAddr_gfx12<0x058, "GLOBAL_LOAD_TR_B64_w32", "global_load_tr_b64">; +} + +let WaveSizePredicate = isWave64, DecoderNamespace = "GFX12W64" in { + defm GLOBAL_LOAD_TR_B128_w64 : VGLOBAL_Real_AllAddr_gfx12<0x057, "GLOBAL_LOAD_TR_B128_w64", "global_load_tr_b128">; + defm GLOBAL_LOAD_TR_B64_w64 : VGLOBAL_Real_AllAddr_gfx12<0x058, "GLOBAL_LOAD_TR_B64_w64", "global_load_tr_b64">; +} + defm GLOBAL_ATOMIC_ORDERED_ADD_B64 : VGLOBAL_Real_Atomics_gfx12<0x073>; defm GLOBAL_ATOMIC_PK_ADD_F16 : VGLOBAL_Real_Atomics_gfx12<0x059, "GLOBAL_ATOMIC_PK_ADD_F16", "global_atomic_pk_add_f16">; defm GLOBAL_ATOMIC_PK_ADD_BF16 : VGLOBAL_Real_Atomics_gfx12<0x05a, "GLOBAL_ATOMIC_PK_ADD_BF16", "global_atomic_pk_add_bf16">; diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.global.load.tr-w32.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.global.load.tr-w32.ll new file mode 100644 index 0000000000000..5382b56b92fb1 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.global.load.tr-w32.ll @@ -0,0 +1,140 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py +; RUN: llc -global-isel=0 -march=amdgcn -mcpu=gfx1200 -verify-machineinstrs -mattr=+wavefrontsize32,-wavefrontsize64 < %s | FileCheck -check-prefixes=GFX12-SDAG-W32 %s +; RUN: llc -global-isel=1 -march=amdgcn -mcpu=gfx1200 -verify-machineinstrs -mattr=+wavefrontsize32,-wavefrontsize64 < %s | FileCheck -check-prefixes=GFX12-GISEL-W32 %s + +declare <2 x i32> @llvm.amdgcn.global.load.tr.v2i32.p1(ptr addrspace(1)) +declare <8 x i16> @llvm.amdgcn.global.load.tr.v8i16.p1(ptr addrspace(1)) +declare <8 x half> @llvm.amdgcn.global.load.tr.v8f16.p1(ptr addrspace(1)) +declare <8 x bfloat> @llvm.amdgcn.global.load.tr.v8bf16.p1(ptr addrspace(1)) + +define amdgpu_kernel void @global_load_tr_b64(ptr addrspace(1) %addr, ptr addrspace(1) %use) { +; GFX12-SDAG-W32-LABEL: global_load_tr_b64: +; GFX12-SDAG-W32: ; %bb.0: ; %entry +; GFX12-SDAG-W32-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-SDAG-W32-NEXT: v_mov_b32_e32 v2, 0 +; GFX12-SDAG-W32-NEXT: s_wait_kmcnt 0x0 +; GFX12-SDAG-W32-NEXT: global_load_tr_b64 v[0:1], v2, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-SDAG-W32-NEXT: s_wait_loadcnt 0x0 +; GFX12-SDAG-W32-NEXT: global_inv scope:SCOPE_SYS +; GFX12-SDAG-W32-NEXT: global_store_b64 v2, v[0:1], s[2:3] +; GFX12-SDAG-W32-NEXT: s_nop 0 +; GFX12-SDAG-W32-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-SDAG-W32-NEXT: s_endpgm +; +; GFX12-GISEL-W32-LABEL: global_load_tr_b64: +; GFX12-GISEL-W32: ; %bb.0: ; %entry +; GFX12-GISEL-W32-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-GISEL-W32-NEXT: v_mov_b32_e32 v2, 0 +; GFX12-GISEL-W32-NEXT: s_wait_kmcnt 0x0 +; GFX12-GISEL-W32-NEXT: global_load_tr_b64 v[0:1], v2, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-GISEL-W32-NEXT: s_wait_loadcnt 0x0 +; GFX12-GISEL-W32-NEXT: global_inv scope:SCOPE_SYS +; GFX12-GISEL-W32-NEXT: global_store_b64 v2, v[0:1], s[2:3] +; GFX12-GISEL-W32-NEXT: s_nop 0 +; GFX12-GISEL-W32-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-GISEL-W32-NEXT: s_endpgm +entry: + %gep = getelementptr i64, ptr addrspace(1) %addr, i32 4 + %val = call <2 x i32> @llvm.amdgcn.global.load.tr.v2i32.p1(ptr addrspace(1) %gep) + store <2 x i32> %val, ptr addrspace(1) %use + ret void +} + +define amdgpu_kernel void @global_load_tr_b128_i16(ptr addrspace(1) %addr, ptr addrspace(1) %use) { +; GFX12-SDAG-W32-LABEL: global_load_tr_b128_i16: +; GFX12-SDAG-W32: ; %bb.0: ; %entry +; GFX12-SDAG-W32-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-SDAG-W32-NEXT: v_mov_b32_e32 v4, 0 +; GFX12-SDAG-W32-NEXT: s_wait_kmcnt 0x0 +; GFX12-SDAG-W32-NEXT: global_load_tr_b128 v[0:3], v4, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-SDAG-W32-NEXT: s_wait_loadcnt 0x0 +; GFX12-SDAG-W32-NEXT: global_inv scope:SCOPE_SYS +; GFX12-SDAG-W32-NEXT: global_store_b128 v4, v[0:3], s[2:3] +; GFX12-SDAG-W32-NEXT: s_nop 0 +; GFX12-SDAG-W32-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-SDAG-W32-NEXT: s_endpgm +; +; GFX12-GISEL-W32-LABEL: global_load_tr_b128_i16: +; GFX12-GISEL-W32: ; %bb.0: ; %entry +; GFX12-GISEL-W32-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-GISEL-W32-NEXT: v_mov_b32_e32 v4, 0 +; GFX12-GISEL-W32-NEXT: s_wait_kmcnt 0x0 +; GFX12-GISEL-W32-NEXT: global_load_tr_b128 v[0:3], v4, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-GISEL-W32-NEXT: s_wait_loadcnt 0x0 +; GFX12-GISEL-W32-NEXT: global_inv scope:SCOPE_SYS +; GFX12-GISEL-W32-NEXT: global_store_b128 v4, v[0:3], s[2:3] +; GFX12-GISEL-W32-NEXT: s_nop 0 +; GFX12-GISEL-W32-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-GISEL-W32-NEXT: s_endpgm +entry: + %gep = getelementptr i64, ptr addrspace(1) %addr, i32 4 + %val = call <8 x i16> @llvm.amdgcn.global.load.tr.v8i16.p1(ptr addrspace(1) %gep) + store <8 x i16> %val, ptr addrspace(1) %use + ret void +} + +define amdgpu_kernel void @global_load_tr_b128_half(ptr addrspace(1) %addr, ptr addrspace(1) %use) { +; GFX12-SDAG-W32-LABEL: global_load_tr_b128_half: +; GFX12-SDAG-W32: ; %bb.0: ; %entry +; GFX12-SDAG-W32-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-SDAG-W32-NEXT: v_mov_b32_e32 v4, 0 +; GFX12-SDAG-W32-NEXT: s_wait_kmcnt 0x0 +; GFX12-SDAG-W32-NEXT: global_load_tr_b128 v[0:3], v4, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-SDAG-W32-NEXT: s_wait_loadcnt 0x0 +; GFX12-SDAG-W32-NEXT: global_inv scope:SCOPE_SYS +; GFX12-SDAG-W32-NEXT: global_store_b128 v4, v[0:3], s[2:3] +; GFX12-SDAG-W32-NEXT: s_nop 0 +; GFX12-SDAG-W32-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-SDAG-W32-NEXT: s_endpgm +; +; GFX12-GISEL-W32-LABEL: global_load_tr_b128_half: +; GFX12-GISEL-W32: ; %bb.0: ; %entry +; GFX12-GISEL-W32-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-GISEL-W32-NEXT: v_mov_b32_e32 v4, 0 +; GFX12-GISEL-W32-NEXT: s_wait_kmcnt 0x0 +; GFX12-GISEL-W32-NEXT: global_load_tr_b128 v[0:3], v4, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-GISEL-W32-NEXT: s_wait_loadcnt 0x0 +; GFX12-GISEL-W32-NEXT: global_inv scope:SCOPE_SYS +; GFX12-GISEL-W32-NEXT: global_store_b128 v4, v[0:3], s[2:3] +; GFX12-GISEL-W32-NEXT: s_nop 0 +; GFX12-GISEL-W32-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-GISEL-W32-NEXT: s_endpgm +entry: + %gep = getelementptr i64, ptr addrspace(1) %addr, i32 4 + %val = call <8 x half> @llvm.amdgcn.global.load.tr.v8f16.p1(ptr addrspace(1) %gep) + store <8 x half> %val, ptr addrspace(1) %use + ret void +} + +define amdgpu_kernel void @global_load_tr_b128_bfloat(ptr addrspace(1) %addr, ptr addrspace(1) %use) { +; GFX12-SDAG-W32-LABEL: global_load_tr_b128_bfloat: +; GFX12-SDAG-W32: ; %bb.0: ; %entry +; GFX12-SDAG-W32-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-SDAG-W32-NEXT: v_mov_b32_e32 v4, 0 +; GFX12-SDAG-W32-NEXT: s_wait_kmcnt 0x0 +; GFX12-SDAG-W32-NEXT: global_load_tr_b128 v[0:3], v4, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-SDAG-W32-NEXT: s_wait_loadcnt 0x0 +; GFX12-SDAG-W32-NEXT: global_inv scope:SCOPE_SYS +; GFX12-SDAG-W32-NEXT: global_store_b128 v4, v[0:3], s[2:3] +; GFX12-SDAG-W32-NEXT: s_nop 0 +; GFX12-SDAG-W32-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-SDAG-W32-NEXT: s_endpgm +; +; GFX12-GISEL-W32-LABEL: global_load_tr_b128_bfloat: +; GFX12-GISEL-W32: ; %bb.0: ; %entry +; GFX12-GISEL-W32-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-GISEL-W32-NEXT: v_mov_b32_e32 v4, 0 +; GFX12-GISEL-W32-NEXT: s_wait_kmcnt 0x0 +; GFX12-GISEL-W32-NEXT: global_load_tr_b128 v[0:3], v4, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-GISEL-W32-NEXT: s_wait_loadcnt 0x0 +; GFX12-GISEL-W32-NEXT: global_inv scope:SCOPE_SYS +; GFX12-GISEL-W32-NEXT: global_store_b128 v4, v[0:3], s[2:3] +; GFX12-GISEL-W32-NEXT: s_nop 0 +; GFX12-GISEL-W32-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-GISEL-W32-NEXT: s_endpgm +entry: + %gep = getelementptr i64, ptr addrspace(1) %addr, i32 4 + %val = call <8 x bfloat> @llvm.amdgcn.global.load.tr.v8bf16.p1(ptr addrspace(1) %gep) + store <8 x bfloat> %val, ptr addrspace(1) %use + ret void +} diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.global.load.tr-w64.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.global.load.tr-w64.ll new file mode 100644 index 0000000000000..0936d17563644 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.global.load.tr-w64.ll @@ -0,0 +1,140 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py +; RUN: llc -global-isel=0 -march=amdgcn -mcpu=gfx1200 -verify-machineinstrs -mattr=-wavefrontsize32,+wavefrontsize64 < %s | FileCheck -check-prefixes=GFX12-SDAG-W64 %s +; RUN: llc -global-isel=1 -march=amdgcn -mcpu=gfx1200 -verify-machineinstrs -mattr=-wavefrontsize32,+wavefrontsize64 < %s | FileCheck -check-prefixes=GFX12-GISEL-W64 %s + +declare i32 @llvm.amdgcn.global.load.tr.i32.p1(ptr addrspace(1)) +declare <4 x i16> @llvm.amdgcn.global.load.tr.v4i16.p1(ptr addrspace(1)) +declare <4 x half> @llvm.amdgcn.global.load.tr.v4f16.p1(ptr addrspace(1)) +declare <4 x bfloat> @llvm.amdgcn.global.load.tr.v4bf16.p1(ptr addrspace(1)) + +define amdgpu_kernel void @global_load_tr_b64(ptr addrspace(1) %addr, ptr addrspace(1) %use) { +; GFX12-SDAG-W64-LABEL: global_load_tr_b64: +; GFX12-SDAG-W64: ; %bb.0: ; %entry +; GFX12-SDAG-W64-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-SDAG-W64-NEXT: v_mov_b32_e32 v0, 0 +; GFX12-SDAG-W64-NEXT: s_wait_kmcnt 0x0 +; GFX12-SDAG-W64-NEXT: global_load_tr_b64 v1, v0, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-SDAG-W64-NEXT: s_wait_loadcnt 0x0 +; GFX12-SDAG-W64-NEXT: global_inv scope:SCOPE_SYS +; GFX12-SDAG-W64-NEXT: global_store_b32 v0, v1, s[2:3] +; GFX12-SDAG-W64-NEXT: s_nop 0 +; GFX12-SDAG-W64-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-SDAG-W64-NEXT: s_endpgm +; +; GFX12-GISEL-W64-LABEL: global_load_tr_b64: +; GFX12-GISEL-W64: ; %bb.0: ; %entry +; GFX12-GISEL-W64-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-GISEL-W64-NEXT: v_mov_b32_e32 v0, 0 +; GFX12-GISEL-W64-NEXT: s_wait_kmcnt 0x0 +; GFX12-GISEL-W64-NEXT: global_load_tr_b64 v1, v0, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-GISEL-W64-NEXT: s_wait_loadcnt 0x0 +; GFX12-GISEL-W64-NEXT: global_inv scope:SCOPE_SYS +; GFX12-GISEL-W64-NEXT: global_store_b32 v0, v1, s[2:3] +; GFX12-GISEL-W64-NEXT: s_nop 0 +; GFX12-GISEL-W64-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-GISEL-W64-NEXT: s_endpgm +entry: + %gep = getelementptr i64, ptr addrspace(1) %addr, i32 4 + %val = call i32 @llvm.amdgcn.global.load.tr.i32.p1(ptr addrspace(1) %gep) + store i32 %val, ptr addrspace(1) %use + ret void +} + +define amdgpu_kernel void @global_load_tr_b128_i16(ptr addrspace(1) %addr, ptr addrspace(1) %use) { +; GFX12-SDAG-W64-LABEL: global_load_tr_b128_i16: +; GFX12-SDAG-W64: ; %bb.0: ; %entry +; GFX12-SDAG-W64-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-SDAG-W64-NEXT: v_mov_b32_e32 v2, 0 +; GFX12-SDAG-W64-NEXT: s_wait_kmcnt 0x0 +; GFX12-SDAG-W64-NEXT: global_load_tr_b128 v[0:1], v2, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-SDAG-W64-NEXT: s_wait_loadcnt 0x0 +; GFX12-SDAG-W64-NEXT: global_inv scope:SCOPE_SYS +; GFX12-SDAG-W64-NEXT: global_store_b64 v2, v[0:1], s[2:3] +; GFX12-SDAG-W64-NEXT: s_nop 0 +; GFX12-SDAG-W64-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-SDAG-W64-NEXT: s_endpgm +; +; GFX12-GISEL-W64-LABEL: global_load_tr_b128_i16: +; GFX12-GISEL-W64: ; %bb.0: ; %entry +; GFX12-GISEL-W64-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-GISEL-W64-NEXT: v_mov_b32_e32 v2, 0 +; GFX12-GISEL-W64-NEXT: s_wait_kmcnt 0x0 +; GFX12-GISEL-W64-NEXT: global_load_tr_b128 v[0:1], v2, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-GISEL-W64-NEXT: s_wait_loadcnt 0x0 +; GFX12-GISEL-W64-NEXT: global_inv scope:SCOPE_SYS +; GFX12-GISEL-W64-NEXT: global_store_b64 v2, v[0:1], s[2:3] +; GFX12-GISEL-W64-NEXT: s_nop 0 +; GFX12-GISEL-W64-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-GISEL-W64-NEXT: s_endpgm +entry: + %gep = getelementptr i64, ptr addrspace(1) %addr, i32 4 + %val = call <4 x i16> @llvm.amdgcn.global.load.tr.v4i16.p1(ptr addrspace(1) %gep) + store <4 x i16> %val, ptr addrspace(1) %use + ret void +} + +define amdgpu_kernel void @global_load_tr_b128_half(ptr addrspace(1) %addr, ptr addrspace(1) %use) { +; GFX12-SDAG-W64-LABEL: global_load_tr_b128_half: +; GFX12-SDAG-W64: ; %bb.0: ; %entry +; GFX12-SDAG-W64-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-SDAG-W64-NEXT: v_mov_b32_e32 v2, 0 +; GFX12-SDAG-W64-NEXT: s_wait_kmcnt 0x0 +; GFX12-SDAG-W64-NEXT: global_load_tr_b128 v[0:1], v2, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-SDAG-W64-NEXT: s_wait_loadcnt 0x0 +; GFX12-SDAG-W64-NEXT: global_inv scope:SCOPE_SYS +; GFX12-SDAG-W64-NEXT: global_store_b64 v2, v[0:1], s[2:3] +; GFX12-SDAG-W64-NEXT: s_nop 0 +; GFX12-SDAG-W64-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-SDAG-W64-NEXT: s_endpgm +; +; GFX12-GISEL-W64-LABEL: global_load_tr_b128_half: +; GFX12-GISEL-W64: ; %bb.0: ; %entry +; GFX12-GISEL-W64-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-GISEL-W64-NEXT: v_mov_b32_e32 v2, 0 +; GFX12-GISEL-W64-NEXT: s_wait_kmcnt 0x0 +; GFX12-GISEL-W64-NEXT: global_load_tr_b128 v[0:1], v2, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-GISEL-W64-NEXT: s_wait_loadcnt 0x0 +; GFX12-GISEL-W64-NEXT: global_inv scope:SCOPE_SYS +; GFX12-GISEL-W64-NEXT: global_store_b64 v2, v[0:1], s[2:3] +; GFX12-GISEL-W64-NEXT: s_nop 0 +; GFX12-GISEL-W64-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-GISEL-W64-NEXT: s_endpgm +entry: + %gep = getelementptr i64, ptr addrspace(1) %addr, i32 4 + %val = call <4 x half> @llvm.amdgcn.global.load.tr.v4f16.p1(ptr addrspace(1) %gep) + store <4 x half> %val, ptr addrspace(1) %use + ret void +} + +define amdgpu_kernel void @global_load_tr_b128_bfloat(ptr addrspace(1) %addr, ptr addrspace(1) %use) { +; GFX12-SDAG-W64-LABEL: global_load_tr_b128_bfloat: +; GFX12-SDAG-W64: ; %bb.0: ; %entry +; GFX12-SDAG-W64-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-SDAG-W64-NEXT: v_mov_b32_e32 v2, 0 +; GFX12-SDAG-W64-NEXT: s_wait_kmcnt 0x0 +; GFX12-SDAG-W64-NEXT: global_load_tr_b128 v[0:1], v2, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-SDAG-W64-NEXT: s_wait_loadcnt 0x0 +; GFX12-SDAG-W64-NEXT: global_inv scope:SCOPE_SYS +; GFX12-SDAG-W64-NEXT: global_store_b64 v2, v[0:1], s[2:3] +; GFX12-SDAG-W64-NEXT: s_nop 0 +; GFX12-SDAG-W64-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-SDAG-W64-NEXT: s_endpgm +; +; GFX12-GISEL-W64-LABEL: global_load_tr_b128_bfloat: +; GFX12-GISEL-W64: ; %bb.0: ; %entry +; GFX12-GISEL-W64-NEXT: s_load_b128 s[0:3], s[0:1], 0x24 +; GFX12-GISEL-W64-NEXT: v_mov_b32_e32 v2, 0 +; GFX12-GISEL-W64-NEXT: s_wait_kmcnt 0x0 +; GFX12-GISEL-W64-NEXT: global_load_tr_b128 v[0:1], v2, s[0:1] offset:32 th:TH_LOAD_NT +; GFX12-GISEL-W64-NEXT: s_wait_loadcnt 0x0 +; GFX12-GISEL-W64-NEXT: global_inv scope:SCOPE_SYS +; GFX12-GISEL-W64-NEXT: global_store_b64 v2, v[0:1], s[2:3] +; GFX12-GISEL-W64-NEXT: s_nop 0 +; GFX12-GISEL-W64-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) +; GFX12-GISEL-W64-NEXT: s_endpgm +entry: + %gep = getelementptr i64, ptr addrspace(1) %addr, i32 4 + %val = call <4 x bfloat> @llvm.amdgcn.global.load.tr.v4bf16.p1(ptr addrspace(1) %gep) + store <4 x bfloat> %val, ptr addrspace(1) %use + ret void +} diff --git a/llvm/test/MC/AMDGPU/gfx11_unsupported.s b/llvm/test/MC/AMDGPU/gfx11_unsupported.s index d110360e5bc62..bfca71ae3a01e 100644 --- a/llvm/test/MC/AMDGPU/gfx11_unsupported.s +++ b/llvm/test/MC/AMDGPU/gfx11_unsupported.s @@ -223,6 +223,12 @@ global_load_lds_ubyte v[2:3], off global_load_lds_ushort v[2:3], off // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU +global_load_tr_b128 v[1:4], v5, s[2:3] +// CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + +global_load_tr_b64 v[1:2], v[3:4], off +// CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + image_atomic_fcmpswap v[1:2], v2, s[12:19] dmask:0x3 dim:SQ_RSRC_IMG_1D unorm // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU diff --git a/llvm/test/MC/AMDGPU/gfx12_asm_global_load_tr.s b/llvm/test/MC/AMDGPU/gfx12_asm_global_load_tr.s new file mode 100644 index 0000000000000..597e0d29b43af --- /dev/null +++ b/llvm/test/MC/AMDGPU/gfx12_asm_global_load_tr.s @@ -0,0 +1,103 @@ +// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1200 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=W64 %s +// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1200 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=W32 %s +// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1200 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W64-ERR --implicit-check-not=error: %s +// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1200 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W32-ERR --implicit-check-not=error: %s + +global_load_tr_b128 v[1:4], v0, s[0:1] offset:-64 +// W64-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W32: encoding: [0x00,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0xc0,0xff,0xff] + +global_load_tr_b128 v[1:4], v0, s[0:1] offset:64 +// W64-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W32: encoding: [0x00,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] + +global_load_tr_b128 v[1:4], v5, s[2:3] +// W64-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W32: encoding: [0x02,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x05,0x00,0x00,0x00] + +global_load_tr_b128 v[1:4], v[0:1], off offset:-64 +// W64-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W32: encoding: [0x7c,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0xc0,0xff,0xff] + +global_load_tr_b128 v[1:4], v[0:1], off offset:64 +// W64-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W32: encoding: [0x7c,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] + +global_load_tr_b128 v[1:4], v[5:6], off +// W64-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W32: encoding: [0x7c,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x05,0x00,0x00,0x00] + +global_load_tr_b64 v[1:2], v0, s[0:1] offset:-64 +// W64-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W32: encoding: [0x00,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0xc0,0xff,0xff] + +global_load_tr_b64 v[1:2], v0, s[0:1] offset:64 +// W64-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W32: encoding: [0x00,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] + +global_load_tr_b64 v[1:2], v3, s[2:3] +// W64-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W32: encoding: [0x02,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x03,0x00,0x00,0x00] + +global_load_tr_b64 v[1:2], v[0:1], off offset:-64 +// W64-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W32: encoding: [0x7c,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0xc0,0xff,0xff] + +global_load_tr_b64 v[1:2], v[0:1], off offset:64 +// W64-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W32: encoding: [0x7c,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] + +global_load_tr_b64 v[1:2], v[3:4], off +// W64-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W32: encoding: [0x7c,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x03,0x00,0x00,0x00] + + + +global_load_tr_b128 v[1:2], v0, s[0:1] offset:-64 +// W32-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W64: encoding: [0x00,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0xc0,0xff,0xff] + +global_load_tr_b128 v[1:2], v0, s[0:1] offset:64 +// W32-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W64: encoding: [0x00,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] + +global_load_tr_b128 v[1:2], v5, s[2:3] +// W32-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W64: encoding: [0x02,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x05,0x00,0x00,0x00] + +global_load_tr_b128 v[1:2], v[0:1], off offset:-64 +// W32-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W64: encoding: [0x7c,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0xc0,0xff,0xff] + +global_load_tr_b128 v[1:2], v[0:1], off offset:64 +// W32-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W64: encoding: [0x7c,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] + +global_load_tr_b128 v[1:2], v[5:6], off +// W32-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W64: encoding: [0x7c,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x05,0x00,0x00,0x00] + +global_load_tr_b64 v1, v0, s[0:1] offset:-64 +// W32-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W64: encoding: [0x00,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0xc0,0xff,0xff] + +global_load_tr_b64 v1, v0, s[0:1] offset:64 +// W32-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W64: encoding: [0x00,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] + +global_load_tr_b64 v1, v3, s[2:3] +// W32-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W64: encoding: [0x02,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x03,0x00,0x00,0x00] + +global_load_tr_b64 v1, v[0:1], off offset:-64 +// W32-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W64: encoding: [0x7c,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0xc0,0xff,0xff] + +global_load_tr_b64 v1, v[0:1], off offset:64 +// W32-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W64: encoding: [0x7c,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] + +global_load_tr_b64 v1, v[3:4], off +// W32-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +// W64: encoding: [0x7c,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x03,0x00,0x00,0x00] + diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_global_load_tr.txt b/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_global_load_tr.txt new file mode 100644 index 0000000000000..e8498d4aef0a1 --- /dev/null +++ b/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_global_load_tr.txt @@ -0,0 +1,34 @@ +# RUN: llvm-mc -disassemble -arch=amdgcn -mcpu=gfx1200 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=W64 %s +# RUN: llvm-mc -disassemble -arch=amdgcn -mcpu=gfx1200 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=W32 %s + +# W32: global_load_tr_b128 v[1:4], v0, s[0:1] offset:64 ; encoding: [0x00,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] +# W64: global_load_tr_b128 v[1:2], v0, s[0:1] offset:64 ; encoding: [0x00,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] +0x00,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00 + +# W32: global_load_tr_b128 v[1:4], v5, s[2:3] ; encoding: [0x02,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x05,0x00,0x00,0x00] +# W64: global_load_tr_b128 v[1:2], v5, s[2:3] ; encoding: [0x02,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x05,0x00,0x00,0x00] +0x02,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x05,0x00,0x00,0x00 + +# W32: global_load_tr_b128 v[1:4], v[0:1], off offset:64 ; encoding: [0x7c,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] +# W64: global_load_tr_b128 v[1:2], v[0:1], off offset:64 ; encoding: [0x7c,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] +0x7c,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00 + +# W32: global_load_tr_b128 v[1:4], v[5:6], off ; encoding: [0x7c,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x05,0x00,0x00,0x00] +# W64: global_load_tr_b128 v[1:2], v[5:6], off ; encoding: [0x7c,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x05,0x00,0x00,0x00] +0x7c,0xc0,0x15,0xee,0x01,0x00,0x00,0x00,0x05,0x00,0x00,0x00 + +# W32: global_load_tr_b64 v[1:2], v0, s[0:1] offset:64 ; encoding: [0x00,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] +# W64: global_load_tr_b64 v1, v0, s[0:1] offset:64 ; encoding: [0x00,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] +0x00,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00 + +# W32: global_load_tr_b64 v[1:2], v3, s[2:3] ; encoding: [0x02,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x03,0x00,0x00,0x00] +# W64: global_load_tr_b64 v1, v3, s[2:3] ; encoding: [0x02,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x03,0x00,0x00,0x00] +0x02,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x03,0x00,0x00,0x00 + +# W32: global_load_tr_b64 v[1:2], v[0:1], off offset:64 ; encoding: [0x7c,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] +# W64: global_load_tr_b64 v1, v[0:1], off offset:64 ; encoding: [0x7c,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00] +0x7c,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x00,0x40,0x00,0x00 + +# W32: global_load_tr_b64 v[1:2], v[3:4], off ; encoding: [0x7c,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x03,0x00,0x00,0x00] +# W64: global_load_tr_b64 v1, v[3:4], off ; encoding: [0x7c,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x03,0x00,0x00,0x00] +0x7c,0x00,0x16,0xee,0x01,0x00,0x00,0x00,0x03,0x00,0x00,0x00