139 lines
5.3 KiB
Common Lisp
139 lines
5.3 KiB
Common Lisp
// License Agreement
|
|
// For Open Source Computer Vision Library
|
|
//
|
|
// 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.
|
|
// Third party copyrights are property of their respective owners.
|
|
//
|
|
// @Authors
|
|
// Rock Li, Rock.li@amd.com
|
|
//
|
|
// Redistribution and use in source and binary forms, with or without modification,
|
|
// are permitted provided that the following conditions are met:
|
|
//
|
|
// * Redistribution's of source code must retain the above copyright notice,
|
|
// this list of conditions and the following disclaimer.
|
|
//
|
|
// * Redistribution's in binary form must reproduce the above copyright notice,
|
|
// this list of conditions and the following disclaimer in the documentation
|
|
// and/or other GpuMaterials provided with the distribution.
|
|
//
|
|
// * The name of the copyright holders may not be used to endorse or promote products
|
|
// derived from this software without specific prior written permission.
|
|
//
|
|
// This software is provided by the copyright holders and contributors as is and
|
|
// any express or implied warranties, including, but not limited to, the implied
|
|
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
|
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
|
// indirect, incidental, special, exemplary, or consequential damages
|
|
// (including, but not limited to, procurement of substitute goods or services;
|
|
// loss of use, data, or profits; or business interruption) however caused
|
|
// and on any theory of liability, whether in contract, strict liability,
|
|
// or tort (including negligence or otherwise) arising in any way out of
|
|
// the use of this software, even if advised of the possibility of such damage.
|
|
|
|
__kernel void bilateral_C1_D0(__global uchar *dst,
|
|
__global const uchar *src,
|
|
const int dst_rows,
|
|
const int dst_cols,
|
|
const int maxk,
|
|
const int radius,
|
|
const int dst_step,
|
|
const int dst_offset,
|
|
const int src_step,
|
|
const int src_rows,
|
|
const int src_cols,
|
|
__constant float *color_weight,
|
|
__constant float *space_weight,
|
|
__constant int *space_ofs)
|
|
{
|
|
int gidx = get_global_id(0);
|
|
int gidy = get_global_id(1);
|
|
if((gidy<dst_rows) && (gidx<dst_cols))
|
|
{
|
|
int src_addr = mad24(gidy+radius,src_step,gidx+radius);
|
|
int dst_addr = mad24(gidy,dst_step,gidx+dst_offset);
|
|
float sum = 0.f, wsum = 0.f;
|
|
|
|
int val0 = (int)src[src_addr];
|
|
for(int k = 0; k < maxk; k++ )
|
|
{
|
|
int val = (int)src[src_addr + space_ofs[k]];
|
|
float w = space_weight[k]*color_weight[abs(val - val0)];
|
|
sum += (float)(val)*w;
|
|
wsum += w;
|
|
}
|
|
dst[dst_addr] = convert_uchar_rtz(sum/wsum+0.5f);
|
|
}
|
|
}
|
|
__kernel void bilateral2_C1_D0(__global uchar *dst,
|
|
__global const uchar *src,
|
|
const int dst_rows,
|
|
const int dst_cols,
|
|
const int maxk,
|
|
const int radius,
|
|
const int dst_step,
|
|
const int dst_offset,
|
|
const int src_step,
|
|
const int src_rows,
|
|
const int src_cols,
|
|
__constant float *color_weight,
|
|
__constant float *space_weight,
|
|
__constant int *space_ofs)
|
|
{
|
|
int gidx = get_global_id(0)<<2;
|
|
int gidy = get_global_id(1);
|
|
if((gidy<dst_rows) && (gidx<dst_cols))
|
|
{
|
|
int src_addr = mad24(gidy+radius,src_step,gidx+radius);
|
|
int dst_addr = mad24(gidy,dst_step,gidx+dst_offset);
|
|
float4 sum = (float4)(0.f), wsum = (float4)(0.f);
|
|
|
|
int4 val0 = convert_int4(vload4(0,src+src_addr));
|
|
for(int k = 0; k < maxk; k++ )
|
|
{
|
|
int4 val = convert_int4(vload4(0,src+src_addr + space_ofs[k]));
|
|
float4 w = (float4)(space_weight[k])*(float4)(color_weight[abs(val.x - val0.x)],color_weight[abs(val.y - val0.y)],color_weight[abs(val.z - val0.z)],color_weight[abs(val.w - val0.w)]);
|
|
sum += convert_float4(val)*w;
|
|
wsum += w;
|
|
}
|
|
*(__global uchar4*)(dst+dst_addr) = convert_uchar4_rtz(sum/wsum+0.5f);
|
|
}
|
|
}
|
|
__kernel void bilateral_C4_D0(__global uchar4 *dst,
|
|
__global const uchar4 *src,
|
|
const int dst_rows,
|
|
const int dst_cols,
|
|
const int maxk,
|
|
const int radius,
|
|
const int dst_step,
|
|
const int dst_offset,
|
|
const int src_step,
|
|
const int src_rows,
|
|
const int src_cols,
|
|
__constant float *color_weight,
|
|
__constant float *space_weight,
|
|
__constant int *space_ofs)
|
|
{
|
|
int gidx = get_global_id(0);
|
|
int gidy = get_global_id(1);
|
|
if((gidy<dst_rows) && (gidx<dst_cols))
|
|
{
|
|
int src_addr = mad24(gidy+radius,src_step,gidx+radius);
|
|
int dst_addr = mad24(gidy,dst_step,gidx+dst_offset);
|
|
float4 sum = (float4)0.f;
|
|
float wsum = 0.f;
|
|
|
|
int4 val0 = convert_int4(src[src_addr]);
|
|
for(int k = 0; k < maxk; k++ )
|
|
{
|
|
int4 val = convert_int4(src[src_addr + space_ofs[k]]);
|
|
float w = space_weight[k]*color_weight[abs(val.x - val0.x)+abs(val.y - val0.y)+abs(val.z - val0.z)];
|
|
sum += convert_float4(val)*(float4)w;
|
|
wsum += w;
|
|
}
|
|
wsum=1.f/wsum;
|
|
dst[dst_addr] = convert_uchar4_rtz(sum*(float4)wsum+(float4)0.5f);
|
|
}
|
|
}
|