OCL_TEST_CYCLE() cv::integral(src, dst, ddepth);
- SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE);
+ SANITY_CHECK(dst, 2e-6, ERROR_RELATIVE);
}
OCL_PERF_TEST_P(IntegralFixture, Integral2, ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32S, CV_32F)))
checkDeviceMaxMemoryAllocSize(srcSize, ddepth);
UMat src(srcSize, CV_8UC1), sum(srcSize + Size(1, 1), ddepth), sqsum(srcSize + Size(1, 1), CV_32F);
- declare.in(src, WARMUP_RNG).out(sum).out(sqsum);
+ declare.in(src, WARMUP_RNG).out(sum, sqsum);
OCL_TEST_CYCLE() cv::integral(src, sum, sqsum, ddepth, CV_32F);
- SANITY_CHECK(sum, 1e-6, ERROR_RELATIVE);
+ SANITY_CHECK(sum, 2e-4, ERROR_RELATIVE);
SANITY_CHECK(sqsum, 5e-5, ERROR_RELATIVE);
}
+++ /dev/null
-/*M///////////////////////////////////////////////////////////////////////////////////////
-//
-// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
-//
-// By downloading, copying, installing or using the software you agree to this license.
-// If you do not agree to this license, do not download, install,
-// copy or use the software.
-//
-//
-// License Agreement
-// For Open Source Computer Vision Library
-//
-// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
-// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// @Authors
-// Shengen Yan,yanshengen@gmail.com
-//
-// Redistribution and use in source and binary forms, with or without modification,
-// are permitted provided that the following conditions are met:
-//
-// * Redistribution's of source code must retain the above copyright notice,
-// this list of conditions and the following disclaimer.
-//
-// * Redistribution's in binary form must reproduce the above copyright notice,
-// this list of conditions and the following disclaimer in the documentation
-// and/or other materials provided with the distribution.
-//
-// * The name of the copyright holders may not be used to endorse or promote products
-// derived from this software without specific prior written permission.
-//
-// This software is provided by the copyright holders and contributors as is and
-// any express or implied warranties, including, but not limited to, the implied
-// warranties of merchantability and fitness for a particular purpose are disclaimed.
-// In no event shall the Intel Corporation or contributors be liable for any direct,
-// indirect, incidental, special, exemplary, or consequential damages
-// (including, but not limited to, procurement of substitute goods or services;
-// loss of use, data, or profits; or business interruption) however caused
-// and on any theory of liability, whether in contract, strict liability,
-// or tort (including negligence or otherwise) arising in any way out of
-// the use of this software, even if advised of the possibility of such damage.
-//
-//M*/
-
-#ifdef DOUBLE_SUPPORT
-#ifdef cl_amd_fp64
-#pragma OPENCL EXTENSION cl_amd_fp64:enable
-#elif defined (cl_khr_fp64)
-#pragma OPENCL EXTENSION cl_khr_fp64:enable
-#endif
-#endif
-
-#if sqdepth == 6
-#define CONVERT(step) ((step)>>1)
-#else
-#define CONVERT(step) ((step))
-#endif
-
-#define LSIZE 256
-#define LSIZE_1 255
-#define LSIZE_2 254
-#define HF_LSIZE 128
-#define LOG_LSIZE 8
-#define LOG_NUM_BANKS 5
-#define NUM_BANKS 32
-#define GET_CONFLICT_OFFSET(lid) ((lid) >> LOG_NUM_BANKS)
-
-#define noconvert
-
-#if sdepth == 4
-
-kernel void integral_cols(__global uchar4 *src, __global int *sum, __global TYPE *sqsum,
- int src_offset, int pre_invalid, int rows, int cols, int src_step, int dst_step, int dst1_step)
-{
- int lid = get_local_id(0);
- int gid = get_group_id(0);
- int4 src_t[2], sum_t[2];
- TYPE4 sqsum_t[2];
- __local int4 lm_sum[2][LSIZE + LOG_LSIZE];
- __local TYPE4 lm_sqsum[2][LSIZE + LOG_LSIZE];
- __local int* sum_p;
- __local TYPE* sqsum_p;
- src_step = src_step >> 2;
- gid = gid << 1;
- for(int i = 0; i < rows; i =i + LSIZE_1)
- {
- src_t[0] = (i + lid < rows ? convert_int4(src[src_offset + (lid+i) * src_step + min(gid, cols - 1)]) : 0);
- src_t[1] = (i + lid < rows ? convert_int4(src[src_offset + (lid+i) * src_step + min(gid + 1, cols - 1)]) : 0);
-
- sum_t[0] = (i == 0 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]);
- sqsum_t[0] = (i == 0 ? (TYPE4)0 : lm_sqsum[0][LSIZE_2 + LOG_LSIZE]);
- sum_t[1] = (i == 0 ? 0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]);
- sqsum_t[1] = (i == 0 ? (TYPE4)0 : lm_sqsum[1][LSIZE_2 + LOG_LSIZE]);
- barrier(CLK_LOCAL_MEM_FENCE);
-
- int bf_loc = lid + GET_CONFLICT_OFFSET(lid);
- lm_sum[0][bf_loc] = src_t[0];
- lm_sqsum[0][bf_loc] = convert_TYPE4(src_t[0] * src_t[0]);
-
- lm_sum[1][bf_loc] = src_t[1];
- lm_sqsum[1][bf_loc] = convert_TYPE4(src_t[1] * src_t[1]);
-
- int offset = 1;
- for(int d = LSIZE >> 1 ; d > 0; d>>=1)
- {
- barrier(CLK_LOCAL_MEM_FENCE);
- int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset;
- ai += GET_CONFLICT_OFFSET(ai);
- bi += GET_CONFLICT_OFFSET(bi);
-
- if((lid & 127) < d)
- {
- lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
- lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai];
- }
- offset <<= 1;
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- if(lid < 2)
- {
- lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0;
- lm_sqsum[lid][LSIZE_2 + LOG_LSIZE] = 0;
- }
- for(int d = 1; d < LSIZE; d <<= 1)
- {
- barrier(CLK_LOCAL_MEM_FENCE);
- offset >>= 1;
- int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset;
- ai += GET_CONFLICT_OFFSET(ai);
- bi += GET_CONFLICT_OFFSET(bi);
-
- if((lid & 127) < d)
- {
- lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
- lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai];
-
- lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai];
- lm_sqsum[lid >> 7][ai] = lm_sqsum[lid >> 7][bi] - lm_sqsum[lid >> 7][ai];
- }
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- int loc_s0 = gid * dst_step + i + lid - 1 - pre_invalid * dst_step /4, loc_s1 = loc_s0 + dst_step ;
- int loc_sq0 = gid * CONVERT(dst1_step) + i + lid - 1 - pre_invalid * dst1_step / sizeof(TYPE),loc_sq1 = loc_sq0 + CONVERT(dst1_step);
- if(lid > 0 && (i+lid) <= rows)
- {
- lm_sum[0][bf_loc] += sum_t[0];
- lm_sum[1][bf_loc] += sum_t[1];
- lm_sqsum[0][bf_loc] += sqsum_t[0];
- lm_sqsum[1][bf_loc] += sqsum_t[1];
- sum_p = (__local int*)(&(lm_sum[0][bf_loc]));
- sqsum_p = (__local TYPE*)(&(lm_sqsum[0][bf_loc]));
- for(int k = 0; k < 4; k++)
- {
- if(gid * 4 + k >= cols + pre_invalid || gid * 4 + k < pre_invalid) continue;
- sum[loc_s0 + k * dst_step / 4] = sum_p[k];
- sqsum[loc_sq0 + k * dst1_step / sizeof(TYPE)] = sqsum_p[k];
- }
- sum_p = (__local int*)(&(lm_sum[1][bf_loc]));
- sqsum_p = (__local TYPE*)(&(lm_sqsum[1][bf_loc]));
- for(int k = 0; k < 4; k++)
- {
- if(gid * 4 + k + 4 >= cols + pre_invalid) break;
- sum[loc_s1 + k * dst_step / 4] = sum_p[k];
- sqsum[loc_sq1 + k * dst1_step / sizeof(TYPE)] = sqsum_p[k];
- }
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- }
-}
-
-kernel void integral_rows(__global int4 *srcsum, __global TYPE4 * srcsqsum,__global int *sum,
- __global TYPE *sqsum, int rows, int cols, int src_step, int src1_step, int sum_step,
- int sqsum_step, int sum_offset, int sqsum_offset)
-{
- int lid = get_local_id(0);
- int gid = get_group_id(0);
- int4 src_t[2], sum_t[2];
- TYPE4 sqsrc_t[2],sqsum_t[2];
- __local int4 lm_sum[2][LSIZE + LOG_LSIZE];
- __local TYPE4 lm_sqsum[2][LSIZE + LOG_LSIZE];
- __local int *sum_p;
- __local TYPE *sqsum_p;
- src_step = src_step >> 4;
- src1_step = (src1_step / sizeof(TYPE)) >> 2 ;
- gid <<= 1;
- for(int i = 0; i < rows; i =i + LSIZE_1)
- {
- src_t[0] = i + lid < rows ? srcsum[(lid+i) * src_step + gid ] : (int4)0;
- sqsrc_t[0] = i + lid < rows ? srcsqsum[(lid+i) * src1_step + gid ] : (TYPE4)0;
- src_t[1] = i + lid < rows ? srcsum[(lid+i) * src_step + gid + 1] : (int4)0;
- sqsrc_t[1] = i + lid < rows ? srcsqsum[(lid+i) * src1_step + gid + 1] : (TYPE4)0;
-
- sum_t[0] = (i == 0 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]);
- sqsum_t[0] = (i == 0 ? (TYPE4)0 : lm_sqsum[0][LSIZE_2 + LOG_LSIZE]);
- sum_t[1] = (i == 0 ? 0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]);
- sqsum_t[1] = (i == 0 ? (TYPE4)0 : lm_sqsum[1][LSIZE_2 + LOG_LSIZE]);
- barrier(CLK_LOCAL_MEM_FENCE);
-
- int bf_loc = lid + GET_CONFLICT_OFFSET(lid);
- lm_sum[0][bf_loc] = src_t[0];
- lm_sqsum[0][bf_loc] = sqsrc_t[0];
-
- lm_sum[1][bf_loc] = src_t[1];
- lm_sqsum[1][bf_loc] = sqsrc_t[1];
-
- int offset = 1;
- for(int d = LSIZE >> 1 ; d > 0; d>>=1)
- {
- barrier(CLK_LOCAL_MEM_FENCE);
- int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset;
- ai += GET_CONFLICT_OFFSET(ai);
- bi += GET_CONFLICT_OFFSET(bi);
-
- if((lid & 127) < d)
- {
- lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
- lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai];
- }
- offset <<= 1;
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- if(lid < 2)
- {
- lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0;
- lm_sqsum[lid][LSIZE_2 + LOG_LSIZE] = 0;
- }
- for(int d = 1; d < LSIZE; d <<= 1)
- {
- barrier(CLK_LOCAL_MEM_FENCE);
- offset >>= 1;
- int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset;
- ai += GET_CONFLICT_OFFSET(ai);
- bi += GET_CONFLICT_OFFSET(bi);
-
- if((lid & 127) < d)
- {
- lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
- lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai];
-
- lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai];
- lm_sqsum[lid >> 7][ai] = lm_sqsum[lid >> 7][bi] - lm_sqsum[lid >> 7][ai];
- }
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- if(gid == 0 && (i + lid) <= rows)
- {
- sum[sum_offset + i + lid] = 0;
- sqsum[sqsum_offset + i + lid] = 0;
- }
- if(i + lid == 0)
- {
- int loc0 = gid * sum_step;
- int loc1 = gid * CONVERT(sqsum_step);
- for(int k = 1; k <= 8; k++)
- {
- if(gid * 4 + k > cols) break;
- sum[sum_offset + loc0 + k * sum_step / 4] = 0;
- sqsum[sqsum_offset + loc1 + k * sqsum_step / sizeof(TYPE)] = 0;
- }
- }
- int loc_s0 = sum_offset + gid * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ;
- int loc_sq0 = sqsum_offset + gid * CONVERT(sqsum_step) + sqsum_step / sizeof(TYPE) + i + lid, loc_sq1 = loc_sq0 + CONVERT(sqsum_step) ;
-
- if(lid > 0 && (i+lid) <= rows)
- {
- lm_sum[0][bf_loc] += sum_t[0];
- lm_sum[1][bf_loc] += sum_t[1];
- lm_sqsum[0][bf_loc] += sqsum_t[0];
- lm_sqsum[1][bf_loc] += sqsum_t[1];
- sum_p = (__local int*)(&(lm_sum[0][bf_loc]));
- sqsum_p = (__local TYPE*)(&(lm_sqsum[0][bf_loc]));
- for(int k = 0; k < 4; k++)
- {
- if(gid * 4 + k >= cols) break;
- sum[loc_s0 + k * sum_step / 4] = sum_p[k];
- sqsum[loc_sq0 + k * sqsum_step / sizeof(TYPE)] = sqsum_p[k];
- }
- sum_p = (__local int*)(&(lm_sum[1][bf_loc]));
- sqsum_p = (__local TYPE*)(&(lm_sqsum[1][bf_loc]));
- for(int k = 0; k < 4; k++)
- {
- if(gid * 4 + 4 + k >= cols) break;
- sum[loc_s1 + k * sum_step / 4] = sum_p[k];
- sqsum[loc_sq1 + k * sqsum_step / sizeof(TYPE)] = sqsum_p[k];
- }
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- }
-}
-
-#elif sdepth == 5
-
-kernel void integral_cols(__global uchar4 *src, __global float *sum, __global TYPE *sqsum,
- int src_offset, int pre_invalid, int rows, int cols, int src_step, int dst_step, int dst1_step)
-{
- int lid = get_local_id(0);
- int gid = get_group_id(0);
- float4 src_t[2], sum_t[2];
- TYPE4 sqsum_t[2];
- __local float4 lm_sum[2][LSIZE + LOG_LSIZE];
- __local TYPE4 lm_sqsum[2][LSIZE + LOG_LSIZE];
- __local float* sum_p;
- __local TYPE* sqsum_p;
- src_step = src_step >> 2;
- gid = gid << 1;
- for(int i = 0; i < rows; i =i + LSIZE_1)
- {
- src_t[0] = (i + lid < rows ? convert_float4(src[src_offset + (lid+i) * src_step + min(gid, cols - 1)]) : (float4)0);
- src_t[1] = (i + lid < rows ? convert_float4(src[src_offset + (lid+i) * src_step + min(gid + 1, cols - 1)]) : (float4)0);
-
- sum_t[0] = (i == 0 ? (float4)0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]);
- sqsum_t[0] = (i == 0 ? (TYPE4)0 : lm_sqsum[0][LSIZE_2 + LOG_LSIZE]);
- sum_t[1] = (i == 0 ? (float4)0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]);
- sqsum_t[1] = (i == 0 ? (TYPE4)0 : lm_sqsum[1][LSIZE_2 + LOG_LSIZE]);
- barrier(CLK_LOCAL_MEM_FENCE);
-
- int bf_loc = lid + GET_CONFLICT_OFFSET(lid);
- lm_sum[0][bf_loc] = src_t[0];
- lm_sqsum[0][bf_loc] = convert_TYPE4(src_t[0] * src_t[0]);
-// printf("%f\n", src_t[0].s0);
-
- lm_sum[1][bf_loc] = src_t[1];
- lm_sqsum[1][bf_loc] = convert_TYPE4(src_t[1] * src_t[1]);
-
- int offset = 1;
- for(int d = LSIZE >> 1 ; d > 0; d>>=1)
- {
- barrier(CLK_LOCAL_MEM_FENCE);
- int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset;
- ai += GET_CONFLICT_OFFSET(ai);
- bi += GET_CONFLICT_OFFSET(bi);
-
- if((lid & 127) < d)
- {
- lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
- lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai];
- }
- offset <<= 1;
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- if(lid < 2)
- {
- lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0;
- lm_sqsum[lid][LSIZE_2 + LOG_LSIZE] = 0;
- }
- for(int d = 1; d < LSIZE; d <<= 1)
- {
- barrier(CLK_LOCAL_MEM_FENCE);
- offset >>= 1;
- int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset;
- ai += GET_CONFLICT_OFFSET(ai);
- bi += GET_CONFLICT_OFFSET(bi);
-
- if((lid & 127) < d)
- {
- lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
- lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai];
-
- lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai];
- lm_sqsum[lid >> 7][ai] = lm_sqsum[lid >> 7][bi] - lm_sqsum[lid >> 7][ai];
- }
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- int loc_s0 = gid * dst_step + i + lid - 1 - pre_invalid * dst_step / 4, loc_s1 = loc_s0 + dst_step ;
- int loc_sq0 = gid * CONVERT(dst1_step) + i + lid - 1 - pre_invalid * dst1_step / sizeof(TYPE), loc_sq1 = loc_sq0 + CONVERT(dst1_step);
- if(lid > 0 && (i+lid) <= rows)
- {
- lm_sum[0][bf_loc] += sum_t[0];
- lm_sum[1][bf_loc] += sum_t[1];
- lm_sqsum[0][bf_loc] += sqsum_t[0];
- lm_sqsum[1][bf_loc] += sqsum_t[1];
- sum_p = (__local float*)(&(lm_sum[0][bf_loc]));
- sqsum_p = (__local TYPE*)(&(lm_sqsum[0][bf_loc]));
- for(int k = 0; k < 4; k++)
- {
- if(gid * 4 + k >= cols + pre_invalid || gid * 4 + k < pre_invalid) continue;
- sum[loc_s0 + k * dst_step / 4] = sum_p[k];
- sqsum[loc_sq0 + k * dst1_step / sizeof(TYPE)] = sqsum_p[k];
- }
- sum_p = (__local float*)(&(lm_sum[1][bf_loc]));
- sqsum_p = (__local TYPE*)(&(lm_sqsum[1][bf_loc]));
- for(int k = 0; k < 4; k++)
- {
- if(gid * 4 + k + 4 >= cols + pre_invalid) break;
- sum[loc_s1 + k * dst_step / 4] = sum_p[k];
- sqsum[loc_sq1 + k * dst1_step / sizeof(TYPE)] = sqsum_p[k];
- }
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- }
-}
-
-kernel void integral_rows(__global float4 *srcsum, __global TYPE4 * srcsqsum, __global float *sum ,
- __global TYPE *sqsum, int rows, int cols, int src_step, int src1_step, int sum_step,
- int sqsum_step, int sum_offset, int sqsum_offset)
-{
- int lid = get_local_id(0);
- int gid = get_group_id(0);
- float4 src_t[2], sum_t[2];
- TYPE4 sqsrc_t[2],sqsum_t[2];
- __local float4 lm_sum[2][LSIZE + LOG_LSIZE];
- __local TYPE4 lm_sqsum[2][LSIZE + LOG_LSIZE];
- __local float *sum_p;
- __local TYPE *sqsum_p;
- src_step = src_step >> 4;
- src1_step = (src1_step / sizeof(TYPE)) >> 2;
- for(int i = 0; i < rows; i =i + LSIZE_1)
- {
- src_t[0] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2] : (float4)0;
- sqsrc_t[0] = i + lid < rows ? srcsqsum[(lid+i) * src1_step + gid * 2] : (TYPE4)0;
- src_t[1] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2 + 1] : (float4)0;
- sqsrc_t[1] = i + lid < rows ? srcsqsum[(lid+i) * src1_step + gid * 2 + 1] : (TYPE4)0;
-
- sum_t[0] = (i == 0 ? (float4)0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]);
- sqsum_t[0] = (i == 0 ? (TYPE4)0 : lm_sqsum[0][LSIZE_2 + LOG_LSIZE]);
- sum_t[1] = (i == 0 ? (float4)0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]);
- sqsum_t[1] = (i == 0 ? (TYPE4)0 : lm_sqsum[1][LSIZE_2 + LOG_LSIZE]);
- barrier(CLK_LOCAL_MEM_FENCE);
-
- int bf_loc = lid + GET_CONFLICT_OFFSET(lid);
- lm_sum[0][bf_loc] = src_t[0];
- lm_sqsum[0][bf_loc] = sqsrc_t[0];
-
- lm_sum[1][bf_loc] = src_t[1];
- lm_sqsum[1][bf_loc] = sqsrc_t[1];
-
- int offset = 1;
- for(int d = LSIZE >> 1 ; d > 0; d>>=1)
- {
- barrier(CLK_LOCAL_MEM_FENCE);
- int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset;
- ai += GET_CONFLICT_OFFSET(ai);
- bi += GET_CONFLICT_OFFSET(bi);
-
- if((lid & 127) < d)
- {
- lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
- lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai];
- }
- offset <<= 1;
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- if(lid < 2)
- {
- lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0;
- lm_sqsum[lid][LSIZE_2 + LOG_LSIZE] = 0;
- }
- for(int d = 1; d < LSIZE; d <<= 1)
- {
- barrier(CLK_LOCAL_MEM_FENCE);
- offset >>= 1;
- int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset;
- ai += GET_CONFLICT_OFFSET(ai);
- bi += GET_CONFLICT_OFFSET(bi);
-
- if((lid & 127) < d)
- {
- lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
- lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai];
-
- lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai];
- lm_sqsum[lid >> 7][ai] = lm_sqsum[lid >> 7][bi] - lm_sqsum[lid >> 7][ai];
- }
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- if(gid == 0 && (i + lid) <= rows)
- {
- sum[sum_offset + i + lid] = 0;
- sqsum[sqsum_offset + i + lid] = 0;
- }
- if(i + lid == 0)
- {
- int loc0 = gid * 2 * sum_step;
- int loc1 = gid * 2 * CONVERT(sqsum_step);
- for(int k = 1; k <= 8; k++)
- {
- if(gid * 8 + k > cols) break;
- sum[sum_offset + loc0 + k * sum_step / 4] = 0;
- sqsum[sqsum_offset + loc1 + k * sqsum_step / sizeof(TYPE)] = 0;
- }
- }
- int loc_s0 = sum_offset + gid * 2 * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ;
- int loc_sq0 = sqsum_offset + gid * 2 * CONVERT(sqsum_step) + sqsum_step / sizeof(TYPE) + i + lid, loc_sq1 = loc_sq0 + CONVERT(sqsum_step) ;
- if(lid > 0 && (i+lid) <= rows)
- {
- lm_sum[0][bf_loc] += sum_t[0];
- lm_sum[1][bf_loc] += sum_t[1];
- lm_sqsum[0][bf_loc] += sqsum_t[0];
- lm_sqsum[1][bf_loc] += sqsum_t[1];
- sum_p = (__local float*)(&(lm_sum[0][bf_loc]));
- sqsum_p = (__local TYPE*)(&(lm_sqsum[0][bf_loc]));
- for(int k = 0; k < 4; k++)
- {
- if(gid * 8 + k >= cols) break;
- sum[loc_s0 + k * sum_step / 4] = sum_p[k];
- sqsum[loc_sq0 + k * sqsum_step / sizeof(TYPE)] = sqsum_p[k];
- }
- sum_p = (__local float*)(&(lm_sum[1][bf_loc]));
- sqsum_p = (__local TYPE*)(&(lm_sqsum[1][bf_loc]));
- for(int k = 0; k < 4; k++)
- {
- if(gid * 8 + 4 + k >= cols) break;
- sum[loc_s1 + k * sum_step / 4] = sum_p[k];
- sqsum[loc_sq1 + k * sqsum_step / sizeof(TYPE)] = sqsum_p[k];
- }
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- }
-}
-
-#endif
/*M///////////////////////////////////////////////////////////////////////////////////////
-//
-// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
-//
-// By downloading, copying, installing or using the software you agree to this license.
-// If you do not agree to this license, do not download, install,
-// copy or use the software.
-//
-//
-// License Agreement
-// For Open Source Computer Vision Library
-//
-// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
-// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+// Copyright (C) 2014, Itseez, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
-//
-// @Authors
-// Shengen Yan,yanshengen@gmail.com
-//
-// Redistribution and use in source and binary forms, with or without modification,
-// are permitted provided that the following conditions are met:
-//
-// * Redistribution's of source code must retain the above copyright notice,
-// this list of conditions and the following disclaimer.
-//
-// * Redistribution's in binary form must reproduce the above copyright notice,
-// this list of conditions and the following disclaimer in the documentation
-// and/or other materials provided with the distribution.
-//
-// * The name of the copyright holders may not be used to endorse or promote products
-// derived from this software without specific prior written permission.
-//
-// This software is provided by the copyright holders and contributors as is and
-// any express or implied warranties, including, but not limited to, the implied
-// warranties of merchantability and fitness for a particular purpose are disclaimed.
-// In no event shall the Intel Corporation or contributors be liable for any direct,
-// indirect, incidental, special, exemplary, or consequential damages
-// (including, but not limited to, procurement of substitute goods or services;
-// loss of use, data, or profits; or business interruption) however caused
-// and on any theory of liability, whether in contract, strict liability,
-// or tort (including negligence or otherwise) arising in any way out of
-// the use of this software, even if advised of the possibility of such damage.
-//
//M*/
#ifdef DOUBLE_SUPPORT
#endif
#endif
-#define LSIZE 256
-#define LSIZE_1 255
-#define LSIZE_2 254
-#define HF_LSIZE 128
-#define LOG_LSIZE 8
-#define LOG_NUM_BANKS 5
-#define NUM_BANKS 32
-#define GET_CONFLICT_OFFSET(lid) ((lid) >> LOG_NUM_BANKS)
-
-#if sdepth == 4
-#define sumT int
-#define vecSumT int4
-#define convertToSum4 convert_int4
-#elif sdepth == 5
-#define sumT float
-#define vecSumT float4
-#define convertToSum4 convert_float4
+#ifndef LOCAL_SUM_SIZE
+#define LOCAL_SUM_SIZE 16
#endif
+#define LOCAL_SUM_STRIDE (LOCAL_SUM_SIZE + 1)
+
-kernel void integral_sum_cols(__global const uchar4 *src, __global uchar *sum_ptr,
- int src_offset, int rows, int cols, int src_step, int dst_step)
+kernel void integral_sum_cols(__global const uchar *src_ptr, int src_step, int src_offset, int rows, int cols,
+ __global uchar *buf_ptr, int buf_step, int buf_offset
+#ifdef SUM_SQUARE
+ ,__global uchar *buf_sq_ptr, int buf_sq_step, int buf_sq_offset
+#endif
+ )
{
- __global sumT *sum = (__global sumT *)sum_ptr;
+ __local sumT lm_sum[LOCAL_SUM_STRIDE * LOCAL_SUM_SIZE];
+#ifdef SUM_SQUARE
+ __local sumSQT lm_sum_sq[LOCAL_SUM_STRIDE * LOCAL_SUM_SIZE];
+#endif
int lid = get_local_id(0);
int gid = get_group_id(0);
- vecSumT src_t[2], sum_t[2];
- __local vecSumT lm_sum[2][LSIZE + LOG_LSIZE];
- __local sumT* sum_p;
- src_step = src_step >> 2;
- gid = gid << 1;
- int lid_prim = ((lid & 127) << 1) + 1;
- for (int i = 0; i < rows; i += LSIZE_1)
- {
- if (i + lid < rows)
- {
- int src_index = mad24((lid+i), src_step, gid + src_offset);
- src_t[0] = convertToSum4(src[src_index]);
- src_t[1] = convertToSum4(src[src_index + 1]);
- }
- else
- {
- src_t[0] = (vecSumT)0;
- src_t[1] = (vecSumT)0;
- }
- if (i == 0)
- {
- sum_t[0] = (vecSumT)0;
- sum_t[1] = (vecSumT)0;
- }
- else
- {
- sum_t[0] = lm_sum[0][LSIZE_2 + LOG_LSIZE];
- sum_t[1] = lm_sum[1][LSIZE_2 + LOG_LSIZE];
- }
- barrier(CLK_LOCAL_MEM_FENCE);
-
- int bf_loc = lid + GET_CONFLICT_OFFSET(lid);
-
- lm_sum[0][bf_loc] = src_t[0];
- lm_sum[1][bf_loc] = src_t[1];
+ int x = get_global_id(0);
+ int src_index = x + src_offset;
- int offset = 1;
- for (int d = LSIZE >> 1 ; d > 0; d>>=1)
+ sumT accum = 0;
+#ifdef SUM_SQUARE
+ sumSQT accum_sq = 0;
+#endif
+ for (int y = 0; y < rows; y += LOCAL_SUM_SIZE)
+ {
+ int lsum_index = lid;
+ #pragma unroll
+ for (int yin = 0; yin < LOCAL_SUM_SIZE; yin++, src_index+=src_step, lsum_index += LOCAL_SUM_STRIDE)
{
- barrier(CLK_LOCAL_MEM_FENCE);
- int ai = offset * lid_prim - 1,bi = ai + offset;
- ai += GET_CONFLICT_OFFSET(ai);
- bi += GET_CONFLICT_OFFSET(bi);
-
- if((lid & 127) < d)
+ if ((x < cols) && (y + yin < rows))
{
- lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
+ __global const uchar *src = src_ptr + src_index;
+ accum += src[0];
+#ifdef SUM_SQUARE
+ sumSQT temp = src[0] * src[0];
+ accum_sq += temp;
+#endif
}
- offset <<= 1;
+ lm_sum[lsum_index] = accum;
+#ifdef SUM_SQUARE
+ lm_sum_sq[lsum_index] = accum_sq;
+#endif
}
barrier(CLK_LOCAL_MEM_FENCE);
- if (lid < 2)
- {
- lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0;
- }
- for (int d = 1; d < LSIZE; d <<= 1)
- {
- barrier(CLK_LOCAL_MEM_FENCE);
- offset >>= 1;
- int ai = offset * lid_prim - 1,bi = ai + offset;
- ai += GET_CONFLICT_OFFSET(ai);
- bi += GET_CONFLICT_OFFSET(bi);
- if((lid & 127) < d)
- {
- lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
- lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai];
- }
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- if (lid > 0 && (i+lid) <= rows)
- {
- int loc_s0 = mad24(gid, dst_step, i + lid - 1), loc_s1 = loc_s0 + dst_step;
- lm_sum[0][bf_loc] += sum_t[0];
- lm_sum[1][bf_loc] += sum_t[1];
- sum_p = (__local sumT*)(&(lm_sum[0][bf_loc]));
- for (int k = 0; k < 4; k++)
- {
- if (gid * 4 + k >= cols)
- break;
- sum[loc_s0 + k * dst_step / 4] = sum_p[k];
- }
- sum_p = (__local sumT*)(&(lm_sum[1][bf_loc]));
- for (int k = 0; k < 4; k++)
- {
- if (gid * 4 + k + 4 >= cols)
- break;
- sum[loc_s1 + k * dst_step / 4] = sum_p[k];
- }
+ //int buf_index = buf_offset + buf_step * LOCAL_SUM_COLS * gid + sizeof(sumT) * y + sizeof(sumT) * lid;
+ int buf_index = mad24(buf_step, LOCAL_SUM_SIZE * gid, mad24((int)sizeof(sumT), y + lid, buf_offset));
+#ifdef SUM_SQUARE
+ int buf_sq_index = mad24(buf_sq_step, LOCAL_SUM_SIZE * gid, mad24((int)sizeof(sumSQT), y + lid, buf_sq_offset));
+#endif
+
+ lsum_index = LOCAL_SUM_STRIDE * lid;
+ #pragma unroll
+ for (int yin = 0; yin < LOCAL_SUM_SIZE; yin++, lsum_index ++)
+ {
+ __global sumT *buf = (__global sumT *)(buf_ptr + buf_index);
+ buf[0] = lm_sum[lsum_index];
+ buf_index += buf_step;
+#ifdef SUM_SQUARE
+ __global sumSQT *bufsq = (__global sumSQT *)(buf_sq_ptr + buf_sq_index);
+ bufsq[0] = lm_sum_sq[lsum_index];
+ buf_sq_index += buf_sq_step;
+#endif
}
barrier(CLK_LOCAL_MEM_FENCE);
}
}
-
-kernel void integral_sum_rows(__global const uchar *srcsum_ptr, __global uchar *sum_ptr,
- int rows, int cols, int src_step, int sum_step, int sum_offset)
+kernel void integral_sum_rows(__global const uchar *buf_ptr, int buf_step, int buf_offset,
+#ifdef SUM_SQUARE
+ __global uchar *buf_sq_ptr, int buf_sq_step, int buf_sq_offset,
+#endif
+ __global uchar *dst_ptr, int dst_step, int dst_offset, int rows, int cols
+#ifdef SUM_SQUARE
+ ,__global uchar *dst_sq_ptr, int dst_sq_step, int dst_sq_offset
+#endif
+ )
{
- __global const vecSumT *srcsum = (__global const vecSumT *)srcsum_ptr;
- __global sumT *sum = (__global sumT *)sum_ptr;
+ __local sumT lm_sum[LOCAL_SUM_STRIDE * LOCAL_SUM_SIZE];
+#ifdef SUM_SQUARE
+ __local sumSQT lm_sum_sq[LOCAL_SUM_STRIDE * LOCAL_SUM_SIZE];
+#endif
int lid = get_local_id(0);
int gid = get_group_id(0);
- vecSumT src_t[2], sum_t[2];
- __local vecSumT lm_sum[2][LSIZE + LOG_LSIZE];
- __local sumT *sum_p;
- src_step = src_step >> 4;
- int lid_prim = ((lid & 127) << 1) + 1;
- for (int i = 0; i < rows; i += LSIZE_1)
+
+ int gs = get_global_size(0);
+
+ int x = get_global_id(0);
+
+ __global sumT *dst = (__global sumT *)(dst_ptr + dst_offset);
+ for (int xin = x; xin < cols; xin += gs)
{
- if (i + lid < rows)
- {
- int sum_idx = mad24(lid + i, src_step, gid * 2);
- src_t[0] = srcsum[sum_idx];
- src_t[1] = srcsum[sum_idx + 1];
- }
- else
- {
- src_t[0] = 0;
- src_t[1] = 0;
- }
- if (i == 0)
- {
- sum_t[0] = 0;
- sum_t[1] = 0;
- }
- else
- {
- sum_t[0] = lm_sum[0][LSIZE_2 + LOG_LSIZE];
- sum_t[1] = lm_sum[1][LSIZE_2 + LOG_LSIZE];
- }
- barrier(CLK_LOCAL_MEM_FENCE);
+ dst[xin] = 0;
+ }
+ dst_offset += dst_step;
- int bf_loc = lid + GET_CONFLICT_OFFSET(lid);
+ if (x < rows - 1)
+ {
+ dst = (__global sumT *)(dst_ptr + mad24(x, dst_step, dst_offset));
+ dst[0] = 0;
+ }
- lm_sum[0][bf_loc] = src_t[0];
- lm_sum[1][bf_loc] = src_t[1];
+ int buf_index = mad24((int)sizeof(sumT), x, buf_offset);
+ sumT accum = 0;
- int offset = 1;
- for (int d = LSIZE >> 1 ; d > 0; d>>=1)
- {
- barrier(CLK_LOCAL_MEM_FENCE);
- int ai = offset * lid_prim - 1, bi = ai + offset;
- ai += GET_CONFLICT_OFFSET(ai);
- bi += GET_CONFLICT_OFFSET(bi);
+#ifdef SUM_SQUARE
+ __global sumSQT *dst_sq = (__global sumT *)(dst_sq_ptr + dst_sq_offset);
+ for (int xin = x; xin < cols; xin += gs)
+ {
+ dst_sq[xin] = 0;
+ }
+ dst_sq_offset += dst_sq_step;
- if((lid & 127) < d)
- {
- lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
- }
- offset <<= 1;
- }
- barrier(CLK_LOCAL_MEM_FENCE);
- if (lid < 2)
- {
- lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0;
- }
- for (int d = 1; d < LSIZE; d <<= 1)
- {
- barrier(CLK_LOCAL_MEM_FENCE);
- offset >>= 1;
- int ai = offset * lid_prim - 1,bi = ai + offset;
- ai += GET_CONFLICT_OFFSET(ai);
- bi += GET_CONFLICT_OFFSET(bi);
+ dst_sq = (__global sumSQT *)(dst_sq_ptr + mad24(x, dst_sq_step, dst_sq_offset));
+ dst_sq[0] = 0;
- if ((lid & 127) < d)
- {
- lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
- lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai];
- }
+ int buf_sq_index = mad24((int)sizeof(sumSQT), x, buf_sq_offset);
+ sumSQT accum_sq = 0;
+#endif
+
+ for (int y = 1; y < cols; y += LOCAL_SUM_SIZE)
+ {
+ int lsum_index = lid;
+ #pragma unroll
+ for (int yin = 0; yin < LOCAL_SUM_SIZE; yin++, lsum_index += LOCAL_SUM_STRIDE)
+ {
+ __global const sumT *buf = (__global const sumT *)(buf_ptr + buf_index);
+ accum += buf[0];
+ lm_sum[lsum_index] = accum;
+ buf_index += buf_step;
+#ifdef SUM_SQUARE
+ __global const sumSQT *buf_sq = (__global const sumSQT *)(buf_sq_ptr + buf_sq_index);
+ accum_sq += buf_sq[0];
+ lm_sum_sq[lsum_index] = accum_sq;
+ buf_sq_index += buf_sq_step;
+#endif
}
barrier(CLK_LOCAL_MEM_FENCE);
- if (gid == 0 && (i + lid) <= rows)
- {
- sum[sum_offset + i + lid] = 0;
- }
- if (i + lid == 0)
- {
- int loc0 = gid * 2 * sum_step;
- for(int k = 1; k <= 8; k++)
- {
- if (gid * 8 + k > cols)
- break;
- sum[sum_offset + loc0 + k * sum_step / 4] = 0;
- }
- }
- if (lid > 0 && (i+lid) <= rows)
+ if (y + lid < cols)
{
- int loc_s0 = sum_offset + gid * 2 * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ;
- lm_sum[0][bf_loc] += sum_t[0];
- lm_sum[1][bf_loc] += sum_t[1];
- sum_p = (__local sumT*)(&(lm_sum[0][bf_loc]));
- for(int k = 0; k < 4; k++)
- {
- if (gid * 8 + k >= cols)
- break;
- sum[loc_s0 + k * sum_step / 4] = sum_p[k];
- }
- sum_p = (__local sumT*)(&(lm_sum[1][bf_loc]));
- for(int k = 0; k < 4; k++)
+ //int dst_index = dst_offset + dst_step * LOCAL_SUM_COLS * gid + sizeof(sumT) * y + sizeof(sumT) * lid;
+ int dst_index = mad24(dst_step, LOCAL_SUM_SIZE * gid, mad24((int)sizeof(sumT), y + lid, dst_offset));
+#ifdef SUM_SQUARE
+ int dst_sq_index = mad24(dst_sq_step, LOCAL_SUM_SIZE * gid, mad24((int)sizeof(sumSQT), y + lid, dst_sq_offset));
+#endif
+ lsum_index = LOCAL_SUM_STRIDE * lid;
+ int yin_max = min(rows - 1 - LOCAL_SUM_SIZE * gid, LOCAL_SUM_SIZE);
+ #pragma unroll
+ for (int yin = 0; yin < yin_max; yin++, lsum_index++)
{
- if (gid * 8 + 4 + k >= cols)
- break;
- sum[loc_s1 + k * sum_step / 4] = sum_p[k];
+ dst = (__global sumT *)(dst_ptr + dst_index);
+ dst[0] = lm_sum[lsum_index];
+ dst_index += dst_step;
+#ifdef SUM_SQUARE
+ dst_sq = (__global sumSQT *)(dst_sq_ptr + dst_sq_index);
+ dst_sq[0] = lm_sum_sq[lsum_index];
+ dst_sq_index += dst_sq_step;
+#endif
}
}
barrier(CLK_LOCAL_MEM_FENCE);
#ifdef HAVE_OPENCL
-enum { vlen = 4 };
-
static bool ocl_integral( InputArray _src, OutputArray _sum, int sdepth )
{
- if ( _src.type() != CV_8UC1 || _src.step() % vlen != 0 || _src.offset() % vlen != 0 ||
- !(sdepth == CV_32S || sdepth == CV_32F) )
- return false;
+ bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
- ocl::Kernel k1("integral_sum_cols", ocl::imgproc::integral_sum_oclsrc,
- format("-D sdepth=%d", sdepth));
- if (k1.empty())
+ if ( (_src.type() != CV_8UC1) ||
+ !(sdepth == CV_32S || sdepth == CV_32F || (doubleSupport && sdepth == CV_64F)))
return false;
- Size size = _src.size(), t_size = Size(((size.height + vlen - 1) / vlen) * vlen, size.width),
- ssize(size.width + 1, size.height + 1);
- _sum.create(ssize, sdepth);
- UMat src = _src.getUMat(), t_sum(t_size, sdepth), sum = _sum.getUMat();
- t_sum = t_sum(Range::all(), Range(0, size.height));
+ static const int tileSize = 16;
+
+ String build_opt = format("-D sumT=%s -D LOCAL_SUM_SIZE=%d%s",
+ ocl::typeToStr(sdepth), tileSize,
+ doubleSupport ? " -D DOUBLE_SUPPORT" : "");
+
+ ocl::Kernel kcols("integral_sum_cols", ocl::imgproc::integral_sum_oclsrc, build_opt);
+ if (kcols.empty())
+ return false;
- int offset = (int)src.offset / vlen;
- int vcols = (src.cols + vlen - 1) / vlen;
- int sum_offset = (int)sum.offset / vlen;
+ UMat src = _src.getUMat();
+ Size src_size = src.size();
+ Size bufsize(((src_size.height + tileSize - 1) / tileSize) * tileSize, ((src_size.width + tileSize - 1) / tileSize) * tileSize);
+ UMat buf(bufsize, sdepth);
+ kcols.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnlyNoSize(buf));
+ size_t gt = src.cols, lt = tileSize;
+ if (!kcols.run(1, >, <, false))
+ return false;
- k1.args(ocl::KernelArg::PtrReadOnly(src), ocl::KernelArg::PtrWriteOnly(t_sum),
- offset, src.rows, src.cols, (int)src.step, (int)t_sum.step);
- size_t gt = ((vcols + 1) / 2) * 256, lt = 256;
- if (!k1.run(1, >, <, false))
+ ocl::Kernel krows("integral_sum_rows", ocl::imgproc::integral_sum_oclsrc, build_opt);
+ if (krows.empty())
return false;
- ocl::Kernel k2("integral_sum_rows", ocl::imgproc::integral_sum_oclsrc,
- format("-D sdepth=%d", sdepth));
- k2.args(ocl::KernelArg::PtrReadOnly(t_sum), ocl::KernelArg::PtrWriteOnly(sum),
- t_sum.rows, t_sum.cols, (int)t_sum.step, (int)sum.step, sum_offset);
+ Size sumsize(src_size.width + 1, src_size.height + 1);
+ _sum.create(sumsize, sdepth);
+ UMat sum = _sum.getUMat();
- size_t gt2 = t_sum.cols * 32, lt2 = 256;
- return k2.run(1, >2, <2, false);
+ krows.args(ocl::KernelArg::ReadOnlyNoSize(buf), ocl::KernelArg::WriteOnly(sum));
+ gt = src.rows;
+ return krows.run(1, >, <, false);
}
static bool ocl_integral( InputArray _src, OutputArray _sum, OutputArray _sqsum, int sdepth, int sqdepth )
{
bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
- if ( _src.type() != CV_8UC1 || _src.step() % vlen != 0 || _src.offset() % vlen != 0 ||
- (!doubleSupport && (sdepth == CV_64F || sqdepth == CV_64F)) )
- return false;
-
- char cvt[40];
- String opts = format("-D sdepth=%d -D sqdepth=%d -D TYPE=%s -D TYPE4=%s4 -D convert_TYPE4=%s%s",
- sdepth, sqdepth, ocl::typeToStr(sqdepth), ocl::typeToStr(sqdepth),
- ocl::convertTypeStr(sdepth, sqdepth, 4, cvt),
- doubleSupport ? " -D DOUBLE_SUPPORT" : "");
-
- ocl::Kernel k1("integral_cols", ocl::imgproc::integral_sqrsum_oclsrc, opts);
- if (k1.empty())
+ if ( _src.type() != CV_8UC1 || (!doubleSupport && (sdepth == CV_64F || sqdepth == CV_64F)) )
return false;
- Size size = _src.size(), dsize = Size(size.width + 1, size.height + 1),
- t_size = Size(((size.height + vlen - 1) / vlen) * vlen, size.width);
- UMat src = _src.getUMat(), t_sum(t_size, sdepth), t_sqsum(t_size, sqdepth);
- t_sum = t_sum(Range::all(), Range(0, size.height));
- t_sqsum = t_sqsum(Range::all(), Range(0, size.height));
-
- _sum.create(dsize, sdepth);
- _sqsum.create(dsize, sqdepth);
- UMat sum = _sum.getUMat(), sqsum = _sqsum.getUMat();
+ static const int tileSize = 16;
- int offset = (int)src.offset / vlen;
- int pre_invalid = src.offset % vlen;
- int vcols = (pre_invalid + src.cols + vlen - 1) / vlen;
- int sum_offset = (int)(sum.offset / sum.elemSize());
- int sqsum_offset = (int)(sqsum.offset / sqsum.elemSize());
+ String build_opt = format("-D SUM_SQUARE -D sumT=%s -D sumSQT=%s -D LOCAL_SUM_SIZE=%d%s",
+ ocl::typeToStr(sdepth), ocl::typeToStr(sqdepth),
+ tileSize,
+ doubleSupport ? " -D DOUBLE_SUPPORT" : "");
- k1.args(ocl::KernelArg::PtrReadOnly(src), ocl::KernelArg::PtrWriteOnly(t_sum),
- ocl::KernelArg::PtrWriteOnly(t_sqsum), offset, pre_invalid, src.rows,
- src.cols, (int)src.step, (int)t_sum.step, (int)t_sqsum.step);
+ ocl::Kernel kcols("integral_sum_cols", ocl::imgproc::integral_sum_oclsrc, build_opt);
+ if (kcols.empty())
+ return false;
- size_t gt = ((vcols + 1) / 2) * 256, lt = 256;
- if (!k1.run(1, >, <, false))
+ UMat src = _src.getUMat();
+ Size src_size = src.size();
+ Size bufsize(((src_size.height + tileSize - 1) / tileSize) * tileSize, ((src_size.width + tileSize - 1) / tileSize) * tileSize);
+ UMat buf(bufsize, sdepth);
+ UMat buf_sq(bufsize, sqdepth);
+ kcols.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnlyNoSize(buf), ocl::KernelArg::WriteOnlyNoSize(buf_sq));
+ size_t gt = src.cols, lt = tileSize;
+ if (!kcols.run(1, >, <, false))
return false;
- ocl::Kernel k2("integral_rows", ocl::imgproc::integral_sqrsum_oclsrc, opts);
- if (k2.empty())
+ ocl::Kernel krows("integral_sum_rows", ocl::imgproc::integral_sum_oclsrc, build_opt);
+ if (krows.empty())
return false;
- k2.args(ocl::KernelArg::PtrReadOnly(t_sum), ocl::KernelArg::PtrReadOnly(t_sqsum),
- ocl::KernelArg::PtrWriteOnly(sum), ocl::KernelArg::PtrWriteOnly(sqsum),
- t_sum.rows, t_sum.cols, (int)t_sum.step, (int)t_sqsum.step,
- (int)sum.step, (int)sqsum.step, sum_offset, sqsum_offset);
+ Size sumsize(src_size.width + 1, src_size.height + 1);
+ _sum.create(sumsize, sdepth);
+ UMat sum = _sum.getUMat();
+ _sqsum.create(sumsize, sqdepth);
+ UMat sum_sq = _sqsum.getUMat();
- size_t gt2 = t_sum.cols * 32, lt2 = 256;
- return k2.run(1, >2, <2, false);
+ krows.args(ocl::KernelArg::ReadOnlyNoSize(buf), ocl::KernelArg::ReadOnlyNoSize(buf_sq), ocl::KernelArg::WriteOnly(sum), ocl::KernelArg::WriteOnlyNoSize(sum_sq));
+ gt = src.rows;
+ return krows.run(1, >, <, false);
}
#endif
projects / profile / ivi / opencv.git / commitdiff