[llvm] r248532 - AMDGPU: Add cache invalidation instructions.
Matt Arsenault via llvm-commits
llvm-commits at lists.llvm.org
Thu Sep 24 12:52:21 PDT 2015
Author: arsenm
Date: Thu Sep 24 14:52:21 2015
New Revision: 248532
URL: http://llvm.org/viewvc/llvm-project?rev=248532&view=rev
Log:
AMDGPU: Add cache invalidation instructions.
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.
Added:
llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.ll
llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.sc.ll
llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.vol.ll
llvm/trunk/test/MC/AMDGPU/buffer_wbinv1l_vol_vi.s
Modified:
llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/trunk/lib/Target/AMDGPU/CIInstructions.td
llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td
llvm/trunk/lib/Target/AMDGPU/SIInstructions.td
llvm/trunk/test/MC/AMDGPU/mubuf.s
Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=248532&r1=248531&r2=248532&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Thu Sep 24 14:52:21 2015
@@ -83,3 +83,21 @@ def int_AMDGPU_read_workdim : AMDGPURead
"__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<[], [], []>;
+
+}
Modified: llvm/trunk/lib/Target/AMDGPU/CIInstructions.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/CIInstructions.td?rev=248532&r1=248531&r2=248532&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/CIInstructions.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/CIInstructions.td Thu Sep 24 14:52:21 2015
@@ -100,6 +100,14 @@ defm DS_WRAP_RTN_F32 : DS_1A1D_RET <0x34
// 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
//===----------------------------------------------------------------------===//
Modified: llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td?rev=248532&r1=248531&r2=248532&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td Thu Sep 24 14:52:21 2015
@@ -2455,6 +2455,23 @@ multiclass MUBUF_Store_Helper <mubuf op,
} // 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),
Modified: llvm/trunk/lib/Target/AMDGPU/SIInstructions.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstructions.td?rev=248532&r1=248531&r2=248532&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIInstructions.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIInstructions.td Thu Sep 24 14:52:21 2015
@@ -30,7 +30,9 @@ def isGCN : Predicate<"Subtarget->getGen
">= 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
Added: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.ll?rev=248532&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.ll (added)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.ll Thu Sep 24 14:52:21 2015
@@ -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 }
Added: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.sc.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.sc.ll?rev=248532&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.sc.ll (added)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.sc.ll Thu Sep 24 14:52:21 2015
@@ -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 }
Added: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.vol.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.vol.ll?rev=248532&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.vol.ll (added)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.vol.ll Thu Sep 24 14:52:21 2015
@@ -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 }
Added: llvm/trunk/test/MC/AMDGPU/buffer_wbinv1l_vol_vi.s
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/MC/AMDGPU/buffer_wbinv1l_vol_vi.s?rev=248532&view=auto
==============================================================================
--- llvm/trunk/test/MC/AMDGPU/buffer_wbinv1l_vol_vi.s (added)
+++ llvm/trunk/test/MC/AMDGPU/buffer_wbinv1l_vol_vi.s Thu Sep 24 14:52:21 2015
@@ -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]
Modified: llvm/trunk/test/MC/AMDGPU/mubuf.s
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/MC/AMDGPU/mubuf.s?rev=248532&r1=248531&r2=248532&view=diff
==============================================================================
--- llvm/trunk/test/MC/AMDGPU/mubuf.s (original)
+++ llvm/trunk/test/MC/AMDGPU/mubuf.s Thu Sep 24 14:52:21 2015
@@ -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
More information about the llvm-commits
mailing list