RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-11-23 11:13:28 +01:00
Code Issues Projects Releases Wiki Activity

AMDGPU: Add cache invalidation instructions.

Browse Source 
These are necessary for implementing mem_fence for
OpenCL 2.0.

The VI assembler tests are disabled since it seems to be
using the wrong encoding or opcode.

llvm-svn: 248532
This commit is contained in:
Matt Arsenault 2015-09-24 19:52:21 +00:00
parent e10636a9fb
commit 304779755d
9 changed files with 127 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

18
include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -83,3 +83,21 @@ def int_AMDGPU_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
"__builtin_amdgpu_read_workdim">;
} // End TargetPrefix = "AMDGPU"
let TargetPrefix = "amdgcn" in {
// SI only
def int_amdgcn_buffer_wbinvl1_sc :
GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
Intrinsic<[], [], []>;
// On CI+
def int_amdgcn_buffer_wbinvl1_vol :
GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
Intrinsic<[], [], []>;
def int_amdgcn_buffer_wbinvl1 :
GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
Intrinsic<[], [], []>;
}

8
lib/Target/AMDGPU/CIInstructions.td
View File

@ -99,6 +99,14 @@ defm DS_WRAP_RTN_F32 : DS_1A1D_RET <0x34, "ds_wrap_rtn_f32", VGPR_32, "ds_wrap_f
// DS_CONDXCHG32_RTN_B64
// DS_CONDXCHG32_RTN_B128
//===----------------------------------------------------------------------===//
// MUBUF Instructions
//===----------------------------------------------------------------------===//
defm BUFFER_WBINVL1_VOL : MUBUF_Invalidate <mubuf<0x70, 0x3f>,
"buffer_wbinvl1_vol", int_amdgcn_buffer_wbinvl1_vol
>;
//===----------------------------------------------------------------------===//
// Flat Instructions
//===----------------------------------------------------------------------===//

17
lib/Target/AMDGPU/SIInstrInfo.td
View File

@ -2455,6 +2455,23 @@ multiclass MUBUF_Store_Helper <mubuf op, string name, RegisterClass vdataClass,
} // End mayLoad = 0, mayStore = 1
}
// For cache invalidation instructions.
multiclass MUBUF_Invalidate <mubuf op, string opName, SDPatternOperator node> {
let hasSideEffects = 1, mayStore = 1, AsmMatchConverter = "" in {
def "" : MUBUF_Pseudo <opName, (outs), (ins), [(node)]>;
// Set everything to 0.
let offset = 0, offen = 0, idxen = 0, glc = 0, vaddr = 0,
vdata = 0, srsrc = 0, slc = 0, tfe = 0, soffset = 0 in {
let addr64 = 0 in {
def _si : MUBUF_Real_si <op, opName, (outs), (ins), opName>;
}
def _vi : MUBUF_Real_vi <op, opName, (outs), (ins), opName>;
}
} // End hasSideEffects = 1, mayStore = 1, AsmMatchConverter = ""
}
class FLAT_Load_Helper <bits<7> op, string asm, RegisterClass regClass> :
FLAT <op, (outs regClass:$vdst),
(ins VReg_64:$addr, glc_flat:$glc, slc_flat:$slc, tfe_flat:$tfe),

13
lib/Target/AMDGPU/SIInstructions.td
View File

@ -30,7 +30,9 @@ def isGCN : Predicate<"Subtarget->getGeneration() "
">= AMDGPUSubtarget::SOUTHERN_ISLANDS">,
AssemblerPredicate<"FeatureGCN">;
def isSI : Predicate<"Subtarget->getGeneration() "
"== AMDGPUSubtarget::SOUTHERN_ISLANDS">;
"== AMDGPUSubtarget::SOUTHERN_ISLANDS">,
AssemblerPredicate<"FeatureSouthernIslands">;
def has16BankLDS : Predicate<"Subtarget->getLDSBankCount() == 16">;
def has32BankLDS : Predicate<"Subtarget->getLDSBankCount() == 32">;
@ -1028,9 +1030,12 @@ defm BUFFER_ATOMIC_XOR : MUBUF_Atomic <
//def BUFFER_ATOMIC_FCMPSWAP_X2 : MUBUF_X2 <mubuf<0x5e>, "buffer_atomic_fcmpswap_x2", []>; // isn't on VI
//def BUFFER_ATOMIC_FMIN_X2 : MUBUF_X2 <mubuf<0x5f>, "buffer_atomic_fmin_x2", []>; // isn't on VI
//def BUFFER_ATOMIC_FMAX_X2 : MUBUF_X2 <mubuf<0x60>, "buffer_atomic_fmax_x2", []>; // isn't on VI
//def BUFFER_WBINVL1_SC : MUBUF_WBINVL1 <mubuf<0x70>, "buffer_wbinvl1_sc", []>; // isn't on CI & VI
//def BUFFER_WBINVL1_VOL : MUBUF_WBINVL1 <mubuf<0x70, 0x3f>, "buffer_wbinvl1_vol", []>; // isn't on SI
//def BUFFER_WBINVL1 : MUBUF_WBINVL1 <mubuf<0x71, 0x3e>, "buffer_wbinvl1", []>;
let SubtargetPredicate = isSI in {
defm BUFFER_WBINVL1_SC : MUBUF_Invalidate <mubuf<0x70>, "buffer_wbinvl1_sc", int_amdgcn_buffer_wbinvl1_sc>; // isn't on CI & VI
}
defm BUFFER_WBINVL1 : MUBUF_Invalidate <mubuf<0x71, 0x3e>, "buffer_wbinvl1", int_amdgcn_buffer_wbinvl1>;
//===----------------------------------------------------------------------===//
// MTBUF Instructions

16
test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.ll Normal file
View File

@ -0,0 +1,16 @@
; RUN: llc -march=amdgcn -mcpu=tahiti -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
; RUN: llc -march=amdgcn -mcpu=fiji -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
declare void @llvm.amdgcn.buffer.wbinvl1() #0
; GCN-LABEL: {{^}}test_buffer_wbinvl1:
; GCN-NEXT: ; BB#0:
; SI-NEXT: buffer_wbinvl1 ; encoding: [0x00,0x00,0xc4,0xe1,0x00,0x00,0x00,0x00]
; VI-NEXT: buffer_wbinvl1 ; encoding: [0x00,0x00,0xf8,0xe0,0x00,0x00,0x00,0x00]
; GCN-NEXT: s_endpgm
define void @test_buffer_wbinvl1() #0 {
call void @llvm.amdgcn.buffer.wbinvl1()
ret void
}
attributes #0 = { nounwind }

14
test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.sc.ll Normal file
View File

@ -0,0 +1,14 @@
; RUN: llc -march=amdgcn -mcpu=tahiti -show-mc-encoding < %s | FileCheck -check-prefix=SI %s
declare void @llvm.amdgcn.buffer.wbinvl1.sc() #0
; SI-LABEL: {{^}}test_buffer_wbinvl1_sc:
; SI-NEXT: ; BB#0:
; SI-NEXT: buffer_wbinvl1_sc ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
; SI-NEXT: s_endpgm
define void @test_buffer_wbinvl1_sc() #0 {
call void @llvm.amdgcn.buffer.wbinvl1.sc()
ret void
}
attributes #0 = { nounwind }

16
test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.vol.ll Normal file
View File

@ -0,0 +1,16 @@
; RUN: llc -march=amdgcn -mcpu=bonaire -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=CI %s
; RUN: llc -march=amdgcn -mcpu=tonga -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
declare void @llvm.amdgcn.buffer.wbinvl1.vol() #0
; GCN-LABEL: {{^}}test_buffer_wbinvl1_vol:
; GCN-NEXT: ; BB#0:
; CI-NEXT: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
; VI-NEXT: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xfc,0xe0,0x00,0x00,0x00,0x00]
; GCN-NEXT: s_endpgm
define void @test_buffer_wbinvl1_vol() #0 {
call void @llvm.amdgcn.buffer.wbinvl1.vol()
ret void
}
attributes #0 = { nounwind }

7
test/MC/AMDGPU/buffer_wbinv1l_vol_vi.s Normal file
View File

@ -0,0 +1,7 @@
// XFAIL: *
// RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck -check-prefix=VI %s
; When assembled, this emits a different encoding value than codegen for the intrinsic
buffer_wbinvl1_vol
// VI: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xfc,0xe0,0x00,0x00,0x00,0x00]

24
test/MC/AMDGPU/mubuf.s
View File

@ -1,5 +1,9 @@
// RUN: llvm-mc -arch=amdgcn -mcpu=tahiti -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=SI -check-prefix=SICI %s
// RUN: llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=CI -check-prefix=SICI %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=SI -check-prefix=SICI %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=CI -check-prefix=SICI %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti %s 2>&1 | FileCheck -check-prefix=NOSI %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck -check-prefix=NOCI %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=tonga %s 2>&1 | FileCheck -check-prefix=NOVI %s
//===----------------------------------------------------------------------===//
// Test for different operand combinations
@ -349,4 +353,20 @@ buffer_store_dwordx2 v[1:2], s[4:7], s1
buffer_store_dwordx4 v[1:4], s[4:7], s1
// SICI: buffer_store_dwordx4 v[1:4], s[4:7], s1 ; encoding: [0x00,0x00,0x78,0xe0,0x00,0x01,0x01,0x01]
//===----------------------------------------------------------------------===//
// Cache invalidation
//===----------------------------------------------------------------------===//
buffer_wbinvl1
// SICI: buffer_wbinvl1 ; encoding: [0x00,0x00,0xc4,0xe1,0x00,0x00,0x00,0x00]
buffer_wbinvl1_sc
// SI: buffer_wbinvl1_sc ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
// NOCI: error: instruction not supported on this GPU
// NOVI: error: instruction not supported on this GPU
buffer_wbinvl1_vol
// CI: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
// NOSI: error: instruction not supported on this GPU
// TODO: Atomics