Add super resolution's OpenCL implementation.
Accuracy and performance tests are also provided.
This commit is contained in:
peng xiao
261
modules/superres/src/opencl/superres_btvl1.cl
Normal file
261
modules/superres/src/opencl/superres_btvl1.cl
Normal file
@@ -0,0 +1,261 @@
|
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Jin Ma jin@multicorewareinc.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 oclMaterials 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.
|
||||
//
|
||||
//M*/
|
||||
|
||||
__kernel void buildMotionMapsKernel(__global float* forwardMotionX,
|
||||
__global float* forwardMotionY,
|
||||
__global float* backwardMotionX,
|
||||
__global float* backwardMotionY,
|
||||
__global float* forwardMapX,
|
||||
__global float* forwardMapY,
|
||||
__global float* backwardMapX,
|
||||
__global float* backwardMapY,
|
||||
int forwardMotionX_row,
|
||||
int forwardMotionX_col,
|
||||
int forwardMotionX_step,
|
||||
int forwardMotionY_step,
|
||||
int backwardMotionX_step,
|
||||
int backwardMotionY_step,
|
||||
int forwardMapX_step,
|
||||
int forwardMapY_step,
|
||||
int backwardMapX_step,
|
||||
int backwardMapY_step
|
||||
)
|
||||
{
|
||||
int x = get_global_id(0);
|
||||
int y = get_global_id(1);
|
||||
|
||||
if(x < forwardMotionX_col && y < forwardMotionX_row)
|
||||
{
|
||||
float fx = forwardMotionX[y * forwardMotionX_step + x];
|
||||
float fy = forwardMotionY[y * forwardMotionY_step + x];
|
||||
|
||||
float bx = backwardMotionX[y * backwardMotionX_step + x];
|
||||
float by = backwardMotionY[y * backwardMotionY_step + x];
|
||||
|
||||
forwardMapX[y * forwardMapX_step + x] = x + bx;
|
||||
forwardMapY[y * forwardMapY_step + x] = y + by;
|
||||
|
||||
backwardMapX[y * backwardMapX_step + x] = x + fx;
|
||||
backwardMapY[y * backwardMapY_step + x] = y + fy;
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void upscaleKernel(__global float* src,
|
||||
__global float* dst,
|
||||
int src_step,
|
||||
int dst_step,
|
||||
int src_row,
|
||||
int src_col,
|
||||
int scale,
|
||||
int channels
|
||||
)
|
||||
{
|
||||
int x = get_global_id(0);
|
||||
int y = get_global_id(1);
|
||||
|
||||
if(x < src_col && y < src_row)
|
||||
{
|
||||
if(channels == 1)
|
||||
{
|
||||
dst[y * scale * dst_step + x * scale] = src[y * src_step + x];
|
||||
}else if(channels == 3)
|
||||
{
|
||||
dst[y * channels * scale * dst_step + 3 * x * scale + 0] = src[y * channels * src_step + 3 * x + 0];
|
||||
dst[y * channels * scale * dst_step + 3 * x * scale + 1] = src[y * channels * src_step + 3 * x + 1];
|
||||
dst[y * channels * scale * dst_step + 3 * x * scale + 2] = src[y * channels * src_step + 3 * x + 2];
|
||||
}else
|
||||
{
|
||||
dst[y * channels * scale * dst_step + 4 * x * scale + 0] = src[y * channels * src_step + 4 * x + 0];
|
||||
dst[y * channels * scale * dst_step + 4 * x * scale + 1] = src[y * channels * src_step + 4 * x + 1];
|
||||
dst[y * channels * scale * dst_step + 4 * x * scale + 2] = src[y * channels * src_step + 4 * x + 2];
|
||||
dst[y * channels * scale * dst_step + 4 * x * scale + 3] = src[y * channels * src_step + 4 * x + 3];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
float diffSign(float a, float b)
|
||||
{
|
||||
return a > b ? 1.0f : a < b ? -1.0f : 0.0f;
|
||||
}
|
||||
|
||||
float3 diffSign3(float3 a, float3 b)
|
||||
{
|
||||
float3 pos;
|
||||
pos.x = a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f;
|
||||
pos.y = a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f;
|
||||
pos.z = a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f;
|
||||
return pos;
|
||||
}
|
||||
|
||||
float4 diffSign4(float4 a, float4 b)
|
||||
{
|
||||
float4 pos;
|
||||
pos.x = a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f;
|
||||
pos.y = a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f;
|
||||
pos.z = a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f;
|
||||
pos.w = 0.0f;
|
||||
return pos;
|
||||
}
|
||||
|
||||
__kernel void diffSignKernel(__global float* src1,
|
||||
__global float* src2,
|
||||
__global float* dst,
|
||||
int src1_row,
|
||||
int src1_col,
|
||||
int dst_step,
|
||||
int src1_step,
|
||||
int src2_step)
|
||||
{
|
||||
int x = get_global_id(0);
|
||||
int y = get_global_id(1);
|
||||
|
||||
if(x < src1_col && y < src1_row)
|
||||
{
|
||||
dst[y * dst_step + x] = diffSign(src1[y * src1_step + x], src2[y * src2_step + x]);
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
__kernel void calcBtvRegularizationKernel(__global float* src,
|
||||
__global float* dst,
|
||||
int src_step,
|
||||
int dst_step,
|
||||
int src_row,
|
||||
int src_col,
|
||||
int ksize,
|
||||
int channels,
|
||||
__global float* c_btvRegWeights
|
||||
)
|
||||
{
|
||||
int x = get_global_id(0) + ksize;
|
||||
int y = get_global_id(1) + ksize;
|
||||
|
||||
if ((y < src_row - ksize) && (x < src_col - ksize))
|
||||
{
|
||||
if(channels == 1)
|
||||
{
|
||||
const float srcVal = src[y * src_step + x];
|
||||
float dstVal = 0.0f;
|
||||
|
||||
for (int m = 0, count = 0; m <= ksize; ++m)
|
||||
{
|
||||
for (int l = ksize; l + m >= 0; --l, ++count)
|
||||
dstVal = dstVal + c_btvRegWeights[count] * (diffSign(srcVal, src[(y + m) * src_step + (x + l)]) - diffSign(src[(y - m) * src_step + (x - l)], srcVal));
|
||||
}
|
||||
dst[y * dst_step + x] = dstVal;
|
||||
}else if(channels == 3)
|
||||
{
|
||||
float3 srcVal;
|
||||
srcVal.x = src[y * src_step + 3 * x + 0];
|
||||
srcVal.y = src[y * src_step + 3 * x + 1];
|
||||
srcVal.z = src[y * src_step + 3 * x + 2];
|
||||
|
||||
float3 dstVal;
|
||||
dstVal.x = 0.0f;
|
||||
dstVal.y = 0.0f;
|
||||
dstVal.z = 0.0f;
|
||||
|
||||
for (int m = 0, count = 0; m <= ksize; ++m)
|
||||
{
|
||||
for (int l = ksize; l + m >= 0; --l, ++count)
|
||||
{
|
||||
float3 src1;
|
||||
src1.x = src[(y + m) * src_step + 3 * (x + l) + 0];
|
||||
src1.y = src[(y + m) * src_step + 3 * (x + l) + 1];
|
||||
src1.z = src[(y + m) * src_step + 3 * (x + l) + 2];
|
||||
|
||||
float3 src2;
|
||||
src2.x = src[(y - m) * src_step + 3 * (x - l) + 0];
|
||||
src2.y = src[(y - m) * src_step + 3 * (x - l) + 1];
|
||||
src2.z = src[(y - m) * src_step + 3 * (x - l) + 2];
|
||||
|
||||
dstVal = dstVal + c_btvRegWeights[count] * (diffSign3(srcVal, src1) - diffSign3(src2, srcVal));
|
||||
}
|
||||
}
|
||||
dst[y * dst_step + 3 * x + 0] = dstVal.x;
|
||||
dst[y * dst_step + 3 * x + 1] = dstVal.y;
|
||||
dst[y * dst_step + 3 * x + 2] = dstVal.z;
|
||||
}else
|
||||
{
|
||||
float4 srcVal;
|
||||
srcVal.x = src[y * src_step + 4 * x + 0];//r type =float
|
||||
srcVal.y = src[y * src_step + 4 * x + 1];//g
|
||||
srcVal.z = src[y * src_step + 4 * x + 2];//b
|
||||
srcVal.w = src[y * src_step + 4 * x + 3];//a
|
||||
|
||||
float4 dstVal;
|
||||
dstVal.x = 0.0f;
|
||||
dstVal.y = 0.0f;
|
||||
dstVal.z = 0.0f;
|
||||
dstVal.w = 0.0f;
|
||||
|
||||
for (int m = 0, count = 0; m <= ksize; ++m)
|
||||
{
|
||||
for (int l = ksize; l + m >= 0; --l, ++count)
|
||||
{
|
||||
float4 src1;
|
||||
src1.x = src[(y + m) * src_step + 4 * (x + l) + 0];
|
||||
src1.y = src[(y + m) * src_step + 4 * (x + l) + 1];
|
||||
src1.z = src[(y + m) * src_step + 4 * (x + l) + 2];
|
||||
src1.w = src[(y + m) * src_step + 4 * (x + l) + 3];
|
||||
|
||||
float4 src2;
|
||||
src2.x = src[(y - m) * src_step + 4 * (x - l) + 0];
|
||||
src2.y = src[(y - m) * src_step + 4 * (x - l) + 1];
|
||||
src2.z = src[(y - m) * src_step + 4 * (x - l) + 2];
|
||||
src2.w = src[(y - m) * src_step + 4 * (x - l) + 3];
|
||||
|
||||
dstVal = dstVal + c_btvRegWeights[count] * (diffSign4(srcVal, src1) - diffSign4(src2, srcVal));
|
||||
|
||||
}
|
||||
}
|
||||
dst[y * dst_step + 4 * x + 0] = dstVal.x;
|
||||
dst[y * dst_step + 4 * x + 1] = dstVal.y;
|
||||
dst[y * dst_step + 4 * x + 2] = dstVal.z;
|
||||
dst[y * dst_step + 4 * x + 3] = dstVal.w;
|
||||
}
|
||||
}
|
||||
}
|
||||
Reference in New Issue
Block a user