From 6fae02c05dc1dbeb3b799f51ab19a72a7ef5c37a Mon Sep 17 00:00:00 2001 From: peng xiao Date: Tue, 28 May 2013 11:12:05 +0800 Subject: [PATCH] Fix some OpenCL kernel file build errors on Mac. --- modules/ocl/src/filtering.cpp | 3 +-- modules/ocl/src/mcwutil.cpp | 34 ++++++++++++++++++------ modules/ocl/src/opencl/arithm_add.cl | 10 +++---- modules/ocl/src/opencl/arithm_add_scalar_mask.cl | 6 ++--- modules/ocl/src/opencl/filtering_morph.cl | 2 +- modules/ocl/src/opencl/imgproc_threshold.cl | 2 +- modules/ocl/src/precomp.hpp | 1 + 7 files changed, 38 insertions(+), 20 deletions(-) diff --git a/modules/ocl/src/filtering.cpp b/modules/ocl/src/filtering.cpp index cc07209..56a70ae 100644 --- a/modules/ocl/src/filtering.cpp +++ b/modules/ocl/src/filtering.cpp @@ -356,8 +356,7 @@ static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel, char compile_option[128]; sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D DILATE %s %s", anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1], - rectKernel?"-D RECTKERNEL":"", - s); + s, rectKernel?"-D RECTKERNEL":""); vector< pair > args; args.push_back(make_pair(sizeof(cl_mem), (void *)&src.data)); args.push_back(make_pair(sizeof(cl_mem), (void *)&dst.data)); diff --git a/modules/ocl/src/mcwutil.cpp b/modules/ocl/src/mcwutil.cpp index 3bcb870..79800c5 100644 --- a/modules/ocl/src/mcwutil.cpp +++ b/modules/ocl/src/mcwutil.cpp @@ -43,9 +43,28 @@ // //M*/ -#define CL_USE_DEPRECATED_OPENCL_1_1_APIS #include "precomp.hpp" +#ifdef __GNUC__ +#if ((__GNUC__ * 100) + __GNUC_MINOR__) >= 402 +#define GCC_DIAG_STR(s) #s +#define GCC_DIAG_JOINSTR(x,y) GCC_DIAG_STR(x ## y) +# define GCC_DIAG_DO_PRAGMA(x) _Pragma (#x) +# define GCC_DIAG_PRAGMA(x) GCC_DIAG_DO_PRAGMA(GCC diagnostic x) +# if ((__GNUC__ * 100) + __GNUC_MINOR__) >= 406 +# define GCC_DIAG_OFF(x) GCC_DIAG_PRAGMA(push) \ +GCC_DIAG_PRAGMA(ignored GCC_DIAG_JOINSTR(-W,x)) +# define GCC_DIAG_ON(x) GCC_DIAG_PRAGMA(pop) +# else +# define GCC_DIAG_OFF(x) GCC_DIAG_PRAGMA(ignored GCC_DIAG_JOINSTR(-W,x)) +# define GCC_DIAG_ON(x) GCC_DIAG_PRAGMA(warning GCC_DIAG_JOINSTR(-W,x)) +# endif +#else +# define GCC_DIAG_OFF(x) +# define GCC_DIAG_ON(x) +#endif +#endif /* __GNUC__ */ + using namespace std; namespace cv @@ -121,6 +140,9 @@ namespace cv build_options, finish_mode); } +#ifdef __GNUC__ + GCC_DIAG_OFF(deprecated-declarations) +#endif cl_mem bindTexture(const oclMat &mat) { cl_mem texture; @@ -180,10 +202,6 @@ namespace cv else #endif { -#ifdef __GNUC__ -#pragma GCC diagnostic push -#pragma GCC diagnostic ignored "-Wdeprecated-declarations" -#endif texture = clCreateImage2D( (cl_context)mat.clCxt->oclContext(), CL_MEM_READ_WRITE, @@ -193,9 +211,6 @@ namespace cv 0, NULL, &err); -#ifdef __GNUC__ -#pragma GCC diagnostic pop -#endif } size_t origin[] = { 0, 0, 0 }; size_t region[] = { mat.cols, mat.rows, 1 }; @@ -225,6 +240,9 @@ namespace cv openCLSafeCall(err); return texture; } +#ifdef __GNUC__ + GCC_DIAG_ON(deprecated-declarations) +#endif void releaseTexture(cl_mem& texture) { openCLFree(texture); diff --git a/modules/ocl/src/opencl/arithm_add.cl b/modules/ocl/src/opencl/arithm_add.cl index 7d4b0a7..070ced4 100644 --- a/modules/ocl/src/opencl/arithm_add.cl +++ b/modules/ocl/src/opencl/arithm_add.cl @@ -127,7 +127,7 @@ __kernel void arithm_add_D2 (__global ushort *src1, int src1_step, int src1_offs #ifdef dst_align #undef dst_align #endif -#define dst_align ((dst_offset >> 1) & 3) +#define dst_align ((dst_offset / 2) & 3) int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); int src2_index = mad24(y, src2_step, (x << 1) + src2_offset - (dst_align << 1)); @@ -165,7 +165,7 @@ __kernel void arithm_add_D3 (__global short *src1, int src1_step, int src1_offse #ifdef dst_align #undef dst_align #endif -#define dst_align ((dst_offset >> 1) & 3) +#define dst_align ((dst_offset / 2) & 3) int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); int src2_index = mad24(y, src2_step, (x << 1) + src2_offset - (dst_align << 1)); @@ -335,7 +335,7 @@ __kernel void arithm_add_with_mask_C1_D2 (__global ushort *src1, int src1_step, #ifdef dst_align #undef dst_align #endif -#define dst_align ((dst_offset >> 1) & 1) +#define dst_align ((dst_offset / 2) & 1) int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); int src2_index = mad24(y, src2_step, (x << 1) + src2_offset - (dst_align << 1)); int mask_index = mad24(y, mask_step, x + mask_offset - dst_align); @@ -375,7 +375,7 @@ __kernel void arithm_add_with_mask_C1_D3 (__global short *src1, int src1_step, i #ifdef dst_align #undef dst_align #endif -#define dst_align ((dst_offset >> 1) & 1) +#define dst_align ((dst_offset / 2) & 1) int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); int src2_index = mad24(y, src2_step, (x << 1) + src2_offset - (dst_align << 1)); int mask_index = mad24(y, mask_step, x + mask_offset - dst_align); @@ -507,7 +507,7 @@ __kernel void arithm_add_with_mask_C2_D0 (__global uchar *src1, int src1_step, i #ifdef dst_align #undef dst_align #endif -#define dst_align ((dst_offset >> 1) & 1) +#define dst_align ((dst_offset / 2) & 1) int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); int src2_index = mad24(y, src2_step, (x << 1) + src2_offset - (dst_align << 1)); int mask_index = mad24(y, mask_step, x + mask_offset - dst_align); diff --git a/modules/ocl/src/opencl/arithm_add_scalar_mask.cl b/modules/ocl/src/opencl/arithm_add_scalar_mask.cl index fdf6592..3dbd376 100644 --- a/modules/ocl/src/opencl/arithm_add_scalar_mask.cl +++ b/modules/ocl/src/opencl/arithm_add_scalar_mask.cl @@ -126,7 +126,7 @@ __kernel void arithm_s_add_with_mask_C1_D2 (__global ushort *src1, int src1_st #ifdef dst_align #undef dst_align #endif -#define dst_align ((dst_offset >> 1) & 1) +#define dst_align ((dst_offset / 2) & 1) int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); int mask_index = mad24(y, mask_step, x + mask_offset - dst_align); @@ -164,7 +164,7 @@ __kernel void arithm_s_add_with_mask_C1_D3 (__global short *src1, int src1_ste #ifdef dst_align #undef dst_align #endif -#define dst_align ((dst_offset >> 1) & 1) +#define dst_align ((dst_offset / 2) & 1) int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); int mask_index = mad24(y, mask_step, x + mask_offset - dst_align); @@ -288,7 +288,7 @@ __kernel void arithm_s_add_with_mask_C2_D0 (__global uchar *src1, int src1_ste #ifdef dst_align #undef dst_align #endif -#define dst_align ((dst_offset >> 1) & 1) +#define dst_align ((dst_offset / 2) & 1) int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); int mask_index = mad24(y, mask_step, x + mask_offset - dst_align); diff --git a/modules/ocl/src/opencl/filtering_morph.cl b/modules/ocl/src/opencl/filtering_morph.cl index 4964000..e659a59 100644 --- a/modules/ocl/src/opencl/filtering_morph.cl +++ b/modules/ocl/src/opencl/filtering_morph.cl @@ -120,7 +120,7 @@ __kernel void morph_C1_D0(__global const uchar * restrict src, int gidy = get_global_id(1); int out_addr = mad24(gidy,dst_step_in_pixel,gidx+dst_offset_in_pixel); - if(gidx+3= 0 && dpos < dst_cols; - ddata = convert_float4(con) != 0 ? ddata : dVal; + ddata = convert_float4(con) != (float4)(0) ? ddata : dVal; if(dstart < dst_cols) { *(__global float4*)(dst+dst_offset+gy*dst_step+dstart) = ddata; diff --git a/modules/ocl/src/precomp.hpp b/modules/ocl/src/precomp.hpp index b2a3e41..4f93eac 100644 --- a/modules/ocl/src/precomp.hpp +++ b/modules/ocl/src/precomp.hpp @@ -78,6 +78,7 @@ #if defined (HAVE_OPENCL) +#define CL_USE_DEPRECATED_OPENCL_1_1_APIS #include "opencv2/ocl/private/util.hpp" #include "safe_call.hpp" -- 2.7.4