From 2b198dcbab2303ebdb9892df1a0e2ea4f56cc235 Mon Sep 17 00:00:00 2001 From: Tom Stellard Date: Wed, 26 Jun 2013 18:18:59 +0000 Subject: r600: Initial support 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 --- configure.py | 2 +- r600/include/clc/synchronization/barrier.h | 2 ++ r600/include/clc/workitem/get_global_id.h | 1 + r600/include/clc/workitem/get_global_size.h | 3 +++ r600/include/clc/workitem/get_group_id.h | 3 +++ r600/include/clc/workitem/get_local_id.h | 3 +++ r600/include/clc/workitem/get_local_size.h | 3 +++ r600/include/clc/workitem/get_num_groups.h | 3 +++ r600/lib/SOURCES | 1 + r600/lib/workitem/get_global_id.cl | 10 ++++++++++ 10 files changed, 30 insertions(+), 1 deletion(-) create mode 100644 r600/include/clc/synchronization/barrier.h create mode 100644 r600/include/clc/workitem/get_global_id.h create mode 100644 r600/include/clc/workitem/get_global_size.h create mode 100644 r600/include/clc/workitem/get_group_id.h create mode 100644 r600/include/clc/workitem/get_local_id.h create mode 100644 r600/include/clc/workitem/get_local_size.h create mode 100644 r600/include/clc/workitem/get_num_groups.h create mode 100644 r600/lib/SOURCES create mode 100644 r600/lib/workitem/get_global_id.cl diff --git a/configure.py b/configure.py index 9ae49b7..4f63c5b 100755 --- a/configure.py +++ b/configure.py @@ -43,7 +43,7 @@ llvm_clang = os.path.join(llvm_bindir, 'clang') llvm_link = os.path.join(llvm_bindir, 'llvm-link') llvm_opt = os.path.join(llvm_bindir, 'opt') -default_targets = ['nvptx--nvidiacl', 'nvptx64--nvidiacl'] +default_targets = ['nvptx--nvidiacl', 'nvptx64--nvidiacl', 'r600--'] targets = args if not targets: diff --git a/r600/include/clc/synchronization/barrier.h b/r600/include/clc/synchronization/barrier.h new file mode 100644 index 0000000..7f150d4 --- /dev/null +++ b/r600/include/clc/synchronization/barrier.h @@ -0,0 +1,2 @@ +_CLC_INLINE void barrier(cl_mem_fence_flags flags) { +} diff --git a/r600/include/clc/workitem/get_global_id.h b/r600/include/clc/workitem/get_global_id.h new file mode 100644 index 0000000..b61450f --- /dev/null +++ b/r600/include/clc/workitem/get_global_id.h @@ -0,0 +1 @@ +size_t get_global_id(uint dim); diff --git a/r600/include/clc/workitem/get_global_size.h b/r600/include/clc/workitem/get_global_size.h new file mode 100644 index 0000000..afd9ae1 --- /dev/null +++ b/r600/include/clc/workitem/get_global_size.h @@ -0,0 +1,3 @@ +_CLC_INLINE size_t get_global_size(uint dim) { + return 0; +} diff --git a/r600/include/clc/workitem/get_group_id.h b/r600/include/clc/workitem/get_group_id.h new file mode 100644 index 0000000..6862dba --- /dev/null +++ b/r600/include/clc/workitem/get_group_id.h @@ -0,0 +1,3 @@ +_CLC_INLINE size_t get_group_id(uint dim) { + return 0; +} diff --git a/r600/include/clc/workitem/get_local_id.h b/r600/include/clc/workitem/get_local_id.h new file mode 100644 index 0000000..22749cd --- /dev/null +++ b/r600/include/clc/workitem/get_local_id.h @@ -0,0 +1,3 @@ +_CLC_INLINE size_t get_local_id(uint dim) { + return 0; +} diff --git a/r600/include/clc/workitem/get_local_size.h b/r600/include/clc/workitem/get_local_size.h new file mode 100644 index 0000000..51d9762 --- /dev/null +++ b/r600/include/clc/workitem/get_local_size.h @@ -0,0 +1,3 @@ +_CLC_INLINE size_t get_local_size(uint dim) { + return 0; +} diff --git a/r600/include/clc/workitem/get_num_groups.h b/r600/include/clc/workitem/get_num_groups.h new file mode 100644 index 0000000..fe1f343 --- /dev/null +++ b/r600/include/clc/workitem/get_num_groups.h @@ -0,0 +1,3 @@ +_CLC_INLINE size_t get_num_groups(uint dim) { + return 0; +} diff --git a/r600/lib/SOURCES b/r600/lib/SOURCES new file mode 100644 index 0000000..0844030 --- /dev/null +++ b/r600/lib/SOURCES @@ -0,0 +1 @@ +workitem/get_global_id.cl diff --git a/r600/lib/workitem/get_global_id.cl b/r600/lib/workitem/get_global_id.cl new file mode 100644 index 0000000..9b0bd94 --- /dev/null +++ b/r600/lib/workitem/get_global_id.cl @@ -0,0 +1,10 @@ +#include + +_CLC_DEF size_t get_global_id(uint dim) { + switch (dim) { + case 0: return __builtin_r600_read_tgid_x()*__builtin_r600_read_ngroups_x()+__builtin_r600_read_tidig_x(); + case 1: return __builtin_r600_read_tgid_y()*__builtin_r600_read_ngroups_y()+__builtin_r600_read_tidig_y(); + case 2: return __builtin_r600_read_tgid_z()*__builtin_r600_read_ngroups_z()+__builtin_r600_read_tidig_z(); + default: return 0; + } +} -- cgit v1.2.3