//
// 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
// 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:
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;
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);
}
//
// 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:
//
__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);
}
/*