diff --git a/libclc/riscv32/lib/SOURCES b/libclc/riscv32/lib/SOURCES index 1c7fb41f156e..22aa6b61a205 100644 --- a/libclc/riscv32/lib/SOURCES +++ b/libclc/riscv32/lib/SOURCES @@ -38,4 +38,6 @@ workgroup/wg.h workgroup/wganyall.cl workgroup/wgbarrier.cl workgroup/wgbcast.cl +workgroup/wgreduce.cl +workgroup/wgscan.cl workgroup/wgscratch.cl diff --git a/libclc/riscv32/lib/workgroup/wg.h b/libclc/riscv32/lib/workgroup/wg.h index aa26a7cff229..fedba2722cdd 100644 --- a/libclc/riscv32/lib/workgroup/wg.h +++ b/libclc/riscv32/lib/workgroup/wg.h @@ -1,3 +1,5 @@ #define MAX_WORKGROUP 32 +#define MAX_THREAD_PER_WG 1024 extern __global int __wg_scratch[MAX_WORKGROUP]; +extern __global int __wi_scratch[MAX_THREAD_PER_WG * MAX_WORKGROUP]; diff --git a/libclc/riscv32/lib/workgroup/wgreduce.cl b/libclc/riscv32/lib/workgroup/wgreduce.cl new file mode 100644 index 000000000000..3f19a1388ca8 --- /dev/null +++ b/libclc/riscv32/lib/workgroup/wgreduce.cl @@ -0,0 +1,89 @@ +#if __OPENCL_C_VERSION__ >= 200 +#include "wg.h" +#include + +#define GEN_REDUCE(TYPE) \ +__attribute__((overloadable,weak,always_inline)) _CLC_DEF _CLC_OVERLOAD TYPE \ +work_group_reduce_add(TYPE a) \ +{ \ + uint n = get_local_size(0); \ + if (n == 1) \ + return a; \ + \ + int gid = get_group_id(0); \ + int i = get_local_id(0); \ + __global TYPE *p = (__global TYPE *)&__wi_scratch[gid * MAX_THREAD_PER_WG]; \ +\ + p[i] = a; \ +\ + work_group_barrier(CLK_LOCAL_MEM_FENCE); \ + if (i == 0) { \ + TYPE res = 0; \ + for (int j = 0; j < n; j++) \ + res += p[j]; \ + p[0] = res; \ + } \ + work_group_barrier(CLK_LOCAL_MEM_FENCE); \ + a = p[0]; \ + work_group_barrier(CLK_LOCAL_MEM_FENCE); \ + return a; \ +} \ +\ +__attribute__((overloadable,weak,always_inline)) _CLC_DEF _CLC_OVERLOAD TYPE \ +work_group_reduce_max(TYPE a) \ +{ \ + uint n = get_local_size(0); \ + if (n == 1) \ + return a; \ + \ + int gid = get_group_id(0); \ + int i = get_local_id(0); \ + __global TYPE *p = (__global TYPE *)&__wi_scratch[gid * MAX_THREAD_PER_WG]; \ +\ + p[i] = a; \ +\ + work_group_barrier(CLK_LOCAL_MEM_FENCE); \ + if (i == 0) { \ + TYPE res = p[0]; \ + for (int j = 0; j < n; j++) \ + res = p[j] > res ? p[j] : res; \ + p[0] = res; \ + } \ + work_group_barrier(CLK_LOCAL_MEM_FENCE); \ + a = p[0]; \ + work_group_barrier(CLK_LOCAL_MEM_FENCE); \ + return a; \ +} \ +\ +__attribute__((overloadable,weak,always_inline)) _CLC_DEF _CLC_OVERLOAD TYPE \ +work_group_reduce_min(TYPE a) \ +{ \ + uint n = get_local_size(0); \ + if (n == 1) \ + return a; \ + \ + int gid = get_group_id(0); \ + int i = get_local_id(0); \ + __global TYPE *p = (__global TYPE *)&__wi_scratch[gid * MAX_THREAD_PER_WG]; \ +\ + p[i] = a; \ +\ + work_group_barrier(CLK_LOCAL_MEM_FENCE); \ + if (i == 0) { \ + TYPE res = p[0]; \ + for (int j = 0; j < n; j++) \ + res = p[j] < res ? p[j] : res; \ + p[0] = res; \ + } \ + work_group_barrier(CLK_LOCAL_MEM_FENCE); \ + a = p[0]; \ + work_group_barrier(CLK_LOCAL_MEM_FENCE); \ + return a; \ +} + +GEN_REDUCE(int) +GEN_REDUCE(uint) +GEN_REDUCE(float) + +#endif + diff --git a/libclc/riscv32/lib/workgroup/wgscan.cl b/libclc/riscv32/lib/workgroup/wgscan.cl new file mode 100644 index 000000000000..e4b1192ea70b --- /dev/null +++ b/libclc/riscv32/lib/workgroup/wgscan.cl @@ -0,0 +1,169 @@ +#if __OPENCL_C_VERSION__ >= 200 +#include "wg.h" +#include + +#define GEN_SCAN_INCLUSIVE(TYPE) \ +__attribute__((overloadable,weak,always_inline)) _CLC_DEF _CLC_OVERLOAD TYPE \ +work_group_scan_inclusive_add(TYPE a) \ +{ \ + uint n = get_local_size(0); \ + if (n == 1) \ + return a; \ + \ + int gid = get_group_id(0); \ + int i = get_local_id(0); \ + __global TYPE *p = (__global TYPE *)&__wi_scratch[gid * MAX_THREAD_PER_WG]; \ +\ + p[i] = a; \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ +\ + if (i == 0) { \ + for (int j = 1; j < n; j++) \ + p[j] = p[j-1] + p[j]; \ + } \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ + a = p[i]; \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ + return a; \ +} \ +\ +__attribute__((overloadable,weak,always_inline)) _CLC_DEF _CLC_OVERLOAD TYPE \ +work_group_scan_inclusive_max(TYPE a) \ +{ \ + uint n = get_local_size(0); \ + if (n == 1) \ + return a; \ + \ + int gid = get_group_id(0); \ + int i = get_local_id(0); \ + __global TYPE *p = (__global TYPE *)&__wi_scratch[gid * MAX_THREAD_PER_WG]; \ +\ + p[i] = a; \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ +\ + if (i == 0) { \ + for (int j = 1; j < n; j++) \ + p[j] = p[j-1] > p[j] ? p[j-1] : p[j]; \ + } \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ + a = p[i]; \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ + return a; \ +} \ +\ +__attribute__((overloadable,weak,always_inline)) _CLC_DEF _CLC_OVERLOAD TYPE \ +work_group_scan_inclusive_min(TYPE a) \ +{ \ + uint n = get_local_size(0); \ + if (n == 1) \ + return a; \ + \ + int gid = get_group_id(0); \ + int i = get_local_id(0); \ + __global TYPE *p = (__global TYPE *)&__wi_scratch[gid * MAX_THREAD_PER_WG]; \ +\ + p[i] = a; \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ +\ + if (i == 0) { \ + for (int j = 1; j < n; j++) \ + p[j] = p[j-1] < p[j] ? p[j-1] : p[j]; \ + } \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ + a = p[i]; \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ + return a; \ +} + +GEN_SCAN_INCLUSIVE(int) +GEN_SCAN_INCLUSIVE(uint) +GEN_SCAN_INCLUSIVE(float) + +#define GEN_SCAN_EXCLUSIVE(TYPE, LB, UB) \ +__attribute__((overloadable,weak,always_inline)) _CLC_DEF _CLC_OVERLOAD TYPE \ +work_group_scan_exclusive_add(TYPE a) \ +{ \ + uint n = get_local_size(0); \ + if (n == 1) \ + return a; \ + \ + int gid = get_group_id(0); \ + int i = get_local_id(0); \ + __global TYPE *p = (__global TYPE *)&__wi_scratch[gid * MAX_THREAD_PER_WG]; \ +\ + p[i] = a; \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ +\ + if (i == 0) { \ + for (int j = 1; j < n; j++) \ + p[j] = p[j-1] + p[j]; \ + } \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ + if (i == 0) \ + a = 0; \ + else \ + a = p[i-1]; \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ + return a; \ +} \ +\ +__attribute__((overloadable,weak,always_inline)) _CLC_DEF _CLC_OVERLOAD TYPE \ +work_group_scan_exclusive_max(TYPE a) \ +{ \ + uint n = get_local_size(0); \ + if (n == 1) \ + return a; \ + \ + int gid = get_group_id(0); \ + int i = get_local_id(0); \ + __global TYPE *p = (__global TYPE *)&__wi_scratch[gid * MAX_THREAD_PER_WG]; \ +\ + p[i] = a; \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ +\ + if (i == 0) { \ + for (int j = 1; j < n; j++) \ + p[j] = p[j-1] > p[j] ? p[j-1] : p[j]; \ + } \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ + if (i == 0) \ + a = LB; \ + else \ + a = p[i-1]; \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ + return a; \ +} \ +\ +__attribute__((overloadable,weak,always_inline)) _CLC_DEF _CLC_OVERLOAD TYPE \ +work_group_scan_exclusive_min(TYPE a) \ +{ \ + uint n = get_local_size(0); \ + if (n == 1) \ + return a; \ + \ + int gid = get_group_id(0); \ + int i = get_local_id(0); \ + __global TYPE *p = (__global TYPE *)&__wi_scratch[gid * MAX_THREAD_PER_WG]; \ +\ + p[i] = a; \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ +\ + if (i == 0) { \ + for (int j = 1; j < n; j++) \ + p[j] = p[j-1] < p[j] ? p[j-1] : p[j]; \ + } \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ + if (i == 0) \ + a = UB; \ + else \ + a = p[i-1]; \ + work_group_barrier(CLK_GLOBAL_MEM_FENCE); \ + return a; \ +} + +GEN_SCAN_EXCLUSIVE(int, INT_MIN, INT_MAX) +GEN_SCAN_EXCLUSIVE(uint, 0U, UINT_MAX) +GEN_SCAN_EXCLUSIVE(float, -INFINITY, INFINITY) + +#endif + diff --git a/libclc/riscv32/lib/workgroup/wgscratch.cl b/libclc/riscv32/lib/workgroup/wgscratch.cl index e6097d89d3db..28a8c499aab1 100644 --- a/libclc/riscv32/lib/workgroup/wgscratch.cl +++ b/libclc/riscv32/lib/workgroup/wgscratch.cl @@ -1,3 +1,4 @@ #include "wg.h" __global int __wg_scratch[MAX_WORKGROUP]; +__global int __wi_scratch[MAX_THREAD_PER_WG * MAX_WORKGROUP];