replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity
263,234 Commits 2 Branches 0 Tags 1.8 GiB
Commit Graph

29 Commits

Author SHA1 Message Date
Reid Kleckner 581a6c5d56 Revert "[AMDGPU] add __builtin_amdgcn_s_getpc" 
This reverts commit r303861, the LLVM intrinsic was reverted.

llvm-svn: 303908
 2017-05-25 20:28:26 +00:00
Tim Corringham 702fe45bcd [AMDGPU] add __builtin_amdgcn_s_getpc 
Summary: Added the builtin corresponding to the s_getpc intrinsic added in llvm D32862

Subscribers: kzhuravl, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye

Differential Revision: https://reviews.llvm.org/D33276

llvm-svn: 303861
 2017-05-25 14:16:11 +00:00
Yaxun Liu af3d4db64b [AMDGPU] Do not require opencl triple environment for OpenCL 
A recent change requires opencl triple environment for compiling OpenCL
program, which causes regressions in libclc.

This patch fixes that. Instead of deducing language based on triple
environment, it checks LangOptions.

Differential Revision: https://reviews.llvm.org/D33445

llvm-svn: 303644
 2017-05-23 16:15:53 +00:00
Yaxun Liu 6d96f16347 CodeGen: Cast alloca to expected address space 
Alloca always returns a pointer in alloca address space, which may
be different from the type defined by the language. For example,
in C++ the auto variables are in the default address space. Therefore
cast alloca to the expected address space when necessary.

Differential Revision: https://reviews.llvm.org/D32248

llvm-svn: 303370
 2017-05-18 18:51:09 +00:00
Yaxun Liu 4d86799219 [AMDGPU] Add builtin functions readlane ds_permute mov_dpp 
Differential Revision: https://reviews.llvm.org/D30551

llvm-svn: 297436
 2017-03-10 01:30:46 +00:00
Jan Vesely 9488560bb8 AMDGPU: export s_sendmsg{halt} instrinsics 
Differential Revision: https://reviews.llvm.org/D30366

llvm-svn: 296241
 2017-02-25 04:20:24 +00:00
Jan Vesely c255097517 AMDGPU: export l1 cache invalidation intrinsics 
Differential Revision: https://reviews.llvm.org/D30360

llvm-svn: 296240
 2017-02-25 04:20:22 +00:00
Jan Vesely d26dbb389f AMDGPU: export s_waitcnt builtin 
Differential Revision: https://reviews.llvm.org/D30359

llvm-svn: 296239
 2017-02-25 04:20:20 +00:00
Matt Arsenault a274b209f5 AMDGPU: Add builtin for fmed3 intrinsic 
llvm-svn: 293600
 2017-01-31 03:42:07 +00:00
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