diff options
| author | Tom Stellard <thomas.stellard@amd.com> | 2016-02-13 00:29:57 +0000 |
|---|---|---|
| committer | Tom Stellard <thomas.stellard@amd.com> | 2016-02-13 00:29:57 +0000 |
| commit | 1debba54a71804331f995f4d7dae2a7d77dd9b26 (patch) | |
| tree | 360c283eab00e360f4e97081cc1bc7b019ac47e3 | |
| parent | 9fb970e422e8efef841b7027c2195a738585ab16 (diff) | |
| download | bcm5719-llvm-1debba54a71804331f995f4d7dae2a7d77dd9b26.tar.gz bcm5719-llvm-1debba54a71804331f995f4d7dae2a7d77dd9b26.zip | |
AMDGPU/SI: Organize intrinsics by subtarget
Reviewers: arsenm
Subscribers: llvm-commits
Differential Revision: http://reviews.llvm.org/D17210
llvm-svn: 260771
| -rw-r--r-- | llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 45 |
1 files changed, 25 insertions, 20 deletions
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index e8d1698cc9e..712dcb03538 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -139,11 +139,6 @@ 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<[], [], []>; @@ -152,21 +147,6 @@ def int_amdgcn_s_dcache_inv : GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">, Intrinsic<[], [], []>; -// CI+ -def int_amdgcn_s_dcache_inv_vol : - GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">, - Intrinsic<[], [], []>; - -// VI -def int_amdgcn_s_dcache_wb : - GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">, - Intrinsic<[], [], []>; - -// VI -def int_amdgcn_s_dcache_wb_vol : - GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">, - Intrinsic<[], [], []>; - def int_amdgcn_dispatch_ptr : GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">, Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; @@ -194,4 +174,29 @@ def int_amdgcn_mbcnt_lo : def int_amdgcn_mbcnt_hi : GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>; + +//===----------------------------------------------------------------------===// +// CI+ Intrinsics +//===----------------------------------------------------------------------===// + +def int_amdgcn_s_dcache_inv_vol : + GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">, + Intrinsic<[], [], []>; + +def int_amdgcn_buffer_wbinvl1_vol : + GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">, + Intrinsic<[], [], []>; + +//===----------------------------------------------------------------------===// +// VI Intrinsics +//===----------------------------------------------------------------------===// + +def int_amdgcn_s_dcache_wb : + GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">, + Intrinsic<[], [], []>; + +def int_amdgcn_s_dcache_wb_vol : + GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">, + Intrinsic<[], [], []>; + } |
