summaryrefslogtreecommitdiffstats
path: root/libclc/r600/lib/SOURCES
Commit message (Collapse)AuthorAgeFilesLines
* R600: Set the noduplicate attribute on barrier() intrinsicsTom Stellard2013-10-311-1/+0
| | | | | | | | 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). llvm-svn: 193753
* Implement nextafter() builtinTom Stellard2013-10-101-0/+1
| | | | | | | | | | | | | | 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. llvm-svn: 192383
* Add atomic_inc and atomic_add builtinsAaron Watry2013-09-051-0/+1
| | | | | Reviewed-by: Aaron Watry <awatry@gmail.com> llvm-svn: 190058
* Added get_num_groupsAaron Watry2013-07-241-0/+1
| | | | | | | | | 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> llvm-svn: 187059
* Fix and re-enable R600 vload/vstore assemblyAaron Watry2013-07-161-0/+2
| | | | | | | | | | | | | | | | | | | 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 llvm-svn: 186416
* Implement barrier() builtinTom Stellard2013-07-081-0/+2
| | | | | | Reviewed and Tested-by: Aaron Watry <awatry@gmail.com> llvm-svn: 185837
* R600: Replace cl implementations with LLVM IR implementationTom Stellard2013-06-261-2/+4
| | | | | | This allows libclc to be built for R600 with upstream clang and LLVM. llvm-svn: 184980
* r600: Add get_global_size() implementationTom Stellard2013-06-261-0/+1
| | | | llvm-svn: 184977
* r600: Initial supportTom Stellard2013-06-261-0/+1
This includes a get_global_id() implementation and function stubs for the other workitem and synchronization functions. llvm-svn: 184975
OpenPOWER on IntegriCloud