[libclc] r184976 - r600: Fix get_global_id implementation

[Tom Stellard]

Author: tstellar
Date: Wed Jun 26 13:19:39 2013
New Revision: 184976

URL: http://llvm.org/viewvc/llvm-project?rev=184976&view=rev
Log:
r600: Fix get_global_id implementation

Modified:
    libclc/trunk/r600/lib/workitem/get_global_id.cl

Modified: 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=184976&r1=184975&r2=184976&view=diff
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_global_id.cl (original)
+++ libclc/trunk/r600/lib/workitem/get_global_id.cl Wed Jun 26 13:19:39 2013
@@ -2,9 +2,9 @@
 
 _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();
+  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;
   }
 }