[arrayfire] 45/408: FEAT: intl/uintl support for all reduce functions
Ghislain Vaillant
ghisvail-guest at moszumanska.debian.org
Mon Sep 21 19:11:13 UTC 2015
This is an automated email from the git hooks/post-receive script.
ghisvail-guest pushed a commit to branch debian/sid
in repository arrayfire.
commit 7883d325a00c0b91f8211b3eab955e593024dedb
Author: Umar Arshad <umar at arrayfire.com>
Date: Thu Jun 25 20:23:45 2015 -0400
FEAT: intl/uintl support for all reduce functions
---
src/api/c/reduce.cpp | 16 ++++++++++++++++
src/backend/cpu/ireduce.cpp | 4 ++++
src/backend/cpu/reduce.cpp | 14 ++++++++++++++
src/backend/cuda/all.cu | 2 ++
src/backend/cuda/any.cu | 2 ++
src/backend/cuda/count.cu | 2 ++
src/backend/cuda/ireduce.cu | 4 ++++
src/backend/cuda/math.hpp | 3 +++
src/backend/cuda/max.cu | 2 ++
src/backend/cuda/min.cu | 2 ++
src/backend/cuda/product.cu | 2 ++
src/backend/cuda/sum.cu | 2 ++
src/backend/opencl/all.cpp | 2 ++
src/backend/opencl/any.cpp | 2 ++
src/backend/opencl/count.cpp | 2 ++
src/backend/opencl/ireduce.cpp | 4 ++++
src/backend/opencl/max.cpp | 2 ++
src/backend/opencl/min.cpp | 2 ++
src/backend/opencl/product.cpp | 2 ++
src/backend/opencl/sum.cpp | 2 ++
20 files changed, 73 insertions(+)
diff --git a/src/api/c/reduce.cpp b/src/api/c/reduce.cpp
index 082c36b..f4d5e61 100644
--- a/src/api/c/reduce.cpp
+++ b/src/api/c/reduce.cpp
@@ -53,6 +53,8 @@ static af_err reduce_type(af_array *out, const af_array in, const int dim)
case c64: res = reduce<op, cdouble, To>(in, dim); break;
case u32: res = reduce<op, uint , To>(in, dim); break;
case s32: res = reduce<op, int , To>(in, dim); break;
+ case u64: res = reduce<op, uintl , To>(in, dim); break;
+ case s64: res = reduce<op, intl , To>(in, dim); break;
case b8: res = reduce<op, char , To>(in, dim); break;
case u8: res = reduce<op, uchar , To>(in, dim); break;
default: TYPE_ERROR(1, type);
@@ -90,6 +92,8 @@ static af_err reduce_common(af_array *out, const af_array in, const int dim)
case c64: res = reduce<op, cdouble, cdouble>(in, dim); break;
case u32: res = reduce<op, uint , uint >(in, dim); break;
case s32: res = reduce<op, int , int >(in, dim); break;
+ case u64: res = reduce<op, uintl , uintl >(in, dim); break;
+ case s64: res = reduce<op, intl , intl >(in, dim); break;
case b8: res = reduce<op, char , char >(in, dim); break;
case u8: res = reduce<op, uchar , uchar >(in, dim); break;
default: TYPE_ERROR(1, type);
@@ -127,6 +131,8 @@ static af_err reduce_promote(af_array *out, const af_array in, const int dim)
case c64: res = reduce<op, cdouble, cdouble>(in, dim); break;
case u32: res = reduce<op, uint , uint >(in, dim); break;
case s32: res = reduce<op, int , int >(in, dim); break;
+ case u64: res = reduce<op, uintl , uintl >(in, dim); break;
+ case s64: res = reduce<op, intl , intl >(in, dim); break;
case u8: res = reduce<op, uchar , uint >(in, dim); break;
// Make sure you are adding only "1" for every non zero value, even if op == af_add_t
case b8: res = reduce<af_notzero_t, char , uint >(in, dim); break;
@@ -199,6 +205,8 @@ static af_err reduce_all_type(double *real, double *imag, const af_array in)
case c64: *real = (double)reduce_all<op, cdouble, To>(in); break;
case u32: *real = (double)reduce_all<op, uint , To>(in); break;
case s32: *real = (double)reduce_all<op, int , To>(in); break;
+ case u64: *real = (double)reduce_all<op, uintl , To>(in); break;
+ case s64: *real = (double)reduce_all<op, intl , To>(in); break;
case b8: *real = (double)reduce_all<op, char , To>(in); break;
case u8: *real = (double)reduce_all<op, uchar , To>(in); break;
default: TYPE_ERROR(1, type);
@@ -230,6 +238,8 @@ static af_err reduce_all_common(double *real_val, double *imag_val, const af_arr
case f64: *real_val = (double)reduce_all<op, double , double >(in); break;
case u32: *real_val = (double)reduce_all<op, uint , uint >(in); break;
case s32: *real_val = (double)reduce_all<op, int , int >(in); break;
+ case u64: *real_val = (double)reduce_all<op, uintl , uintl >(in); break;
+ case s64: *real_val = (double)reduce_all<op, intl , intl >(in); break;
case b8: *real_val = (double)reduce_all<op, char , char >(in); break;
case u8: *real_val = (double)reduce_all<op, uchar , uchar >(in); break;
@@ -276,6 +286,8 @@ static af_err reduce_all_promote(double *real_val, double *imag_val, const af_ar
case f64: *real_val = (double)reduce_all<op, double , double >(in); break;
case u32: *real_val = (double)reduce_all<op, uint , uint >(in); break;
case s32: *real_val = (double)reduce_all<op, int , int >(in); break;
+ case u64: *real_val = (double)reduce_all<op, uintl , uintl >(in); break;
+ case s64: *real_val = (double)reduce_all<op, intl , intl >(in); break;
case u8: *real_val = (double)reduce_all<op, uchar , uint >(in); break;
// Make sure you are adding only "1" for every non zero value, even if op == af_add_t
case b8: *real_val = (double)reduce_all<af_notzero_t, char , uint >(in); break;
@@ -378,6 +390,8 @@ static af_err ireduce_common(af_array *val, af_array *idx, const af_array in, co
case c64: ireduce<op, cdouble>(&res, &loc, in, dim); break;
case u32: ireduce<op, uint >(&res, &loc, in, dim); break;
case s32: ireduce<op, int >(&res, &loc, in, dim); break;
+ case u64: ireduce<op, uintl >(&res, &loc, in, dim); break;
+ case s64: ireduce<op, intl >(&res, &loc, in, dim); break;
case b8: ireduce<op, char >(&res, &loc, in, dim); break;
case u8: ireduce<op, uchar >(&res, &loc, in, dim); break;
default: TYPE_ERROR(1, type);
@@ -428,6 +442,8 @@ static af_err ireduce_all_common(double *real_val, double *imag_val,
case f64: *real_val = (double)ireduce_all<op, double>(loc, in); break;
case u32: *real_val = (double)ireduce_all<op, uint >(loc, in); break;
case s32: *real_val = (double)ireduce_all<op, int >(loc, in); break;
+ case u64: *real_val = (double)ireduce_all<op, uintl >(loc, in); break;
+ case s64: *real_val = (double)ireduce_all<op, intl >(loc, in); break;
case b8: *real_val = (double)ireduce_all<op, char >(loc, in); break;
case u8: *real_val = (double)ireduce_all<op, uchar >(loc, in); break;
diff --git a/src/backend/cpu/ireduce.cpp b/src/backend/cpu/ireduce.cpp
index 9819eec..199a0be 100644
--- a/src/backend/cpu/ireduce.cpp
+++ b/src/backend/cpu/ireduce.cpp
@@ -181,6 +181,8 @@ namespace cpu
INSTANTIATE(af_min_t, cdouble)
INSTANTIATE(af_min_t, int )
INSTANTIATE(af_min_t, uint )
+ INSTANTIATE(af_min_t, intl )
+ INSTANTIATE(af_min_t, uintl )
INSTANTIATE(af_min_t, char )
INSTANTIATE(af_min_t, uchar )
@@ -191,6 +193,8 @@ namespace cpu
INSTANTIATE(af_max_t, cdouble)
INSTANTIATE(af_max_t, int )
INSTANTIATE(af_max_t, uint )
+ INSTANTIATE(af_max_t, intl )
+ INSTANTIATE(af_max_t, uintl )
INSTANTIATE(af_max_t, char )
INSTANTIATE(af_max_t, uchar )
}
diff --git a/src/backend/cpu/reduce.cpp b/src/backend/cpu/reduce.cpp
index 79b4860..428e5d9 100644
--- a/src/backend/cpu/reduce.cpp
+++ b/src/backend/cpu/reduce.cpp
@@ -130,6 +130,8 @@ namespace cpu
INSTANTIATE(af_min_t, cdouble, cdouble)
INSTANTIATE(af_min_t, int , int )
INSTANTIATE(af_min_t, uint , uint )
+ INSTANTIATE(af_min_t, intl , intl )
+ INSTANTIATE(af_min_t, uintl , uintl )
INSTANTIATE(af_min_t, char , char )
INSTANTIATE(af_min_t, uchar , uchar )
@@ -140,6 +142,8 @@ namespace cpu
INSTANTIATE(af_max_t, cdouble, cdouble)
INSTANTIATE(af_max_t, int , int )
INSTANTIATE(af_max_t, uint , uint )
+ INSTANTIATE(af_max_t, intl , intl )
+ INSTANTIATE(af_max_t, uintl , uintl )
INSTANTIATE(af_max_t, char , char )
INSTANTIATE(af_max_t, uchar , uchar )
@@ -150,6 +154,8 @@ namespace cpu
INSTANTIATE(af_add_t, cdouble, cdouble)
INSTANTIATE(af_add_t, int , int )
INSTANTIATE(af_add_t, uint , uint )
+ INSTANTIATE(af_add_t, intl , intl )
+ INSTANTIATE(af_add_t, uintl , uintl )
INSTANTIATE(af_add_t, char , int )
INSTANTIATE(af_add_t, uchar , uint )
@@ -160,6 +166,8 @@ namespace cpu
INSTANTIATE(af_mul_t, cdouble, cdouble)
INSTANTIATE(af_mul_t, int , int )
INSTANTIATE(af_mul_t, uint , uint )
+ INSTANTIATE(af_mul_t, intl , intl )
+ INSTANTIATE(af_mul_t, uintl , uintl )
INSTANTIATE(af_mul_t, char , int )
INSTANTIATE(af_mul_t, uchar , uint )
@@ -170,6 +178,8 @@ namespace cpu
INSTANTIATE(af_notzero_t, cdouble, uint)
INSTANTIATE(af_notzero_t, int , uint)
INSTANTIATE(af_notzero_t, uint , uint)
+ INSTANTIATE(af_notzero_t, intl , uint)
+ INSTANTIATE(af_notzero_t, uintl , uint)
INSTANTIATE(af_notzero_t, char , uint)
INSTANTIATE(af_notzero_t, uchar , uint)
@@ -180,6 +190,8 @@ namespace cpu
INSTANTIATE(af_or_t, cdouble, char)
INSTANTIATE(af_or_t, int , char)
INSTANTIATE(af_or_t, uint , char)
+ INSTANTIATE(af_or_t, intl , char)
+ INSTANTIATE(af_or_t, uintl , char)
INSTANTIATE(af_or_t, char , char)
INSTANTIATE(af_or_t, uchar , char)
@@ -190,6 +202,8 @@ namespace cpu
INSTANTIATE(af_and_t, cdouble, char)
INSTANTIATE(af_and_t, int , char)
INSTANTIATE(af_and_t, uint , char)
+ INSTANTIATE(af_and_t, intl , char)
+ INSTANTIATE(af_and_t, uintl , char)
INSTANTIATE(af_and_t, char , char)
INSTANTIATE(af_and_t, uchar , char)
}
diff --git a/src/backend/cuda/all.cu b/src/backend/cuda/all.cu
index 5d79de4..bfc070a 100644
--- a/src/backend/cuda/all.cu
+++ b/src/backend/cuda/all.cu
@@ -18,6 +18,8 @@ namespace cuda
INSTANTIATE(af_and_t, cdouble, char)
INSTANTIATE(af_and_t, int , char)
INSTANTIATE(af_and_t, uint , char)
+ INSTANTIATE(af_and_t, intl , char)
+ INSTANTIATE(af_and_t, uintl , char)
INSTANTIATE(af_and_t, char , char)
INSTANTIATE(af_and_t, uchar , char)
}
diff --git a/src/backend/cuda/any.cu b/src/backend/cuda/any.cu
index 994db2f..836970e 100644
--- a/src/backend/cuda/any.cu
+++ b/src/backend/cuda/any.cu
@@ -18,6 +18,8 @@ namespace cuda
INSTANTIATE(af_or_t, cdouble, char)
INSTANTIATE(af_or_t, int , char)
INSTANTIATE(af_or_t, uint , char)
+ INSTANTIATE(af_or_t, intl , char)
+ INSTANTIATE(af_or_t, uintl , char)
INSTANTIATE(af_or_t, char , char)
INSTANTIATE(af_or_t, uchar , char)
}
diff --git a/src/backend/cuda/count.cu b/src/backend/cuda/count.cu
index 5d58b7a..d624141 100644
--- a/src/backend/cuda/count.cu
+++ b/src/backend/cuda/count.cu
@@ -18,6 +18,8 @@ namespace cuda
INSTANTIATE(af_notzero_t, cdouble, uint)
INSTANTIATE(af_notzero_t, int , uint)
INSTANTIATE(af_notzero_t, uint , uint)
+ INSTANTIATE(af_notzero_t, intl , uint)
+ INSTANTIATE(af_notzero_t, uintl , uint)
INSTANTIATE(af_notzero_t, char , uint)
INSTANTIATE(af_notzero_t, uchar , uint)
}
diff --git a/src/backend/cuda/ireduce.cu b/src/backend/cuda/ireduce.cu
index 79fffd0..0c14a01 100644
--- a/src/backend/cuda/ireduce.cu
+++ b/src/backend/cuda/ireduce.cu
@@ -49,6 +49,8 @@ namespace cuda
INSTANTIATE(af_min_t, cdouble)
INSTANTIATE(af_min_t, int )
INSTANTIATE(af_min_t, uint )
+ INSTANTIATE(af_min_t, intl )
+ INSTANTIATE(af_min_t, uintl )
INSTANTIATE(af_min_t, char )
INSTANTIATE(af_min_t, uchar )
@@ -59,6 +61,8 @@ namespace cuda
INSTANTIATE(af_max_t, cdouble)
INSTANTIATE(af_max_t, int )
INSTANTIATE(af_max_t, uint )
+ INSTANTIATE(af_max_t, intl )
+ INSTANTIATE(af_max_t, uintl )
INSTANTIATE(af_max_t, char )
INSTANTIATE(af_max_t, uchar )
}
diff --git a/src/backend/cuda/math.hpp b/src/backend/cuda/math.hpp
index 3f7dcd4..577db84 100644
--- a/src/backend/cuda/math.hpp
+++ b/src/backend/cuda/math.hpp
@@ -99,6 +99,9 @@ namespace cuda
template<> __device__ int limit_max<int>() { return 0x7fffffff; }
template<> __device__ int limit_min<int>() { return 0x80000000; }
+ template<> __device__ intl limit_max<intl>() { return 0x7fffffffffffffff; }
+ template<> __device__ intl limit_min<intl>() { return 0x8000000000000000; }
+ template<> __device__ uintl limit_max<uintl>() { return 1ULL << (8 * sizeof(uintl) - 1); }
template<> __device__ char limit_max<char>() { return 0x7f; }
template<> __device__ char limit_min<char>() { return 0x80; }
template<> __device__ float limit_max<float>() { return CUDART_INF_F; }
diff --git a/src/backend/cuda/max.cu b/src/backend/cuda/max.cu
index 68efee5..7841422 100644
--- a/src/backend/cuda/max.cu
+++ b/src/backend/cuda/max.cu
@@ -18,6 +18,8 @@ namespace cuda
INSTANTIATE(af_max_t, cdouble, cdouble)
INSTANTIATE(af_max_t, int , int )
INSTANTIATE(af_max_t, uint , uint )
+ INSTANTIATE(af_max_t, intl , intl )
+ INSTANTIATE(af_max_t, uintl , uintl )
INSTANTIATE(af_max_t, char , char )
INSTANTIATE(af_max_t, uchar , uchar )
}
diff --git a/src/backend/cuda/min.cu b/src/backend/cuda/min.cu
index 13b3596..0251414 100644
--- a/src/backend/cuda/min.cu
+++ b/src/backend/cuda/min.cu
@@ -18,6 +18,8 @@ namespace cuda
INSTANTIATE(af_min_t, cdouble, cdouble)
INSTANTIATE(af_min_t, int , int )
INSTANTIATE(af_min_t, uint , uint )
+ INSTANTIATE(af_min_t, intl , intl )
+ INSTANTIATE(af_min_t, uintl , uintl )
INSTANTIATE(af_min_t, char , char )
INSTANTIATE(af_min_t, uchar , uchar )
}
diff --git a/src/backend/cuda/product.cu b/src/backend/cuda/product.cu
index d8463f9..abc5c1f 100644
--- a/src/backend/cuda/product.cu
+++ b/src/backend/cuda/product.cu
@@ -18,6 +18,8 @@ namespace cuda
INSTANTIATE(af_mul_t, cdouble, cdouble)
INSTANTIATE(af_mul_t, int , int )
INSTANTIATE(af_mul_t, uint , uint )
+ INSTANTIATE(af_mul_t, intl , intl )
+ INSTANTIATE(af_mul_t, uintl , uintl )
INSTANTIATE(af_mul_t, char , int )
INSTANTIATE(af_mul_t, uchar , uint )
}
diff --git a/src/backend/cuda/sum.cu b/src/backend/cuda/sum.cu
index 4f66874..407cc98 100644
--- a/src/backend/cuda/sum.cu
+++ b/src/backend/cuda/sum.cu
@@ -18,6 +18,8 @@ namespace cuda
INSTANTIATE(af_add_t, cdouble, cdouble)
INSTANTIATE(af_add_t, int , int )
INSTANTIATE(af_add_t, uint , uint )
+ INSTANTIATE(af_add_t, intl , intl )
+ INSTANTIATE(af_add_t, uintl , uintl )
INSTANTIATE(af_add_t, char , int )
INSTANTIATE(af_add_t, uchar , uint )
}
diff --git a/src/backend/opencl/all.cpp b/src/backend/opencl/all.cpp
index 8f224c9..4f5c131 100644
--- a/src/backend/opencl/all.cpp
+++ b/src/backend/opencl/all.cpp
@@ -18,6 +18,8 @@ namespace opencl
INSTANTIATE(af_and_t, cdouble, char)
INSTANTIATE(af_and_t, int , char)
INSTANTIATE(af_and_t, uint , char)
+ INSTANTIATE(af_and_t, intl , char)
+ INSTANTIATE(af_and_t, uintl , char)
INSTANTIATE(af_and_t, char , char)
INSTANTIATE(af_and_t, uchar , char)
}
diff --git a/src/backend/opencl/any.cpp b/src/backend/opencl/any.cpp
index e58c5a7..ee8599d 100644
--- a/src/backend/opencl/any.cpp
+++ b/src/backend/opencl/any.cpp
@@ -18,6 +18,8 @@ namespace opencl
INSTANTIATE(af_or_t, cdouble, char)
INSTANTIATE(af_or_t, int , char)
INSTANTIATE(af_or_t, uint , char)
+ INSTANTIATE(af_or_t, intl , char)
+ INSTANTIATE(af_or_t, uintl , char)
INSTANTIATE(af_or_t, char , char)
INSTANTIATE(af_or_t, uchar , char)
}
diff --git a/src/backend/opencl/count.cpp b/src/backend/opencl/count.cpp
index f8ef0b8..e5ad4bf 100644
--- a/src/backend/opencl/count.cpp
+++ b/src/backend/opencl/count.cpp
@@ -18,6 +18,8 @@ namespace opencl
INSTANTIATE(af_notzero_t, cdouble, uint)
INSTANTIATE(af_notzero_t, int , uint)
INSTANTIATE(af_notzero_t, uint , uint)
+ INSTANTIATE(af_notzero_t, intl , uint)
+ INSTANTIATE(af_notzero_t, uintl , uint)
INSTANTIATE(af_notzero_t, char , uint)
INSTANTIATE(af_notzero_t, uchar , uint)
}
diff --git a/src/backend/opencl/ireduce.cpp b/src/backend/opencl/ireduce.cpp
index 2b735c8..698137c 100644
--- a/src/backend/opencl/ireduce.cpp
+++ b/src/backend/opencl/ireduce.cpp
@@ -47,6 +47,8 @@ namespace opencl
INSTANTIATE(af_min_t, cdouble)
INSTANTIATE(af_min_t, int )
INSTANTIATE(af_min_t, uint )
+ INSTANTIATE(af_min_t, intl )
+ INSTANTIATE(af_min_t, uintl )
INSTANTIATE(af_min_t, char )
INSTANTIATE(af_min_t, uchar )
@@ -57,6 +59,8 @@ namespace opencl
INSTANTIATE(af_max_t, cdouble)
INSTANTIATE(af_max_t, int )
INSTANTIATE(af_max_t, uint )
+ INSTANTIATE(af_max_t, intl )
+ INSTANTIATE(af_max_t, uintl )
INSTANTIATE(af_max_t, char )
INSTANTIATE(af_max_t, uchar )
}
diff --git a/src/backend/opencl/max.cpp b/src/backend/opencl/max.cpp
index f4615c7..d3bee0e 100644
--- a/src/backend/opencl/max.cpp
+++ b/src/backend/opencl/max.cpp
@@ -18,6 +18,8 @@ namespace opencl
INSTANTIATE(af_max_t, cdouble, cdouble)
INSTANTIATE(af_max_t, int , int )
INSTANTIATE(af_max_t, uint , uint )
+ INSTANTIATE(af_max_t, intl , intl )
+ INSTANTIATE(af_max_t, uintl , uintl )
INSTANTIATE(af_max_t, char , char )
INSTANTIATE(af_max_t, uchar , uchar )
}
diff --git a/src/backend/opencl/min.cpp b/src/backend/opencl/min.cpp
index 442b0d6..9962fdb 100644
--- a/src/backend/opencl/min.cpp
+++ b/src/backend/opencl/min.cpp
@@ -18,6 +18,8 @@ namespace opencl
INSTANTIATE(af_min_t, cdouble, cdouble)
INSTANTIATE(af_min_t, int , int )
INSTANTIATE(af_min_t, uint , uint )
+ INSTANTIATE(af_min_t, intl , intl )
+ INSTANTIATE(af_min_t, uintl , uintl )
INSTANTIATE(af_min_t, char , char )
INSTANTIATE(af_min_t, uchar , uchar )
}
diff --git a/src/backend/opencl/product.cpp b/src/backend/opencl/product.cpp
index b3a9825..3f32cae 100644
--- a/src/backend/opencl/product.cpp
+++ b/src/backend/opencl/product.cpp
@@ -18,6 +18,8 @@ namespace opencl
INSTANTIATE(af_mul_t, cdouble, cdouble)
INSTANTIATE(af_mul_t, int , int )
INSTANTIATE(af_mul_t, uint , uint )
+ INSTANTIATE(af_mul_t, intl , intl )
+ INSTANTIATE(af_mul_t, uintl , uintl )
INSTANTIATE(af_mul_t, char , int )
INSTANTIATE(af_mul_t, uchar , uint )
}
diff --git a/src/backend/opencl/sum.cpp b/src/backend/opencl/sum.cpp
index b2e3d47..cbe3c5f 100644
--- a/src/backend/opencl/sum.cpp
+++ b/src/backend/opencl/sum.cpp
@@ -18,6 +18,8 @@ namespace opencl
INSTANTIATE(af_add_t, cdouble, cdouble)
INSTANTIATE(af_add_t, int , int )
INSTANTIATE(af_add_t, uint , uint )
+ INSTANTIATE(af_add_t, intl , intl )
+ INSTANTIATE(af_add_t, uintl , uintl )
INSTANTIATE(af_add_t, char , int )
INSTANTIATE(af_add_t, uchar , uint )
}
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/debian-science/packages/arrayfire.git
More information about the debian-science-commits
mailing list