Matt Arsenault
24b5ae4497
AMDGPU: Add builtin for getreg intrinsic
...
llvm-svn: 292636
2017-01-20 19:24:22 +00:00
Konstantin Zhuravlyov
62ae8f671c
[AMDGPU] Change frexp.exp builtin to return i16 for f16 input
...
Differential Revision: https://reviews.llvm.org/D26863
llvm-svn: 287390
2016-11-18 22:31:51 +00:00
Stanislav Mekhanoshin
cd433d2811
[AMDGPU] Add wave barrier builtin
...
The wave barrier represents the discardable barrier. Its main purpose is to
carry convergent attribute, thus preventing illegal CFG optimizations. All lanes
in a wave come to convergence point simultaneously with SIMT, thus no special
instruction is needed in the ISA. The barrier is discarded during code generation.
Differential Revision: https://reviews.llvm.org/D26584
llvm-svn: 287006
2016-11-15 18:58:03 +00:00
Valery Pykhtin
4b5d9d16d3
[AMDGPU] add s_incperflevel/s_decperflevel builtins
...
Differential revision: https://reviews.llvm.org/D23668
llvm-svn: 279235
2016-08-19 12:54:31 +00:00
Changpeng Fang
03bdd8f797
AMDGPU: Add clang builtin for ds_swizzle.
...
Summary:
int __builtin_amdgcn_ds_swizzle (int a, int imm);
while imm is a constant.
Differential Revision:
http://reviews.llvm.org/D23682
llvm-svn: 279165
2016-08-18 22:04:54 +00:00
Wei Ding
91c8450967
AMDGPU : Add Clang builtin intrinsics for compare with the full
...
wavefront result.
Differential Revision: http://reviews.llvm.org/D22934
llvm-svn: 277824
2016-08-05 15:38:46 +00:00
Matt Arsenault
c7536a5d60
AMDGPU: Remove legacy ldexp builtin
...
llvm-svn: 275623
2016-07-15 21:33:06 +00:00
Matt Arsenault
c86671da09
AMDGPU: Update for rsq intrinsic changes
...
llvm-svn: 275622
2016-07-15 21:33:02 +00:00
Wei Ding
ea41f356bb
AMDGPU: Add Clang Builtin for v_lerp_u8
...
Differential Revision: http://reviews.llvm.org/D22380
llvm-svn: 275577
2016-07-15 16:43:03 +00:00
Jan Vesely
d7e03a5bd9
AMDGPU: Export workitem builtins
...
Reviewers: tstellardAMD
Differential Revision: http://reviews.llvm.org/D20299
llvm-svn: 275030
2016-07-10 22:38:04 +00:00
Matt Arsenault
64665bc50d
AMDGPU: Add builtin to read exec mask
...
llvm-svn: 273965
2016-06-28 00:13:17 +00:00
Matt Arsenault
250024f905
AMDGPU: Verify subtarget specific builtins
...
Cleanup setup of subtarget features.
llvm-svn: 272091
2016-06-08 01:56:42 +00:00
Matt Arsenault
2d51059ebb
AMDGPU: Add fract builtin
...
llvm-svn: 271080
2016-05-28 00:43:27 +00:00
Matt Arsenault
3fb963389e
AMDGPU: Add frexp_mant + frexp_exp builtins
...
llvm-svn: 264960
2016-03-30 22:57:40 +00:00
Matt Arsenault
39edcd0e1d
AMDGPU: Add builtins for recently added intrinsics
...
llvm-svn: 262126
2016-02-27 09:54:43 +00:00
Matt Arsenault
9b277b4ad4
AMDGPU: Add sin/cos builtins
...
llvm-svn: 260783
2016-02-13 01:21:09 +00:00
Matt Arsenault
f5c1f47181
AMDGPU: Update builtin for intrinsic change
...
llvm-svn: 260781
2016-02-13 01:03:09 +00:00
Matt Arsenault
cf70cb9d00
AMDGPU: Add amdgcn cube builtins
...
llvm-svn: 258794
2016-01-26 06:37:54 +00:00
Matt Arsenault
721d21b821
AMDGPU: Add barrier builtin
...
llvm-svn: 258564
2016-01-22 21:56:30 +00:00
Matt Arsenault
8a4078c741
AMDGPU: Rename builtins to use amdgcn prefix
...
Keep the ones still used by libclc around for now.
Emit the new amdgcn intrinsic name if not targeting r600,
in which case the old AMDGPU name is still used.
llvm-svn: 258560
2016-01-22 21:30:53 +00:00