Commit Graph

49 Commits

Author SHA1 Message Date
Matt Arsenault 97f3f0bab0 AMDGPU: Add intrinsic for s_setreg
This will be more useful with fenv access implemented.
2020-05-28 14:26:38 -04:00
Yaxun (Sam) Liu b72fce1ffd Fix __builtin_amdgcn_workgroup_size_x/y/z return type
https://reviews.llvm.org/D77390
2020-04-03 09:56:30 -04:00
Yaxun (Sam) Liu 369e26ca9e [AMDGPU] Add __builtin_amdgcn_workgroup_size_x/y/z
The main purpose of introducing these builtins is to add a range
metadata [1, 1025) on the work group size loaded from dispatch
ptr, which cannot be done by source code.

Differential Revision: https://reviews.llvm.org/D76772
2020-03-28 01:03:20 -04:00
Matt Arsenault 3f533006ba AMDGPU: Emit llvm.fshr for __builtin_amdgcn_alignbit
These are equivalent. The generic rotate builtins do not directly map
to the fshr intrinsic.
2020-03-23 16:51:25 -04:00
Yaxun (Sam) Liu a57d9652a0 Make __builtin_amdgcn_dispatch_ptr dereferenceable and align at 4
Differential Revision: https://reviews.llvm.org/D75028
2020-02-25 13:58:20 -05:00
Matt Arsenault 64d7af09f5 AMDGPU: Add missing builtin declarations
llvm-svn: 367431
2019-07-31 14:03:05 +00:00
Matt Arsenault e56865d40c AMDGPU: Add some missing builtins
llvm-svn: 366286
2019-07-17 00:01:03 +00:00
Matt Arsenault 5495f78165 AMDGPU: Fix missing declaration for mbcnt builtins
llvm-svn: 364251
2019-06-24 23:34:06 +00:00
Matt Arsenault bcdbc9a115 AMDGPU: Add DS GWS sema builtins
llvm-svn: 363986
2019-06-20 21:33:57 +00:00
Matt Arsenault f46f41411b Reapply "r363684: AMDGPU: Add GWS instruction builtins"
llvm-svn: 363871
2019-06-19 19:55:49 +00:00
Simon Pilgrim 6828bc5614 Revert rL363684 : AMDGPU: Add GWS instruction builtins
........
Depends on rL363678 which was reverted at rL363797

llvm-svn: 363824
2019-06-19 15:35:45 +00:00
Matt Arsenault 2acc717627 AMDGPU: Add GWS instruction builtins
llvm-svn: 363684
2019-06-18 14:10:01 +00:00
Stanislav Mekhanoshin 8a8131a3f6 [AMDGPU] gfx1010 wave32 clang support
Differential Revision: https://reviews.llvm.org/D63209

llvm-svn: 363341
2019-06-13 23:47:59 +00:00
Matt Arsenault b72888647b AMDGPU: Add ds append/consume builtins
llvm-svn: 352443
2019-01-28 23:59:18 +00:00
Matt Arsenault e3d81572c1 AMDGPU: Fix missing declaration of queue ptr builtin
llvm-svn: 338754
2018-08-02 18:24:55 +00:00
Matt Arsenault c65f966d76 Try to make builtin address space declarations not useless
The way address space declarations for builtins currently work
is nearly useless. The code assumes the address spaces used for
builtins is a confusingly named "target address space" from user
code using __attribute__((address_space(N))) that matches
the builtin declaration. There's no way to use this to declare
a builtin that returns a language specific address space.
The terminology used is highly cofusing since it has nothing
to do with the the address space selected by the target to use
for a language address space.

This feature is essentially unused as-is. AMDGPU and NVPTX
are the only in-tree targets attempting to use this. The AMDGPU
builtins certainly do not behave as intended (i.e. all of the
builtins returning pointers can never compile because the numbered
address space never matches the expected named address space).

The NVPTX builtins are missing tests for some, and the others
seem to rely on an implicit addrspacecast.

Change the used address space for builtins based on a target
hook to allow using a language address space for a builtin.
This allows the same builtin declaration to be used for multiple
languages with similarly purposed address spaces (e.g. the same
AMDGPU builtin can be used in OpenCL and CUDA even though the
constant address spaces are arbitarily different).

This breaks the possibility of using arbitrary numbered
address spaces alongside the named address spaces for builtins.
If this is an issue we probably need to introduce another builtin
declaration character to distinguish language address spaces from
so-called "target address spaces".

llvm-svn: 338707
2018-08-02 12:14:28 +00:00
Yaxun Liu 651bd73c02 [AMDGPU] Change constant addr space to 4
Differential Revision: https://reviews.llvm.org/D43171

llvm-svn: 325031
2018-02-13 18:01:21 +00:00
Matt Arsenault f12e3b848a AMDGPU: Add read_exec_lo/hi builtins
llvm-svn: 315238
2017-10-09 20:06:37 +00:00
Matt Arsenault cbe0dd13d2 AMDGPU: Fix missing declaration for __builtin_amdgcn_dispatch_ptr
llvm-svn: 315219
2017-10-09 17:44:18 +00:00
Konstantin Zhuravlyov 1f144a18ff Resubmit r303861.
[AMDGPU] add __builtin_amdgcn_s_getpc

Patch by Tim Corringham

llvm-svn: 304033
2017-05-26 21:08:20 +00:00
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