summaryrefslogtreecommitdiff
path: root/r600
AgeCommit message (Collapse)AuthorFilesLines
2016-02-17amdgcn: Use new workitem intrinsicsMatt Arsenault3-0/+62
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@261042 91177308-0d34-0410-b5e6-96231b3b80d8
2016-02-13Split sources for amdgcn and r600Matt Arsenault28-649/+11
Most files remain in a common amdgpu directory. Also switches barriers to to use convergent, and use llvm.amdgcn.s.barrier. This now requires 3.9/trunk to build amdgcn. git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@260777 91177308-0d34-0410-b5e6-96231b3b80d8
2015-09-21r600: Add image writing builtins.Tom Stellard5-0/+83
Patch by: Zoltan Gilian git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@248161 91177308-0d34-0410-b5e6-96231b3b80d8
2015-09-21r600: Add image reading builtins.Tom Stellard5-0/+110
Patch by: Zoltan Gilian git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@248160 91177308-0d34-0410-b5e6-96231b3b80d8
2015-09-21Add image attribute getter builtinsTom Stellard7-0/+153
Added get_image_* OpenCL builtins to the headers. Added implementation to the r600 target. Patch by: Zoltan Gilian git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@248159 91177308-0d34-0410-b5e6-96231b3b80d8
2015-07-10R600: Implement accurate double precision sqrt v2Tom Stellard2-0/+60
v2: - Use same implementation for R600 and gcn. git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@241907 91177308-0d34-0410-b5e6-96231b3b80d8
2015-05-06r600: Use __clc_ldexp on asics that don't implement the intructionJan Vesely1-1/+10
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu> Reviewed-by: Tom Stellard <thomas.stellard@amd.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@236649 91177308-0d34-0410-b5e6-96231b3b80d8
2015-05-06math: Add ldexp implementationTom Stellard2-30/+1
Signed-off-by: Aaron Watry <awatry@gmail.com> Tom Stellard: - Add denormal handling. - Share vectorization code with r600 implementation. Patch By: Aaron Watry git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@236639 91177308-0d34-0410-b5e6-96231b3b80d8
2015-05-06Implement ldexp for R600/SITom Stellard3-0/+68
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@236638 91177308-0d34-0410-b5e6-96231b3b80d8
2014-12-31r600: get_work_dim: Update metadata syntax for LLVM 3.6Tom Stellard1-1/+1
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@225042 91177308-0d34-0410-b5e6-96231b3b80d8
2014-10-22r600: Fix get_work_dim range metadataJan Vesely1-1/+1
Reviewed-by: Matt Arsenault <Matthew.Arsenault@amd.com> Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@220388 91177308-0d34-0410-b5e6-96231b3b80d8
2014-10-15r600: Use llvm intrinsic to read work dimension informationJan Vesely2-0/+9
v2: Fix function declaration Add range metadata to r600 implementation v3: change prefix to AMDGPU Reviewed-by: Tom Stellard <tom@stellard.net> Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@219793 91177308-0d34-0410-b5e6-96231b3b80d8
2014-09-16R600: Map Address spaces for atomic_cmpxchgAaron Watry1-0/+19
Signed-off-by: Aaron Watry <awatry@gmail.com> Reviewed-by: Tom Stellard <thomas.stellard@amd.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@217925 91177308-0d34-0410-b5e6-96231b3b80d8
2014-09-16R600: Map address spaces for atomic_xchgAaron Watry1-0/+1
Signed-off-by: Aaron Watry <awatry@gmail.com> Reviewed-by: Tom Stellard <thomas.stellard@amd.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@217924 91177308-0d34-0410-b5e6-96231b3b80d8
2014-09-16R600: Map address spaces for atomic_minAaron Watry1-0/+10
Signed-off-by: Aaron Watry <awatry@gmail.com> Reviewed-by: Tom Stellard <thomas.stellard@amd.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@217923 91177308-0d34-0410-b5e6-96231b3b80d8
2014-09-16R600: Map address spaces for atomic_xorAaron Watry1-0/+1
Signed-off-by: Aaron Watry <awatry@gmail.com> Reviewed-by: Tom Stellard <thomas.stellard@amd.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@217922 91177308-0d34-0410-b5e6-96231b3b80d8
2014-09-16R600: Map addr spaces and use atomic_maxAaron Watry1-5/+16
Signed-off-by: Aaron Watry <awatry@gmail.com> Reviewed-by: Tom Stellard <thomas.stellard@amd.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@217921 91177308-0d34-0410-b5e6-96231b3b80d8
2014-09-16R600: Map address spaces for atomic_orAaron Watry1-0/+1
Signed-off-by: Aaron Watry <awatry@gmail.com> Reviewed-by: Tom Stellard <thomas.stellard@amd.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@217920 91177308-0d34-0410-b5e6-96231b3b80d8
2014-09-16R600: Map atomic_and address spacesAaron Watry1-0/+1
Signed-off-by: Aaron Watry <awatry@gmail.com> Reviewed-by: Tom Stellard <thomas.stellard@amd.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@217919 91177308-0d34-0410-b5e6-96231b3b80d8
2014-08-20vload/vstore: Use casts instead of scalarizing everything in CLC versionAaron Watry3-189/+0
This generates bitcode which is indistinguishable from what was hand-written for int32 types in v[load|store]_impl.ll. v4: Use vec2+scalar for vec3 load/stores to prevent corruption (per Tom) v3: Also remove unused generic/lib/shared/v[load|store]_impl.ll v2: (Per Matt Arsenault) Fix alignment issues with vector load stores Signed-off-by: Aaron Watry <awatry@gmail.com> Reviewed-by: Tom Stellard <thomas.stellard@amd.com> CC: Matt Arsenault <Matthew.Arsenault@amd.com> CC: Tom Stellard <thomas.stellard@amd.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@216069 91177308-0d34-0410-b5e6-96231b3b80d8
2014-06-24Move clcmacro.h to avoid cluttering user namespace v2Jeroen Ketema1-0/+1
v2: - use quotes instead of <> - add include to r600/lib/math/nextafter.c changed Reviewed-by: Tom Stellard <tom@stellard.net> Reviewed-by: Aaron Watry <awatry@gmail.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@211576 91177308-0d34-0410-b5e6-96231b3b80d8
2013-10-31R600: Set the noduplicate attribute on barrier() intrinsicsTom Stellard3-19/+30
This will prevent LLVM optimization passes from creating illegal uses of the barrier() intrinsic (e.g. calling barrier() from a conditional that is not executed by all threads). git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@193753 91177308-0d34-0410-b5e6-96231b3b80d8
2013-10-10Implement nextafter() builtinTom Stellard2-0/+4
There are two implementations of nextafter(): 1. Using clang's __builtin_nextafter. Clang replaces this builtin with a call to nextafter which is part of libm. Therefore, this implementation will only work for targets with an implementation of libm (e.g. most CPU targets). 2. The other implementation is written in OpenCL C. This function is known internally as __clc_nextafter and can be used by targets that don't have access to libm. git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@192383 91177308-0d34-0410-b5e6-96231b3b80d8
2013-09-06Add atomic_sub and atomic_dec builtin functionsAaron Watry1-0/+1
Reviewed-by: Tom Stellard <thomas.stellard@amd.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@190201 91177308-0d34-0410-b5e6-96231b3b80d8
2013-09-05Add atomic_inc and atomic_add builtinsAaron Watry2-0/+21
Reviewed-by: Aaron Watry <awatry@gmail.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@190058 91177308-0d34-0410-b5e6-96231b3b80d8
2013-08-12Enable assembly vload3 int/uint constant/global for R600Aaron Watry1-16/+2
It's supported by the R600 LLVM back-end now, at least for evergreen. Signed-off-by: Aaron Watry <awatry@gmail.com> Reviewed-by: Tom Stellard <thomas.stellard@amd.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@188180 91177308-0d34-0410-b5e6-96231b3b80d8
2013-08-12Add vload* for addrspace(2) and use as constant load for R600Aaron Watry1-2/+8
Signed-off-by: Aaron Watry <awatry@gmail.com> Reviewed-by: Tom Stellard <thomas.stellard@amd.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@188179 91177308-0d34-0410-b5e6-96231b3b80d8
2013-07-24Added get_num_groupsAaron Watry2-0/+19
The get_num_groups function was missing for r600g. I did the same thing as the other workitem functions. Reviewed-by: Tom Stellard <thomas.stellard@amd.com> Reviewed-by: Aaron Watry <awatry@gmail.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@187059 91177308-0d34-0410-b5e6-96231b3b80d8
2013-07-16Fix and re-enable R600 vload/vstore assemblyAaron Watry3-0/+198
The assembly optimizations were making unsafe assumptions about which address spaces had which identifiers. Also, fix vload/vstore with 64-bit pointers. This was broken previously on Radeon SI. This version still only has assembly versions of int/uint 2/4/8/16 for global loads and stores on R600, but it does it in a way that would be very easily extended to private/local/constant and could also be handled easily on other architectures. v2: 1) Leave v[load|store]_impl.ll in generic/lib 2) Remove vload_if.ll and vstore_if.ll interfaces 3) Fix address+offset calculations 3) Remove offset from assembly arg list git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@186416 91177308-0d34-0410-b5e6-96231b3b80d8
2013-07-08Implement barrier() builtinTom Stellard3-0/+29
Reviewed and Tested-by: Aaron Watry <awatry@gmail.com> git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@185837 91177308-0d34-0410-b5e6-96231b3b80d8
2013-06-26r600: Fix implementations of get_group_id.ll and get_local_size.llTom Stellard2-12/+12
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@185005 91177308-0d34-0410-b5e6-96231b3b80d8
2013-06-26r600: Add overrides fileTom Stellard1-0/+2
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184983 91177308-0d34-0410-b5e6-96231b3b80d8
2013-06-26R600: Replace cl implementations with LLVM IR implementationTom Stellard7-22/+76
This allows libclc to be built for R600 with upstream clang and LLVM. git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184980 91177308-0d34-0410-b5e6-96231b3b80d8
2013-06-26Move R600 headers into generic directoryTom Stellard7-16/+0
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184978 91177308-0d34-0410-b5e6-96231b3b80d8
2013-06-26r600: Add get_global_size() implementationTom Stellard3-3/+12
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184977 91177308-0d34-0410-b5e6-96231b3b80d8
2013-06-26r600: Fix get_global_id implementationTom Stellard1-3/+3
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184976 91177308-0d34-0410-b5e6-96231b3b80d8
2013-06-26r600: Initial supportTom Stellard9-0/+29
This includes a get_global_id() implementation and function stubs for the other workitem and synchronization functions. git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184975 91177308-0d34-0410-b5e6-96231b3b80d8
generated by cgit v1.2.3 (git 2.39.1) at 2024-06-10 01:09:43 +0000