[libclc] r184980 - R600: Replace cl implementations with LLVM IR implementation

[Tom Stellard]

Author: tstellar
Date: Wed Jun 26 13:20:00 2013
New Revision: 184980

URL: http://llvm.org/viewvc/llvm-project?rev=184980&view=rev
Log:
R600: Replace cl implementations with LLVM IR implementation

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

Added:
    libclc/trunk/r600/lib/workitem/get_global_size.ll
    libclc/trunk/r600/lib/workitem/get_group_id.ll
    libclc/trunk/r600/lib/workitem/get_local_id.ll
    libclc/trunk/r600/lib/workitem/get_local_size.ll
Removed:
    libclc/trunk/r600/lib/workitem/get_global_id.cl
    libclc/trunk/r600/lib/workitem/get_global_size.cl
Modified:
    libclc/trunk/r600/lib/SOURCES

Modified: libclc/trunk/r600/lib/SOURCES
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/SOURCES?rev=184980&r1=184979&r2=184980&view=diff
==============================================================================
--- libclc/trunk/r600/lib/SOURCES (original)
+++ libclc/trunk/r600/lib/SOURCES Wed Jun 26 13:20:00 2013
@@ -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

Removed: libclc/trunk/r600/lib/workitem/get_global_id.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_global_id.cl?rev=184979&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_global_id.cl (original)
+++ libclc/trunk/r600/lib/workitem/get_global_id.cl (removed)
@@ -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;
-  }
-}

Removed: libclc/trunk/r600/lib/workitem/get_global_size.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_global_size.cl?rev=184979&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_global_size.cl (original)
+++ libclc/trunk/r600/lib/workitem/get_global_size.cl (removed)
@@ -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;
-  }
-}

Added: libclc/trunk/r600/lib/workitem/get_global_size.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_global_size.ll?rev=184980&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_global_size.ll (added)
+++ libclc/trunk/r600/lib/workitem/get_global_size.ll Wed Jun 26 13:20:00 2013
@@ -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
+}

Added: libclc/trunk/r600/lib/workitem/get_group_id.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_group_id.ll?rev=184980&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_group_id.ll (added)
+++ libclc/trunk/r600/lib/workitem/get_group_id.ll Wed Jun 26 13:20:00 2013
@@ -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
+}

Added: libclc/trunk/r600/lib/workitem/get_local_id.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_local_id.ll?rev=184980&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_local_id.ll (added)
+++ libclc/trunk/r600/lib/workitem/get_local_id.ll Wed Jun 26 13:20:00 2013
@@ -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
+}

Added: libclc/trunk/r600/lib/workitem/get_local_size.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_local_size.ll?rev=184980&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_local_size.ll (added)
+++ libclc/trunk/r600/lib/workitem/get_local_size.ll Wed Jun 26 13:20:00 2013
@@ -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
+}