From 65b7e2018c07ae5c3b6af81c3b629da968db22f4 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Thu, 21 Nov 2013 14:21:43 +0400 Subject: [PATCH] fixed kernel compilation --- modules/core/src/opencl/arithm.cl | 30 +++++++++++++++--------------- modules/core/src/opencl/copyset.cl | 8 ++++---- 2 files changed, 19 insertions(+), 19 deletions(-) diff --git a/modules/core/src/opencl/arithm.cl b/modules/core/src/opencl/arithm.cl index baba41a01..114abe840 100644 --- a/modules/core/src/opencl/arithm.cl +++ b/modules/core/src/opencl/arithm.cl @@ -68,7 +68,7 @@ #define CV_32S 4 #define CV_32F 5 -#define dstelem *(dstT*)(dstptr + dst_index) +#define dstelem *(__global dstT*)(dstptr + dst_index) #define noconvert(x) x #ifndef workT @@ -76,14 +76,14 @@ #define srcT1 dstT #define srcT2 dstT #define workT dstT - #define srcelem1 *(dstT*)(srcptr1 + src1_index) - #define srcelem2 *(dstT*)(srcptr2 + src2_index) + #define srcelem1 *(__global dstT*)(srcptr1 + src1_index) + #define srcelem2 *(__global dstT*)(srcptr2 + src2_index) #define convertToDT noconvert #else - #define srcelem1 convertToWT1(*(srcT1*)(srcptr1 + src1_index)) - #define srcelem2 convertToWT2(*(srcT2*)(srcptr2 + src2_index)) + #define srcelem1 convertToWT1(*(__global srcT1*)(srcptr1 + src1_index)) + #define srcelem2 convertToWT2(*(__global srcT2*)(srcptr2 + src2_index)) #endif @@ -221,9 +221,9 @@ __kernel void KF(__global const uchar* srcptr1, int srcstep1, int srcoffset1, if (x < cols && y < rows) { - int src1_index = mad24(y, srcstep1, x*sizeof(srcT1) + srcoffset1); - int src2_index = mad24(y, srcstep2, x*sizeof(srcT2) + srcoffset2); - int dst_index = mad24(y, dststep, x*sizeof(dstT) + dstoffset); + int src1_index = mad24(y, srcstep1, x*(int)sizeof(srcT1) + srcoffset1); + int src2_index = mad24(y, srcstep2, x*(int)sizeof(srcT2) + srcoffset2); + int dst_index = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset); PROCESS_ELEM; //printf("(x=%d, y=%d). %d, %d, %d\n", x, y, (int)srcelem1, (int)srcelem2, (int)dstelem); @@ -246,9 +246,9 @@ __kernel void KF(__global const uchar* srcptr1, int srcstep1, int srcoffset1, int mask_index = mad24(y, maskstep, x + maskoffset); if( mask[mask_index] ) { - int src1_index = mad24(y, srcstep1, x*sizeof(srcT1) + srcoffset1); - int src2_index = mad24(y, srcstep2, x*sizeof(srcT2) + srcoffset2); - int dst_index = mad24(y, dststep, x*sizeof(dstT) + dstoffset); + int src1_index = mad24(y, srcstep1, x*(int)sizeof(srcT1) + srcoffset1); + int src2_index = mad24(y, srcstep2, x*(int)sizeof(srcT2) + srcoffset2); + int dst_index = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset); PROCESS_ELEM; } @@ -266,8 +266,8 @@ __kernel void KF(__global const uchar* srcptr1, int srcstep1, int srcoffset1, if (x < cols && y < rows) { - int src1_index = mad24(y, srcstep1, x*sizeof(srcT1) + srcoffset1); - int dst_index = mad24(y, dststep, x*sizeof(dstT) + dstoffset); + int src1_index = mad24(y, srcstep1, x*(int)sizeof(srcT1) + srcoffset1); + int dst_index = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset); PROCESS_ELEM; } @@ -288,8 +288,8 @@ __kernel void KF(__global const uchar* srcptr1, int srcstep1, int srcoffset1, int mask_index = mad24(y, maskstep, x + maskoffset); if( mask[mask_index] ) { - int src1_index = mad24(y, srcstep1, x*sizeof(srcT1) + srcoffset1); - int dst_index = mad24(y, dststep, x*sizeof(dstT) + dstoffset); + int src1_index = mad24(y, srcstep1, x*(int)sizeof(srcT1) + srcoffset1); + int dst_index = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset); PROCESS_ELEM; } diff --git a/modules/core/src/opencl/copyset.cl b/modules/core/src/opencl/copyset.cl index df5bdf588..8fb5a00cf 100644 --- a/modules/core/src/opencl/copyset.cl +++ b/modules/core/src/opencl/copyset.cl @@ -53,8 +53,8 @@ __kernel void setMask(__global const uchar* mask, int maskstep, int maskoffset, int mask_index = mad24(y, maskstep, x + maskoffset); if( mask[mask_index] ) { - int dst_index = mad24(y, dststep, x*sizeof(dstT) + dstoffset); - *(dstT*)(dstptr + dst_index) = value; + int dst_index = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset); + *(__global dstT*)(dstptr + dst_index) = value; } } } @@ -67,7 +67,7 @@ __kernel void set(__global uchar* dstptr, int dststep, int dstoffset, if (x < cols && y < rows) { - int dst_index = mad24(y, dststep, x*sizeof(dstT) + dstoffset); - *(dstT*)(dstptr + dst_index) = value; + int dst_index = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset); + *(__global dstT*)(dstptr + dst_index) = value; } }