[VENTUS][#119]Complete workgroup function implementation

Implement work_group_reduce_<op> functions in wgreduce.cl .
Implement work_group_scan_inclusive_<op> work_group_scan_exclusive_<op> functions in wgscan.cl .
Passed corresponding OPENCL-CTS tests.

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

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];

libclc/riscv32/lib/workgroup/wgreduce.cl

@@ -0,0 +1,89 @@
+#if __OPENCL_C_VERSION__ >= 200
+#include "wg.h"
+#include <clc/clc.h>
+
+#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
+

libclc/riscv32/lib/workgroup/wgscan.cl

@@ -0,0 +1,169 @@
+#if __OPENCL_C_VERSION__ >= 200
+#include "wg.h"
+#include <clc/clc.h>
+
+#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
+

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];