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

[VENTUS][#119]Add work_group_broadcast implementation 

Passed corresponding OPENCL-CTS test.
This commit is contained in:
ziliangzl 2024-05-13 11:22:56 +08:00
parent d63caa094a
commit 4789f2096b
2 changed files with 58 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
libclc/riscv32/lib/SOURCES
View File

@ -37,4 +37,5 @@ compiler-rt/udivmodti4.cl
workgroup/wg.h
workgroup/wganyall.cl
workgroup/wgbarrier.cl
workgroup/wgbcast.cl
workgroup/wgscratch.cl

57
libclc/riscv32/lib/workgroup/wgbcast.cl Normal file
View File

@ -0,0 +1,57 @@
#if __OPENCL_C_VERSION__ >= 200
#include "wg.h"
#include <clc/clc.h>
#define GEN_BROADCAST(TYPE) \
__attribute__((overloadable,weak,always_inline)) _CLC_DEF _CLC_OVERLOAD TYPE \
work_group_broadcast(TYPE a, uint local_id_x) \
{ \
__global TYPE *p = (__global TYPE *)__wg_scratch; \
int gid = get_group_id(0); \
if (get_local_id(0) == local_id_x) { \
p[gid] = a; \
} \
work_group_barrier(CLK_GLOBAL_MEM_FENCE); \
a = p[gid]; \
work_group_barrier(CLK_GLOBAL_MEM_FENCE); \
return a; \
} \
\
__attribute__((overloadable,weak,always_inline)) _CLC_DEF _CLC_OVERLOAD TYPE \
work_group_broadcast(TYPE a, size_t local_id_x, size_t local_id_y) \
{ \
__global TYPE *p = (__global TYPE *)__wg_scratch; \
int gid = get_group_id(0); \
if (get_local_id(0) == local_id_x && get_local_id(1) == local_id_y) { \
p[gid] = a; \
} \
work_group_barrier(CLK_GLOBAL_MEM_FENCE); \
a = p[gid]; \
work_group_barrier(CLK_GLOBAL_MEM_FENCE); \
return a; \
} \
\
__attribute__((overloadable,weak,always_inline)) _CLC_DEF _CLC_OVERLOAD TYPE \
work_group_broadcast(TYPE a, size_t local_id_x,size_t local_id_y, \
size_t local_id_z) \
{ \
__global TYPE *p = (__global TYPE *)__wg_scratch; \
int gid = get_group_id(0); \
if (get_local_id(0) == local_id_x && \
get_local_id(1) == local_id_y && \
get_local_id(2) == local_id_z) { \
p[gid] = a; \
} \
work_group_barrier(CLK_GLOBAL_MEM_FENCE); \
a = p[gid]; \
work_group_barrier(CLK_GLOBAL_MEM_FENCE); \
return a; \
}
GEN_BROADCAST(uint)
GEN_BROADCAST(int)
GEN_BROADCAST(float)
#endif