Commit Graph

17 Commits

Author SHA1 Message Date
Matt Arsenault 958fce3192 amdgcn: Fix return type of get_num_groups
llvm-svn: 279723
2016-08-25 07:31:40 +00:00
Matt Arsenault 26d9c41ff6 amdgcn: Fix return type for get_global_size
llvm-svn: 279644
2016-08-24 17:52:04 +00:00
Matt Arsenault 314364cbd2 amdgpu: Fix default case value for get_local_size
llvm-svn: 279359
2016-08-20 04:17:17 +00:00
Matt Arsenault 220268d177 amdgcn: Fix get_local_size IR return type
llvm-svn: 279350
2016-08-20 00:01:21 +00:00
Jan Vesely a82e080b57 AMDGPU: Implement get_global_offset builtin
Also fix get_global_id to consider offset
No idea how to add this for ptx, so they are stuck with the old get_global_id
implementation.

v2: split to a separate patch

v3: Switch R600 to use implictarg.ptr

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 276443
2016-07-22 17:24:24 +00:00
Jan Vesely 74f02db922 AMDGPU: Use clang intrinsics for workitem builtins
v2: split into 2 patches
    use clang builtins for other intrinsics as well

v3: Fix warnings
    Switch r600 to use implictarg.ptr

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 276442
2016-07-22 17:24:20 +00:00
Matt Arsenault 45e6eaaa05 amdgcn: Use new workitem intrinsics
llvm-svn: 261042
2016-02-17 00:27:27 +00:00
Matt Arsenault a48e15c6cb Split sources for amdgcn and r600
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.

llvm-svn: 260777
2016-02-13 01:01:59 +00:00
Tom Stellard 24ea64e050 r600: get_work_dim: Update metadata syntax for LLVM 3.6
llvm-svn: 225042
2014-12-31 15:27:59 +00:00
Jan Vesely ae50c89589 r600: Fix get_work_dim range metadata
Reviewed-by: Matt Arsenault <Matthew.Arsenault@amd.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 220388
2014-10-22 14:32:53 +00:00
Jan Vesely 260827caa2 r600: Use llvm intrinsic to read work dimension information
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>
llvm-svn: 219793
2014-10-15 15:08:06 +00:00
Aaron Watry bde11213e7 Added get_num_groups
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
2013-07-24 18:03:38 +00:00
Tom Stellard c0af47de00 r600: Fix implementations of get_group_id.ll and get_local_size.ll
llvm-svn: 185005
2013-06-26 18:22:00 +00:00
Tom Stellard f2f5a86620 R600: Replace cl implementations with LLVM IR implementation
This allows libclc to be built for R600 with upstream clang and LLVM.

llvm-svn: 184980
2013-06-26 18:20:00 +00:00
Tom Stellard 38f0ac9d5e r600: Add get_global_size() implementation
llvm-svn: 184977
2013-06-26 18:19:44 +00:00
Tom Stellard ac14c4e878 r600: Fix get_global_id implementation
llvm-svn: 184976
2013-06-26 18:19:39 +00:00
Tom Stellard 879327fcdc r600: Initial support
This includes a get_global_id() implementation and function stubs for
the other workitem and synchronization functions.

llvm-svn: 184975
2013-06-26 18:18:59 +00:00