From 09bcc061dd32361b4864c3b1e62ef5d7121571e5 Mon Sep 17 00:00:00 2001 From: vbystricky Date: Thu, 19 Jun 2014 14:39:49 +0400 Subject: [PATCH] Change kernel for optimization. Remove restriction to align data Fix kernel compilation errors on AMD system Fix licanse information in cl file Support CV_64F destination type Change build options of the kernel Optimize sum of square Remove separate kernel for integral square Increase epsilon for perfomance tests Increase epsilon for perfomance tests Test double support on AMD devices Fix some issues Try to fix problems with AMD device Try to solve problem with AMD device Fix error of destination size in kernel Fix warnings --- modules/imgproc/perf/opencl/perf_imgproc.cpp | 6 +- modules/imgproc/src/opencl/integral_sqrsum.cl | 512 ------------------ modules/imgproc/src/opencl/integral_sum.cl | 374 +++++-------- modules/imgproc/src/sumpixels.cpp | 120 ++-- 4 files changed, 193 insertions(+), 819 deletions(-) delete mode 100644 modules/imgproc/src/opencl/integral_sqrsum.cl diff --git a/modules/imgproc/perf/opencl/perf_imgproc.cpp b/modules/imgproc/perf/opencl/perf_imgproc.cpp index 1b3ba7f19e..7f0770853d 100644 --- a/modules/imgproc/perf/opencl/perf_imgproc.cpp +++ b/modules/imgproc/perf/opencl/perf_imgproc.cpp @@ -231,7 +231,7 @@ OCL_PERF_TEST_P(IntegralFixture, Integral1, ::testing::Combine(OCL_TEST_SIZES, O 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))) @@ -243,11 +243,11 @@ OCL_PERF_TEST_P(IntegralFixture, Integral2, ::testing::Combine(OCL_TEST_SIZES, O 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); } diff --git a/modules/imgproc/src/opencl/integral_sqrsum.cl b/modules/imgproc/src/opencl/integral_sqrsum.cl deleted file mode 100644 index 8b5d2452b8..0000000000 --- a/modules/imgproc/src/opencl/integral_sqrsum.cl +++ /dev/null @@ -1,512 +0,0 @@ -/*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 diff --git a/modules/imgproc/src/opencl/integral_sum.cl b/modules/imgproc/src/opencl/integral_sum.cl index 333c7121cb..49a3bde955 100644 --- a/modules/imgproc/src/opencl/integral_sum.cl +++ b/modules/imgproc/src/opencl/integral_sum.cl @@ -1,46 +1,9 @@ /*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 @@ -51,237 +14,170 @@ #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); diff --git a/modules/imgproc/src/sumpixels.cpp b/modules/imgproc/src/sumpixels.cpp index 1d246ec7bc..e7694b01a9 100755 --- a/modules/imgproc/src/sumpixels.cpp +++ b/modules/imgproc/src/sumpixels.cpp @@ -235,97 +235,87 @@ typedef void (*IntegralFunc)(const uchar* src, size_t srcstep, uchar* sum, size_ #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