diff --git a/src/api/c/histogram.cpp b/src/api/c/histogram.cpp index ed9472cc83..f04f4a23df 100644 --- 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(in, nbins, minval, maxval, info.isLinear()); break; + case f16: + output = histogram(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 index 19ef3a9728..2b044efd02 100644 --- a/src/backend/cpu/histogram.cpp +++ b/src/backend/cpu/histogram.cpp @@ -8,6 +8,7 @@ ********************************************************/ #include +#include #include #include #include @@ -15,6 +16,7 @@ #include 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 index 903f2d2204..4b18f94b5b 100644 --- a/src/backend/cpu/kernel/histogram.hpp +++ b/src/backend/cpu/kernel/histogram.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include namespace cpu { namespace kernel { @@ -23,6 +24,7 @@ void histogram(Param out, CParam in, const unsigned nbins, dim4 const oStrides = out.strides(); dim_t const nElems = inDims[0] * inDims[1]; + auto minValT = compute_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 out, CParam 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(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 index 4488c17873..4c320ed6bc 100644 --- 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 index e9f8ce50b5..a2680de686 100644 --- a/src/backend/cuda/histogram.cpp +++ b/src/backend/cuda/histogram.cpp @@ -8,12 +8,14 @@ ********************************************************/ #include +#include #include #include #include #include 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 index 8c1ed0c128..3cd68a1485 100644 --- a/src/backend/cuda/kernel/histogram.cuh +++ b/src/backend/cuda/kernel/histogram.cuh @@ -10,6 +10,7 @@ #include #include #include +#include namespace cuda { @@ -21,12 +22,13 @@ __global__ void histogram(Param out, CParam 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 *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 minvalT(minval); // If nbins > max shared memory allocated, then just use atomicAdd on global // memory @@ -43,7 +45,7 @@ __global__ void histogram(Param out, CParam 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(compute_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 index 929daf67e8..7963d07d3c 100644 --- a/src/backend/opencl/histogram.cpp +++ b/src/backend/opencl/histogram.cpp @@ -8,12 +8,14 @@ ********************************************************/ #include +#include #include #include #include #include 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 index cf0d12b0b9..e9dee59789 100644 --- 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 index c13c329a43..826eebd506 100644 --- 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 +typedef ::testing::Types TestTypes; // register the type list @@ -48,7 +48,7 @@ void histTest(string pTestFile, unsigned nbins, double minval, double maxval) { vector > in; vector > tests; - readTests(pTestFile, numDims, in, tests); + readTests(pTestFile, numDims, in, tests); dim4 dims = numDims[0]; af_array outArray = 0;