You can not select more than 25 topics
Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
3127 lines
102 KiB
3127 lines
102 KiB
/*
|
|
Copyright 1999-2021 ImageMagick Studio LLC, a non-profit organization
|
|
dedicated to making software imaging solutions freely available.
|
|
|
|
You may not use this file except in compliance with the License. You may
|
|
obtain a copy of the License at
|
|
|
|
https://imagemagick.org/script/license.php
|
|
|
|
Unless required by applicable law or agreed to in writing, software
|
|
distributed under the License is distributed on an "AS IS" BASIS,
|
|
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
See the License for the specific language governing permissions and
|
|
limitations under the License.
|
|
|
|
MagickCore private kernels for accelerated functions.
|
|
*/
|
|
|
|
#ifndef MAGICKCORE_ACCELERATE_KERNELS_PRIVATE_H
|
|
#define MAGICKCORE_ACCELERATE_KERNELS_PRIVATE_H
|
|
|
|
#if defined(__cplusplus) || defined(c_plusplus)
|
|
extern "C" {
|
|
#endif
|
|
|
|
#if defined(MAGICKCORE_OPENCL_SUPPORT)
|
|
|
|
/*
|
|
Define declarations.
|
|
*/
|
|
#define OPENCL_DEFINE(VAR,...) "\n #""define " #VAR " " #__VA_ARGS__ " \n"
|
|
#define OPENCL_ELIF(...) "\n #""elif " #__VA_ARGS__ " \n"
|
|
#define OPENCL_ELSE() "\n #""else " " \n"
|
|
#define OPENCL_ENDIF() "\n #""endif " " \n"
|
|
#define OPENCL_IF(...) "\n #""if " #__VA_ARGS__ " \n"
|
|
#define STRINGIFY(...) #__VA_ARGS__ "\n"
|
|
|
|
const char *accelerateKernels =
|
|
|
|
/*
|
|
Define declarations.
|
|
*/
|
|
OPENCL_DEFINE(SigmaUniform, (attenuate*0.015625f))
|
|
OPENCL_DEFINE(SigmaGaussian, (attenuate*0.015625f))
|
|
OPENCL_DEFINE(SigmaImpulse, (attenuate*0.1f))
|
|
OPENCL_DEFINE(SigmaLaplacian, (attenuate*0.0390625f))
|
|
OPENCL_DEFINE(SigmaMultiplicativeGaussian, (attenuate*0.5f))
|
|
OPENCL_DEFINE(SigmaPoisson, (attenuate*12.5f))
|
|
OPENCL_DEFINE(SigmaRandom, (attenuate))
|
|
OPENCL_DEFINE(TauGaussian, (attenuate*0.078125f))
|
|
OPENCL_DEFINE(MagickMax(x,y), (((x) > (y)) ? (x) : (y)))
|
|
OPENCL_DEFINE(MagickMin(x,y), (((x) < (y)) ? (x) : (y)))
|
|
|
|
/*
|
|
Typedef declarations.
|
|
*/
|
|
STRINGIFY(
|
|
typedef enum
|
|
{
|
|
UndefinedColorspace,
|
|
CMYColorspace, /* negated linear RGB colorspace */
|
|
CMYKColorspace, /* CMY with Black separation */
|
|
GRAYColorspace, /* Single Channel greyscale (non-linear) image */
|
|
HCLColorspace,
|
|
HCLpColorspace,
|
|
HSBColorspace,
|
|
HSIColorspace,
|
|
HSLColorspace,
|
|
HSVColorspace, /* alias for HSB */
|
|
HWBColorspace,
|
|
LabColorspace,
|
|
LCHColorspace, /* alias for LCHuv */
|
|
LCHabColorspace, /* Cylindrical (Polar) Lab */
|
|
LCHuvColorspace, /* Cylindrical (Polar) Luv */
|
|
LogColorspace,
|
|
LMSColorspace,
|
|
LuvColorspace,
|
|
OHTAColorspace,
|
|
Rec601YCbCrColorspace,
|
|
Rec709YCbCrColorspace,
|
|
RGBColorspace, /* Linear RGB colorspace */
|
|
scRGBColorspace, /* ??? */
|
|
sRGBColorspace, /* Default: non-linear sRGB colorspace */
|
|
TransparentColorspace,
|
|
xyYColorspace,
|
|
XYZColorspace, /* IEEE Color Reference colorspace */
|
|
YCbCrColorspace,
|
|
YCCColorspace,
|
|
YDbDrColorspace,
|
|
YIQColorspace,
|
|
YPbPrColorspace,
|
|
YUVColorspace,
|
|
LinearGRAYColorspace /* Single Channel greyscale (linear) image */
|
|
} ColorspaceType;
|
|
)
|
|
|
|
STRINGIFY(
|
|
typedef enum
|
|
{
|
|
UndefinedCompositeOp,
|
|
AlphaCompositeOp,
|
|
AtopCompositeOp,
|
|
BlendCompositeOp,
|
|
BlurCompositeOp,
|
|
BumpmapCompositeOp,
|
|
ChangeMaskCompositeOp,
|
|
ClearCompositeOp,
|
|
ColorBurnCompositeOp,
|
|
ColorDodgeCompositeOp,
|
|
ColorizeCompositeOp,
|
|
CopyBlackCompositeOp,
|
|
CopyBlueCompositeOp,
|
|
CopyCompositeOp,
|
|
CopyCyanCompositeOp,
|
|
CopyGreenCompositeOp,
|
|
CopyMagentaCompositeOp,
|
|
CopyAlphaCompositeOp,
|
|
CopyRedCompositeOp,
|
|
CopyYellowCompositeOp,
|
|
DarkenCompositeOp,
|
|
DarkenIntensityCompositeOp,
|
|
DifferenceCompositeOp,
|
|
DisplaceCompositeOp,
|
|
DissolveCompositeOp,
|
|
DistortCompositeOp,
|
|
DivideDstCompositeOp,
|
|
DivideSrcCompositeOp,
|
|
DstAtopCompositeOp,
|
|
DstCompositeOp,
|
|
DstInCompositeOp,
|
|
DstOutCompositeOp,
|
|
DstOverCompositeOp,
|
|
ExclusionCompositeOp,
|
|
HardLightCompositeOp,
|
|
HardMixCompositeOp,
|
|
HueCompositeOp,
|
|
InCompositeOp,
|
|
IntensityCompositeOp,
|
|
LightenCompositeOp,
|
|
LightenIntensityCompositeOp,
|
|
LinearBurnCompositeOp,
|
|
LinearDodgeCompositeOp,
|
|
LinearLightCompositeOp,
|
|
LuminizeCompositeOp,
|
|
MathematicsCompositeOp,
|
|
MinusDstCompositeOp,
|
|
MinusSrcCompositeOp,
|
|
ModulateCompositeOp,
|
|
ModulusAddCompositeOp,
|
|
ModulusSubtractCompositeOp,
|
|
MultiplyCompositeOp,
|
|
NoCompositeOp,
|
|
OutCompositeOp,
|
|
OverCompositeOp,
|
|
OverlayCompositeOp,
|
|
PegtopLightCompositeOp,
|
|
PinLightCompositeOp,
|
|
PlusCompositeOp,
|
|
ReplaceCompositeOp,
|
|
SaturateCompositeOp,
|
|
ScreenCompositeOp,
|
|
SoftLightCompositeOp,
|
|
SrcAtopCompositeOp,
|
|
SrcCompositeOp,
|
|
SrcInCompositeOp,
|
|
SrcOutCompositeOp,
|
|
SrcOverCompositeOp,
|
|
ThresholdCompositeOp,
|
|
VividLightCompositeOp,
|
|
XorCompositeOp,
|
|
StereoCompositeOp
|
|
} CompositeOperator;
|
|
)
|
|
|
|
STRINGIFY(
|
|
typedef enum
|
|
{
|
|
UndefinedFunction,
|
|
ArcsinFunction,
|
|
ArctanFunction,
|
|
PolynomialFunction,
|
|
SinusoidFunction
|
|
} MagickFunction;
|
|
)
|
|
|
|
STRINGIFY(
|
|
typedef enum
|
|
{
|
|
UndefinedNoise,
|
|
UniformNoise,
|
|
GaussianNoise,
|
|
MultiplicativeGaussianNoise,
|
|
ImpulseNoise,
|
|
LaplacianNoise,
|
|
PoissonNoise,
|
|
RandomNoise
|
|
} NoiseType;
|
|
)
|
|
|
|
STRINGIFY(
|
|
typedef enum
|
|
{
|
|
UndefinedPixelIntensityMethod = 0,
|
|
AveragePixelIntensityMethod,
|
|
BrightnessPixelIntensityMethod,
|
|
LightnessPixelIntensityMethod,
|
|
MSPixelIntensityMethod,
|
|
Rec601LumaPixelIntensityMethod,
|
|
Rec601LuminancePixelIntensityMethod,
|
|
Rec709LumaPixelIntensityMethod,
|
|
Rec709LuminancePixelIntensityMethod,
|
|
RMSPixelIntensityMethod
|
|
} PixelIntensityMethod;
|
|
)
|
|
|
|
STRINGIFY(
|
|
typedef enum
|
|
{
|
|
BoxWeightingFunction = 0,
|
|
TriangleWeightingFunction,
|
|
CubicBCWeightingFunction,
|
|
HannWeightingFunction,
|
|
HammingWeightingFunction,
|
|
BlackmanWeightingFunction,
|
|
GaussianWeightingFunction,
|
|
QuadraticWeightingFunction,
|
|
JincWeightingFunction,
|
|
SincWeightingFunction,
|
|
SincFastWeightingFunction,
|
|
KaiserWeightingFunction,
|
|
WelchWeightingFunction,
|
|
BohmanWeightingFunction,
|
|
LagrangeWeightingFunction,
|
|
CosineWeightingFunction
|
|
} ResizeWeightingFunctionType;
|
|
)
|
|
|
|
STRINGIFY(
|
|
typedef enum
|
|
{
|
|
UndefinedChannel = 0x0000,
|
|
RedChannel = 0x0001,
|
|
GrayChannel = 0x0001,
|
|
CyanChannel = 0x0001,
|
|
GreenChannel = 0x0002,
|
|
MagentaChannel = 0x0002,
|
|
BlueChannel = 0x0004,
|
|
YellowChannel = 0x0004,
|
|
BlackChannel = 0x0008,
|
|
AlphaChannel = 0x0010,
|
|
OpacityChannel = 0x0010,
|
|
IndexChannel = 0x0020, /* Color Index Table? */
|
|
ReadMaskChannel = 0x0040, /* Pixel is Not Readable? */
|
|
WriteMaskChannel = 0x0080, /* Pixel is Write Protected? */
|
|
MetaChannel = 0x0100, /* ???? */
|
|
CompositeChannels = 0x001F,
|
|
AllChannels = 0x7ffffff,
|
|
/*
|
|
Special purpose channel types.
|
|
FUTURE: are these needed any more - they are more like hacks
|
|
SyncChannels for example is NOT a real channel but a 'flag'
|
|
It really says -- "User has not defined channels"
|
|
Though it does have extra meaning in the "-auto-level" operator
|
|
*/
|
|
TrueAlphaChannel = 0x0100, /* extract actual alpha channel from opacity */
|
|
RGBChannels = 0x0200, /* set alpha from grayscale mask in RGB */
|
|
GrayChannels = 0x0400,
|
|
SyncChannels = 0x20000, /* channels modified as a single unit */
|
|
DefaultChannels = AllChannels
|
|
} ChannelType; /* must correspond to PixelChannel */
|
|
)
|
|
|
|
/*
|
|
Helper functions.
|
|
*/
|
|
|
|
OPENCL_IF((MAGICKCORE_QUANTUM_DEPTH == 8))
|
|
|
|
STRINGIFY(
|
|
static inline CLQuantum ScaleCharToQuantum(const unsigned char value)
|
|
{
|
|
return((CLQuantum) value);
|
|
}
|
|
)
|
|
|
|
OPENCL_ELIF((MAGICKCORE_QUANTUM_DEPTH == 16))
|
|
|
|
STRINGIFY(
|
|
static inline CLQuantum ScaleCharToQuantum(const unsigned char value)
|
|
{
|
|
return((CLQuantum) (257.0f*value));
|
|
}
|
|
)
|
|
|
|
OPENCL_ELIF((MAGICKCORE_QUANTUM_DEPTH == 32))
|
|
|
|
STRINGIFY(
|
|
static inline CLQuantum ScaleCharToQuantum(const unsigned char value)
|
|
{
|
|
return((CLQuantum) (16843009.0*value));
|
|
}
|
|
)
|
|
|
|
OPENCL_ENDIF()
|
|
|
|
OPENCL_IF((MAGICKCORE_HDRI_SUPPORT == 1))
|
|
|
|
STRINGIFY(
|
|
static inline CLQuantum ClampToQuantum(const float value)
|
|
{
|
|
return (CLQuantum) clamp(value, 0.0f, QuantumRange);
|
|
}
|
|
)
|
|
|
|
OPENCL_ELSE()
|
|
|
|
STRINGIFY(
|
|
static inline CLQuantum ClampToQuantum(const float value)
|
|
{
|
|
return (CLQuantum) (clamp(value, 0.0f, QuantumRange) + 0.5f);
|
|
}
|
|
)
|
|
|
|
OPENCL_ENDIF()
|
|
|
|
STRINGIFY(
|
|
static inline int ClampToCanvas(const int offset,const int range)
|
|
{
|
|
return clamp(offset, (int)0, range-1);
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
static inline uint ScaleQuantumToMap(CLQuantum value)
|
|
{
|
|
if (value >= (CLQuantum) MaxMap)
|
|
return ((uint)MaxMap);
|
|
else
|
|
return ((uint)value);
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
static inline float PerceptibleReciprocal(const float x)
|
|
{
|
|
float sign = x < (float) 0.0 ? (float) -1.0 : (float) 1.0;
|
|
return((sign*x) >= MagickEpsilon ? (float) 1.0/x : sign*((float) 1.0/MagickEpsilon));
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
|
|
static inline unsigned int getPixelIndex(const unsigned int number_channels,
|
|
const unsigned int columns, const unsigned int x, const unsigned int y)
|
|
{
|
|
return (x * number_channels) + (y * columns * number_channels);
|
|
}
|
|
|
|
static inline float getPixelRed(const __global CLQuantum *p) { return (float)*p; }
|
|
static inline float getPixelGreen(const __global CLQuantum *p) { return (float)*(p+1); }
|
|
static inline float getPixelBlue(const __global CLQuantum *p) { return (float)*(p+2); }
|
|
static inline float getPixelAlpha(const __global CLQuantum *p,const unsigned int number_channels) { return (float)*(p+number_channels-1); }
|
|
|
|
static inline void setPixelRed(__global CLQuantum *p,const CLQuantum value) { *p=value; }
|
|
static inline void setPixelGreen(__global CLQuantum *p,const CLQuantum value) { *(p+1)=value; }
|
|
static inline void setPixelBlue(__global CLQuantum *p,const CLQuantum value) { *(p+2)=value; }
|
|
static inline void setPixelAlpha(__global CLQuantum *p,const unsigned int number_channels,const CLQuantum value) { *(p+number_channels-1)=value; }
|
|
|
|
static inline CLQuantum getBlue(CLPixelType p) { return p.x; }
|
|
static inline void setBlue(CLPixelType* p, CLQuantum value) { (*p).x = value; }
|
|
static inline float getBlueF4(float4 p) { return p.x; }
|
|
static inline void setBlueF4(float4* p, float value) { (*p).x = value; }
|
|
|
|
static inline CLQuantum getGreen(CLPixelType p) { return p.y; }
|
|
static inline void setGreen(CLPixelType* p, CLQuantum value) { (*p).y = value; }
|
|
static inline float getGreenF4(float4 p) { return p.y; }
|
|
static inline void setGreenF4(float4* p, float value) { (*p).y = value; }
|
|
|
|
static inline CLQuantum getRed(CLPixelType p) { return p.z; }
|
|
static inline void setRed(CLPixelType* p, CLQuantum value) { (*p).z = value; }
|
|
static inline float getRedF4(float4 p) { return p.z; }
|
|
static inline void setRedF4(float4* p, float value) { (*p).z = value; }
|
|
|
|
static inline CLQuantum getAlpha(CLPixelType p) { return p.w; }
|
|
static inline void setAlpha(CLPixelType* p, CLQuantum value) { (*p).w = value; }
|
|
static inline float getAlphaF4(float4 p) { return p.w; }
|
|
static inline void setAlphaF4(float4* p, float value) { (*p).w = value; }
|
|
|
|
static inline void ReadChannels(const __global CLQuantum *p, const unsigned int number_channels,
|
|
const ChannelType channel, float *red, float *green, float *blue, float *alpha)
|
|
{
|
|
if ((channel & RedChannel) != 0)
|
|
*red=getPixelRed(p);
|
|
|
|
if (number_channels > 2)
|
|
{
|
|
if ((channel & GreenChannel) != 0)
|
|
*green=getPixelGreen(p);
|
|
|
|
if ((channel & BlueChannel) != 0)
|
|
*blue=getPixelBlue(p);
|
|
}
|
|
|
|
if (((number_channels == 4) || (number_channels == 2)) &&
|
|
((channel & AlphaChannel) != 0))
|
|
*alpha=getPixelAlpha(p,number_channels);
|
|
}
|
|
|
|
static inline float4 ReadAllChannels(const __global CLQuantum *image, const unsigned int number_channels,
|
|
const unsigned int columns, const unsigned int x, const unsigned int y)
|
|
{
|
|
const __global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y);
|
|
|
|
float4 pixel;
|
|
|
|
pixel.x=getPixelRed(p);
|
|
|
|
if (number_channels > 2)
|
|
{
|
|
pixel.y=getPixelGreen(p);
|
|
pixel.z=getPixelBlue(p);
|
|
}
|
|
|
|
if ((number_channels == 4) || (number_channels == 2))
|
|
pixel.w=getPixelAlpha(p,number_channels);
|
|
return(pixel);
|
|
}
|
|
|
|
static inline float4 ReadFloat4(const __global CLQuantum *image, const unsigned int number_channels,
|
|
const unsigned int columns, const unsigned int x, const unsigned int y, const ChannelType channel)
|
|
{
|
|
const __global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y);
|
|
|
|
float red = 0.0f;
|
|
float green = 0.0f;
|
|
float blue = 0.0f;
|
|
float alpha = 0.0f;
|
|
|
|
ReadChannels(p, number_channels, channel, &red, &green, &blue, &alpha);
|
|
return (float4)(red, green, blue, alpha);
|
|
}
|
|
|
|
static inline void WriteChannels(__global CLQuantum *p, const unsigned int number_channels,
|
|
const ChannelType channel, float red, float green, float blue, float alpha)
|
|
{
|
|
if ((channel & RedChannel) != 0)
|
|
setPixelRed(p,ClampToQuantum(red));
|
|
|
|
if (number_channels > 2)
|
|
{
|
|
if ((channel & GreenChannel) != 0)
|
|
setPixelGreen(p,ClampToQuantum(green));
|
|
|
|
if ((channel & BlueChannel) != 0)
|
|
setPixelBlue(p,ClampToQuantum(blue));
|
|
}
|
|
|
|
if (((number_channels == 4) || (number_channels == 2)) &&
|
|
((channel & AlphaChannel) != 0))
|
|
setPixelAlpha(p,number_channels,ClampToQuantum(alpha));
|
|
}
|
|
|
|
static inline void WriteAllChannels(__global CLQuantum *image, const unsigned int number_channels,
|
|
const unsigned int columns, const unsigned int x, const unsigned int y, float4 pixel)
|
|
{
|
|
__global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y);
|
|
|
|
setPixelRed(p,ClampToQuantum(pixel.x));
|
|
|
|
if (number_channels > 2)
|
|
{
|
|
setPixelGreen(p,ClampToQuantum(pixel.y));
|
|
setPixelBlue(p,ClampToQuantum(pixel.z));
|
|
}
|
|
|
|
if ((number_channels == 4) || (number_channels == 2))
|
|
setPixelAlpha(p,number_channels,ClampToQuantum(pixel.w));
|
|
}
|
|
|
|
static inline void WriteFloat4(__global CLQuantum *image, const unsigned int number_channels,
|
|
const unsigned int columns, const unsigned int x, const unsigned int y, const ChannelType channel,
|
|
float4 pixel)
|
|
{
|
|
__global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y);
|
|
WriteChannels(p, number_channels, channel, pixel.x, pixel.y, pixel.z, pixel.w);
|
|
}
|
|
|
|
static inline float GetPixelIntensity(const unsigned int colorspace,
|
|
const unsigned int method,float red,float green,float blue)
|
|
{
|
|
float intensity;
|
|
|
|
if ((colorspace == GRAYColorspace) || (colorspace == LinearGRAYColorspace))
|
|
return red;
|
|
|
|
switch (method)
|
|
{
|
|
case AveragePixelIntensityMethod:
|
|
{
|
|
intensity=(red+green+blue)/3.0;
|
|
break;
|
|
}
|
|
case BrightnessPixelIntensityMethod:
|
|
{
|
|
intensity=MagickMax(MagickMax(red,green),blue);
|
|
break;
|
|
}
|
|
case LightnessPixelIntensityMethod:
|
|
{
|
|
intensity=(MagickMin(MagickMin(red,green),blue)+
|
|
MagickMax(MagickMax(red,green),blue))/2.0;
|
|
break;
|
|
}
|
|
case MSPixelIntensityMethod:
|
|
{
|
|
intensity=(float) (((float) red*red+green*green+blue*blue)/
|
|
(3.0*QuantumRange));
|
|
break;
|
|
}
|
|
case Rec601LumaPixelIntensityMethod:
|
|
{
|
|
/*
|
|
if (image->colorspace == RGBColorspace)
|
|
{
|
|
red=EncodePixelGamma(red);
|
|
green=EncodePixelGamma(green);
|
|
blue=EncodePixelGamma(blue);
|
|
}
|
|
*/
|
|
intensity=0.298839*red+0.586811*green+0.114350*blue;
|
|
break;
|
|
}
|
|
case Rec601LuminancePixelIntensityMethod:
|
|
{
|
|
/*
|
|
if (image->colorspace == sRGBColorspace)
|
|
{
|
|
red=DecodePixelGamma(red);
|
|
green=DecodePixelGamma(green);
|
|
blue=DecodePixelGamma(blue);
|
|
}
|
|
*/
|
|
intensity=0.298839*red+0.586811*green+0.114350*blue;
|
|
break;
|
|
}
|
|
case Rec709LumaPixelIntensityMethod:
|
|
default:
|
|
{
|
|
/*
|
|
if (image->colorspace == RGBColorspace)
|
|
{
|
|
red=EncodePixelGamma(red);
|
|
green=EncodePixelGamma(green);
|
|
blue=EncodePixelGamma(blue);
|
|
}
|
|
*/
|
|
intensity=0.212656*red+0.715158*green+0.072186*blue;
|
|
break;
|
|
}
|
|
case Rec709LuminancePixelIntensityMethod:
|
|
{
|
|
/*
|
|
if (image->colorspace == sRGBColorspace)
|
|
{
|
|
red=DecodePixelGamma(red);
|
|
green=DecodePixelGamma(green);
|
|
blue=DecodePixelGamma(blue);
|
|
}
|
|
*/
|
|
intensity=0.212656*red+0.715158*green+0.072186*blue;
|
|
break;
|
|
}
|
|
case RMSPixelIntensityMethod:
|
|
{
|
|
intensity=(float) (sqrt((float) red*red+green*green+blue*blue)/
|
|
sqrt(3.0));
|
|
break;
|
|
}
|
|
}
|
|
|
|
return intensity;
|
|
}
|
|
|
|
static inline int mirrorBottom(int value)
|
|
{
|
|
return (value < 0) ? - (value) : value;
|
|
}
|
|
|
|
static inline int mirrorTop(int value, int width)
|
|
{
|
|
return (value >= width) ? (2 * width - value - 1) : value;
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% A d d N o i s e %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
/*
|
|
Part of MWC64X by David Thomas, dt10@imperial.ac.uk
|
|
This is provided under BSD, full license is with the main package.
|
|
See http://www.doc.ic.ac.uk/~dt10/research
|
|
*/
|
|
|
|
// Pre: a<M, b<M
|
|
// Post: r=(a+b) mod M
|
|
ulong MWC_AddMod64(ulong a, ulong b, ulong M)
|
|
{
|
|
ulong v=a+b;
|
|
//if( (v>=M) || (v<a) )
|
|
if( (v>=M) || (convert_float(v) < convert_float(a)) ) // workaround for what appears to be an optimizer bug.
|
|
v=v-M;
|
|
return v;
|
|
}
|
|
|
|
// Pre: a<M,b<M
|
|
// Post: r=(a*b) mod M
|
|
// This could be done more efficently, but it is portable, and should
|
|
// be easy to understand. It can be replaced with any of the better
|
|
// modular multiplication algorithms (for example if you know you have
|
|
// double precision available or something).
|
|
ulong MWC_MulMod64(ulong a, ulong b, ulong M)
|
|
{
|
|
ulong r=0;
|
|
while(a!=0){
|
|
if(a&1)
|
|
r=MWC_AddMod64(r,b,M);
|
|
b=MWC_AddMod64(b,b,M);
|
|
a=a>>1;
|
|
}
|
|
return r;
|
|
}
|
|
|
|
// Pre: a<M, e>=0
|
|
// Post: r=(a^b) mod M
|
|
// This takes at most ~64^2 modular additions, so probably about 2^15 or so instructions on
|
|
// most architectures
|
|
ulong MWC_PowMod64(ulong a, ulong e, ulong M)
|
|
{
|
|
ulong sqr=a, acc=1;
|
|
while(e!=0){
|
|
if(e&1)
|
|
acc=MWC_MulMod64(acc,sqr,M);
|
|
sqr=MWC_MulMod64(sqr,sqr,M);
|
|
e=e>>1;
|
|
}
|
|
return acc;
|
|
}
|
|
|
|
uint2 MWC_SkipImpl_Mod64(uint2 curr, ulong A, ulong M, ulong distance)
|
|
{
|
|
ulong m=MWC_PowMod64(A, distance, M);
|
|
ulong x=curr.x*(ulong)A+curr.y;
|
|
x=MWC_MulMod64(x, m, M);
|
|
return (uint2)((uint)(x/A), (uint)(x%A));
|
|
}
|
|
|
|
uint2 MWC_SeedImpl_Mod64(ulong A, ulong M, uint vecSize, uint vecOffset, ulong streamBase, ulong streamGap)
|
|
{
|
|
// This is an arbitrary constant for starting LCG jumping from. I didn't
|
|
// want to start from 1, as then you end up with the two or three first values
|
|
// being a bit poor in ones - once you've decided that, one constant is as
|
|
// good as any another. There is no deep mathematical reason for it, I just
|
|
// generated a random number.
|
|
enum{ MWC_BASEID = 4077358422479273989UL };
|
|
|
|
ulong dist=streamBase + (get_global_id(0)*vecSize+vecOffset)*streamGap;
|
|
ulong m=MWC_PowMod64(A, dist, M);
|
|
|
|
ulong x=MWC_MulMod64(MWC_BASEID, m, M);
|
|
return (uint2)((uint)(x/A), (uint)(x%A));
|
|
}
|
|
|
|
//! Represents the state of a particular generator
|
|
typedef struct{ uint x; uint c; uint seed0; ulong seed1; } mwc64x_state_t;
|
|
|
|
void MWC64X_Step(mwc64x_state_t *s)
|
|
{
|
|
uint X=s->x, C=s->c;
|
|
|
|
uint Xn=s->seed0*X+C;
|
|
uint carry=(uint)(Xn<C); // The (Xn<C) will be zero or one for scalar
|
|
uint Cn=mad_hi(s->seed0,X,carry);
|
|
|
|
s->x=Xn;
|
|
s->c=Cn;
|
|
}
|
|
|
|
void MWC64X_Skip(mwc64x_state_t *s, ulong distance)
|
|
{
|
|
uint2 tmp=MWC_SkipImpl_Mod64((uint2)(s->x,s->c), s->seed0, s->seed1, distance);
|
|
s->x=tmp.x;
|
|
s->c=tmp.y;
|
|
}
|
|
|
|
void MWC64X_SeedStreams(mwc64x_state_t *s, ulong baseOffset, ulong perStreamOffset)
|
|
{
|
|
uint2 tmp=MWC_SeedImpl_Mod64(s->seed0, s->seed1, 1, 0, baseOffset, perStreamOffset);
|
|
s->x=tmp.x;
|
|
s->c=tmp.y;
|
|
}
|
|
|
|
//! Return a 32-bit integer in the range [0..2^32)
|
|
uint MWC64X_NextUint(mwc64x_state_t *s)
|
|
{
|
|
uint res=s->x ^ s->c;
|
|
MWC64X_Step(s);
|
|
return res;
|
|
}
|
|
|
|
//
|
|
// End of MWC64X excerpt
|
|
//
|
|
|
|
float mwcReadPseudoRandomValue(mwc64x_state_t* rng)
|
|
{
|
|
return (1.0f * MWC64X_NextUint(rng)) / (float)(0xffffffff); // normalized to 1.0
|
|
}
|
|
|
|
float mwcGenerateDifferentialNoise(mwc64x_state_t* r, float pixel, NoiseType noise_type, float attenuate)
|
|
{
|
|
float
|
|
alpha,
|
|
beta,
|
|
noise,
|
|
sigma;
|
|
|
|
noise = 0.0f;
|
|
alpha=mwcReadPseudoRandomValue(r);
|
|
switch(noise_type)
|
|
{
|
|
case UniformNoise:
|
|
default:
|
|
{
|
|
noise=(pixel+QuantumRange*SigmaUniform*(alpha-0.5f));
|
|
break;
|
|
}
|
|
case GaussianNoise:
|
|
{
|
|
float
|
|
gamma,
|
|
tau;
|
|
|
|
if (alpha == 0.0f)
|
|
alpha=1.0f;
|
|
beta=mwcReadPseudoRandomValue(r);
|
|
gamma=sqrt(-2.0f*log(alpha));
|
|
sigma=gamma*cospi((2.0f*beta));
|
|
tau=gamma*sinpi((2.0f*beta));
|
|
noise=pixel+sqrt(pixel)*SigmaGaussian*sigma+QuantumRange*TauGaussian*tau;
|
|
break;
|
|
}
|
|
case ImpulseNoise:
|
|
{
|
|
if (alpha < (SigmaImpulse/2.0f))
|
|
noise=0.0f;
|
|
else
|
|
if (alpha >= (1.0f-(SigmaImpulse/2.0f)))
|
|
noise=QuantumRange;
|
|
else
|
|
noise=pixel;
|
|
break;
|
|
}
|
|
case LaplacianNoise:
|
|
{
|
|
if (alpha <= 0.5f)
|
|
{
|
|
if (alpha <= MagickEpsilon)
|
|
noise=(pixel-QuantumRange);
|
|
else
|
|
noise=(pixel+QuantumRange*SigmaLaplacian*log(2.0f*alpha)+
|
|
0.5f);
|
|
break;
|
|
}
|
|
beta=1.0f-alpha;
|
|
if (beta <= (0.5f*MagickEpsilon))
|
|
noise=(pixel+QuantumRange);
|
|
else
|
|
noise=(pixel-QuantumRange*SigmaLaplacian*log(2.0f*beta)+0.5f);
|
|
break;
|
|
}
|
|
case MultiplicativeGaussianNoise:
|
|
{
|
|
sigma=1.0f;
|
|
if (alpha > MagickEpsilon)
|
|
sigma=sqrt(-2.0f*log(alpha));
|
|
beta=mwcReadPseudoRandomValue(r);
|
|
noise=(pixel+pixel*SigmaMultiplicativeGaussian*sigma*
|
|
cospi((2.0f*beta))/2.0f);
|
|
break;
|
|
}
|
|
case PoissonNoise:
|
|
{
|
|
float
|
|
poisson;
|
|
unsigned int i;
|
|
poisson=exp(-SigmaPoisson*QuantumScale*pixel);
|
|
for (i=0; alpha > poisson; i++)
|
|
{
|
|
beta=mwcReadPseudoRandomValue(r);
|
|
alpha*=beta;
|
|
}
|
|
noise=(QuantumRange*i*PerceptibleReciprocal(SigmaPoisson));
|
|
break;
|
|
}
|
|
case RandomNoise:
|
|
{
|
|
noise=(QuantumRange*SigmaRandom*alpha);
|
|
break;
|
|
}
|
|
}
|
|
return noise;
|
|
}
|
|
|
|
__kernel
|
|
void AddNoise(const __global CLQuantum *image,
|
|
const unsigned int number_channels,const ChannelType channel,
|
|
const unsigned int length,const unsigned int pixelsPerWorkItem,
|
|
const NoiseType noise_type,const float attenuate,const unsigned int seed0,
|
|
const unsigned int seed1,const unsigned int numRandomNumbersPerPixel,
|
|
__global CLQuantum *filteredImage)
|
|
{
|
|
mwc64x_state_t rng;
|
|
rng.seed0 = seed0;
|
|
rng.seed1 = seed1;
|
|
|
|
uint span = pixelsPerWorkItem * numRandomNumbersPerPixel; // length of RNG substream each workitem will use
|
|
uint offset = span * get_local_size(0) * get_group_id(0); // offset of this workgroup's RNG substream (in master stream);
|
|
MWC64X_SeedStreams(&rng, offset, span); // Seed the RNG streams
|
|
|
|
uint pos = get_group_id(0) * get_local_size(0) * pixelsPerWorkItem * number_channels + (get_local_id(0) * number_channels);
|
|
uint count = pixelsPerWorkItem;
|
|
|
|
while (count > 0 && pos < length)
|
|
{
|
|
const __global CLQuantum *p = image + pos;
|
|
__global CLQuantum *q = filteredImage + pos;
|
|
|
|
float red;
|
|
float green;
|
|
float blue;
|
|
float alpha;
|
|
|
|
ReadChannels(p, number_channels, channel, &red, &green, &blue, &alpha);
|
|
|
|
if ((channel & RedChannel) != 0)
|
|
red=mwcGenerateDifferentialNoise(&rng,red,noise_type,attenuate);
|
|
|
|
if (number_channels > 2)
|
|
{
|
|
if ((channel & GreenChannel) != 0)
|
|
green=mwcGenerateDifferentialNoise(&rng,green,noise_type,attenuate);
|
|
|
|
if ((channel & BlueChannel) != 0)
|
|
blue=mwcGenerateDifferentialNoise(&rng,blue,noise_type,attenuate);
|
|
}
|
|
|
|
if (((number_channels == 4) || (number_channels == 2)) &&
|
|
((channel & AlphaChannel) != 0))
|
|
alpha=mwcGenerateDifferentialNoise(&rng,alpha,noise_type,attenuate);
|
|
|
|
WriteChannels(q, number_channels, channel, red, green, blue, alpha);
|
|
|
|
pos += (get_local_size(0) * number_channels);
|
|
count--;
|
|
}
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% B l u r %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
/*
|
|
Reduce image noise and reduce detail levels by row
|
|
*/
|
|
__kernel void BlurRow(const __global CLQuantum *image,
|
|
const unsigned int number_channels,const ChannelType channel,
|
|
__constant float *filter,const unsigned int width,
|
|
const unsigned int imageColumns,const unsigned int imageRows,
|
|
__local float4 *temp,__global float4 *tempImage)
|
|
{
|
|
const int x = get_global_id(0);
|
|
const int y = get_global_id(1);
|
|
|
|
const int columns = imageColumns;
|
|
|
|
const unsigned int radius = (width-1)/2;
|
|
const int wsize = get_local_size(0);
|
|
const unsigned int loadSize = wsize+width;
|
|
|
|
//group coordinate
|
|
const int groupX=get_local_size(0)*get_group_id(0);
|
|
|
|
//parallel load and clamp
|
|
for (int i=get_local_id(0); i < loadSize; i=i+get_local_size(0))
|
|
{
|
|
int cx = ClampToCanvas(i + groupX - radius, columns);
|
|
temp[i] = ReadFloat4(image, number_channels, columns, cx, y, channel);
|
|
}
|
|
|
|
// barrier
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
// only do the work if this is not a patched item
|
|
if (get_global_id(0) < columns)
|
|
{
|
|
// compute
|
|
float4 result = (float4) 0;
|
|
|
|
int i = 0;
|
|
|
|
for ( ; i+7 < width; )
|
|
{
|
|
for (int j=0; j < 8; j++)
|
|
result+=filter[i+j]*temp[i+j+get_local_id(0)];
|
|
i+=8;
|
|
}
|
|
|
|
for ( ; i < width; i++)
|
|
result+=filter[i]*temp[i+get_local_id(0)];
|
|
|
|
// write back to global
|
|
tempImage[y*columns+x] = result;
|
|
}
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
/*
|
|
Reduce image noise and reduce detail levels by line
|
|
*/
|
|
__kernel void BlurColumn(const __global float4 *blurRowData,
|
|
const unsigned int number_channels,const ChannelType channel,
|
|
__constant float *filter,const unsigned int width,
|
|
const unsigned int imageColumns,const unsigned int imageRows,
|
|
__local float4 *temp,__global CLQuantum *filteredImage)
|
|
{
|
|
const int x = get_global_id(0);
|
|
const int y = get_global_id(1);
|
|
|
|
const int columns = imageColumns;
|
|
const int rows = imageRows;
|
|
|
|
unsigned int radius = (width-1)/2;
|
|
const int wsize = get_local_size(1);
|
|
const unsigned int loadSize = wsize+width;
|
|
|
|
//group coordinate
|
|
const int groupX=get_local_size(0)*get_group_id(0);
|
|
const int groupY=get_local_size(1)*get_group_id(1);
|
|
//notice that get_local_size(0) is 1, so
|
|
//groupX=get_group_id(0);
|
|
|
|
//parallel load and clamp
|
|
for (int i = get_local_id(1); i < loadSize; i=i+get_local_size(1))
|
|
temp[i] = blurRowData[ClampToCanvas(i+groupY-radius, rows) * columns + groupX];
|
|
|
|
// barrier
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
// only do the work if this is not a patched item
|
|
if (get_global_id(1) < rows)
|
|
{
|
|
// compute
|
|
float4 result = (float4) 0;
|
|
|
|
int i = 0;
|
|
|
|
for ( ; i+7 < width; )
|
|
{
|
|
for (int j=0; j < 8; j++)
|
|
result+=filter[i+j]*temp[i+j+get_local_id(1)];
|
|
i+=8;
|
|
}
|
|
|
|
for ( ; i < width; i++)
|
|
result+=filter[i]*temp[i+get_local_id(1)];
|
|
|
|
// write back to global
|
|
WriteFloat4(filteredImage, number_channels, columns, x, y, channel, result);
|
|
}
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% C o n t r a s t %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
|
|
static inline float4 ConvertRGBToHSB(const float4 pixel)
|
|
{
|
|
float4 result=0.0f;
|
|
result.w=pixel.w;
|
|
float tmax=MagickMax(MagickMax(pixel.x,pixel.y),pixel.z);
|
|
if (tmax != 0.0f)
|
|
{
|
|
float tmin=MagickMin(MagickMin(pixel.x,pixel.y),pixel.z);
|
|
float delta=tmax-tmin;
|
|
|
|
result.y=delta/tmax;
|
|
result.z=QuantumScale*tmax;
|
|
if (delta != 0.0f)
|
|
{
|
|
result.x =((pixel.x == tmax) ? 0.0f : ((pixel.y == tmax) ? 2.0f : 4.0f));
|
|
result.x+=((pixel.x == tmax) ? (pixel.y-pixel.z) : ((pixel.y == tmax) ?
|
|
(pixel.z-pixel.x) : (pixel.x-pixel.y)))/delta;
|
|
result.x/=6.0f;
|
|
result.x+=(result.x < 0.0f) ? 0.0f : 1.0f;
|
|
}
|
|
}
|
|
return(result);
|
|
}
|
|
|
|
static inline float4 ConvertHSBToRGB(const float4 pixel)
|
|
{
|
|
float hue=pixel.x;
|
|
float saturation=pixel.y;
|
|
float brightness=pixel.z;
|
|
|
|
float4 result=pixel;
|
|
|
|
if (saturation == 0.0f)
|
|
{
|
|
result.x=result.y=result.z=ClampToQuantum(QuantumRange*brightness);
|
|
}
|
|
else
|
|
{
|
|
float h=6.0f*(hue-floor(hue));
|
|
float f=h-floor(h);
|
|
float p=brightness*(1.0f-saturation);
|
|
float q=brightness*(1.0f-saturation*f);
|
|
float t=brightness*(1.0f-(saturation*(1.0f-f)));
|
|
int ih = (int)h;
|
|
|
|
if (ih == 1)
|
|
{
|
|
result.x=ClampToQuantum(QuantumRange*q);
|
|
result.y=ClampToQuantum(QuantumRange*brightness);
|
|
result.z=ClampToQuantum(QuantumRange*p);
|
|
}
|
|
else if (ih == 2)
|
|
{
|
|
result.x=ClampToQuantum(QuantumRange*p);
|
|
result.y=ClampToQuantum(QuantumRange*brightness);
|
|
result.z=ClampToQuantum(QuantumRange*t);
|
|
}
|
|
else if (ih == 3)
|
|
{
|
|
result.x=ClampToQuantum(QuantumRange*p);
|
|
result.y=ClampToQuantum(QuantumRange*q);
|
|
result.z=ClampToQuantum(QuantumRange*brightness);
|
|
}
|
|
else if (ih == 4)
|
|
{
|
|
result.x=ClampToQuantum(QuantumRange*t);
|
|
result.y=ClampToQuantum(QuantumRange*p);
|
|
result.z=ClampToQuantum(QuantumRange*brightness);
|
|
}
|
|
else if (ih == 5)
|
|
{
|
|
result.x=ClampToQuantum(QuantumRange*brightness);
|
|
result.y=ClampToQuantum(QuantumRange*p);
|
|
result.z=ClampToQuantum(QuantumRange*q);
|
|
}
|
|
else
|
|
{
|
|
result.x=ClampToQuantum(QuantumRange*brightness);
|
|
result.y=ClampToQuantum(QuantumRange*t);
|
|
result.z=ClampToQuantum(QuantumRange*p);
|
|
}
|
|
}
|
|
return(result);
|
|
}
|
|
|
|
__kernel void Contrast(__global CLQuantum *image,
|
|
const unsigned int number_channels,const int sign)
|
|
{
|
|
const int x=get_global_id(0);
|
|
const int y=get_global_id(1);
|
|
const unsigned int columns=get_global_size(0);
|
|
|
|
float4 pixel=ReadAllChannels(image,number_channels,columns,x,y);
|
|
if (number_channels < 3)
|
|
pixel.y=pixel.z=pixel.x;
|
|
|
|
pixel=ConvertRGBToHSB(pixel);
|
|
float brightness=pixel.z;
|
|
brightness+=0.5f*sign*(0.5f*(sinpi(brightness-0.5f)+1.0f)-brightness);
|
|
brightness=clamp(brightness,0.0f,1.0f);
|
|
pixel.z=brightness;
|
|
pixel=ConvertHSBToRGB(pixel);
|
|
|
|
WriteAllChannels(image,number_channels,columns,x,y,pixel);
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% C o n t r a s t S t r e t c h %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
/*
|
|
*/
|
|
__kernel void Histogram(__global CLPixelType * restrict im,
|
|
const ChannelType channel,
|
|
const unsigned int colorspace,
|
|
const unsigned int method,
|
|
__global uint4 * restrict histogram)
|
|
{
|
|
const int x = get_global_id(0);
|
|
const int y = get_global_id(1);
|
|
const int columns = get_global_size(0);
|
|
const int c = x + y * columns;
|
|
if ((channel & SyncChannels) != 0)
|
|
{
|
|
float red=(float)getRed(im[c]);
|
|
float green=(float)getGreen(im[c]);
|
|
float blue=(float)getBlue(im[c]);
|
|
|
|
float intensity = GetPixelIntensity(colorspace, method, red, green, blue);
|
|
uint pos = ScaleQuantumToMap(ClampToQuantum(intensity));
|
|
atomic_inc((__global uint *)(&(histogram[pos]))+2); //red position
|
|
}
|
|
else
|
|
{
|
|
// for equalizing, we always need all channels?
|
|
// otherwise something more
|
|
}
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
/*
|
|
*/
|
|
__kernel void ContrastStretch(__global CLPixelType * restrict im,
|
|
const ChannelType channel,
|
|
__global CLPixelType * restrict stretch_map,
|
|
const float4 white, const float4 black)
|
|
{
|
|
const int x = get_global_id(0);
|
|
const int y = get_global_id(1);
|
|
const int columns = get_global_size(0);
|
|
const int c = x + y * columns;
|
|
|
|
uint ePos;
|
|
CLPixelType oValue, eValue;
|
|
CLQuantum red, green, blue, alpha;
|
|
|
|
//read from global
|
|
oValue=im[c];
|
|
|
|
if ((channel & RedChannel) != 0)
|
|
{
|
|
if (getRedF4(white) != getRedF4(black))
|
|
{
|
|
ePos = ScaleQuantumToMap(getRed(oValue));
|
|
eValue = stretch_map[ePos];
|
|
red = getRed(eValue);
|
|
}
|
|
}
|
|
|
|
if ((channel & GreenChannel) != 0)
|
|
{
|
|
if (getGreenF4(white) != getGreenF4(black))
|
|
{
|
|
ePos = ScaleQuantumToMap(getGreen(oValue));
|
|
eValue = stretch_map[ePos];
|
|
green = getGreen(eValue);
|
|
}
|
|
}
|
|
|
|
if ((channel & BlueChannel) != 0)
|
|
{
|
|
if (getBlueF4(white) != getBlueF4(black))
|
|
{
|
|
ePos = ScaleQuantumToMap(getBlue(oValue));
|
|
eValue = stretch_map[ePos];
|
|
blue = getBlue(eValue);
|
|
}
|
|
}
|
|
|
|
if ((channel & AlphaChannel) != 0)
|
|
{
|
|
if (getAlphaF4(white) != getAlphaF4(black))
|
|
{
|
|
ePos = ScaleQuantumToMap(getAlpha(oValue));
|
|
eValue = stretch_map[ePos];
|
|
alpha = getAlpha(eValue);
|
|
}
|
|
}
|
|
|
|
//write back
|
|
im[c]=(CLPixelType)(blue, green, red, alpha);
|
|
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% C o n v o l v e %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
__kernel
|
|
void ConvolveOptimized(const __global CLPixelType *input, __global CLPixelType *output,
|
|
const unsigned int imageWidth, const unsigned int imageHeight,
|
|
__constant float *filter, const unsigned int filterWidth, const unsigned int filterHeight,
|
|
const uint matte, const ChannelType channel, __local CLPixelType *pixelLocalCache, __local float* filterCache) {
|
|
|
|
int2 blockID;
|
|
blockID.x = get_global_id(0) / get_local_size(0);
|
|
blockID.y = get_global_id(1) / get_local_size(1);
|
|
|
|
// image area processed by this workgroup
|
|
int2 imageAreaOrg;
|
|
imageAreaOrg.x = blockID.x * get_local_size(0);
|
|
imageAreaOrg.y = blockID.y * get_local_size(1);
|
|
|
|
int2 midFilterDimen;
|
|
midFilterDimen.x = (filterWidth-1)/2;
|
|
midFilterDimen.y = (filterHeight-1)/2;
|
|
|
|
int2 cachedAreaOrg = imageAreaOrg - midFilterDimen;
|
|
|
|
// dimension of the local cache
|
|
int2 cachedAreaDimen;
|
|
cachedAreaDimen.x = get_local_size(0) + filterWidth - 1;
|
|
cachedAreaDimen.y = get_local_size(1) + filterHeight - 1;
|
|
|
|
// cache the pixels accessed by this workgroup in local memory
|
|
int localID = get_local_id(1)*get_local_size(0)+get_local_id(0);
|
|
int cachedAreaNumPixels = cachedAreaDimen.x * cachedAreaDimen.y;
|
|
int groupSize = get_local_size(0) * get_local_size(1);
|
|
for (int i = localID; i < cachedAreaNumPixels; i+=groupSize) {
|
|
|
|
int2 cachedAreaIndex;
|
|
cachedAreaIndex.x = i % cachedAreaDimen.x;
|
|
cachedAreaIndex.y = i / cachedAreaDimen.x;
|
|
|
|
int2 imagePixelIndex;
|
|
imagePixelIndex = cachedAreaOrg + cachedAreaIndex;
|
|
|
|
// only support EdgeVirtualPixelMethod through ClampToCanvas
|
|
// TODO: implement other virtual pixel method
|
|
imagePixelIndex.x = ClampToCanvas(imagePixelIndex.x, imageWidth);
|
|
imagePixelIndex.y = ClampToCanvas(imagePixelIndex.y, imageHeight);
|
|
|
|
pixelLocalCache[i] = input[imagePixelIndex.y * imageWidth + imagePixelIndex.x];
|
|
}
|
|
|
|
// cache the filter
|
|
for (int i = localID; i < filterHeight*filterWidth; i+=groupSize) {
|
|
filterCache[i] = filter[i];
|
|
}
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
|
|
int2 imageIndex;
|
|
imageIndex.x = imageAreaOrg.x + get_local_id(0);
|
|
imageIndex.y = imageAreaOrg.y + get_local_id(1);
|
|
|
|
// if out-of-range, stops here and quit
|
|
if (imageIndex.x >= imageWidth
|
|
|| imageIndex.y >= imageHeight) {
|
|
return;
|
|
}
|
|
|
|
int filterIndex = 0;
|
|
float4 sum = (float4)0.0f;
|
|
float gamma = 0.0f;
|
|
if (((channel & AlphaChannel) == 0) || (matte == 0)) {
|
|
int cacheIndexY = get_local_id(1);
|
|
for (int j = 0; j < filterHeight; j++) {
|
|
int cacheIndexX = get_local_id(0);
|
|
for (int i = 0; i < filterWidth; i++) {
|
|
CLPixelType p = pixelLocalCache[cacheIndexY*cachedAreaDimen.x + cacheIndexX];
|
|
float f = filterCache[filterIndex];
|
|
|
|
sum.x += f * p.x;
|
|
sum.y += f * p.y;
|
|
sum.z += f * p.z;
|
|
sum.w += f * p.w;
|
|
|
|
gamma += f;
|
|
filterIndex++;
|
|
cacheIndexX++;
|
|
}
|
|
cacheIndexY++;
|
|
}
|
|
}
|
|
else {
|
|
int cacheIndexY = get_local_id(1);
|
|
for (int j = 0; j < filterHeight; j++) {
|
|
int cacheIndexX = get_local_id(0);
|
|
for (int i = 0; i < filterWidth; i++) {
|
|
|
|
CLPixelType p = pixelLocalCache[cacheIndexY*cachedAreaDimen.x + cacheIndexX];
|
|
float alpha = QuantumScale*p.w;
|
|
float f = filterCache[filterIndex];
|
|
float g = alpha * f;
|
|
|
|
sum.x += g*p.x;
|
|
sum.y += g*p.y;
|
|
sum.z += g*p.z;
|
|
sum.w += f*p.w;
|
|
|
|
gamma += g;
|
|
filterIndex++;
|
|
cacheIndexX++;
|
|
}
|
|
cacheIndexY++;
|
|
}
|
|
gamma = PerceptibleReciprocal(gamma);
|
|
sum.xyz = gamma*sum.xyz;
|
|
}
|
|
CLPixelType outputPixel;
|
|
outputPixel.x = ClampToQuantum(sum.x);
|
|
outputPixel.y = ClampToQuantum(sum.y);
|
|
outputPixel.z = ClampToQuantum(sum.z);
|
|
outputPixel.w = ((channel & AlphaChannel)!=0)?ClampToQuantum(sum.w):input[imageIndex.y * imageWidth + imageIndex.x].w;
|
|
|
|
output[imageIndex.y * imageWidth + imageIndex.x] = outputPixel;
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
__kernel
|
|
void Convolve(const __global CLPixelType *input, __global CLPixelType *output,
|
|
const uint imageWidth, const uint imageHeight,
|
|
__constant float *filter, const unsigned int filterWidth, const unsigned int filterHeight,
|
|
const uint matte, const ChannelType channel) {
|
|
|
|
int2 imageIndex;
|
|
imageIndex.x = get_global_id(0);
|
|
imageIndex.y = get_global_id(1);
|
|
|
|
/*
|
|
unsigned int imageWidth = get_global_size(0);
|
|
unsigned int imageHeight = get_global_size(1);
|
|
*/
|
|
if (imageIndex.x >= imageWidth
|
|
|| imageIndex.y >= imageHeight)
|
|
return;
|
|
|
|
int2 midFilterDimen;
|
|
midFilterDimen.x = (filterWidth-1)/2;
|
|
midFilterDimen.y = (filterHeight-1)/2;
|
|
|
|
int filterIndex = 0;
|
|
float4 sum = (float4)0.0f;
|
|
float gamma = 0.0f;
|
|
if (((channel & AlphaChannel) == 0) || (matte == 0)) {
|
|
for (int j = 0; j < filterHeight; j++) {
|
|
int2 inputPixelIndex;
|
|
inputPixelIndex.y = imageIndex.y - midFilterDimen.y + j;
|
|
inputPixelIndex.y = ClampToCanvas(inputPixelIndex.y, imageHeight);
|
|
for (int i = 0; i < filterWidth; i++) {
|
|
inputPixelIndex.x = imageIndex.x - midFilterDimen.x + i;
|
|
inputPixelIndex.x = ClampToCanvas(inputPixelIndex.x, imageWidth);
|
|
|
|
CLPixelType p = input[inputPixelIndex.y * imageWidth + inputPixelIndex.x];
|
|
float f = filter[filterIndex];
|
|
|
|
sum.x += f * p.x;
|
|
sum.y += f * p.y;
|
|
sum.z += f * p.z;
|
|
sum.w += f * p.w;
|
|
|
|
gamma += f;
|
|
|
|
filterIndex++;
|
|
}
|
|
}
|
|
}
|
|
else {
|
|
|
|
for (int j = 0; j < filterHeight; j++) {
|
|
int2 inputPixelIndex;
|
|
inputPixelIndex.y = imageIndex.y - midFilterDimen.y + j;
|
|
inputPixelIndex.y = ClampToCanvas(inputPixelIndex.y, imageHeight);
|
|
for (int i = 0; i < filterWidth; i++) {
|
|
inputPixelIndex.x = imageIndex.x - midFilterDimen.x + i;
|
|
inputPixelIndex.x = ClampToCanvas(inputPixelIndex.x, imageWidth);
|
|
|
|
CLPixelType p = input[inputPixelIndex.y * imageWidth + inputPixelIndex.x];
|
|
float alpha = QuantumScale*p.w;
|
|
float f = filter[filterIndex];
|
|
float g = alpha * f;
|
|
|
|
sum.x += g*p.x;
|
|
sum.y += g*p.y;
|
|
sum.z += g*p.z;
|
|
sum.w += f*p.w;
|
|
|
|
gamma += g;
|
|
|
|
|
|
filterIndex++;
|
|
}
|
|
}
|
|
gamma = PerceptibleReciprocal(gamma);
|
|
sum.xyz = gamma*sum.xyz;
|
|
}
|
|
|
|
CLPixelType outputPixel;
|
|
outputPixel.x = ClampToQuantum(sum.x);
|
|
outputPixel.y = ClampToQuantum(sum.y);
|
|
outputPixel.z = ClampToQuantum(sum.z);
|
|
outputPixel.w = ((channel & AlphaChannel)!=0)?ClampToQuantum(sum.w):input[imageIndex.y * imageWidth + imageIndex.x].w;
|
|
|
|
output[imageIndex.y * imageWidth + imageIndex.x] = outputPixel;
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% D e s p e c k l e %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
|
|
__kernel void HullPass1(const __global CLPixelType *inputImage, __global CLPixelType *outputImage
|
|
, const unsigned int imageWidth, const unsigned int imageHeight
|
|
, const int2 offset, const int polarity, const int matte) {
|
|
|
|
int x = get_global_id(0);
|
|
int y = get_global_id(1);
|
|
|
|
CLPixelType v = inputImage[y*imageWidth+x];
|
|
|
|
int2 neighbor;
|
|
neighbor.y = y + offset.y;
|
|
neighbor.x = x + offset.x;
|
|
|
|
int2 clampedNeighbor;
|
|
clampedNeighbor.x = ClampToCanvas(neighbor.x, imageWidth);
|
|
clampedNeighbor.y = ClampToCanvas(neighbor.y, imageHeight);
|
|
|
|
CLPixelType r = (clampedNeighbor.x == neighbor.x
|
|
&& clampedNeighbor.y == neighbor.y)?inputImage[clampedNeighbor.y*imageWidth+clampedNeighbor.x]
|
|
:(CLPixelType)0;
|
|
|
|
int sv[4];
|
|
sv[0] = (int)v.x;
|
|
sv[1] = (int)v.y;
|
|
sv[2] = (int)v.z;
|
|
sv[3] = (int)v.w;
|
|
|
|
int sr[4];
|
|
sr[0] = (int)r.x;
|
|
sr[1] = (int)r.y;
|
|
sr[2] = (int)r.z;
|
|
sr[3] = (int)r.w;
|
|
|
|
if (polarity > 0) {
|
|
\n #pragma unroll 4\n
|
|
for (unsigned int i = 0; i < 4; i++) {
|
|
sv[i] = (sr[i] >= (sv[i]+ScaleCharToQuantum(2)))?(sv[i]+ScaleCharToQuantum(1)):sv[i];
|
|
}
|
|
}
|
|
else {
|
|
\n #pragma unroll 4\n
|
|
for (unsigned int i = 0; i < 4; i++) {
|
|
sv[i] = (sr[i] <= (sv[i]-ScaleCharToQuantum(2)))?(sv[i]-ScaleCharToQuantum(1)):sv[i];
|
|
}
|
|
|
|
}
|
|
|
|
v.x = (CLQuantum)sv[0];
|
|
v.y = (CLQuantum)sv[1];
|
|
v.z = (CLQuantum)sv[2];
|
|
|
|
if (matte!=0)
|
|
v.w = (CLQuantum)sv[3];
|
|
|
|
outputImage[y*imageWidth+x] = v;
|
|
|
|
}
|
|
|
|
|
|
)
|
|
|
|
STRINGIFY(
|
|
|
|
__kernel void HullPass2(const __global CLPixelType *inputImage, __global CLPixelType *outputImage
|
|
, const unsigned int imageWidth, const unsigned int imageHeight
|
|
, const int2 offset, const int polarity, const int matte) {
|
|
|
|
int x = get_global_id(0);
|
|
int y = get_global_id(1);
|
|
|
|
CLPixelType v = inputImage[y*imageWidth+x];
|
|
|
|
int2 neighbor, clampedNeighbor;
|
|
|
|
neighbor.y = y + offset.y;
|
|
neighbor.x = x + offset.x;
|
|
clampedNeighbor.x = ClampToCanvas(neighbor.x, imageWidth);
|
|
clampedNeighbor.y = ClampToCanvas(neighbor.y, imageHeight);
|
|
|
|
CLPixelType r = (clampedNeighbor.x == neighbor.x
|
|
&& clampedNeighbor.y == neighbor.y)?inputImage[clampedNeighbor.y*imageWidth+clampedNeighbor.x]
|
|
:(CLPixelType)0;
|
|
|
|
|
|
neighbor.y = y - offset.y;
|
|
neighbor.x = x - offset.x;
|
|
clampedNeighbor.x = ClampToCanvas(neighbor.x, imageWidth);
|
|
clampedNeighbor.y = ClampToCanvas(neighbor.y, imageHeight);
|
|
|
|
CLPixelType s = (clampedNeighbor.x == neighbor.x
|
|
&& clampedNeighbor.y == neighbor.y)?inputImage[clampedNeighbor.y*imageWidth+clampedNeighbor.x]
|
|
:(CLPixelType)0;
|
|
|
|
|
|
int sv[4];
|
|
sv[0] = (int)v.x;
|
|
sv[1] = (int)v.y;
|
|
sv[2] = (int)v.z;
|
|
sv[3] = (int)v.w;
|
|
|
|
int sr[4];
|
|
sr[0] = (int)r.x;
|
|
sr[1] = (int)r.y;
|
|
sr[2] = (int)r.z;
|
|
sr[3] = (int)r.w;
|
|
|
|
int ss[4];
|
|
ss[0] = (int)s.x;
|
|
ss[1] = (int)s.y;
|
|
ss[2] = (int)s.z;
|
|
ss[3] = (int)s.w;
|
|
|
|
if (polarity > 0) {
|
|
\n #pragma unroll 4\n
|
|
for (unsigned int i = 0; i < 4; i++) {
|
|
//sv[i] = (ss[i] >= (sv[i]+ScaleCharToQuantum(2)) && sr[i] > sv[i] ) ? (sv[i]+ScaleCharToQuantum(1)):sv[i];
|
|
//
|
|
//sv[i] =(!( (int)(ss[i] >= (sv[i]+ScaleCharToQuantum(2))) && (int) (sr[i] > sv[i] ) )) ? sv[i]:(sv[i]+ScaleCharToQuantum(1));
|
|
//sv[i] =(( (int)( ss[i] < (sv[i]+ScaleCharToQuantum(2))) || (int) ( sr[i] <= sv[i] ) )) ? sv[i]:(sv[i]+ScaleCharToQuantum(1));
|
|
sv[i] =(( (int)( ss[i] < (sv[i]+ScaleCharToQuantum(2))) + (int) ( sr[i] <= sv[i] ) ) !=0) ? sv[i]:(sv[i]+ScaleCharToQuantum(1));
|
|
}
|
|
}
|
|
else {
|
|
\n #pragma unroll 4\n
|
|
for (unsigned int i = 0; i < 4; i++) {
|
|
//sv[i] = (ss[i] <= (sv[i]-ScaleCharToQuantum(2)) && sr[i] < sv[i] ) ? (sv[i]-ScaleCharToQuantum(1)):sv[i];
|
|
//
|
|
//sv[i] = ( (int)(ss[i] <= (sv[i]-ScaleCharToQuantum(2)) ) + (int)( sr[i] < sv[i] ) ==0) ? sv[i]:(sv[i]-ScaleCharToQuantum(1));
|
|
sv[i] = (( (int)(ss[i] > (sv[i]-ScaleCharToQuantum(2))) + (int)( sr[i] >= sv[i] )) !=0) ? sv[i]:(sv[i]-ScaleCharToQuantum(1));
|
|
}
|
|
}
|
|
|
|
v.x = (CLQuantum)sv[0];
|
|
v.y = (CLQuantum)sv[1];
|
|
v.z = (CLQuantum)sv[2];
|
|
|
|
if (matte!=0)
|
|
v.w = (CLQuantum)sv[3];
|
|
|
|
outputImage[y*imageWidth+x] = v;
|
|
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% E q u a l i z e %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
/*
|
|
*/
|
|
__kernel void Equalize(__global CLPixelType * restrict im,
|
|
const ChannelType channel,
|
|
__global CLPixelType * restrict equalize_map,
|
|
const float4 white, const float4 black)
|
|
{
|
|
const int x = get_global_id(0);
|
|
const int y = get_global_id(1);
|
|
const int columns = get_global_size(0);
|
|
const int c = x + y * columns;
|
|
|
|
uint ePos;
|
|
CLPixelType oValue, eValue;
|
|
CLQuantum red, green, blue, alpha;
|
|
|
|
//read from global
|
|
oValue=im[c];
|
|
|
|
if ((channel & SyncChannels) != 0)
|
|
{
|
|
if (getRedF4(white) != getRedF4(black))
|
|
{
|
|
ePos = ScaleQuantumToMap(getRed(oValue));
|
|
eValue = equalize_map[ePos];
|
|
red = getRed(eValue);
|
|
ePos = ScaleQuantumToMap(getGreen(oValue));
|
|
eValue = equalize_map[ePos];
|
|
green = getRed(eValue);
|
|
ePos = ScaleQuantumToMap(getBlue(oValue));
|
|
eValue = equalize_map[ePos];
|
|
blue = getRed(eValue);
|
|
ePos = ScaleQuantumToMap(getAlpha(oValue));
|
|
eValue = equalize_map[ePos];
|
|
alpha = getRed(eValue);
|
|
|
|
//write back
|
|
im[c]=(CLPixelType)(blue, green, red, alpha);
|
|
}
|
|
|
|
}
|
|
|
|
// for equalizing, we always need all channels?
|
|
// otherwise something more
|
|
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% F u n c t i o n %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
/*
|
|
apply FunctionImageChannel(braightness-contrast)
|
|
*/
|
|
CLQuantum ApplyFunction(float pixel,const MagickFunction function,
|
|
const unsigned int number_parameters,__constant float *parameters)
|
|
{
|
|
float result = 0.0f;
|
|
|
|
switch (function)
|
|
{
|
|
case PolynomialFunction:
|
|
{
|
|
for (unsigned int i=0; i < number_parameters; i++)
|
|
result = result*QuantumScale*pixel + parameters[i];
|
|
result *= QuantumRange;
|
|
break;
|
|
}
|
|
case SinusoidFunction:
|
|
{
|
|
float freq,phase,ampl,bias;
|
|
freq = ( number_parameters >= 1 ) ? parameters[0] : 1.0f;
|
|
phase = ( number_parameters >= 2 ) ? parameters[1] : 0.0f;
|
|
ampl = ( number_parameters >= 3 ) ? parameters[2] : 0.5f;
|
|
bias = ( number_parameters >= 4 ) ? parameters[3] : 0.5f;
|
|
result = QuantumRange*(ampl*sin(2.0f*MagickPI*
|
|
(freq*QuantumScale*pixel + phase/360.0f)) + bias);
|
|
break;
|
|
}
|
|
case ArcsinFunction:
|
|
{
|
|
float width,range,center,bias;
|
|
width = ( number_parameters >= 1 ) ? parameters[0] : 1.0f;
|
|
center = ( number_parameters >= 2 ) ? parameters[1] : 0.5f;
|
|
range = ( number_parameters >= 3 ) ? parameters[2] : 1.0f;
|
|
bias = ( number_parameters >= 4 ) ? parameters[3] : 0.5f;
|
|
|
|
result = 2.0f/width*(QuantumScale*pixel - center);
|
|
result = range/MagickPI*asin(result)+bias;
|
|
result = ( result <= -1.0f ) ? bias - range/2.0f : result;
|
|
result = ( result >= 1.0f ) ? bias + range/2.0f : result;
|
|
result *= QuantumRange;
|
|
break;
|
|
}
|
|
case ArctanFunction:
|
|
{
|
|
float slope,range,center,bias;
|
|
slope = ( number_parameters >= 1 ) ? parameters[0] : 1.0f;
|
|
center = ( number_parameters >= 2 ) ? parameters[1] : 0.5f;
|
|
range = ( number_parameters >= 3 ) ? parameters[2] : 1.0f;
|
|
bias = ( number_parameters >= 4 ) ? parameters[3] : 0.5f;
|
|
result = MagickPI*slope*(QuantumScale*pixel-center);
|
|
result = QuantumRange*(range/MagickPI*atan(result) + bias);
|
|
break;
|
|
}
|
|
case UndefinedFunction:
|
|
break;
|
|
}
|
|
return(ClampToQuantum(result));
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
/*
|
|
Improve brightness / contrast of the image
|
|
channel : define which channel is improved
|
|
function : the function called to enchance the brightness contrast
|
|
number_parameters : numbers of parameters
|
|
parameters : the parameter
|
|
*/
|
|
__kernel void ComputeFunction(__global CLQuantum *image,const unsigned int number_channels,
|
|
const ChannelType channel,const MagickFunction function,const unsigned int number_parameters,
|
|
__constant float *parameters)
|
|
{
|
|
const unsigned int x = get_global_id(0);
|
|
const unsigned int y = get_global_id(1);
|
|
const unsigned int columns = get_global_size(0);
|
|
__global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y);
|
|
|
|
float red;
|
|
float green;
|
|
float blue;
|
|
float alpha;
|
|
|
|
ReadChannels(p, number_channels, channel, &red, &green, &blue, &alpha);
|
|
|
|
if ((channel & RedChannel) != 0)
|
|
red=ApplyFunction(red, function, number_parameters, parameters);
|
|
|
|
if (number_channels > 2)
|
|
{
|
|
if ((channel & GreenChannel) != 0)
|
|
green=ApplyFunction(green, function, number_parameters, parameters);
|
|
|
|
if ((channel & BlueChannel) != 0)
|
|
blue=ApplyFunction(blue, function, number_parameters, parameters);
|
|
}
|
|
|
|
if (((number_channels == 4) || (number_channels == 2)) &&
|
|
((channel & AlphaChannel) != 0))
|
|
alpha=ApplyFunction(alpha, function, number_parameters, parameters);
|
|
|
|
WriteChannels(p, number_channels, channel, red, green, blue, alpha);
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% G r a y s c a l e %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
__kernel void Grayscale(__global CLQuantum *image,const int number_channels,
|
|
const unsigned int colorspace,const unsigned int method)
|
|
{
|
|
const unsigned int x = get_global_id(0);
|
|
const unsigned int y = get_global_id(1);
|
|
const unsigned int columns = get_global_size(0);
|
|
__global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y);
|
|
|
|
float
|
|
blue,
|
|
green,
|
|
red;
|
|
|
|
red=getPixelRed(p);
|
|
green=getPixelGreen(p);
|
|
blue=getPixelBlue(p);
|
|
|
|
CLQuantum intensity=ClampToQuantum(GetPixelIntensity(colorspace, method, red, green, blue));
|
|
|
|
setPixelRed(p,intensity);
|
|
setPixelGreen(p,intensity);
|
|
setPixelBlue(p,intensity);
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% L o c a l C o n t r a s t %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
|
|
__kernel void LocalContrastBlurRow(__global CLPixelType *srcImage, __global CLPixelType *dstImage, __global float *tmpImage,
|
|
const int radius,
|
|
const int imageWidth,
|
|
const int imageHeight)
|
|
{
|
|
const float4 RGB = ((float4)(0.2126f, 0.7152f, 0.0722f, 0.0f));
|
|
|
|
int x = get_local_id(0);
|
|
int y = get_global_id(1);
|
|
|
|
if ((x >= imageWidth) || (y >= imageHeight))
|
|
return;
|
|
|
|
global CLPixelType *src = srcImage + y * imageWidth;
|
|
|
|
for (int i = x; i < imageWidth; i += get_local_size(0)) {
|
|
float sum = 0.0f;
|
|
float weight = 1.0f;
|
|
|
|
int j = i - radius;
|
|
while ((j + 7) < i) {
|
|
for (int k = 0; k < 8; ++k) // Unroll 8x
|
|
sum += (weight + k) * dot(RGB, convert_float4(src[mirrorBottom(j+k)]));
|
|
weight += 8.0f;
|
|
j+=8;
|
|
}
|
|
while (j < i) {
|
|
sum += weight * dot(RGB, convert_float4(src[mirrorBottom(j)]));
|
|
weight += 1.0f;
|
|
++j;
|
|
}
|
|
|
|
while ((j + 7) < radius + i) {
|
|
for (int k = 0; k < 8; ++k) // Unroll 8x
|
|
sum += (weight - k) * dot(RGB, convert_float4(src[mirrorTop(j + k, imageWidth)]));
|
|
weight -= 8.0f;
|
|
j+=8;
|
|
}
|
|
while (j < radius + i) {
|
|
sum += weight * dot(RGB, convert_float4(src[mirrorTop(j, imageWidth)]));
|
|
weight -= 1.0f;
|
|
++j;
|
|
}
|
|
|
|
tmpImage[i + y * imageWidth] = sum / ((radius + 1) * (radius + 1));
|
|
}
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
__kernel void LocalContrastBlurApplyColumn(__global CLPixelType *srcImage, __global CLPixelType *dstImage, __global float *blurImage,
|
|
const int radius,
|
|
const float strength,
|
|
const int imageWidth,
|
|
const int imageHeight)
|
|
{
|
|
const float4 RGB = (float4)(0.2126f, 0.7152f, 0.0722f, 0.0f);
|
|
|
|
int x = get_global_id(0);
|
|
int y = get_global_id(1);
|
|
|
|
if ((x >= imageWidth) || (y >= imageHeight))
|
|
return;
|
|
|
|
global float *src = blurImage + x;
|
|
|
|
float sum = 0.0f;
|
|
float weight = 1.0f;
|
|
|
|
int j = y - radius;
|
|
while ((j + 7) < y) {
|
|
for (int k = 0; k < 8; ++k) // Unroll 8x
|
|
sum += (weight + k) * src[mirrorBottom(j+k) * imageWidth];
|
|
weight += 8.0f;
|
|
j+=8;
|
|
}
|
|
while (j < y) {
|
|
sum += weight * src[mirrorBottom(j) * imageWidth];
|
|
weight += 1.0f;
|
|
++j;
|
|
}
|
|
|
|
while ((j + 7) < radius + y) {
|
|
for (int k = 0; k < 8; ++k) // Unroll 8x
|
|
sum += (weight - k) * src[mirrorTop(j + k, imageHeight) * imageWidth];
|
|
weight -= 8.0f;
|
|
j+=8;
|
|
}
|
|
while (j < radius + y) {
|
|
sum += weight * src[mirrorTop(j, imageHeight) * imageWidth];
|
|
weight -= 1.0f;
|
|
++j;
|
|
}
|
|
|
|
CLPixelType pixel = srcImage[x + y * imageWidth];
|
|
float srcVal = dot(RGB, convert_float4(pixel));
|
|
float mult = (srcVal - (sum / ((radius + 1) * (radius + 1)))) * (strength / 100.0f);
|
|
mult = (srcVal + mult) / srcVal;
|
|
|
|
pixel.x = ClampToQuantum(pixel.x * mult);
|
|
pixel.y = ClampToQuantum(pixel.y * mult);
|
|
pixel.z = ClampToQuantum(pixel.z * mult);
|
|
|
|
dstImage[x + y * imageWidth] = pixel;
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% M o d u l a t e %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
|
|
static inline void ConvertRGBToHSL(const CLQuantum red,const CLQuantum green, const CLQuantum blue,
|
|
float *hue, float *saturation, float *lightness)
|
|
{
|
|
float
|
|
c,
|
|
tmax,
|
|
tmin;
|
|
|
|
/*
|
|
Convert RGB to HSL colorspace.
|
|
*/
|
|
tmax=MagickMax(QuantumScale*red,MagickMax(QuantumScale*green, QuantumScale*blue));
|
|
tmin=MagickMin(QuantumScale*red,MagickMin(QuantumScale*green, QuantumScale*blue));
|
|
|
|
c=tmax-tmin;
|
|
|
|
*lightness=(tmax+tmin)/2.0;
|
|
if (c <= 0.0)
|
|
{
|
|
*hue=0.0;
|
|
*saturation=0.0;
|
|
return;
|
|
}
|
|
|
|
if (tmax == (QuantumScale*red))
|
|
{
|
|
*hue=(QuantumScale*green-QuantumScale*blue)/c;
|
|
if ((QuantumScale*green) < (QuantumScale*blue))
|
|
*hue+=6.0;
|
|
}
|
|
else
|
|
if (tmax == (QuantumScale*green))
|
|
*hue=2.0+(QuantumScale*blue-QuantumScale*red)/c;
|
|
else
|
|
*hue=4.0+(QuantumScale*red-QuantumScale*green)/c;
|
|
|
|
*hue*=60.0/360.0;
|
|
if (*lightness <= 0.5)
|
|
*saturation=c/(2.0*(*lightness));
|
|
else
|
|
*saturation=c/(2.0-2.0*(*lightness));
|
|
}
|
|
|
|
static inline void ConvertHSLToRGB(const float hue,const float saturation, const float lightness,
|
|
CLQuantum *red,CLQuantum *green,CLQuantum *blue)
|
|
{
|
|
float
|
|
b,
|
|
c,
|
|
g,
|
|
h,
|
|
tmin,
|
|
r,
|
|
x;
|
|
|
|
/*
|
|
Convert HSL to RGB colorspace.
|
|
*/
|
|
h=hue*360.0;
|
|
if (lightness <= 0.5)
|
|
c=2.0*lightness*saturation;
|
|
else
|
|
c=(2.0-2.0*lightness)*saturation;
|
|
tmin=lightness-0.5*c;
|
|
h-=360.0*floor(h/360.0);
|
|
h/=60.0;
|
|
x=c*(1.0-fabs(h-2.0*floor(h/2.0)-1.0));
|
|
switch ((int) floor(h) % 6)
|
|
{
|
|
case 0:
|
|
{
|
|
r=tmin+c;
|
|
g=tmin+x;
|
|
b=tmin;
|
|
break;
|
|
}
|
|
case 1:
|
|
{
|
|
r=tmin+x;
|
|
g=tmin+c;
|
|
b=tmin;
|
|
break;
|
|
}
|
|
case 2:
|
|
{
|
|
r=tmin;
|
|
g=tmin+c;
|
|
b=tmin+x;
|
|
break;
|
|
}
|
|
case 3:
|
|
{
|
|
r=tmin;
|
|
g=tmin+x;
|
|
b=tmin+c;
|
|
break;
|
|
}
|
|
case 4:
|
|
{
|
|
r=tmin+x;
|
|
g=tmin;
|
|
b=tmin+c;
|
|
break;
|
|
}
|
|
case 5:
|
|
{
|
|
r=tmin+c;
|
|
g=tmin;
|
|
b=tmin+x;
|
|
break;
|
|
}
|
|
default:
|
|
{
|
|
r=0.0;
|
|
g=0.0;
|
|
b=0.0;
|
|
}
|
|
}
|
|
*red=ClampToQuantum(QuantumRange*r);
|
|
*green=ClampToQuantum(QuantumRange*g);
|
|
*blue=ClampToQuantum(QuantumRange*b);
|
|
}
|
|
|
|
static inline void ModulateHSL(const float percent_hue, const float percent_saturation,const float percent_lightness,
|
|
CLQuantum *red,CLQuantum *green,CLQuantum *blue)
|
|
{
|
|
float
|
|
hue,
|
|
lightness,
|
|
saturation;
|
|
|
|
/*
|
|
Increase or decrease color lightness, saturation, or hue.
|
|
*/
|
|
ConvertRGBToHSL(*red,*green,*blue,&hue,&saturation,&lightness);
|
|
hue+=0.5*(0.01*percent_hue-1.0);
|
|
while (hue < 0.0)
|
|
hue+=1.0;
|
|
while (hue >= 1.0)
|
|
hue-=1.0;
|
|
saturation*=0.01*percent_saturation;
|
|
lightness*=0.01*percent_lightness;
|
|
ConvertHSLToRGB(hue,saturation,lightness,red,green,blue);
|
|
}
|
|
|
|
__kernel void Modulate(__global CLPixelType *im,
|
|
const float percent_brightness,
|
|
const float percent_hue,
|
|
const float percent_saturation,
|
|
const int colorspace)
|
|
{
|
|
|
|
const int x = get_global_id(0);
|
|
const int y = get_global_id(1);
|
|
const int columns = get_global_size(0);
|
|
const int c = x + y * columns;
|
|
|
|
CLPixelType pixel = im[c];
|
|
|
|
CLQuantum
|
|
blue,
|
|
green,
|
|
red;
|
|
|
|
red=getRed(pixel);
|
|
green=getGreen(pixel);
|
|
blue=getBlue(pixel);
|
|
|
|
switch (colorspace)
|
|
{
|
|
case HSLColorspace:
|
|
default:
|
|
{
|
|
ModulateHSL(percent_hue, percent_saturation, percent_brightness,
|
|
&red, &green, &blue);
|
|
}
|
|
|
|
}
|
|
|
|
CLPixelType filteredPixel;
|
|
|
|
setRed(&filteredPixel, red);
|
|
setGreen(&filteredPixel, green);
|
|
setBlue(&filteredPixel, blue);
|
|
filteredPixel.w = pixel.w;
|
|
|
|
im[c] = filteredPixel;
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% M o t i o n B l u r %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
__kernel
|
|
void MotionBlur(const __global CLPixelType *input, __global CLPixelType *output,
|
|
const unsigned int imageWidth, const unsigned int imageHeight,
|
|
const __global float *filter, const unsigned int width, const __global int2* offset,
|
|
const float4 bias,
|
|
const ChannelType channel, const unsigned int matte) {
|
|
|
|
int2 currentPixel;
|
|
currentPixel.x = get_global_id(0);
|
|
currentPixel.y = get_global_id(1);
|
|
|
|
if (currentPixel.x >= imageWidth
|
|
|| currentPixel.y >= imageHeight)
|
|
return;
|
|
|
|
float4 pixel;
|
|
pixel.x = (float)bias.x;
|
|
pixel.y = (float)bias.y;
|
|
pixel.z = (float)bias.z;
|
|
pixel.w = (float)bias.w;
|
|
|
|
if (((channel & AlphaChannel) == 0) || (matte == 0)) {
|
|
|
|
for (int i = 0; i < width; i++) {
|
|
// only support EdgeVirtualPixelMethod through ClampToCanvas
|
|
// TODO: implement other virtual pixel method
|
|
int2 samplePixel = currentPixel + offset[i];
|
|
samplePixel.x = ClampToCanvas(samplePixel.x, imageWidth);
|
|
samplePixel.y = ClampToCanvas(samplePixel.y, imageHeight);
|
|
CLPixelType samplePixelValue = input[ samplePixel.y * imageWidth + samplePixel.x];
|
|
|
|
pixel.x += (filter[i] * (float)samplePixelValue.x);
|
|
pixel.y += (filter[i] * (float)samplePixelValue.y);
|
|
pixel.z += (filter[i] * (float)samplePixelValue.z);
|
|
pixel.w += (filter[i] * (float)samplePixelValue.w);
|
|
}
|
|
|
|
CLPixelType outputPixel;
|
|
outputPixel.x = ClampToQuantum(pixel.x);
|
|
outputPixel.y = ClampToQuantum(pixel.y);
|
|
outputPixel.z = ClampToQuantum(pixel.z);
|
|
outputPixel.w = ClampToQuantum(pixel.w);
|
|
output[currentPixel.y * imageWidth + currentPixel.x] = outputPixel;
|
|
}
|
|
else {
|
|
|
|
float gamma = 0.0f;
|
|
for (int i = 0; i < width; i++) {
|
|
// only support EdgeVirtualPixelMethod through ClampToCanvas
|
|
// TODO: implement other virtual pixel method
|
|
int2 samplePixel = currentPixel + offset[i];
|
|
samplePixel.x = ClampToCanvas(samplePixel.x, imageWidth);
|
|
samplePixel.y = ClampToCanvas(samplePixel.y, imageHeight);
|
|
|
|
CLPixelType samplePixelValue = input[ samplePixel.y * imageWidth + samplePixel.x];
|
|
|
|
float alpha = QuantumScale*samplePixelValue.w;
|
|
float k = filter[i];
|
|
pixel.x = pixel.x + k * alpha * samplePixelValue.x;
|
|
pixel.y = pixel.y + k * alpha * samplePixelValue.y;
|
|
pixel.z = pixel.z + k * alpha * samplePixelValue.z;
|
|
|
|
pixel.w += k * alpha * samplePixelValue.w;
|
|
|
|
gamma+=k*alpha;
|
|
}
|
|
gamma = PerceptibleReciprocal(gamma);
|
|
pixel.xyz = gamma*pixel.xyz;
|
|
|
|
CLPixelType outputPixel;
|
|
outputPixel.x = ClampToQuantum(pixel.x);
|
|
outputPixel.y = ClampToQuantum(pixel.y);
|
|
outputPixel.z = ClampToQuantum(pixel.z);
|
|
outputPixel.w = ClampToQuantum(pixel.w);
|
|
output[currentPixel.y * imageWidth + currentPixel.x] = outputPixel;
|
|
}
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% R e s i z e %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
// Based on Box from resize.c
|
|
float BoxResizeFilter(const float x)
|
|
{
|
|
return 1.0f;
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
// Based on CubicBC from resize.c
|
|
float CubicBC(const float x,const __global float* resizeFilterCoefficients)
|
|
{
|
|
/*
|
|
Cubic Filters using B,C determined values:
|
|
Mitchell-Netravali B = 1/3 C = 1/3 "Balanced" cubic spline filter
|
|
Catmull-Rom B = 0 C = 1/2 Interpolatory and exact on linears
|
|
Spline B = 1 C = 0 B-Spline Gaussian approximation
|
|
Hermite B = 0 C = 0 B-Spline interpolator
|
|
|
|
See paper by Mitchell and Netravali, Reconstruction Filters in Computer
|
|
Graphics Computer Graphics, Volume 22, Number 4, August 1988
|
|
http://www.cs.utexas.edu/users/fussell/courses/cs384g/lectures/mitchell/
|
|
Mitchell.pdf.
|
|
|
|
Coefficents are determined from B,C values:
|
|
P0 = ( 6 - 2*B )/6 = coeff[0]
|
|
P1 = 0
|
|
P2 = (-18 +12*B + 6*C )/6 = coeff[1]
|
|
P3 = ( 12 - 9*B - 6*C )/6 = coeff[2]
|
|
Q0 = ( 8*B +24*C )/6 = coeff[3]
|
|
Q1 = ( -12*B -48*C )/6 = coeff[4]
|
|
Q2 = ( 6*B +30*C )/6 = coeff[5]
|
|
Q3 = ( - 1*B - 6*C )/6 = coeff[6]
|
|
|
|
which are used to define the filter:
|
|
|
|
P0 + P1*x + P2*x^2 + P3*x^3 0 <= x < 1
|
|
Q0 + Q1*x + Q2*x^2 + Q3*x^3 1 <= x < 2
|
|
|
|
which ensures function is continuous in value and derivative (slope).
|
|
*/
|
|
if (x < 1.0)
|
|
return(resizeFilterCoefficients[0]+x*(x*
|
|
(resizeFilterCoefficients[1]+x*resizeFilterCoefficients[2])));
|
|
if (x < 2.0)
|
|
return(resizeFilterCoefficients[3]+x*(resizeFilterCoefficients[4]+x*
|
|
(resizeFilterCoefficients[5]+x*resizeFilterCoefficients[6])));
|
|
return(0.0);
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
float Sinc(const float x)
|
|
{
|
|
if (x != 0.0f)
|
|
{
|
|
const float alpha=(float) (MagickPI*x);
|
|
return sinpi(x)/alpha;
|
|
}
|
|
return(1.0f);
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
float Triangle(const float x)
|
|
{
|
|
/*
|
|
1st order (linear) B-Spline, bilinear interpolation, Tent 1D filter, or
|
|
a Bartlett 2D Cone filter. Also used as a Bartlett Windowing function
|
|
for Sinc().
|
|
*/
|
|
return ((x<1.0f)?(1.0f-x):0.0f);
|
|
}
|
|
)
|
|
|
|
|
|
STRINGIFY(
|
|
float Hann(const float x)
|
|
{
|
|
/*
|
|
Cosine window function:
|
|
0.5+0.5*cos(pi*x).
|
|
*/
|
|
const float cosine=cos((MagickPI*x));
|
|
return(0.5f+0.5f*cosine);
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
float Hamming(const float x)
|
|
{
|
|
/*
|
|
Offset cosine window function:
|
|
.54 + .46 cos(pi x).
|
|
*/
|
|
const float cosine=cos((MagickPI*x));
|
|
return(0.54f+0.46f*cosine);
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
float Blackman(const float x)
|
|
{
|
|
/*
|
|
Blackman: 2nd order cosine windowing function:
|
|
0.42 + 0.5 cos(pi x) + 0.08 cos(2pi x)
|
|
|
|
Refactored by Chantal Racette and Nicolas Robidoux to one trig call and
|
|
five flops.
|
|
*/
|
|
const float cosine=cos((MagickPI*x));
|
|
return(0.34f+cosine*(0.5f+cosine*0.16f));
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
static inline float applyResizeFilter(const float x, const ResizeWeightingFunctionType filterType, const __global float* filterCoefficients)
|
|
{
|
|
switch (filterType)
|
|
{
|
|
/* Call Sinc even for SincFast to get better precision on GPU
|
|
and to avoid thread divergence. Sinc is pretty fast on GPU anyway...*/
|
|
case SincWeightingFunction:
|
|
case SincFastWeightingFunction:
|
|
return Sinc(x);
|
|
case CubicBCWeightingFunction:
|
|
return CubicBC(x,filterCoefficients);
|
|
case BoxWeightingFunction:
|
|
return BoxResizeFilter(x);
|
|
case TriangleWeightingFunction:
|
|
return Triangle(x);
|
|
case HannWeightingFunction:
|
|
return Hann(x);
|
|
case HammingWeightingFunction:
|
|
return Hamming(x);
|
|
case BlackmanWeightingFunction:
|
|
return Blackman(x);
|
|
|
|
default:
|
|
return 0.0f;
|
|
}
|
|
}
|
|
)
|
|
|
|
|
|
STRINGIFY(
|
|
static inline float getResizeFilterWeight(const __global float* resizeFilterCubicCoefficients, const ResizeWeightingFunctionType resizeFilterType
|
|
, const ResizeWeightingFunctionType resizeWindowType
|
|
, const float resizeFilterScale, const float resizeWindowSupport, const float resizeFilterBlur, const float x)
|
|
{
|
|
float scale;
|
|
float xBlur = fabs(x/resizeFilterBlur);
|
|
if (resizeWindowSupport < MagickEpsilon
|
|
|| resizeWindowType == BoxWeightingFunction)
|
|
{
|
|
scale = 1.0f;
|
|
}
|
|
else
|
|
{
|
|
scale = resizeFilterScale;
|
|
scale = applyResizeFilter(xBlur*scale, resizeWindowType, resizeFilterCubicCoefficients);
|
|
}
|
|
float weight = scale * applyResizeFilter(xBlur, resizeFilterType, resizeFilterCubicCoefficients);
|
|
return weight;
|
|
}
|
|
|
|
)
|
|
|
|
;
|
|
const char *accelerateKernels2 =
|
|
|
|
STRINGIFY(
|
|
|
|
static inline unsigned int getNumWorkItemsPerPixel(const unsigned int pixelPerWorkgroup, const unsigned int numWorkItems) {
|
|
return (numWorkItems/pixelPerWorkgroup);
|
|
}
|
|
|
|
// returns the index of the pixel for the current workitem to compute.
|
|
// returns -1 if this workitem doesn't need to participate in any computation
|
|
static inline int pixelToCompute(const unsigned itemID, const unsigned int pixelPerWorkgroup, const unsigned int numWorkItems) {
|
|
const unsigned int numWorkItemsPerPixel = getNumWorkItemsPerPixel(pixelPerWorkgroup, numWorkItems);
|
|
int pixelIndex = itemID/numWorkItemsPerPixel;
|
|
pixelIndex = (pixelIndex<pixelPerWorkgroup)?pixelIndex:-1;
|
|
return pixelIndex;
|
|
}
|
|
|
|
)
|
|
|
|
STRINGIFY(
|
|
__kernel __attribute__((reqd_work_group_size(256, 1, 1)))
|
|
void ResizeHorizontalFilter(const __global CLQuantum *inputImage, const unsigned int number_channels,
|
|
const unsigned int inputColumns, const unsigned int inputRows, __global CLQuantum *filteredImage,
|
|
const unsigned int filteredColumns, const unsigned int filteredRows, const float xFactor,
|
|
const int resizeFilterType, const int resizeWindowType, const __global float *resizeFilterCubicCoefficients,
|
|
const float resizeFilterScale, const float resizeFilterSupport, const float resizeFilterWindowSupport,
|
|
const float resizeFilterBlur, __local CLQuantum *inputImageCache, const int numCachedPixels,
|
|
const unsigned int pixelPerWorkgroup, const unsigned int pixelChunkSize,
|
|
__local float4 *outputPixelCache, __local float *densityCache, __local float *gammaCache)
|
|
{
|
|
// calculate the range of resized image pixels computed by this workgroup
|
|
const unsigned int startX = get_group_id(0)*pixelPerWorkgroup;
|
|
const unsigned int stopX = MagickMin(startX + pixelPerWorkgroup,filteredColumns);
|
|
const unsigned int actualNumPixelToCompute = stopX - startX;
|
|
|
|
// calculate the range of input image pixels to cache
|
|
float scale = MagickMax(1.0f/xFactor+MagickEpsilon ,1.0f);
|
|
const float support = MagickMax(scale*resizeFilterSupport,0.5f);
|
|
scale = PerceptibleReciprocal(scale);
|
|
|
|
const int cacheRangeStartX = MagickMax((int)((startX+0.5f)/xFactor+MagickEpsilon-support+0.5f),(int)(0));
|
|
const int cacheRangeEndX = MagickMin((int)(cacheRangeStartX + numCachedPixels), (int)inputColumns);
|
|
|
|
// cache the input pixels into local memory
|
|
const unsigned int y = get_global_id(1);
|
|
const unsigned int pos = getPixelIndex(number_channels, inputColumns, cacheRangeStartX, y);
|
|
const unsigned int num_elements = (cacheRangeEndX - cacheRangeStartX) * number_channels;
|
|
event_t e = async_work_group_copy(inputImageCache, inputImage + pos, num_elements, 0);
|
|
wait_group_events(1, &e);
|
|
|
|
unsigned int alpha_index = (number_channels == 4) || (number_channels == 2) ? number_channels - 1 : 0;
|
|
unsigned int totalNumChunks = (actualNumPixelToCompute+pixelChunkSize-1)/pixelChunkSize;
|
|
for (unsigned int chunk = 0; chunk < totalNumChunks; chunk++)
|
|
{
|
|
const unsigned int chunkStartX = startX + chunk*pixelChunkSize;
|
|
const unsigned int chunkStopX = MagickMin(chunkStartX + pixelChunkSize, stopX);
|
|
const unsigned int actualNumPixelInThisChunk = chunkStopX - chunkStartX;
|
|
|
|
// determine which resized pixel computed by this workitem
|
|
const unsigned int itemID = get_local_id(0);
|
|
const unsigned int numItems = getNumWorkItemsPerPixel(actualNumPixelInThisChunk, get_local_size(0));
|
|
|
|
const int pixelIndex = pixelToCompute(itemID, actualNumPixelInThisChunk, get_local_size(0));
|
|
|
|
float4 filteredPixel = (float4)0.0f;
|
|
float density = 0.0f;
|
|
float gamma = 0.0f;
|
|
// -1 means this workitem doesn't participate in the computation
|
|
if (pixelIndex != -1)
|
|
{
|
|
// x coordinated of the resized pixel computed by this workitem
|
|
const int x = chunkStartX + pixelIndex;
|
|
|
|
// calculate how many steps required for this pixel
|
|
const float bisect = (x+0.5)/xFactor+MagickEpsilon;
|
|
const unsigned int start = (unsigned int)MagickMax(bisect-support+0.5f,0.0f);
|
|
const unsigned int stop = (unsigned int)MagickMin(bisect+support+0.5f,(float)inputColumns);
|
|
const unsigned int n = stop - start;
|
|
|
|
// calculate how many steps this workitem will contribute
|
|
unsigned int numStepsPerWorkItem = n / numItems;
|
|
numStepsPerWorkItem += ((numItems*numStepsPerWorkItem)==n?0:1);
|
|
|
|
const unsigned int startStep = (itemID%numItems)*numStepsPerWorkItem;
|
|
if (startStep < n)
|
|
{
|
|
const unsigned int stopStep = MagickMin(startStep+numStepsPerWorkItem, n);
|
|
|
|
unsigned int cacheIndex = start+startStep-cacheRangeStartX;
|
|
for (unsigned int i = startStep; i < stopStep; i++, cacheIndex++)
|
|
{
|
|
float weight = getResizeFilterWeight(resizeFilterCubicCoefficients,
|
|
(ResizeWeightingFunctionType) resizeFilterType,
|
|
(ResizeWeightingFunctionType) resizeWindowType,
|
|
resizeFilterScale, resizeFilterWindowSupport,
|
|
resizeFilterBlur, scale*(start + i - bisect + 0.5));
|
|
|
|
float4 cp = (float4)0.0f;
|
|
|
|
__local CLQuantum *p = inputImageCache + (cacheIndex*number_channels);
|
|
cp.x = (float) *(p);
|
|
if (number_channels > 2)
|
|
{
|
|
cp.y = (float) *(p + 1);
|
|
cp.z = (float) *(p + 2);
|
|
}
|
|
|
|
if (alpha_index != 0)
|
|
{
|
|
cp.w = (float) *(p + alpha_index);
|
|
|
|
float alpha = weight * QuantumScale * cp.w;
|
|
|
|
filteredPixel.x += alpha * cp.x;
|
|
filteredPixel.y += alpha * cp.y;
|
|
filteredPixel.z += alpha * cp.z;
|
|
filteredPixel.w += weight * cp.w;
|
|
gamma += alpha;
|
|
}
|
|
else
|
|
filteredPixel += ((float4) weight)*cp;
|
|
|
|
density += weight;
|
|
}
|
|
}
|
|
}
|
|
|
|
// initialize the accumulators to zero
|
|
if (itemID < actualNumPixelInThisChunk) {
|
|
outputPixelCache[itemID] = (float4)0.0f;
|
|
densityCache[itemID] = 0.0f;
|
|
if (alpha_index != 0)
|
|
gammaCache[itemID] = 0.0f;
|
|
}
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
// accumulatte the filtered pixel value and the density
|
|
for (unsigned int i = 0; i < numItems; i++) {
|
|
if (pixelIndex != -1) {
|
|
if (itemID%numItems == i) {
|
|
outputPixelCache[pixelIndex]+=filteredPixel;
|
|
densityCache[pixelIndex]+=density;
|
|
if (alpha_index != 0)
|
|
gammaCache[pixelIndex]+=gamma;
|
|
}
|
|
}
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
}
|
|
|
|
if (itemID < actualNumPixelInThisChunk)
|
|
{
|
|
float4 filteredPixel = outputPixelCache[itemID];
|
|
|
|
float gamma = 0.0f;
|
|
if (alpha_index != 0)
|
|
gamma = gammaCache[itemID];
|
|
|
|
float density = densityCache[itemID];
|
|
if ((density != 0.0f) && (density != 1.0f))
|
|
{
|
|
density = PerceptibleReciprocal(density);
|
|
filteredPixel *= (float4) density;
|
|
if (alpha_index != 0)
|
|
gamma *= density;
|
|
}
|
|
|
|
if (alpha_index != 0)
|
|
{
|
|
gamma = PerceptibleReciprocal(gamma);
|
|
filteredPixel.x *= gamma;
|
|
filteredPixel.y *= gamma;
|
|
filteredPixel.z *= gamma;
|
|
}
|
|
|
|
WriteAllChannels(filteredImage, number_channels, filteredColumns, chunkStartX + itemID, y, filteredPixel);
|
|
}
|
|
}
|
|
}
|
|
)
|
|
|
|
|
|
STRINGIFY(
|
|
__kernel __attribute__((reqd_work_group_size(1, 256, 1)))
|
|
void ResizeVerticalFilter(const __global CLQuantum *inputImage, const unsigned int number_channels,
|
|
const unsigned int inputColumns, const unsigned int inputRows, __global CLQuantum *filteredImage,
|
|
const unsigned int filteredColumns, const unsigned int filteredRows, const float yFactor,
|
|
const int resizeFilterType, const int resizeWindowType, const __global float *resizeFilterCubicCoefficients,
|
|
const float resizeFilterScale, const float resizeFilterSupport, const float resizeFilterWindowSupport,
|
|
const float resizeFilterBlur, __local CLQuantum *inputImageCache, const int numCachedPixels,
|
|
const unsigned int pixelPerWorkgroup, const unsigned int pixelChunkSize,
|
|
__local float4 *outputPixelCache, __local float *densityCache, __local float *gammaCache)
|
|
{
|
|
// calculate the range of resized image pixels computed by this workgroup
|
|
const unsigned int startY = get_group_id(1)*pixelPerWorkgroup;
|
|
const unsigned int stopY = MagickMin(startY + pixelPerWorkgroup,filteredRows);
|
|
const unsigned int actualNumPixelToCompute = stopY - startY;
|
|
|
|
// calculate the range of input image pixels to cache
|
|
float scale = MagickMax(1.0f/yFactor+MagickEpsilon ,1.0f);
|
|
const float support = MagickMax(scale*resizeFilterSupport,0.5f);
|
|
scale = PerceptibleReciprocal(scale);
|
|
|
|
const int cacheRangeStartY = MagickMax((int)((startY+0.5f)/yFactor+MagickEpsilon-support+0.5f),(int)(0));
|
|
const int cacheRangeEndY = MagickMin((int)(cacheRangeStartY + numCachedPixels), (int)inputRows);
|
|
|
|
// cache the input pixels into local memory
|
|
const unsigned int x = get_global_id(0);
|
|
unsigned int pos = getPixelIndex(number_channels, inputColumns, x, cacheRangeStartY);
|
|
unsigned int rangeLength = cacheRangeEndY-cacheRangeStartY;
|
|
unsigned int stride = inputColumns * number_channels;
|
|
for (unsigned int i = 0; i < number_channels; i++)
|
|
{
|
|
event_t e = async_work_group_strided_copy(inputImageCache + (rangeLength*i), inputImage+pos+i, rangeLength, stride, 0);
|
|
wait_group_events(1,&e);
|
|
}
|
|
|
|
unsigned int alpha_index = (number_channels == 4) || (number_channels == 2) ? number_channels - 1 : 0;
|
|
unsigned int totalNumChunks = (actualNumPixelToCompute+pixelChunkSize-1)/pixelChunkSize;
|
|
for (unsigned int chunk = 0; chunk < totalNumChunks; chunk++)
|
|
{
|
|
const unsigned int chunkStartY = startY + chunk*pixelChunkSize;
|
|
const unsigned int chunkStopY = MagickMin(chunkStartY + pixelChunkSize, stopY);
|
|
const unsigned int actualNumPixelInThisChunk = chunkStopY - chunkStartY;
|
|
|
|
// determine which resized pixel computed by this workitem
|
|
const unsigned int itemID = get_local_id(1);
|
|
const unsigned int numItems = getNumWorkItemsPerPixel(actualNumPixelInThisChunk, get_local_size(1));
|
|
|
|
const int pixelIndex = pixelToCompute(itemID, actualNumPixelInThisChunk, get_local_size(1));
|
|
|
|
float4 filteredPixel = (float4)0.0f;
|
|
float density = 0.0f;
|
|
float gamma = 0.0f;
|
|
// -1 means this workitem doesn't participate in the computation
|
|
if (pixelIndex != -1)
|
|
{
|
|
// x coordinated of the resized pixel computed by this workitem
|
|
const int y = chunkStartY + pixelIndex;
|
|
|
|
// calculate how many steps required for this pixel
|
|
const float bisect = (y+0.5)/yFactor+MagickEpsilon;
|
|
const unsigned int start = (unsigned int)MagickMax(bisect-support+0.5f,0.0f);
|
|
const unsigned int stop = (unsigned int)MagickMin(bisect+support+0.5f,(float)inputRows);
|
|
const unsigned int n = stop - start;
|
|
|
|
// calculate how many steps this workitem will contribute
|
|
unsigned int numStepsPerWorkItem = n / numItems;
|
|
numStepsPerWorkItem += ((numItems*numStepsPerWorkItem)==n?0:1);
|
|
|
|
const unsigned int startStep = (itemID%numItems)*numStepsPerWorkItem;
|
|
if (startStep < n)
|
|
{
|
|
const unsigned int stopStep = MagickMin(startStep+numStepsPerWorkItem, n);
|
|
|
|
unsigned int cacheIndex = start+startStep-cacheRangeStartY;
|
|
for (unsigned int i = startStep; i < stopStep; i++, cacheIndex++)
|
|
{
|
|
float weight = getResizeFilterWeight(resizeFilterCubicCoefficients,
|
|
(ResizeWeightingFunctionType) resizeFilterType,
|
|
(ResizeWeightingFunctionType) resizeWindowType,
|
|
resizeFilterScale, resizeFilterWindowSupport,
|
|
resizeFilterBlur, scale*(start + i - bisect + 0.5));
|
|
|
|
float4 cp = (float4)0.0f;
|
|
|
|
__local CLQuantum *p = inputImageCache + cacheIndex;
|
|
cp.x = (float) *(p);
|
|
if (number_channels > 2)
|
|
{
|
|
cp.y = (float) *(p + rangeLength);
|
|
cp.z = (float) *(p + (rangeLength * 2));
|
|
}
|
|
|
|
if (alpha_index != 0)
|
|
{
|
|
cp.w = (float) *(p + (rangeLength * alpha_index));
|
|
|
|
float alpha = weight * QuantumScale * cp.w;
|
|
|
|
filteredPixel.x += alpha * cp.x;
|
|
filteredPixel.y += alpha * cp.y;
|
|
filteredPixel.z += alpha * cp.z;
|
|
filteredPixel.w += weight * cp.w;
|
|
gamma += alpha;
|
|
}
|
|
else
|
|
filteredPixel += ((float4) weight)*cp;
|
|
|
|
density += weight;
|
|
}
|
|
}
|
|
}
|
|
|
|
// initialize the accumulators to zero
|
|
if (itemID < actualNumPixelInThisChunk) {
|
|
outputPixelCache[itemID] = (float4)0.0f;
|
|
densityCache[itemID] = 0.0f;
|
|
if (alpha_index != 0)
|
|
gammaCache[itemID] = 0.0f;
|
|
}
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
// accumulatte the filtered pixel value and the density
|
|
for (unsigned int i = 0; i < numItems; i++) {
|
|
if (pixelIndex != -1) {
|
|
if (itemID%numItems == i) {
|
|
outputPixelCache[pixelIndex]+=filteredPixel;
|
|
densityCache[pixelIndex]+=density;
|
|
if (alpha_index != 0)
|
|
gammaCache[pixelIndex]+=gamma;
|
|
}
|
|
}
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
}
|
|
|
|
if (itemID < actualNumPixelInThisChunk)
|
|
{
|
|
float4 filteredPixel = outputPixelCache[itemID];
|
|
|
|
float gamma = 0.0f;
|
|
if (alpha_index != 0)
|
|
gamma = gammaCache[itemID];
|
|
|
|
float density = densityCache[itemID];
|
|
if ((density != 0.0f) && (density != 1.0f))
|
|
{
|
|
density = PerceptibleReciprocal(density);
|
|
filteredPixel *= (float4) density;
|
|
if (alpha_index != 0)
|
|
gamma *= density;
|
|
}
|
|
|
|
if (alpha_index != 0)
|
|
{
|
|
gamma = PerceptibleReciprocal(gamma);
|
|
filteredPixel.x *= gamma;
|
|
filteredPixel.y *= gamma;
|
|
filteredPixel.z *= gamma;
|
|
}
|
|
|
|
WriteAllChannels(filteredImage, number_channels, filteredColumns, x, chunkStartY + itemID, filteredPixel);
|
|
}
|
|
}
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% R o t a t i o n a l B l u r %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
__kernel void RotationalBlur(const __global CLQuantum *image,
|
|
const unsigned int number_channels,const unsigned int channel,
|
|
const float2 blurCenter,__constant float *cos_theta,
|
|
__constant float *sin_theta,const unsigned int cossin_theta_size,
|
|
__global CLQuantum *filteredImage)
|
|
{
|
|
const int x = get_global_id(0);
|
|
const int y = get_global_id(1);
|
|
const int columns = get_global_size(0);
|
|
const int rows = get_global_size(1);
|
|
unsigned int step = 1;
|
|
float center_x = (float) x - blurCenter.x;
|
|
float center_y = (float) y - blurCenter.y;
|
|
float radius = hypot(center_x, center_y);
|
|
|
|
float blur_radius = hypot(blurCenter.x, blurCenter.y);
|
|
|
|
if (radius > MagickEpsilon)
|
|
{
|
|
step = (unsigned int) (blur_radius / radius);
|
|
if (step == 0)
|
|
step = 1;
|
|
if (step >= cossin_theta_size)
|
|
step = cossin_theta_size-1;
|
|
}
|
|
|
|
float4 result = 0.0f;
|
|
float normalize = 0.0f;
|
|
float gamma = 0.0f;
|
|
|
|
for (unsigned int i=0; i<cossin_theta_size; i+=step)
|
|
{
|
|
int cx = ClampToCanvas(blurCenter.x+center_x*cos_theta[i]-center_y*sin_theta[i]+0.5f,columns);
|
|
int cy = ClampToCanvas(blurCenter.y+center_x*sin_theta[i]+center_y*cos_theta[i]+0.5f,rows);
|
|
|
|
float4 pixel = ReadAllChannels(image, number_channels, columns, cx, cy);
|
|
|
|
if ((number_channels == 4) || (number_channels == 2))
|
|
{
|
|
float alpha = (float)(QuantumScale*pixel.w);
|
|
|
|
gamma += alpha;
|
|
|
|
result.x += alpha * pixel.x;
|
|
result.y += alpha * pixel.y;
|
|
result.z += alpha * pixel.z;
|
|
result.w += pixel.w;
|
|
}
|
|
else
|
|
result += pixel;
|
|
|
|
normalize += 1.0f;
|
|
}
|
|
|
|
normalize = PerceptibleReciprocal(normalize);
|
|
|
|
if ((number_channels == 4) || (number_channels == 2))
|
|
{
|
|
gamma = PerceptibleReciprocal(gamma);
|
|
result.x *= gamma;
|
|
result.y *= gamma;
|
|
result.z *= gamma;
|
|
result.w *= normalize;
|
|
}
|
|
else
|
|
result *= normalize;
|
|
|
|
WriteFloat4(filteredImage, number_channels, columns, x, y, channel, result);
|
|
}
|
|
)
|
|
|
|
/*
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
% %
|
|
% %
|
|
% %
|
|
% U n s h a r p M a s k %
|
|
% %
|
|
% %
|
|
% %
|
|
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
|
|
*/
|
|
|
|
STRINGIFY(
|
|
__kernel void UnsharpMaskBlurColumn(const __global CLQuantum* image,
|
|
const __global float4 *blurRowData,const unsigned int number_channels,
|
|
const ChannelType channel,const unsigned int columns,
|
|
const unsigned int rows,__local float4* cachedData,
|
|
__local float* cachedFilter,const __global float *filter,
|
|
const unsigned int width,const float gain, const float threshold,
|
|
__global CLQuantum *filteredImage)
|
|
{
|
|
const unsigned int radius = (width-1)/2;
|
|
|
|
// cache the pixel shared by the workgroup
|
|
const int groupX = get_group_id(0);
|
|
const int groupStartY = get_group_id(1)*get_local_size(1) - radius;
|
|
const int groupStopY = (get_group_id(1)+1)*get_local_size(1) + radius;
|
|
|
|
if ((groupStartY >= 0) && (groupStopY < rows))
|
|
{
|
|
event_t e = async_work_group_strided_copy(cachedData,
|
|
blurRowData+groupStartY*columns+groupX,groupStopY-groupStartY,columns,0);
|
|
wait_group_events(1,&e);
|
|
}
|
|
else
|
|
{
|
|
for (int i = get_local_id(1); i < (groupStopY - groupStartY); i+=get_local_size(1))
|
|
cachedData[i] = blurRowData[ClampToCanvas(groupStartY+i,rows)*columns + groupX];
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
}
|
|
// cache the filter as well
|
|
event_t e = async_work_group_copy(cachedFilter,filter,width,0);
|
|
wait_group_events(1,&e);
|
|
|
|
// only do the work if this is not a patched item
|
|
const int cy = get_global_id(1);
|
|
|
|
if (cy < rows)
|
|
{
|
|
float4 blurredPixel = (float4) 0.0f;
|
|
|
|
int i = 0;
|
|
|
|
for ( ; i+7 < width; )
|
|
{
|
|
for (int j=0; j < 8; j++, i++)
|
|
blurredPixel+=cachedFilter[i+j]*cachedData[i+j+get_local_id(1)];
|
|
}
|
|
|
|
for ( ; i < width; i++)
|
|
blurredPixel+=cachedFilter[i]*cachedData[i+get_local_id(1)];
|
|
|
|
float4 inputImagePixel = ReadFloat4(image,number_channels,columns,groupX,cy,channel);
|
|
float4 outputPixel = inputImagePixel - blurredPixel;
|
|
|
|
float quantumThreshold = QuantumRange*threshold;
|
|
|
|
int4 mask = isless(fabs(2.0f*outputPixel), (float4)quantumThreshold);
|
|
outputPixel = select(inputImagePixel + outputPixel * gain, inputImagePixel, mask);
|
|
|
|
//write back
|
|
WriteFloat4(filteredImage,number_channels,columns,groupX,cy,channel,outputPixel);
|
|
}
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
__kernel void UnsharpMask(const __global CLQuantum *image,const unsigned int number_channels,
|
|
const ChannelType channel,__constant float *filter,const unsigned int width,
|
|
const unsigned int columns,const unsigned int rows,__local float4 *pixels,
|
|
const float gain,const float threshold,__global CLQuantum *filteredImage)
|
|
{
|
|
const unsigned int x = get_global_id(0);
|
|
const unsigned int y = get_global_id(1);
|
|
|
|
const unsigned int radius = (width - 1) / 2;
|
|
|
|
int row = y - radius;
|
|
int baseRow = get_group_id(1) * get_local_size(1) - radius;
|
|
int endRow = (get_group_id(1) + 1) * get_local_size(1) + radius;
|
|
|
|
while (row < endRow) {
|
|
int srcy = (row < 0) ? -row : row; // mirror pad
|
|
srcy = (srcy >= rows) ? (2 * rows - srcy - 1) : srcy;
|
|
|
|
float4 value = 0.0f;
|
|
|
|
int ix = x - radius;
|
|
int i = 0;
|
|
|
|
while (i + 7 < width) {
|
|
for (int j = 0; j < 8; ++j) { // unrolled
|
|
int srcx = ix + j;
|
|
srcx = (srcx < 0) ? -srcx : srcx;
|
|
srcx = (srcx >= columns) ? (2 * columns - srcx - 1) : srcx;
|
|
value += filter[i + j] * ReadFloat4(image, number_channels, columns, srcx, srcy, channel);
|
|
}
|
|
ix += 8;
|
|
i += 8;
|
|
}
|
|
|
|
while (i < width) {
|
|
int srcx = (ix < 0) ? -ix : ix; // mirror pad
|
|
srcx = (srcx >= columns) ? (2 * columns - srcx - 1) : srcx;
|
|
value += filter[i] * ReadFloat4(image, number_channels, columns, srcx, srcy, channel);
|
|
++i;
|
|
++ix;
|
|
}
|
|
pixels[(row - baseRow) * get_local_size(0) + get_local_id(0)] = value;
|
|
row += get_local_size(1);
|
|
}
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
const int px = get_local_id(0);
|
|
const int py = get_local_id(1);
|
|
const int prp = get_local_size(0);
|
|
float4 value = (float4)(0.0f);
|
|
|
|
int i = 0;
|
|
while (i + 7 < width) {
|
|
for (int j = 0; j < 8; ++j) // unrolled
|
|
value += (float4)(filter[i]) * pixels[px + (py + i + j) * prp];
|
|
i += 8;
|
|
}
|
|
while (i < width) {
|
|
value += (float4)(filter[i]) * pixels[px + (py + i) * prp];
|
|
++i;
|
|
}
|
|
|
|
if ((x < columns) && (y < rows)) {
|
|
float4 srcPixel = ReadFloat4(image, number_channels, columns, x, y, channel);
|
|
float4 diff = srcPixel - value;
|
|
|
|
float quantumThreshold = QuantumRange*threshold;
|
|
|
|
int4 mask = isless(fabs(2.0f * diff), (float4)quantumThreshold);
|
|
value = select(srcPixel + diff * gain, srcPixel, mask);
|
|
|
|
WriteFloat4(filteredImage, number_channels, columns, x, y, channel, value);
|
|
}
|
|
}
|
|
)
|
|
|
|
STRINGIFY(
|
|
__kernel __attribute__((reqd_work_group_size(64, 4, 1)))
|
|
void WaveletDenoise(__global CLQuantum *srcImage,__global CLQuantum *dstImage,
|
|
const unsigned int number_channels,const unsigned int max_channels,
|
|
const float threshold,const int passes,const unsigned int imageWidth,
|
|
const unsigned int imageHeight)
|
|
{
|
|
const int pad = (1 << (passes - 1));
|
|
const int tileSize = 64;
|
|
const int tileRowPixels = 64;
|
|
const float noise[] = { 0.8002, 0.2735, 0.1202, 0.0585, 0.0291, 0.0152, 0.0080, 0.0044 };
|
|
|
|
CLQuantum stage[48]; // 16 * 3 (we only need 3 channels)
|
|
|
|
local float buffer[64 * 64];
|
|
|
|
int srcx = (get_group_id(0) + get_global_offset(0) / tileSize) * (tileSize - 2 * pad) - pad + get_local_id(0);
|
|
int srcy = (get_group_id(1) + get_global_offset(1) / 4) * (tileSize - 2 * pad) - pad;
|
|
|
|
for (int i = get_local_id(1); i < tileSize; i += get_local_size(1)) {
|
|
int pos = (mirrorTop(mirrorBottom(srcx), imageWidth) * number_channels) +
|
|
(mirrorTop(mirrorBottom(srcy + i), imageHeight)) * imageWidth * number_channels;
|
|
|
|
for (int channel = 0; channel < max_channels; ++channel)
|
|
stage[(i / 4) + (16 * channel)] = srcImage[pos + channel];
|
|
}
|
|
|
|
for (int channel = 0; channel < max_channels; ++channel) {
|
|
// Load LDS
|
|
for (int i = get_local_id(1); i < tileSize; i += get_local_size(1))
|
|
buffer[get_local_id(0) + i * tileRowPixels] = convert_float(stage[(i / 4) + (16 * channel)]);
|
|
|
|
// Process
|
|
|
|
float tmp[16];
|
|
float accum[16];
|
|
float pixel;
|
|
|
|
for (int i = 0; i < 16; i++)
|
|
accum[i]=0.0f;
|
|
|
|
for (int pass = 0; pass < passes; ++pass) {
|
|
const int radius = 1 << pass;
|
|
const int x = get_local_id(0);
|
|
const float thresh = threshold * noise[pass];
|
|
|
|
// Apply horizontal hat
|
|
for (int i = get_local_id(1); i < tileSize; i += get_local_size(1)) {
|
|
const int offset = i * tileRowPixels;
|
|
if (pass == 0)
|
|
tmp[i / 4] = buffer[x + offset]; // snapshot input on first pass
|
|
pixel = 0.5f * tmp[i / 4] + 0.25 * (buffer[mirrorBottom(x - radius) + offset] + buffer[mirrorTop(x + radius, tileSize) + offset]);
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
buffer[x + offset] = pixel;
|
|
}
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
// Apply vertical hat
|
|
for (int i = get_local_id(1); i < tileSize; i += get_local_size(1)) {
|
|
pixel = 0.5f * buffer[x + i * tileRowPixels] + 0.25 * (buffer[x + mirrorBottom(i - radius) * tileRowPixels] + buffer[x + mirrorTop(i + radius, tileRowPixels) * tileRowPixels]);
|
|
float delta = tmp[i / 4] - pixel;
|
|
tmp[i / 4] = pixel; // hold output in tmp until all workitems are done
|
|
if (delta < -thresh)
|
|
delta += thresh;
|
|
else if (delta > thresh)
|
|
delta -= thresh;
|
|
else
|
|
delta = 0;
|
|
accum[i / 4] += delta;
|
|
}
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
if (pass < passes - 1)
|
|
for (int i = get_local_id(1); i < tileSize; i += get_local_size(1))
|
|
buffer[x + i * tileRowPixels] = tmp[i / 4]; // store lowpass for next pass
|
|
else // last pass
|
|
for (int i = get_local_id(1); i < tileSize; i += get_local_size(1))
|
|
accum[i / 4] += tmp[i / 4]; // add the lowpass signal back to output
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
}
|
|
|
|
for (int i = get_local_id(1); i < tileSize; i += get_local_size(1))
|
|
stage[(i / 4) + (16 * channel)] = ClampToQuantum(accum[i / 4]);
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
}
|
|
|
|
// Write from stage to output
|
|
|
|
if ((get_local_id(0) >= pad) && (get_local_id(0) < tileSize - pad) && (srcx >= 0) && (srcx < imageWidth)) {
|
|
for (int i = get_local_id(1); i < tileSize; i += get_local_size(1)) {
|
|
if ((i >= pad) && (i < tileSize - pad) && (srcy + i >= 0) && (srcy + i < imageHeight)) {
|
|
int pos = (srcx * number_channels) + ((srcy + i) * (imageWidth * number_channels));
|
|
for (int channel = 0; channel < max_channels; ++channel) {
|
|
dstImage[pos + channel] = stage[(i / 4) + (16 * channel)];
|
|
}
|
|
}
|
|
}
|
|
}
|
|
}
|
|
)
|
|
|
|
;
|
|
|
|
#endif // MAGICKCORE_OPENCL_SUPPORT
|
|
|
|
#if defined(__cplusplus) || defined(c_plusplus)
|
|
}
|
|
#endif
|
|
|
|
#endif // MAGICKCORE_ACCELERATE_KERNELS_PRIVATE_H
|