replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

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
This commit is contained in:
Jan Vesely 2016-07-22 17:24:24 +00:00
parent 74f02db922
commit a82e080b57
9 changed files with 33 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
libclc/amdgcn/lib/SOURCES
View File

@ -1,5 +1,6 @@
math/ldexp.cl
synchronization/barrier_impl.ll
workitem/get_global_offset.cl
workitem/get_group_id.cl
workitem/get_local_id.cl
workitem/get_work_dim.cl

11
libclc/amdgcn/lib/workitem/get_global_offset.cl Normal file
View File

@ -0,0 +1,11 @@
#include <clc/clc.h>
_CLC_DEF uint get_global_offset(uint dim)
{
__attribute__((address_space(2))) uint * ptr =
(__attribute__((address_space(2))) uint *)
__builtin_amdgcn_implicitarg_ptr();
if (dim < 3)
return ptr[dim + 1];
return 0;
}

1
libclc/generic/include/clc/clc.h
View File

@ -30,6 +30,7 @@
#include <clc/workitem/get_local_id.h>
#include <clc/workitem/get_num_groups.h>
#include <clc/workitem/get_group_id.h>
#include <clc/workitem/get_global_offset.h>
/* 6.11.2 Math Functions */
#include <clc/math/acos.h>

1
libclc/generic/include/clc/workitem/get_global_offset.h Normal file
View File

@ -0,0 +1 @@
_CLC_DECL size_t get_global_offset(uint dim);

2
libclc/generic/lib/workitem/get_global_id.cl
View File

@ -1,5 +1,5 @@
#include <clc/clc.h>
_CLC_DEF size_t get_global_id(uint dim) {
return get_group_id(dim)*get_local_size(dim) + get_local_id(dim);
return get_group_id(dim) * get_local_size(dim) + get_local_id(dim) + get_global_offset(dim);
}

1
libclc/ptx-nvidiacl/lib/SOURCES
View File

@ -1,4 +1,5 @@
synchronization/barrier.cl
workitem/get_global_id.cl
workitem/get_group_id.cl
workitem/get_local_id.cl
workitem/get_local_size.cl

5
libclc/ptx-nvidiacl/lib/workitem/get_global_id.cl Normal file
View File

@ -0,0 +1,5 @@
#include <clc/clc.h>
_CLC_DEF size_t get_global_id(uint dim) {
return get_group_id(dim) * get_local_size(dim) + get_local_id(dim);
}

1
libclc/r600/lib/SOURCES
View File

@ -1,4 +1,5 @@
synchronization/barrier_impl.ll
workitem/get_global_offset.cl
workitem/get_group_id.cl
workitem/get_local_id.cl
workitem/get_work_dim.cl

11
libclc/r600/lib/workitem/get_global_offset.cl Normal file
View File

@ -0,0 +1,11 @@
#include <clc/clc.h>
_CLC_DEF uint get_global_offset(uint dim)
{
__attribute__((address_space(7))) uint * ptr =
(__attribute__((address_space(7))) uint *)
__builtin_r600_implicitarg_ptr();
if (dim < 3)
return ptr[dim + 1];
return 0;
}