OSchip
/
llvm-project
mirror of https://github.com/THU-DSP-LAB/llvm-project.git
13
0
Fork 2
Code Issues Packages Projects Releases Wiki Activity

[#56][fix] Fix the implementation of get_local_linear_id

This commit is contained in:
zhoujingya 2023-11-10 09:26:41 +08:00 committed by zhoujing
parent e04c1a6ec7
commit efd82b9d86
1 changed files with 16 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

19
libclc/riscv32/lib/workitem/get_global_linear_id.cl
View File

@ -1,7 +1,20 @@
#include <clc/clc.h>
extern size_t __builtin_riscv_global_linear_id();
_CLC_DEF _CLC_OVERLOAD size_t get_global_linear_id() {
return __builtin_riscv_global_linear_id();
uint dim = get_work_dim() - 1;
switch (dim) {
case 0:
return get_global_id(0) - get_global_offset(0);
;
case 1:
return (get_global_id(1) - get_global_offset(1)) * get_global_size(0) +
(get_global_id(0) - get_global_offset(0));
case 2:
return ((get_global_id(2) - get_global_offset(2)) * get_global_size(1) +
(get_global_id(1) - get_global_offset(1))) *
get_global_size(0) +
(get_global_id(0) - get_global_offset(0));
default:
return 0;
}
}