llvm.org GIT mirror llvm / 1348e9d
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. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@248532 91177308-0d34-0410-b5e6-96231b3b80d8 Matt Arsenault 5 years ago
9 changed file(s) with 127 addition(s) and 6 deletion(s). Raw diff Collapse all Expand all
8282 "__builtin_amdgpu_read_workdim">;
8383
8484 } // End TargetPrefix = "AMDGPU"
85
86 let TargetPrefix = "amdgcn" in {
87
88 // SI only
89 def int_amdgcn_buffer_wbinvl1_sc :
90 GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
91 Intrinsic<[], [], []>;
92
93 // On CI+
94 def int_amdgcn_buffer_wbinvl1_vol :
95 GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
96 Intrinsic<[], [], []>;
97
98 def int_amdgcn_buffer_wbinvl1 :
99 GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
100 Intrinsic<[], [], []>;
101
102 }
9999 // DS_CONDXCHG32_RTN_B128
100100
101101 //===----------------------------------------------------------------------===//
102 // MUBUF Instructions
103 //===----------------------------------------------------------------------===//
104
105 defm BUFFER_WBINVL1_VOL : MUBUF_Invalidate ,
106 "buffer_wbinvl1_vol", int_amdgcn_buffer_wbinvl1_vol
107 >;
108
109 //===----------------------------------------------------------------------===//
102110 // Flat Instructions
103111 //===----------------------------------------------------------------------===//
104112
24542454 } // End mayLoad = 0, mayStore = 1
24552455 }
24562456
2457 // For cache invalidation instructions.
2458 multiclass MUBUF_Invalidate {
2459 let hasSideEffects = 1, mayStore = 1, AsmMatchConverter = "" in {
2460 def "" : MUBUF_Pseudo ;
2461
2462 // Set everything to 0.
2463 let offset = 0, offen = 0, idxen = 0, glc = 0, vaddr = 0,
2464 vdata = 0, srsrc = 0, slc = 0, tfe = 0, soffset = 0 in {
2465 let addr64 = 0 in {
2466 def _si : MUBUF_Real_si ;
2467 }
2468
2469 def _vi : MUBUF_Real_vi ;
2470 }
2471 } // End hasSideEffects = 1, mayStore = 1, AsmMatchConverter = ""
2472 }
2473
24572474 class FLAT_Load_Helper op, string asm, RegisterClass regClass> :
24582475 FLAT
24592476 (ins VReg_64:$addr, glc_flat:$glc, slc_flat:$slc, tfe_flat:$tfe),
2929 ">= AMDGPUSubtarget::SOUTHERN_ISLANDS">,
3030 AssemblerPredicate<"FeatureGCN">;
3131 def isSI : Predicate<"Subtarget->getGeneration() "
32 "== AMDGPUSubtarget::SOUTHERN_ISLANDS">;
32 "== AMDGPUSubtarget::SOUTHERN_ISLANDS">,
33 AssemblerPredicate<"FeatureSouthernIslands">;
34
3335
3436 def has16BankLDS : Predicate<"Subtarget->getLDSBankCount() == 16">;
3537 def has32BankLDS : Predicate<"Subtarget->getLDSBankCount() == 32">;
10271029 //def BUFFER_ATOMIC_FCMPSWAP_X2 : MUBUF_X2 , "buffer_atomic_fcmpswap_x2", []>; // isn't on VI
10281030 //def BUFFER_ATOMIC_FMIN_X2 : MUBUF_X2 , "buffer_atomic_fmin_x2", []>; // isn't on VI
10291031 //def BUFFER_ATOMIC_FMAX_X2 : MUBUF_X2 , "buffer_atomic_fmax_x2", []>; // isn't on VI
1030 //def BUFFER_WBINVL1_SC : MUBUF_WBINVL1 , "buffer_wbinvl1_sc", []>; // isn't on CI & VI
1031 //def BUFFER_WBINVL1_VOL : MUBUF_WBINVL1 , "buffer_wbinvl1_vol", []>; // isn't on SI
1032 //def BUFFER_WBINVL1 : MUBUF_WBINVL1 , "buffer_wbinvl1", []>;
1032
1033 let SubtargetPredicate = isSI in {
1034 defm BUFFER_WBINVL1_SC : MUBUF_Invalidate , "buffer_wbinvl1_sc", int_amdgcn_buffer_wbinvl1_sc>; // isn't on CI & VI
1035 }
1036
1037 defm BUFFER_WBINVL1 : MUBUF_Invalidate , "buffer_wbinvl1", int_amdgcn_buffer_wbinvl1>;
10331038
10341039 //===----------------------------------------------------------------------===//
10351040 // MTBUF Instructions
0 ; RUN: llc -march=amdgcn -mcpu=tahiti -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
1 ; RUN: llc -march=amdgcn -mcpu=fiji -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
2
3 declare void @llvm.amdgcn.buffer.wbinvl1() #0
4
5 ; GCN-LABEL: {{^}}test_buffer_wbinvl1:
6 ; GCN-NEXT: ; BB#0:
7 ; SI-NEXT: buffer_wbinvl1 ; encoding: [0x00,0x00,0xc4,0xe1,0x00,0x00,0x00,0x00]
8 ; VI-NEXT: buffer_wbinvl1 ; encoding: [0x00,0x00,0xf8,0xe0,0x00,0x00,0x00,0x00]
9 ; GCN-NEXT: s_endpgm
10 define void @test_buffer_wbinvl1() #0 {
11 call void @llvm.amdgcn.buffer.wbinvl1()
12 ret void
13 }
14
15 attributes #0 = { nounwind }
0 ; RUN: llc -march=amdgcn -mcpu=tahiti -show-mc-encoding < %s | FileCheck -check-prefix=SI %s
1
2 declare void @llvm.amdgcn.buffer.wbinvl1.sc() #0
3
4 ; SI-LABEL: {{^}}test_buffer_wbinvl1_sc:
5 ; SI-NEXT: ; BB#0:
6 ; SI-NEXT: buffer_wbinvl1_sc ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
7 ; SI-NEXT: s_endpgm
8 define void @test_buffer_wbinvl1_sc() #0 {
9 call void @llvm.amdgcn.buffer.wbinvl1.sc()
10 ret void
11 }
12
13 attributes #0 = { nounwind }
0 ; RUN: llc -march=amdgcn -mcpu=bonaire -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=CI %s
1 ; RUN: llc -march=amdgcn -mcpu=tonga -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
2
3 declare void @llvm.amdgcn.buffer.wbinvl1.vol() #0
4
5 ; GCN-LABEL: {{^}}test_buffer_wbinvl1_vol:
6 ; GCN-NEXT: ; BB#0:
7 ; CI-NEXT: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
8 ; VI-NEXT: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xfc,0xe0,0x00,0x00,0x00,0x00]
9 ; GCN-NEXT: s_endpgm
10 define void @test_buffer_wbinvl1_vol() #0 {
11 call void @llvm.amdgcn.buffer.wbinvl1.vol()
12 ret void
13 }
14
15 attributes #0 = { nounwind }
0 // XFAIL: *
1 // RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck -check-prefix=VI %s
2
3 ; When assembled, this emits a different encoding value than codegen for the intrinsic
4
5 buffer_wbinvl1_vol
6 // VI: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xfc,0xe0,0x00,0x00,0x00,0x00]
None // RUN: llvm-mc -arch=amdgcn -mcpu=tahiti -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=SI -check-prefix=SICI %s
1 // RUN: llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=CI -check-prefix=SICI %s
0 // RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=SI -check-prefix=SICI %s
1 // RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=CI -check-prefix=SICI %s
2
3 // RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti %s 2>&1 | FileCheck -check-prefix=NOSI %s
4 // RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck -check-prefix=NOCI %s
5 // RUN: not llvm-mc -arch=amdgcn -mcpu=tonga %s 2>&1 | FileCheck -check-prefix=NOVI %s
26
37 //===----------------------------------------------------------------------===//
48 // Test for different operand combinations
348352 buffer_store_dwordx4 v[1:4], s[4:7], s1
349353 // SICI: buffer_store_dwordx4 v[1:4], s[4:7], s1 ; encoding: [0x00,0x00,0x78,0xe0,0x00,0x01,0x01,0x01]
350354
355 //===----------------------------------------------------------------------===//
356 // Cache invalidation
357 //===----------------------------------------------------------------------===//
358
359 buffer_wbinvl1
360 // SICI: buffer_wbinvl1 ; encoding: [0x00,0x00,0xc4,0xe1,0x00,0x00,0x00,0x00]
361
362 buffer_wbinvl1_sc
363 // SI: buffer_wbinvl1_sc ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
364 // NOCI: error: instruction not supported on this GPU
365 // NOVI: error: instruction not supported on this GPU
366
367 buffer_wbinvl1_vol
368 // CI: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
369 // NOSI: error: instruction not supported on this GPU
370
351371 // TODO: Atomics