maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
351,705 Commits 23 Branches 0 Tags 2.2 GiB
Commit Graph

24 Commits

Author SHA1 Message Date
Jan Vesely 35b7ac4c30 r600: Convert get_num_groups to clc 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewer: Aaron Watry
llvm-svn: 346077
 2018-11-04 00:35:12 +00:00
Jan Vesely cc6c2ef3b4 r600: Convert get_global_size to clc 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewer: Aaron Watry
llvm-svn: 346076
 2018-11-04 00:35:08 +00:00
Jan Vesely 5fa4e06e27 r600: Convert get_local_size to clc 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewer: Aaron Watry
llvm-svn: 346075
 2018-11-04 00:35:03 +00:00
Jan Vesely 92357a2336 r600: Update datalayout after LLVM r328656 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 329291
 2018-04-05 14:47:57 +00:00
Jan Vesely 83cd840010 r600: Fix datalayout after clang r324101 
r324101 switched around AS numbering

Acked-by: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325864
 2018-02-23 07:36:51 +00:00
Jan Vesely 66b32ad9ad r600: Add missing datalayout to .ll files 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Acked-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 316238
 2017-10-20 21:00:31 +00:00
Jeroen Ketema fe9fa89854 Let get_work_dim take exactly 0 arguments 
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 314634
 2017-10-01 20:11:46 +00:00
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