maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

R600: Replace cl implementations with LLVM IR implementation 

This allows libclc to be built for R600 with upstream clang and LLVM.

llvm-svn: 184980
This commit is contained in:
Tom Stellard 2013-06-26 18:20:00 +00:00
parent 10b6c22e8d
commit f2f5a86620
7 changed files with 76 additions and 22 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
libclc/r600/lib/SOURCES
View File

@ -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

10
libclc/r600/lib/workitem/get_global_id.cl
View File

@ -1,10 +0,0 @@
#include <clc/clc.h>
_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;
}
}

10
libclc/r600/lib/workitem/get_global_size.cl
View File

@ -1,10 +0,0 @@
#include <clc/clc.h>
_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;
}
}

18
libclc/r600/lib/workitem/get_global_size.ll Normal file
View File

@ -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
}

18
libclc/r600/lib/workitem/get_group_id.ll Normal file
View File

@ -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
}

18
libclc/r600/lib/workitem/get_local_id.ll Normal file
View File

@ -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
}

18
libclc/r600/lib/workitem/get_local_size.ll Normal file
View File

@ -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
}