generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Fix ocl::equalizeHist mismatch

Browse Source
This commit is contained in:
yao 2013-02-22 15:48:31 +08:00
parent 6ebb0e2ad2
commit 7810254198
2 changed files with 14 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
modules/ocl/src/imgproc.cpp
View File

@ -12,6 +12,7 @@
//
// 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.
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
@ -23,6 +24,7 @@
// Zhang Ying, zhangying913@gmail.com
// Xu Pang, pangxu010@163.com
// Wu Zailong, bullet@yeah.net
// Wenju He, wenju@multicorewareinc.com
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
@ -1524,7 +1526,7 @@ namespace cv
mat_dst.create(mat_src.rows, mat_src.cols, CV_8UC1);
oclMat mat_hist(1, 256, CV_32SC1);
//mat_hist.setTo(0);
calcHist(mat_src, mat_hist);
Context *clCxt = mat_src.clCxt;
@ -1533,10 +1535,10 @@ namespace cv
size_t globalThreads[3] = { 256, 1, 1};
oclMat lut(1, 256, CV_8UC1);
vector<pair<size_t , const void *> > args;
float scale = 255.f / (mat_src.rows * mat_src.cols);
int total = mat_src.rows * mat_src.cols;
args.push_back( make_pair( sizeof(cl_mem), (void *)&lut.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_hist.data));
args.push_back( make_pair( sizeof(cl_float), (void *)&scale));
args.push_back( make_pair( sizeof(int), (void *)&total));
openCLExecuteKernel(clCxt, &imgproc_histogram, kernelName, globalThreads, localThreads, args, -1, -1);
LUT(mat_src, lut, mat_dst);
}

13
modules/ocl/src/kernels/imgproc_histogram.cl
View File

@ -3,12 +3,14 @@
//
// 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.
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
// Niko Li, newlife20080214@gmail.com
// Jia Haipeng, jiahaipeng95@gmail.com
// Xu Pang, pangxu010@163.com
// Wenju He, wenju@multicorewareinc.com
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
@ -189,24 +191,27 @@ __kernel __attribute__((reqd_work_group_size(256,1,1)))void merge_hist(__global
__kernel __attribute__((reqd_work_group_size(256,1,1)))void calLUT(
__global uchar * dst,
__constant int * hist,
float scale)
int total)
{
int lid = get_local_id(0);
__local int sumhist[HISTOGRAM256_BIN_COUNT];
//__local uchar lut[HISTOGRAM256_BIN_COUNT+1];
__local int sumhist[HISTOGRAM256_BIN_COUNT+1];
sumhist[lid]=hist[lid];
barrier(CLK_LOCAL_MEM_FENCE);
if(lid==0)
{
int sum = 0;
for(int i=0;i<HISTOGRAM256_BIN_COUNT;i++)
int i = 0;
while (!sumhist[i]) ++i;
sumhist[HISTOGRAM256_BIN_COUNT] = sumhist[i];
for(sumhist[i++] = 0; i<HISTOGRAM256_BIN_COUNT; i++)
{
sum+=sumhist[i];
sumhist[i]=sum;
}
}
barrier(CLK_LOCAL_MEM_FENCE);
float scale = 255.f/(total - sumhist[HISTOGRAM256_BIN_COUNT]);
dst[lid]= lid == 0 ? 0 : convert_uchar_sat(convert_float(sumhist[lid])*scale);
}
/*