Repository: opencv
Issue: opencv_perf_objdetect test "OCL_HOGFixture_HOG.HOG": unsupported call to function reduce_smem in normalize_hists_kernel
Labels: bug,category: ocl

When running "bin/opencv_perf_objdetect" it gives the following warning/error you find below.
Deinlining the function "reduce_smem" like with the following patch fixes that, but i don't know if there is a more proper fix:

diff --git a/modules/objdetect/src/opencl/objdetect_hog.cl b/modules/objdetect/src/opencl/objdetect_hog.cl
index 4fae320..b43fe70 100644
--- a/modules/objdetect/src/opencl/objdetect_hog.cl
+++ b/modules/objdetect/src/opencl/objdetect_hog.cl
@@ -207,7 +207,7 @@ __kernel void normalize_hists_36_kernel(__global float\* block_hists,
 //-------------------------------------------------------------
 //  Normalization of histograms via L2Hys_norm
 //
-inline float reduce_smem(volatile __local float\* smem, int size)
+float reduce_smem(volatile __local float\* smem, int size)
 {
     unsigned int tid = get_local_id(0);
     float sum = smem[tid];

Time compensation is 0
CTEST_FULL_OUTPUT
OpenCV version: 3.1.0-dev
OpenCV VCS version: 3.1.0-315-gecf65a0-dirty
Build type: release
Parallel framework: pthreads
CPU features: mmx sse sse2 sse3
OpenCL Platforms:
    Clover
        iGPU: AMD TAHITI (DRM 2.43.0, LLVM 3.7.1) (OpenCL 1.1 MESA 11.1.2)
Current OpenCL device:
    Type = iGPU
    Name = AMD TAHITI (DRM 2.43.0, LLVM 3.7.1)
    Version = OpenCL 1.1 MESA 11.1.2
    Compute units = 28
    Max work group size = 256
    Local memory size = 32 kB
    Max memory allocation size = 256 MB
    Double support = Yes
    Host unified memory = Yes
    Has AMD Blas = No
    Has AMD Fft = No
    Preferred vector width char = 16
    Preferred vector width short = 8
    Preferred vector width int = 4
    Preferred vector width long = 2
    Preferred vector width float = 4
    Preferred vector width double = 2
[==========] Running 28 tests from 2 test cases.
[----------] Global test environment set-up.
[----------] 1 test from OCL_HOGFixture_HOG
[ RUN      ] OCL_HOGFixture_HOG.HOG
OpenCL program build log:  -D AMD_DEVICE
unsupported call to function reduce_smem in normalize_hists_kernel
.
OpenCL program build log:  -D AMD_DEVICE
unsupported call to function reduce_smem in normalize_hists_kernel
.
OpenCL program build log:  -D AMD_DEVICE
unsupported call to function reduce_smem in normalize_hists_kernel
.
OpenCL program build log:  -D AMD_DEVICE
unsupported call to function reduce_smem in normalize_hists_kernel
.
OpenCL program build log:  -D AMD_DEVICE
unsupported call to function reduce_smem in normalize_hists_kernel
.
OpenCL program build log:  -D AMD_DEVICE
unsupported call to function reduce_smem in normalize_hists_kernel
.
OpenCL program build log:  -D AMD_DEVICE
unsupported call to function reduce_smem in normalize_hists_kernel
.
OpenCL program build log:  -D AMD_DEVICE
unsupported call to function reduce_smem in normalize_hists_kernel
.
OpenCL program build log:  -D AMD_DEVICE
unsupported call to function reduce_smem in normalize_hists_kernel
.
OpenCL program build log:  -D AMD_DEVICE
unsupported call to function reduce_smem in normalize_hists_kernel

[ PERFSTAT ](samples = 13, mean = 15608.84, median = 15618.32, stddev = 72.35 %280.5%%29)
[       OK ] OCL_HOGFixture_HOG.HOG (204664 ms)
[----------] 1 test from OCL_HOGFixture_HOG (204664 ms total)