This is an archive of the discontinued LLVM Phabricator instance.

Differential D127955

[AMDGPU] Add GFX11 llvm.amdgcn.ds.add.gs.reg.rtn / llvm.amdgcn.ds.sub.gs.reg.rtn intrinsics
ClosedPublic

Authored by foad on Jun 16 2022, 5:28 AM.

Details

Reviewers
Joe_Nash
rampitec
critson
piotr
Group Reviewers
Restricted Project
Commits
rG36ec1fcaac8e: [AMDGPU] Add GFX11 llvm.amdgcn.ds.add.gs.reg.rtn / llvm.amdgcn.ds.sub.gs.reg.

Diff Detail

Event Timeline

foad created this revision.Jun 16 2022, 5:28 AM
Herald added a project: Restricted Project. · View Herald TranscriptJun 16 2022, 5:28 AM
Herald added subscribers: kosarev, jsilvanus, hsmhsm and 10 others. · View Herald Transcript
foad requested review of this revision.Jun 16 2022, 5:28 AM
Herald added a project: Restricted Project. · View Herald TranscriptJun 16 2022, 5:28 AM
Herald added subscribers: llvm-commits, wdng. · View Herald Transcript
foad added a reviewer: Restricted Project.Jun 16 2022, 5:28 AM
Harbormaster completed remote builds in B170238: Diff 437511.Jun 16 2022, 6:36 AM
rampitec accepted this revision.Jun 16 2022, 10:20 AM
This revision is now accepted and ready to land.Jun 16 2022, 10:20 AM
This revision was landed with ongoing or failed builds.Jun 16 2022, 10:33 AM
Closed by commit rG36ec1fcaac8e: [AMDGPU] Add GFX11 llvm.amdgcn.ds.add.gs.reg.rtn / llvm.amdgcn.ds.sub.gs.reg. (authored by foad). · Explain Why
This revision was automatically updated to reflect the committed changes.
foad added a commit: rG36ec1fcaac8e: [AMDGPU] Add GFX11 llvm.amdgcn.ds.add.gs.reg.rtn / llvm.amdgcn.ds.sub.gs.reg..
arsenm added inline comments.Jun 16 2022, 4:18 PM
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ds.add.gs.reg.rtn.ll
3

These tests should use explicit -global-isel=0/1 for whenever the default switches

foad added inline comments.Jun 17 2022, 3:53 AM
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ds.add.gs.reg.rtn.ll
3

Done in 03a991032571, but it's only the tip of the iceberg.

Revision Contents

PathSize
llvm/
include/
llvm/
IR/
10 lines
lib/
Target/
AMDGPU/
5 lines
28 lines
test/
CodeGen/
AMDGPU/
70 lines
71 lines

Diff 437511

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Show First 20 Lines • Show All 1,898 Lines • ▼ Show 20 Lines
// GFX11 Intrinsics // GFX11 Intrinsics
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// llvm.amdgcn.permlane64 <src0> // llvm.amdgcn.permlane64 <src0>
def int_amdgcn_permlane64 : def int_amdgcn_permlane64 :
Intrinsic<[llvm_i32_ty], [llvm_i32_ty], Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
[IntrNoMem, IntrConvergent, IntrWillReturn]>; [IntrNoMem, IntrConvergent, IntrWillReturn]>;
def int_amdgcn_ds_add_gs_reg_rtn :
GCCBuiltin<"__builtin_amdgcn_ds_add_gs_reg_rtn">,
Intrinsic<[llvm_anyint_ty], [llvm_i32_ty, llvm_i32_ty],
[ImmArg<ArgIndex<1>>, IntrHasSideEffects, IntrWillReturn]>;
def int_amdgcn_ds_sub_gs_reg_rtn :
GCCBuiltin<"__builtin_amdgcn_ds_sub_gs_reg_rtn">,
Intrinsic<[llvm_anyint_ty], [llvm_i32_ty, llvm_i32_ty],
[ImmArg<ArgIndex<1>>, IntrHasSideEffects, IntrWillReturn]>;
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// Deep learning intrinsics. // Deep learning intrinsics.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// f32 %r = llvm.amdgcn.fdot2(v2f16 %a, v2f16 %b, f32 %c, i1 %clamp) // f32 %r = llvm.amdgcn.fdot2(v2f16 %a, v2f16 %b, f32 %c, i1 %clamp)
// %r = %a[0] * %b[0] + %a[1] * %b[1] + %c // %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
def int_amdgcn_fdot2 : def int_amdgcn_fdot2 :
GCCBuiltin<"__builtin_amdgcn_fdot2">, GCCBuiltin<"__builtin_amdgcn_fdot2">,
▲ Show 20 Lines • Show All 235 Lines • Show Last 20 Lines

llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

Show First 20 Lines • Show All 4,654 Lines • ▼ Show 20 Lines case Intrinsic::amdgcn_ds_gws_sema_release_all: {
OpdsMapping[1] = AMDGPU::getValueMapping(Bank, 32); OpdsMapping[1] = AMDGPU::getValueMapping(Bank, 32);
break; break;
} }
case Intrinsic::amdgcn_global_load_lds: { case Intrinsic::amdgcn_global_load_lds: {
OpdsMapping[1] = getVGPROpMapping(MI.getOperand(1).getReg(), MRI, *TRI); OpdsMapping[1] = getVGPROpMapping(MI.getOperand(1).getReg(), MRI, *TRI);
OpdsMapping[2] = getSGPROpMapping(MI.getOperand(2).getReg(), MRI, *TRI); OpdsMapping[2] = getSGPROpMapping(MI.getOperand(2).getReg(), MRI, *TRI);
break; break;
} }
case Intrinsic::amdgcn_ds_add_gs_reg_rtn:
case Intrinsic::amdgcn_ds_sub_gs_reg_rtn:
OpdsMapping[0] = getVGPROpMapping(MI.getOperand(0).getReg(), MRI, *TRI);
OpdsMapping[2] = getVGPROpMapping(MI.getOperand(2).getReg(), MRI, *TRI);
break;
default: default:
return getInvalidInstructionMapping(); return getInvalidInstructionMapping();
} }
break; break;
} }
case AMDGPU::G_SELECT: { case AMDGPU::G_SELECT: {
unsigned Size = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits(); unsigned Size = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits();
unsigned Op2Bank = getRegBankID(MI.getOperand(2).getReg(), MRI, unsigned Op2Bank = getRegBankID(MI.getOperand(2).getReg(), MRI,
▲ Show 20 Lines • Show All 105 Lines • Show Last 20 Lines

llvm/lib/Target/AMDGPU/DSInstructions.td

Show First 20 Lines • Show All 1,106 Lines • ▼ Show 20 Lines
>; >;
} }
def : Pat < def : Pat <
(SIds_ordered_count i32:$value, i16:$offset), (SIds_ordered_count i32:$value, i16:$offset),
(DS_ORDERED_COUNT $value, (as_i16imm $offset)) (DS_ORDERED_COUNT $value, (as_i16imm $offset))
>; >;
def : GCNPat <
(i64 (int_amdgcn_ds_add_gs_reg_rtn i32:$src, timm:$offset32)),
(DS_ADD_GS_REG_RTN VGPR_32:$src, (as_i32timm $offset32))
>;
def : GCNPat <
(i32 (int_amdgcn_ds_add_gs_reg_rtn i32:$src, timm:$offset32)),
(EXTRACT_SUBREG
(i64 (COPY_TO_REGCLASS
(DS_ADD_GS_REG_RTN VGPR_32:$src, (as_i32timm $offset32)),
VReg_64)),
sub0)
>;
def : GCNPat <
(i64 (int_amdgcn_ds_sub_gs_reg_rtn i32:$src, timm:$offset32)),
(DS_SUB_GS_REG_RTN VGPR_32:$src, (as_i32timm $offset32))
>;
def : GCNPat <
(i32 (int_amdgcn_ds_sub_gs_reg_rtn i32:$src, timm:$offset32)),
(EXTRACT_SUBREG
(i64 (COPY_TO_REGCLASS
(DS_SUB_GS_REG_RTN VGPR_32:$src, (as_i32timm $offset32)),
VReg_64)),
sub0)
>;
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// Target-specific instruction encodings. // Target-specific instruction encodings.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// Base ENC_DS for GFX6, GFX7, GFX10, GFX11. // Base ENC_DS for GFX6, GFX7, GFX10, GFX11.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
▲ Show 20 Lines • Show All 505 Lines • Show Last 20 Lines

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ds.add.gs.reg.rtn.ll

  • This file was added.
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
; RUN: llc -march=amdgcn -mcpu=gfx1100 -verify-machineinstrs < %s | FileCheck %s
; RUN: llc -global-isel -march=amdgcn -mcpu=gfx1100 -verify-machineinstrs < %s | FileCheck %s
arsenmUnsubmitted
Not Done
ReplyInline Actions

These tests should use explicit -global-isel=0/1 for whenever the default switches

arsenm: These tests should use explicit -global-isel=0/1 for whenever the default switches
foadAuthorUnsubmitted
Done
ReplyInline Actions

Done in 03a991032571, but it's only the tip of the iceberg.

foad: Done in 03a991032571, but it's only the tip of the iceberg.
declare i32 @llvm.amdgcn.ds.add.gs.reg.rtn.i32(i32, i32 immarg)
declare i64 @llvm.amdgcn.ds.add.gs.reg.rtn.i64(i32, i32 immarg)
define amdgpu_gs void @test_add_32(i32 %arg) {
; CHECK-LABEL: test_add_32:
; CHECK: ; %bb.0:
; CHECK-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: ds_add_gs_reg_rtn v[0:1], v0 offset:16 gds
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: buffer_gl0_inv
; CHECK-NEXT: buffer_gl1_inv
; CHECK-NEXT: s_endpgm
%unused = call i32 @llvm.amdgcn.ds.add.gs.reg.rtn.i32(i32 %arg, i32 16)
ret void
}
define amdgpu_gs void @test_add_32_use(i32 %arg, i32 addrspace(1)* %out) {
; CHECK-LABEL: test_add_32_use:
; CHECK: ; %bb.0:
; CHECK-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: ds_add_gs_reg_rtn v[3:4], v0 offset:16 gds
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: buffer_gl0_inv
; CHECK-NEXT: buffer_gl1_inv
; CHECK-NEXT: global_store_b32 v[1:2], v3, off
; CHECK-NEXT: s_endpgm
%res = call i32 @llvm.amdgcn.ds.add.gs.reg.rtn.i32(i32 %arg, i32 16)
store i32 %res, i32 addrspace(1)* %out, align 4
ret void
}
define amdgpu_gs void @test_add_64(i32 %arg) {
; CHECK-LABEL: test_add_64:
; CHECK: ; %bb.0:
; CHECK-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: ds_add_gs_reg_rtn v[0:1], v0 offset:32 gds
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: buffer_gl0_inv
; CHECK-NEXT: buffer_gl1_inv
; CHECK-NEXT: s_endpgm
%unused = call i64 @llvm.amdgcn.ds.add.gs.reg.rtn.i64(i32 %arg, i32 32)
ret void
}
define amdgpu_gs void @test_add_64_use(i32 %arg, i64 addrspace(1)* %out) {
; CHECK-LABEL: test_add_64_use:
; CHECK: ; %bb.0:
; CHECK-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: ds_add_gs_reg_rtn v[3:4], v0 offset:32 gds
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: buffer_gl0_inv
; CHECK-NEXT: buffer_gl1_inv
; CHECK-NEXT: global_store_b64 v[1:2], v[3:4], off
; CHECK-NEXT: s_endpgm
%res = call i64 @llvm.amdgcn.ds.add.gs.reg.rtn.i64(i32 %arg, i32 32)
store i64 %res, i64 addrspace(1)* %out, align 4
ret void
}

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ds.sub.gs.reg.rtn.ll

  • This file was added.
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
; RUN: llc -march=amdgcn -mcpu=gfx1100 -verify-machineinstrs < %s | FileCheck %s
; RUN: llc -global-isel -march=amdgcn -mcpu=gfx1100 -verify-machineinstrs < %s | FileCheck %s
declare i32 @llvm.amdgcn.ds.sub.gs.reg.rtn.i32(i32, i32 immarg)
declare i64 @llvm.amdgcn.ds.sub.gs.reg.rtn.i64(i32, i32 immarg)
define amdgpu_gs void @test_sub_32(i32 %arg) {
; CHECK-LABEL: test_sub_32:
; CHECK: ; %bb.0:
; CHECK-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: ds_sub_gs_reg_rtn v[0:1], v0 offset:16 gds
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: buffer_gl0_inv
; CHECK-NEXT: buffer_gl1_inv
; CHECK-NEXT: s_endpgm
%unused = call i32 @llvm.amdgcn.ds.sub.gs.reg.rtn.i32(i32 %arg, i32 16)
ret void
}
define amdgpu_gs void @test_sub_32_use(i32 %arg, i32 addrspace(1)* %out) {
; CHECK-LABEL: test_sub_32_use:
; CHECK: ; %bb.0:
; CHECK-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: ds_sub_gs_reg_rtn v[3:4], v0 offset:16 gds
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: buffer_gl0_inv
; CHECK-NEXT: buffer_gl1_inv
; CHECK-NEXT: global_store_b32 v[1:2], v3, off
; CHECK-NEXT: s_endpgm
%res = call i32 @llvm.amdgcn.ds.sub.gs.reg.rtn.i32(i32 %arg, i32 16)
store i32 %res, i32 addrspace(1)* %out, align 4
ret void
}
define amdgpu_gs void @test_sub_64(i32 %arg) {
; CHECK-LABEL: test_sub_64:
; CHECK: ; %bb.0:
; CHECK-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: ds_sub_gs_reg_rtn v[0:1], v0 offset:32 gds
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: buffer_gl0_inv
; CHECK-NEXT: buffer_gl1_inv
; CHECK-NEXT: s_endpgm
%unused = call i64 @llvm.amdgcn.ds.sub.gs.reg.rtn.i64(i32 %arg, i32 32)
ret void
}
define amdgpu_gs void @test_sub_64_use(i32 %arg, i64 addrspace(1)* %out) {
; CHECK-LABEL: test_sub_64_use:
; CHECK: ; %bb.0:
; CHECK-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: ds_sub_gs_reg_rtn v[3:4], v0 offset:32 gds
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: s_waitcnt_vscnt null, 0x0
; CHECK-NEXT: buffer_gl0_inv
; CHECK-NEXT: buffer_gl1_inv
; CHECK-NEXT: global_store_b64 v[1:2], v[3:4], off
; CHECK-NEXT: s_endpgm
%res = call i64 @llvm.amdgcn.ds.sub.gs.reg.rtn.i64(i32 %arg, i32 32)
store i64 %res, i64 addrspace(1)* %out, align 4
ret void
}