Skip to content

[AMDGPU] Add s_setprio_inc_wg gfx1250 instruction #145152

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: users/rampitec/06-20-_amdgpu_rename_call_instructions_from_b64_to_i64
Choose a base branch
from
+100 −2
Open

[AMDGPU] Add s_setprio_inc_wg gfx1250 instruction #145152

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 6 additions & 0 deletions clang/include/clang/Basic/BuiltinsAMDGPU.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -636,5 +636,11 @@ TARGET_BUILTIN(__builtin_amdgcn_bitop3_b16, "ssssIUi", "nc", "bitop3-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_bf16_f32, "V2yV2yfUiIb", "nc", "f32-to-f16bf16-cvt-sr-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_f16_f32, "V2hV2hfUiIb", "nc", "f32-to-f16bf16-cvt-sr-insts")

//===----------------------------------------------------------------------===//
// GFX1250+ only builtins.
//===----------------------------------------------------------------------===//

TARGET_BUILTIN(__builtin_amdgcn_s_setprio_inc_wg, "vIs", "n", "setprio-inc-wg-inst")
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Missing the sema tests for the supported subtarget, and non-immediate input errors


#undef BUILTIN
#undef TARGET_BUILTIN
7 changes: 7 additions & 0 deletions clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-err.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,7 @@
// REQUIRES: amdgpu-registered-target

// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -verify -emit-llvm -o - %s

void test_setprio_inc_wg(short a) {
__builtin_amdgcn_s_setprio_inc_wg(a); // expected-error {{'__builtin_amdgcn_s_setprio_inc_wg' must be a constant integer}}
}
12 changes: 12 additions & 0 deletions clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,12 @@
// 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 gfx1250 -emit-llvm -o - %s | FileCheck %s
// REQUIRES: amdgpu-registered-target

// CHECK-LABEL: @test_setprio_inc_wg(
// CHECK-NEXT: entry:
// CHECK-NEXT: call void @llvm.amdgcn.s.setprio.inc.wg(i16 10)
// CHECK-NEXT: ret void
//
void test_setprio_inc_wg() {
__builtin_amdgcn_s_setprio_inc_wg(10);
}
5 changes: 5 additions & 0 deletions llvm/include/llvm/IR/IntrinsicsAMDGPU.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2054,6 +2054,11 @@ def int_amdgcn_s_setprio :
DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
IntrHasSideEffects]>;

def int_amdgcn_s_setprio_inc_wg :
ClangBuiltin<"__builtin_amdgcn_s_setprio_inc_wg">,
DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
IntrHasSideEffects]>;

def int_amdgcn_s_ttracedata :
ClangBuiltin<"__builtin_amdgcn_s_ttracedata">,
DefaultAttrsIntrinsic<[], [llvm_i32_ty],
Expand Down
10 changes: 10 additions & 0 deletions llvm/lib/Target/AMDGPU/AMDGPU.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1118,6 +1118,12 @@ def FeatureWaitXcnt : SubtargetFeature<"wait-xcnt",
"Has s_wait_xcnt instruction"
>;

def FeatureSetPrioIncWgInst : SubtargetFeature<"setprio-inc-wg-inst",
"HasSetPrioIncWgInst",
"true",
"Has s_setprio_inc_wg instruction."
>;

//===------------------------------------------------------------===//
// Subtarget Features (options and debugging)
//===------------------------------------------------------------===//
Expand Down Expand Up @@ -1940,6 +1946,7 @@ def FeatureISAVersion12_50 : FeatureSet<
FeatureMemoryAtomicFAddF32DenormalSupport,
FeatureKernargPreload,
FeatureLshlAddU64Inst,
FeatureSetPrioIncWgInst,
]>;

def FeatureISAVersion12_Generic: FeatureSet<
Expand Down Expand Up @@ -2662,6 +2669,9 @@ def HasAshrPkInsts : Predicate<"Subtarget->hasAshrPkInsts()">,
def HasLshlAddU64Inst : Predicate<"Subtarget->hasLshlAddU64Inst()">,
AssemblerPredicate<(all_of FeatureLshlAddU64Inst)>;

def HasSetPrioIncWgInst : Predicate<"Subtarget->hasSetPrioIncWgInst()">,
AssemblerPredicate<(all_of FeatureSetPrioIncWgInst)>;

// Include AMDGPU TD files
include "SISchedule.td"
include "GCNProcessors.td"
Expand Down
6 changes: 6 additions & 0 deletions llvm/lib/Target/AMDGPU/GCNSubtarget.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -262,6 +262,7 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
bool HasMinimum3Maximum3PKF16 = false;
bool HasLshlAddU64Inst = false;
bool HasPointSampleAccel = false;
bool HasSetPrioIncWgInst = false;

bool RequiresCOV6 = false;
bool UseBlockVGPROpsForCSR = false;
Expand Down Expand Up @@ -1465,6 +1466,11 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
/// values.
bool hasSignedScratchOffsets() const { return getGeneration() >= GFX12; }

bool hasGFX1250Insts() const { return GFX1250Insts; }

// \returns true if target has S_SETPRIO_INC_WG instruction.
bool hasSetPrioIncWgInst() const { return HasSetPrioIncWgInst; }

// \returns true if S_GETPC_B64 zero-extends the result from 48 bits instead
// of sign-extending.
bool hasGetPCZeroExtension() const { return GFX12Insts; }
Expand Down
6 changes: 5 additions & 1 deletion llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9575,7 +9575,8 @@ static unsigned subtargetEncodingFamily(const GCNSubtarget &ST) {
case AMDGPUSubtarget::GFX11:
return SIEncodingFamily::GFX11;
case AMDGPUSubtarget::GFX12:
return SIEncodingFamily::GFX12;
return ST.hasGFX1250Insts() ? SIEncodingFamily::GFX1250
: SIEncodingFamily::GFX12;
}
llvm_unreachable("Unknown subtarget generation!");
}
Expand Down Expand Up @@ -9669,6 +9670,9 @@ int SIInstrInfo::pseudoToMCOpcode(int Opcode) const {

int MCOp = AMDGPU::getMCOpcode(Opcode, Gen);

if (MCOp == (uint16_t)-1 && ST.hasGFX1250Insts())
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

So we want a 0xffff instead of a 0xffffffff here? Why is that?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

0xffffffff means it is already a real opcode. 0xffff means it is not found in the subtarget column.

MCOp = AMDGPU::getMCOpcode(Opcode, SIEncodingFamily::GFX12);

// -1 means that Opcode is already a native instruction.
if (MCOp == -1)
return Opcode;
Expand Down
3 changes: 2 additions & 1 deletion llvm/lib/Target/AMDGPU/SIInstrInfo.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3069,7 +3069,8 @@ def getMCOpcodeGen : InstrMapping {
[!cast<string>(SIEncodingFamily.GFX90A)],
[!cast<string>(SIEncodingFamily.GFX940)],
[!cast<string>(SIEncodingFamily.GFX11)],
[!cast<string>(SIEncodingFamily.GFX12)]];
[!cast<string>(SIEncodingFamily.GFX12)],
[!cast<string>(SIEncodingFamily.GFX1250)]];
}

// Get equivalent SOPK instruction.
Expand Down
6 changes: 6 additions & 0 deletions llvm/lib/Target/AMDGPU/SOPInstructions.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1632,6 +1632,11 @@ def S_SETPRIO : SOPP_Pseudo <"s_setprio", (ins i16imm:$simm16), "$simm16",
[(int_amdgcn_s_setprio timm:$simm16)]> {
}

def S_SETPRIO_INC_WG : SOPP_Pseudo <"s_setprio_inc_wg", (ins i16imm:$simm16), "$simm16",
[(int_amdgcn_s_setprio_inc_wg timm:$simm16)]> {
let SubtargetPredicate = HasSetPrioIncWgInst;
}

let Uses = [EXEC, M0] in {
def S_SENDMSG : SOPP_Pseudo <"s_sendmsg" , (ins SendMsg:$simm16), "$simm16",
[(int_amdgcn_s_sendmsg (i32 timm:$simm16), M0)]> {
Expand Down Expand Up @@ -2594,6 +2599,7 @@ defm S_WAIT_STORECNT_DSCNT : SOPP_Real_32_gfx12<0x049>;
//===----------------------------------------------------------------------===//
// SOPP - GFX1250 only.
//===----------------------------------------------------------------------===//
defm S_SETPRIO_INC_WG : SOPP_Real_32_gfx12<0x03e>;
defm S_WAIT_XCNT : SOPP_Real_32_gfx12<0x045>;

//===----------------------------------------------------------------------===//
Expand Down
34 changes: 34 additions & 0 deletions llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.inc.wg.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,34 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
; RUN: llc -march=amdgcn -mcpu=gfx1250 -show-mc-encoding -verify-machineinstrs < %s | FileCheck -check-prefix=GFX1250 %s
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think it is recommended to use -mtriple=amdgcn-amd-amdhsa instead of -march. There was a bulk update in upstream in the past.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Will do. This is a very old test of course.

; RUN: llc -global-isel -march=amdgcn -mcpu=gfx1250 -show-mc-encoding -verify-machineinstrs < %s | FileCheck -check-prefix=GFX1250 %s

declare void @llvm.amdgcn.s.setprio.inc.wg(i16) #0

define void @test_llvm.amdgcn.s.setprio.inc.wg() #0 {
; GFX1250-LABEL: test_llvm.amdgcn.s.setprio.inc.wg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0 ; encoding: [0x00,0x00,0xc8,0xbf]
; GFX1250-NEXT: s_wait_kmcnt 0x0 ; encoding: [0x00,0x00,0xc7,0xbf]
; GFX1250-NEXT: s_setprio_inc_wg 0 ; encoding: [0x00,0x00,0xbe,0xbf]
; GFX1250-NEXT: s_setprio_inc_wg 1 ; encoding: [0x01,0x00,0xbe,0xbf]
; GFX1250-NEXT: s_setprio_inc_wg 2 ; encoding: [0x02,0x00,0xbe,0xbf]
; GFX1250-NEXT: s_setprio_inc_wg 3 ; encoding: [0x03,0x00,0xbe,0xbf]
; GFX1250-NEXT: s_setprio_inc_wg 10 ; encoding: [0x0a,0x00,0xbe,0xbf]
; GFX1250-NEXT: s_setprio_inc_wg -1 ; encoding: [0xff,0xff,0xbe,0xbf]
; GFX1250-NEXT: s_setprio_inc_wg 0 ; encoding: [0x00,0x00,0xbe,0xbf]
; GFX1250-NEXT: s_setprio_inc_wg 1 ; encoding: [0x01,0x00,0xbe,0xbf]
; GFX1250-NEXT: s_setprio_inc_wg -1 ; encoding: [0xff,0xff,0xbe,0xbf]
; GFX1250-NEXT: s_set_pc_i64 s[30:31] ; encoding: [0x1e,0x48,0x80,0xbe]
call void @llvm.amdgcn.s.setprio.inc.wg(i16 0)
call void @llvm.amdgcn.s.setprio.inc.wg(i16 1)
call void @llvm.amdgcn.s.setprio.inc.wg(i16 2)
call void @llvm.amdgcn.s.setprio.inc.wg(i16 3)
call void @llvm.amdgcn.s.setprio.inc.wg(i16 10)
call void @llvm.amdgcn.s.setprio.inc.wg(i16 65535)
call void @llvm.amdgcn.s.setprio.inc.wg(i16 65536)
call void @llvm.amdgcn.s.setprio.inc.wg(i16 65537)
call void @llvm.amdgcn.s.setprio.inc.wg(i16 -1)
ret void
}

attributes #0 = { nounwind }
4 changes: 4 additions & 0 deletions llvm/test/MC/AMDGPU/gfx1250_asm_sopp.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,3 +12,7 @@ s_wait_xcnt 0x7
s_wait_xcnt 0xf
// GFX1250: [0x0f,0x00,0xc5,0xbf]
// GFX12-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU

s_setprio_inc_wg 100
// GFX1250: [0x64,0x00,0xbe,0xbf]
// GFX12-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
3 changes: 3 additions & 0 deletions llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_sopp.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,3 +8,6 @@

# GFX1250: s_wait_xcnt 0xf ; encoding: [0x0f,0x00,0xc5,0xbf]
0x0f,0x00,0xc5,0xbf

# GFX1250: s_setprio_inc_wg 0x64 ; encoding: [0x64,0x00,0xbe,0xbf]
0x64,0x00,0xbe,0xbf
Loading