arrayfire · umar456 · Aug 8, 2020 · Aug 7, 2020
diff --git a/src/api/c/histogram.cpp b/src/api/c/histogram.cpp
@@ -78,6 +78,10 @@ af_err af_histogram(af_array *out, const af_array in, const unsigned nbins,
                 output = histogram<uchar>(in, nbins, minval, maxval,
                                           info.isLinear());
                 break;
+            case f16:
+                output = histogram<common::half>(in, nbins, minval, maxval,
+                                                 info.isLinear());
+                break;
             default: TYPE_ERROR(1, type);
         }
         std::swap(*out, output);

diff --git a/src/backend/cpu/histogram.cpp b/src/backend/cpu/histogram.cpp
@@ -8,13 +8,15 @@
  ********************************************************/
 
 #include <Array.hpp>
+#include <common/half.hpp>
 #include <histogram.hpp>
 #include <kernel/histogram.hpp>
 #include <platform.hpp>
 #include <queue.hpp>
 #include <af/dim4.hpp>
 
 using af::dim4;
+using common::half;
 
 namespace cpu {
 
@@ -50,5 +52,6 @@ INSTANTIATE(short)
 INSTANTIATE(ushort)
 INSTANTIATE(intl)
 INSTANTIATE(uintl)
+INSTANTIATE(half)
 
 }  // namespace cpu
diff --git a/src/backend/cpu/kernel/histogram.hpp b/src/backend/cpu/kernel/histogram.hpp
@@ -9,6 +9,7 @@
 
 #pragma once
 #include <Param.hpp>
+#include <types.hpp>
 
 namespace cpu {
 namespace kernel {
@@ -23,6 +24,7 @@ void histogram(Param<uint> out, CParam<T> in, const unsigned nbins,
     dim4 const oStrides = out.strides();
     dim_t const nElems  = inDims[0] * inDims[1];
 
+    auto minValT = compute_t<T>(minval);
     for (dim_t b3 = 0; b3 < outDims[3]; b3++) {
         uint* outData   = out.get() + b3 * oStrides[3];
         const T* inData = in.get() + b3 * iStrides[3];
@@ -32,7 +34,7 @@ void histogram(Param<uint> out, CParam<T> in, const unsigned nbins,
                     IsLinear
                         ? i
                         : ((i % inDims[0]) + (i / inDims[0]) * iStrides[1]);
-                int bin = (int)((inData[idx] - minval) / step);
+                int bin = (int)((compute_t<T>(inData[idx]) - minValT) / step);
                 bin     = std::max(bin, 0);
                 bin     = std::min(bin, (int)(nbins - 1));
                 outData[bin]++;

diff --git a/src/backend/cuda/CMakeLists.txt b/src/backend/cuda/CMakeLists.txt
@@ -289,7 +289,7 @@ cuda_add_library(af_cuda_static_cuda_library STATIC
 
     OPTIONS
     ${platform_flags} ${cuda_cxx_flags} ${af_cuda_static_flags}
-    -Xcudafe \"--diag_suppress=1427\" -DAFDLL
+    -Xcudafe --display_error_number -Xcudafe \"--diag_suppress=1427\" -DAFDLL
 )
 
 set_target_properties(af_cuda_static_cuda_library
@@ -648,6 +648,7 @@ cuda_add_library(afcuda
     OPTIONS
     ${platform_flags}
     ${cuda_cxx_flags}
+    -Xcudafe --display_error_number
     -Xcudafe \"--diag_suppress=1427\"
   )
 

diff --git a/src/backend/cuda/histogram.cpp b/src/backend/cuda/histogram.cpp
@@ -8,12 +8,14 @@
  ********************************************************/
 
 #include <Array.hpp>
+#include <common/half.hpp>
 #include <err_cuda.hpp>
 #include <histogram.hpp>
 #include <kernel/histogram.hpp>
 #include <af/dim4.hpp>
 
 using af::dim4;
+using common::half;
 
 namespace cuda {
 
@@ -43,5 +45,6 @@ INSTANTIATE(short)
 INSTANTIATE(ushort)
 INSTANTIATE(intl)
 INSTANTIATE(uintl)
+INSTANTIATE(half)
 
 }  // namespace cuda
diff --git a/src/backend/cuda/kernel/histogram.cuh b/src/backend/cuda/kernel/histogram.cuh
@@ -10,6 +10,7 @@
 #include <Param.hpp>
 #include <math.hpp>
 #include <shared.hpp>
+#include <types.hpp>
 
 namespace cuda {
 
@@ -21,12 +22,13 @@ __global__ void histogram(Param<uint> out, CParam<T> in, int len, int nbins,
 
     // offset input and output to account for batch ops
     unsigned b2   = blockIdx.x / nBBS;
-    const T *iptr = in.ptr + b2 * in.strides[2] + blockIdx.y * in.strides[3];
+    const data_t<T> *iptr = in.ptr + b2 * in.strides[2] + blockIdx.y * in.strides[3];
     uint *optr    = out.ptr + b2 * out.strides[2] + blockIdx.y * out.strides[3];
 
     int start = (blockIdx.x - b2 * nBBS) * THRD_LOAD * blockDim.x + threadIdx.x;
     int end   = min((start + THRD_LOAD * blockDim.x), len);
     float step = (maxval - minval) / (float)nbins;
+    compute_t<T> minvalT(minval);
 
     // If nbins > max shared memory allocated, then just use atomicAdd on global
     // memory
@@ -43,7 +45,7 @@ __global__ void histogram(Param<uint> out, CParam<T> in, int len, int nbins,
             isLinear
                 ? row
                 : ((row % in.dims[0]) + (row / in.dims[0]) * in.strides[1]);
-        int bin = (int)((iptr[idx] - minval) / step);
+        int bin = (int)(static_cast<float>(compute_t<T>(iptr[idx]) - minvalT) / step);
         bin     = (bin < 0) ? 0 : bin;
         bin     = (bin >= nbins) ? (nbins - 1) : bin;
 

diff --git a/src/backend/opencl/histogram.cpp b/src/backend/opencl/histogram.cpp
@@ -8,12 +8,14 @@
  ********************************************************/
 
 #include <Array.hpp>
+#include <common/half.hpp>
 #include <err_opencl.hpp>
 #include <histogram.hpp>
 #include <kernel/histogram.hpp>
 #include <af/dim4.hpp>
 
 using af::dim4;
+using common::half;
 
 namespace opencl {
 
@@ -43,5 +45,6 @@ INSTANTIATE(short)
 INSTANTIATE(ushort)
 INSTANTIATE(intl)
 INSTANTIATE(uintl)
+INSTANTIATE(half)
 
 }  // namespace opencl
diff --git a/test/arrayfire_test.cpp b/test/arrayfire_test.cpp
@@ -322,6 +322,7 @@ INSTANTIATE(half_float::half, half_float::half, float);
 
 INSTANTIATE(double, af_cdouble, float);
 INSTANTIATE(float, af_cfloat, float);
+INSTANTIATE(half_float::half, uint, uint);
 
 #undef INSTANTIATE
 

diff --git a/test/histogram.cpp b/test/histogram.cpp
@@ -32,8 +32,8 @@ class Histogram : public ::testing::Test {
 };
 
 // create a list of types to be tested
-typedef ::testing::Types<float, double, int, uint, char, uchar, short, ushort,
-                         intl, uintl>
+typedef ::testing::Types<half_float::half, float, double, int, uint, char,
+                         uchar, short, ushort, intl, uintl>
     TestTypes;
 
 // register the type list
@@ -48,7 +48,7 @@ void histTest(string pTestFile, unsigned nbins, double minval, double maxval) {
 
     vector<vector<inType> > in;
     vector<vector<outType> > tests;
-    readTests<inType, uint, int>(pTestFile, numDims, in, tests);
+    readTests<inType, uint, uint>(pTestFile, numDims, in, tests);
     dim4 dims = numDims[0];
 
     af_array outArray = 0;