summaryrefslogtreecommitdiff
path: root/r600
AgeCommit message (Collapse)AuthorFilesLines
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