Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2025-02-09 10:55:03 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

PTX: move implementations of work-item and synchronisation functions

Browse Source 
to lib, and add header files in generic.  Incorporates a patch by
Tom Stellard!

llvm-svn: 161313
This commit is contained in:
Peter Collingbourne 2012-08-05 22:25:37 +00:00
parent 1e373f07af
commit a385c53413
18 changed files with 38 additions and 21 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
libclc/generic/include/clc/synchronization/barrier.h Normal file
View File

@ -0,0 +1 @@
_CLC_DECL void barrier(cl_mem_fence_flags flags);

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

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

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

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

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

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

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

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

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

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

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

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

2
libclc/generic/lib/SOURCES
View File

@ -12,3 +12,5 @@ integer/sub_sat.ll
integer/sub_sat_impl.ll
math/hypot.cl
math/mad.cl
workitem/get_global_id.cl
workitem/get_global_size.cl

5
libclc/generic/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);
}

5
libclc/generic/lib/workitem/get_global_size.cl Normal file
View File

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

8
libclc/ptx-nvidiacl/include/clc/workitem/get_global_id.h
View File

@ -1,8 +0,0 @@
_CLC_INLINE size_t get_global_id(uint dim) {
switch (dim) {
case 0: return __builtin_ptx_read_ctaid_x()*__builtin_ptx_read_ntid_x()+__builtin_ptx_read_tid_x();
case 1: return __builtin_ptx_read_ctaid_y()*__builtin_ptx_read_ntid_y()+__builtin_ptx_read_tid_y();
case 2: return __builtin_ptx_read_ctaid_z()*__builtin_ptx_read_ntid_z()+__builtin_ptx_read_tid_z();
default: return 0;
}
}

8
libclc/ptx-nvidiacl/include/clc/workitem/get_global_size.h
View File

@ -1,8 +0,0 @@
_CLC_INLINE size_t get_global_size(uint dim) {
switch (dim) {
case 0: return __builtin_ptx_read_nctaid_x()*__builtin_ptx_read_ntid_x();
case 1: return __builtin_ptx_read_nctaid_y()*__builtin_ptx_read_ntid_y();
case 2: return __builtin_ptx_read_nctaid_z()*__builtin_ptx_read_ntid_z();
default: return 0;
}
}

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

@ -0,0 +1,4 @@
workitem/get_group_id.cl
workitem/get_local_id.cl
workitem/get_local_size.cl
workitem/get_num_groups.cl

4
libclc/ptx-nvidiacl/include/clc/synchronization/barrier.h → libclc/ptx-nvidiacl/lib/synchronization/barrier.cl
View File

@ -1,4 +1,6 @@
_CLC_INLINE void barrier(cl_mem_fence_flags flags) {
#include <clc/clc.h>
_CLC_DEF void barrier(cl_mem_fence_flags flags) {
if (flags & CLK_LOCAL_MEM_FENCE) {
__builtin_ptx_bar_sync(0);
}

4
libclc/ptx-nvidiacl/include/clc/workitem/get_group_id.h → libclc/ptx-nvidiacl/lib/workitem/get_group_id.cl
View File

@ -1,4 +1,6 @@
_CLC_INLINE size_t get_group_id(uint dim) {
#include <clc/clc.h>
_CLC_DEF size_t get_group_id(uint dim) {
switch (dim) {
case 0: return __builtin_ptx_read_ctaid_x();
case 1: return __builtin_ptx_read_ctaid_y();

4
libclc/ptx-nvidiacl/include/clc/workitem/get_local_id.h → libclc/ptx-nvidiacl/lib/workitem/get_local_id.cl
View File

@ -1,4 +1,6 @@
_CLC_INLINE size_t get_local_id(uint dim) {
#include <clc/clc.h>
_CLC_DEF size_t get_local_id(uint dim) {
switch (dim) {
case 0: return __builtin_ptx_read_tid_x();
case 1: return __builtin_ptx_read_tid_y();

4
libclc/ptx-nvidiacl/include/clc/workitem/get_local_size.h → libclc/ptx-nvidiacl/lib/workitem/get_local_size.cl
View File

@ -1,4 +1,6 @@
_CLC_INLINE size_t get_local_size(uint dim) {
#include <clc/clc.h>
_CLC_DEF size_t get_local_size(uint dim) {
switch (dim) {
case 0: return __builtin_ptx_read_ntid_x();
case 1: return __builtin_ptx_read_ntid_y();

4
libclc/ptx-nvidiacl/include/clc/workitem/get_num_groups.h → libclc/ptx-nvidiacl/lib/workitem/get_num_groups.cl
View File

@ -1,4 +1,6 @@
_CLC_INLINE size_t get_num_groups(uint dim) {
#include <clc/clc.h>
_CLC_DEF size_t get_num_groups(uint dim) {
switch (dim) {
case 0: return __builtin_ptx_read_nctaid_x();
case 1: return __builtin_ptx_read_nctaid_y();