From f2f5a86620414072509f9f218d0a6a8a9587c231 Mon Sep 17 00:00:00 2001 From: Tom Stellard Date: Wed, 26 Jun 2013 18:20:00 +0000 Subject: R600: Replace cl implementations with LLVM IR implementation This allows libclc to be built for R600 with upstream clang and LLVM. llvm-svn: 184980 --- libclc/r600/lib/SOURCES | 6 ++++-- libclc/r600/lib/workitem/get_global_id.cl | 10 ---------- libclc/r600/lib/workitem/get_global_size.cl | 10 ---------- libclc/r600/lib/workitem/get_global_size.ll | 18 ++++++++++++++++++ libclc/r600/lib/workitem/get_group_id.ll | 18 ++++++++++++++++++ libclc/r600/lib/workitem/get_local_id.ll | 18 ++++++++++++++++++ libclc/r600/lib/workitem/get_local_size.ll | 18 ++++++++++++++++++ 7 files changed, 76 insertions(+), 22 deletions(-) delete mode 100644 libclc/r600/lib/workitem/get_global_id.cl delete mode 100644 libclc/r600/lib/workitem/get_global_size.cl create mode 100644 libclc/r600/lib/workitem/get_global_size.ll create mode 100644 libclc/r600/lib/workitem/get_group_id.ll create mode 100644 libclc/r600/lib/workitem/get_local_id.ll create mode 100644 libclc/r600/lib/workitem/get_local_size.ll (limited to 'libclc') diff --git a/libclc/r600/lib/SOURCES b/libclc/r600/lib/SOURCES index 644d2f39733..af8c8c80822 100644 --- a/libclc/r600/lib/SOURCES +++ b/libclc/r600/lib/SOURCES @@ -1,2 +1,4 @@ -workitem/get_global_id.cl -workitem/get_global_size.cl +workitem/get_group_id.ll +workitem/get_local_size.ll +workitem/get_local_id.ll +workitem/get_global_size.ll diff --git a/libclc/r600/lib/workitem/get_global_id.cl b/libclc/r600/lib/workitem/get_global_id.cl deleted file mode 100644 index 671f657acf2..00000000000 --- a/libclc/r600/lib/workitem/get_global_id.cl +++ /dev/null @@ -1,10 +0,0 @@ -#include - -_CLC_DEF size_t get_global_id(uint dim) { - switch (dim) { - case 0: return __builtin_r600_read_tgid_x()*__builtin_r600_read_local_size_x()+__builtin_r600_read_tidig_x(); - case 1: return __builtin_r600_read_tgid_y()*__builtin_r600_read_local_size_y()+__builtin_r600_read_tidig_y(); - case 2: return __builtin_r600_read_tgid_z()*__builtin_r600_read_local_size_z()+__builtin_r600_read_tidig_z(); - default: return 0; - } -} diff --git a/libclc/r600/lib/workitem/get_global_size.cl b/libclc/r600/lib/workitem/get_global_size.cl deleted file mode 100644 index 4e47bbec916..00000000000 --- a/libclc/r600/lib/workitem/get_global_size.cl +++ /dev/null @@ -1,10 +0,0 @@ -#include - -_CLC_DEF size_t get_global_size(uint dim) { - switch (dim) { - case 0: return __builtin_r600_read_global_size_x(); - case 1: return __builtin_r600_read_global_size_y(); - case 2: return __builtin_r600_read_global_size_z(); - default: return 1; - } -} diff --git a/libclc/r600/lib/workitem/get_global_size.ll b/libclc/r600/lib/workitem/get_global_size.ll new file mode 100644 index 00000000000..ac2d08d8ee1 --- /dev/null +++ b/libclc/r600/lib/workitem/get_global_size.ll @@ -0,0 +1,18 @@ +declare i32 @llvm.r600.read.global.size.x() nounwind readnone +declare i32 @llvm.r600.read.global.size.y() nounwind readnone +declare i32 @llvm.r600.read.global.size.z() nounwind readnone + +define i32 @get_global_size(i32 %dim) nounwind readnone alwaysinline { + switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim] +x_dim: + %x = call i32 @llvm.r600.read.global.size.x() nounwind readnone + ret i32 %x +y_dim: + %y = call i32 @llvm.r600.read.global.size.y() nounwind readnone + ret i32 %y +z_dim: + %z = call i32 @llvm.r600.read.global.size.z() nounwind readnone + ret i32 %z +default: + ret i32 0 +} diff --git a/libclc/r600/lib/workitem/get_group_id.ll b/libclc/r600/lib/workitem/get_group_id.ll new file mode 100644 index 00000000000..51316482e91 --- /dev/null +++ b/libclc/r600/lib/workitem/get_group_id.ll @@ -0,0 +1,18 @@ +declare i32 @llvm.r600.read.local.size.x() nounwind readnone +declare i32 @llvm.r600.read.local.size.y() nounwind readnone +declare i32 @llvm.r600.read.local.size.z() nounwind readnone + +define i32 @get_group_id(i32 %dim) nounwind readnone alwaysinline { + switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim] +x_dim: + %x = call i32 @llvm.r600.read.local.size.x() nounwind readnone + ret i32 %x +y_dim: + %y = call i32 @llvm.r600.read.local.size.y() nounwind readnone + ret i32 %y +z_dim: + %z = call i32 @llvm.r600.read.local.size.z() nounwind readnone + ret i32 %z +default: + ret i32 0 +} diff --git a/libclc/r600/lib/workitem/get_local_id.ll b/libclc/r600/lib/workitem/get_local_id.ll new file mode 100644 index 00000000000..ac5522a7822 --- /dev/null +++ b/libclc/r600/lib/workitem/get_local_id.ll @@ -0,0 +1,18 @@ +declare i32 @llvm.r600.read.tidig.x() nounwind readnone +declare i32 @llvm.r600.read.tidig.y() nounwind readnone +declare i32 @llvm.r600.read.tidig.z() nounwind readnone + +define i32 @get_local_id(i32 %dim) nounwind readnone alwaysinline { + switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim] +x_dim: + %x = call i32 @llvm.r600.read.tidig.x() nounwind readnone + ret i32 %x +y_dim: + %y = call i32 @llvm.r600.read.tidig.y() nounwind readnone + ret i32 %y +z_dim: + %z = call i32 @llvm.r600.read.tidig.z() nounwind readnone + ret i32 %z +default: + ret i32 0 +} diff --git a/libclc/r600/lib/workitem/get_local_size.ll b/libclc/r600/lib/workitem/get_local_size.ll new file mode 100644 index 00000000000..6a71f758ea4 --- /dev/null +++ b/libclc/r600/lib/workitem/get_local_size.ll @@ -0,0 +1,18 @@ +declare i32 @llvm.r600.read.tgid.x() nounwind readnone +declare i32 @llvm.r600.read.tgid.y() nounwind readnone +declare i32 @llvm.r600.read.tgid.z() nounwind readnone + +define i32 @get_local_size(i32 %dim) nounwind readnone alwaysinline { + switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim] +x_dim: + %x = call i32 @llvm.r600.read.tgid.x() nounwind readnone + ret i32 %x +y_dim: + %y = call i32 @llvm.r600.read.tgid.y() nounwind readnone + ret i32 %y +z_dim: + %z = call i32 @llvm.r600.read.tgid.z() nounwind readnone + ret i32 %z +default: + ret i32 0 +} -- cgit v1.2.1