summaryrefslogtreecommitdiff
path: root/kernels
diff options
context:
space:
mode:
authorLuo Xionghu <xionghu.luo@intel.com>2015-01-15 13:22:08 +0800
committerZhigang Gong <zhigang.gong@intel.com>2015-01-15 14:44:13 +0800
commitdb2f53e1798cdb3f716c6cb4811b91483653f9c9 (patch)
treeb979cc237a1fc94e0aafa3fa9e411ead56c56faa /kernels
parent8189d9250ebe3ad562b5ee0810a1ebd190abf899 (diff)
downloadbeignet-db2f53e1798cdb3f716c6cb4811b91483653f9c9.tar.gz
add clz(count leading zero) utest.
this kernl calls the llvm __builtin_clz to generate llvm.clz function then call the gen instruction clz, different from the test compiler_clz_int, which use the fbh to implement. Signed-off-by: Luo Xionghu <xionghu.luo@intel.com> Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Diffstat (limited to 'kernels')
-rw-r--r--kernels/compiler_clz.cl12
1 files changed, 12 insertions, 0 deletions
diff --git a/kernels/compiler_clz.cl b/kernels/compiler_clz.cl
new file mode 100644
index 00000000..7ab62619
--- /dev/null
+++ b/kernels/compiler_clz.cl
@@ -0,0 +1,12 @@
+#define COMPILER_CLZ(TYPE) \
+ kernel void compiler_clz_##TYPE(global TYPE* src, global TYPE* dst) \
+{ \
+ __global TYPE* A = &src[get_global_id(0)]; \
+ __global TYPE* B = &dst[get_global_id(0)]; \
+ *B = __builtin_clz(*A); \
+}
+
+COMPILER_CLZ(uint)
+COMPILER_CLZ(ulong)
+COMPILER_CLZ(ushort)
+COMPILER_CLZ(uchar)