Merge pull request #2167 from ilya-lavrenov:tapi_kern_warn

pull/2158/merge
Andrey Pavlenko 11 years ago committed by OpenCV Buildbot
commit 801d2d0a32
  1. 7
      modules/core/src/opencl/arithm.cl
  2. 4
      modules/imgproc/src/color.cpp
  3. 2
      modules/imgproc/src/imgwarp.cpp
  4. 4
      modules/imgproc/src/opencl/cvtcolor.cl
  5. 2
      modules/imgproc/src/opencl/filterSepCol.cl
  6. 8
      modules/imgproc/src/opencl/filterSepRow.cl
  7. 1
      modules/imgproc/src/opencl/resize.cl

@ -196,12 +196,12 @@
#elif defined OP_PHASE_RADIANS #elif defined OP_PHASE_RADIANS
#define PROCESS_ELEM \ #define PROCESS_ELEM \
workT tmp = atan2(srcelem2, srcelem1); \ workT tmp = atan2(srcelem2, srcelem1); \
if(tmp < 0) tmp += 6.283185307179586232; \ if(tmp < 0) tmp += 6.283185307179586232f; \
dstelem = tmp dstelem = tmp
#elif defined OP_PHASE_DEGREES #elif defined OP_PHASE_DEGREES
#define PROCESS_ELEM \ #define PROCESS_ELEM \
workT tmp = atan2(srcelem2, srcelem1)*57.29577951308232286465; \ workT tmp = atan2(srcelem2, srcelem1)*57.29577951308232286465f; \
if(tmp < 0) tmp += 360; \ if(tmp < 0) tmp += 360; \
dstelem = tmp dstelem = tmp
@ -223,7 +223,6 @@ dstelem = v > (dstT)(0) ? log(v) : log(-v)
#define dstT uchar #define dstT uchar
#define srcT2 srcT1 #define srcT2 srcT1
#define convertToWT1 #define convertToWT1
#define convertToWT2
#define PROCESS_ELEM dstelem = convert_uchar(srcelem1 CMP_OPERATOR srcelem2 ? 255 : 0) #define PROCESS_ELEM dstelem = convert_uchar(srcelem1 CMP_OPERATOR srcelem2 ? 255 : 0)
#elif defined OP_CONVERT_SCALE_ABS #elif defined OP_CONVERT_SCALE_ABS
@ -313,7 +312,9 @@ __kernel void KF(__global const uchar* srcptr1, int srcstep1, int srcoffset1,
if (x < cols && y < rows) if (x < cols && y < rows)
{ {
int src1_index = mad24(y, srcstep1, x*(int)sizeof(srcT1) + srcoffset1); int src1_index = mad24(y, srcstep1, x*(int)sizeof(srcT1) + srcoffset1);
#if !(defined(OP_RECIP_SCALE) || defined(OP_NOT))
int src2_index = mad24(y, srcstep2, x*(int)sizeof(srcT2) + srcoffset2); int src2_index = mad24(y, srcstep2, x*(int)sizeof(srcT2) + srcoffset2);
#endif
int dst_index = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset); int dst_index = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset);
EXTRA_INDEX; EXTRA_INDEX;

@ -2996,7 +2996,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
} }
else else
k.create(kernelName.c_str(), ocl::imgproc::cvtcolor_oclsrc, k.create(kernelName.c_str(), ocl::imgproc::cvtcolor_oclsrc,
format("-D depth=%d -D hscale=%f -D bidx=%d -D scn=%d -D dcn=3", depth, hrange*(1.f/360.f), bidx, scn)); format("-D depth=%d -D hscale=%ff -D bidx=%d -D scn=%d -D dcn=3", depth, hrange*(1.f/360.f), bidx, scn));
break; break;
} }
case COLOR_HSV2BGR: case COLOR_HSV2RGB: case COLOR_HSV2BGR_FULL: case COLOR_HSV2RGB_FULL: case COLOR_HSV2BGR: case COLOR_HSV2RGB: case COLOR_HSV2BGR_FULL: case COLOR_HSV2RGB_FULL:
@ -3014,7 +3014,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
String kernelName = String(is_hsv ? "HSV" : "HLS") + "2RGB"; String kernelName = String(is_hsv ? "HSV" : "HLS") + "2RGB";
k.create(kernelName.c_str(), ocl::imgproc::cvtcolor_oclsrc, k.create(kernelName.c_str(), ocl::imgproc::cvtcolor_oclsrc,
format("-D depth=%d -D dcn=%d -D scn=3 -D bidx=%d -D hrange=%d -D hscale=%f", format("-D depth=%d -D dcn=%d -D scn=3 -D bidx=%d -D hrange=%d -D hscale=%ff",
depth, dcn, bidx, hrange, 6.f/hrange)); depth, dcn, bidx, hrange, 6.f/hrange));
break; break;
} }

@ -2004,7 +2004,7 @@ static bool ocl_resize( InputArray _src, OutputArray _dst, Size dsize,
{ {
int wdepth2 = std::max(CV_32F, depth), wtype2 = CV_MAKE_TYPE(wdepth2, cn); int wdepth2 = std::max(CV_32F, depth), wtype2 = CV_MAKE_TYPE(wdepth2, cn);
buildOption = buildOption + format(" -D convertToT=%s -D WT2V=%s -D convertToWT2V=%s -D INTER_AREA_FAST" buildOption = buildOption + format(" -D convertToT=%s -D WT2V=%s -D convertToWT2V=%s -D INTER_AREA_FAST"
" -D XSCALE=%d -D YSCALE=%d -D SCALE=%f", " -D XSCALE=%d -D YSCALE=%d -D SCALE=%ff",
ocl::convertTypeStr(wdepth2, depth, cn, cvt[0]), ocl::convertTypeStr(wdepth2, depth, cn, cvt[0]),
ocl::typeToStr(wtype2), ocl::convertTypeStr(wdepth, wdepth2, cn, cvt[1]), ocl::typeToStr(wtype2), ocl::convertTypeStr(wdepth, wdepth2, cn, cvt[1]),
iscale_x, iscale_y, 1.0f / (iscale_x * iscale_y)); iscale_x, iscale_y, 1.0f / (iscale_x * iscale_y));

@ -472,7 +472,7 @@ __kernel void RGB(__global const uchar* srcptr, int src_step, int src_offset,
dst[0] = src[2]; dst[0] = src[2];
dst[1] = src[1]; dst[1] = src[1];
dst[2] = src[0]; dst[2] = src[0];
#elif defined ORDER #else
dst[0] = src[0]; dst[0] = src[0];
dst[1] = src[1]; dst[1] = src[1];
dst[2] = src[2]; dst[2] = src[2];
@ -728,7 +728,7 @@ __kernel void RGB2HSV(__global const uchar* srcptr, int src_step, int src_offset
diff = v - vmin; diff = v - vmin;
s = diff/(float)(fabs(v) + FLT_EPSILON); s = diff/(float)(fabs(v) + FLT_EPSILON);
diff = (float)(60./(diff + FLT_EPSILON)); diff = (float)(60.f/(diff + FLT_EPSILON));
if( v == r ) if( v == r )
h = (g - b)*diff; h = (g - b)*diff;
else if( v == g ) else if( v == g )

@ -71,7 +71,7 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter
const int dst_step_in_pixel, const int dst_step_in_pixel,
const int dst_cols, const int dst_cols,
const int dst_rows, const int dst_rows,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSY+1))))) __constant float * mat_kernel)
{ {
int x = get_global_id(0); int x = get_global_id(0);
int y = get_global_id(1); int y = get_global_id(1);

@ -154,7 +154,7 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
int dst_step_in_pixel, int dst_step_in_pixel,
int dst_cols, int dst_rows, int dst_cols, int dst_rows,
int radiusy, int radiusy,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1))))) __constant float * mat_kernel)
{ {
int x = get_global_id(0)<<2; int x = get_global_id(0)<<2;
int y = get_global_id(1); int y = get_global_id(1);
@ -298,7 +298,7 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
int dst_step_in_pixel, int dst_step_in_pixel,
int dst_cols, int dst_rows, int dst_cols, int dst_rows,
int radiusy, int radiusy,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1))))) __constant float * mat_kernel)
{ {
int x = get_global_id(0); int x = get_global_id(0);
int y = get_global_id(1); int y = get_global_id(1);
@ -392,7 +392,7 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
int dst_step_in_pixel, int dst_step_in_pixel,
int dst_cols, int dst_rows, int dst_cols, int dst_rows,
int radiusy, int radiusy,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1))))) __constant float * mat_kernel)
{ {
int x = get_global_id(0); int x = get_global_id(0);
int y = get_global_id(1); int y = get_global_id(1);
@ -485,7 +485,7 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
int dst_step_in_pixel, int dst_step_in_pixel,
int dst_cols, int dst_rows, int dst_cols, int dst_rows,
int radiusy, int radiusy,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1))))) __constant float * mat_kernel)
{ {
int x = get_global_id(0); int x = get_global_id(0);
int y = get_global_id(1); int y = get_global_id(1);

@ -78,7 +78,6 @@ __kernel void resizeLN(__global const uchar* srcptr, int srcstep, int srcoffset,
int y_ = INC(y,srcrows); int y_ = INC(y,srcrows);
int x_ = INC(x,srccols); int x_ = INC(x,srccols);
__global const PIXTYPE* src = (__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE));
#if depth <= 4 #if depth <= 4

Loading…
Cancel
Save