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);
const int columns = imageColumns;
- const int rows = imageRows;
const unsigned int radius = (width-1)/2;
const int wsize = get_local_size(0);
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);
const int columns = imageColumns;
- const int rows = imageRows;
const unsigned int radius = (width-1)/2;
const int wsize = get_local_size(0);
)
STRINGIFY(
- __kernel __kernel __attribute__((reqd_work_group_size(256, 1, 1)))
+ __kernel __attribute__((reqd_work_group_size(256, 1, 1)))
void ResizeHorizontalFilter(const __global CLPixelType* inputImage, const unsigned int inputColumns, const unsigned int inputRows, const unsigned int matte
, const float xFactor, __global CLPixelType* filteredImage, const unsigned int filteredColumns, const unsigned int filteredRows
, const int resizeFilterType, const int resizeWindowType
STRINGIFY(
- __kernel __kernel __attribute__((reqd_work_group_size(1, 256, 1)))
+ __kernel __attribute__((reqd_work_group_size(1, 256, 1)))
void ResizeVerticalFilter(const __global CLPixelType* inputImage, const unsigned int inputColumns, const unsigned int inputRows, const unsigned int matte
, const float yFactor, __global CLPixelType* filteredImage, const unsigned int filteredColumns, const unsigned int filteredRows
, const int resizeFilterType, const int resizeWindowType
STRINGIFY(
- __kernel __kernel __attribute__((reqd_work_group_size(1, 256, 1)))
+ __kernel __attribute__((reqd_work_group_size(1, 256, 1)))
void ResizeVerticalFilterSinc(const __global CLPixelType* inputImage, const unsigned int inputColumns, const unsigned int inputRows, const unsigned int matte
, const float yFactor, __global CLPixelType* filteredImage, const unsigned int filteredColumns, const unsigned int filteredRows
, const int resizeFilterType, const int resizeWindowType
,inputImageCache,numCachedPixels,pixelPerWorkgroup,pixelChunkSize
,outputPixelCache,densityCache,gammaCache);
}
+ )
+
+ STRINGIFY(
+
+
+ __kernel void randomNumberGeneratorKernel(__global uint* seeds, const float normalizeRand
+ , __global float* randomNumbers, const uint init
+ ,const uint numRandomNumbers) {
+
+ unsigned int id = get_global_id(0);
+ unsigned int seed[4];
+
+ if (init!=0) {
+ seed[0] = seeds[id*4];
+ seed[1] = 0x50a7f451;
+ seed[2] = 0x5365417e;
+ seed[3] = 0xc3a4171a;
+ }
+ else {
+ seed[0] = seeds[id*4];
+ seed[1] = seeds[id*4+1];
+ seed[2] = seeds[id*4+2];
+ seed[3] = seeds[id*4+3];
+ }
+
+ unsigned int numRandomNumbersPerItem = (numRandomNumbers+get_global_size(0)-1)/get_global_size(0);
+ for (unsigned int i = 0; i < numRandomNumbersPerItem; i++) {
+ do
+ {
+ unsigned int alpha=(unsigned int) (seed[1] ^ (seed[1] << 11));
+ seed[1]=seed[2];
+ seed[2]=seed[3];
+ seed[3]=seed[0];
+ seed[0]=(seed[0] ^ (seed[0] >> 19)) ^ (alpha ^ (alpha >> 8));
+ } while (seed[0] == ~0UL);
+ unsigned int pos = (get_group_id(0)*get_local_size(0)*numRandomNumbersPerItem)
+ + get_local_size(0) * i + get_local_id(0);
+
+ if (pos >= numRandomNumbers)
+ break;
+ randomNumbers[pos] = normalizeRand*seed[0];
+ }
+
+ /* save the seeds for the time*/
+ seeds[id*4] = seed[0];
+ seeds[id*4+1] = seed[1];
+ seeds[id*4+2] = seed[2];
+ seeds[id*4+3] = seed[3];
+ }
+
+ )
+
+
+ STRINGIFY(
+
+ typedef enum
+ {
+ UndefinedNoise,
+ UniformNoise,
+ GaussianNoise,
+ MultiplicativeGaussianNoise,
+ ImpulseNoise,
+ LaplacianNoise,
+ PoissonNoise,
+ RandomNoise
+ } NoiseType;
+
+ typedef struct {
+ const global float* rns;
+ } RandomNumbers;
+
+
+ float GetPseudoRandomValue(RandomNumbers* r) {
+ float v = *r->rns;
+ r->rns++;
+ return v;
+ }
+ )
+
+ 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))
+
+ STRINGIFY(
+ float GenerateDifferentialNoise(RandomNumbers* r, CLQuantum pixel, NoiseType noise_type, float attenuate) {
+
+ float
+ alpha,
+ beta,
+ noise,
+ sigma;
+
+ noise = 0.0f;
+ alpha=GetPseudoRandomValue(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=GetPseudoRandomValue(r);
+ gamma=sqrt(-2.0f*log(alpha));
+ sigma=gamma*cospi((2.0f*beta));
+ tau=gamma*sinpi((2.0f*beta));
+ noise=(float)(pixel+sqrt((float) 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=(float)QuantumRange;
+ else
+ noise=(float)pixel;
+ break;
+ }
+ case LaplacianNoise:
+ {
+ if (alpha <= 0.5f)
+ {
+ if (alpha <= MagickEpsilon)
+ noise=(float) (pixel-QuantumRange);
+ else
+ noise=(float) (pixel+QuantumRange*SigmaLaplacian*log(2.0f*alpha)+
+ 0.5f);
+ break;
+ }
+ beta=1.0f-alpha;
+ if (beta <= (0.5f*MagickEpsilon))
+ noise=(float) (pixel+QuantumRange);
+ else
+ noise=(float) (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=GetPseudoRandomValue(r);
+ noise=(float) (pixel+pixel*SigmaMultiplicativeGaussian*sigma*
+ cospi((float) (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=GetPseudoRandomValue(r);
+ alpha*=beta;
+ }
+ noise=(float) (QuantumRange*i/SigmaPoisson);
+ break;
+ }
+ case RandomNoise:
+ {
+ noise=(float) (QuantumRange*SigmaRandom*alpha);
+ break;
+ }
+
+ };
+ return noise;
+ }
+
+ __kernel
+ void AddNoiseImage(const __global CLPixelType* inputImage, __global CLPixelType* filteredImage
+ ,const unsigned int inputColumns, const unsigned int inputRows
+ ,const ChannelType channel
+ ,const NoiseType noise_type, const float attenuate
+ ,const __global float* randomNumbers, const unsigned int numRandomNumbersPerPixel
+ ,const unsigned int rowOffset) {
+
+ unsigned int x = get_global_id(0);
+ unsigned int y = get_global_id(1) + rowOffset;
+ RandomNumbers r;
+ r.rns = randomNumbers + (get_global_id(1) * inputColumns + get_global_id(0))*numRandomNumbersPerPixel;
+
+ CLPixelType p = inputImage[y*inputColumns+x];
+ CLPixelType q = filteredImage[y*inputColumns+x];
+
+ if ((channel&RedChannel)!=0) {
+ setRed(&q,ClampToQuantum(GenerateDifferentialNoise(&r,getRed(p),noise_type,attenuate)));
+ }
+
+ if ((channel&GreenChannel)!=0) {
+ setGreen(&q,ClampToQuantum(GenerateDifferentialNoise(&r,getGreen(p),noise_type,attenuate)));
+ }
+
+ if ((channel&BlueChannel)!=0) {
+ setBlue(&q,ClampToQuantum(GenerateDifferentialNoise(&r,getBlue(p),noise_type,attenuate)));
+ }
+
+ if ((channel & OpacityChannel) != 0) {
+ setOpacity(&q,ClampToQuantum(GenerateDifferentialNoise(&r,getOpacity(p),noise_type,attenuate)));
+ }
+
+ filteredImage[y*inputColumns+x] = q;
+ }
+
)
;
+
+
#endif // MAGICKCORE_OPENCL_SUPPORT
#if defined(__cplusplus) || defined(c_plusplus)
#include "MagickCore/pixel-private.h"
#include "MagickCore/prepress.h"
#include "MagickCore/quantize.h"
+#include "MagickCore/random_.h"
+#include "MagickCore/random-private.h"
#include "MagickCore/registry.h"
#include "MagickCore/resize.h"
#include "MagickCore/resize-private.h"
#include "CLPerfMarker.h"
#endif
+#define MAGICK_MAX(x,y) (((x) >= (y))?(x):(y))
+#define MAGICK_MIN(x,y) (((x) <= (y))?(x):(y))
+
#if defined(MAGICKCORE_OPENCL_SUPPORT)
#define ALIGNED(pointer,type) ((((long)(pointer)) & (sizeof(type)-1)) == 0)
}
-
-static inline double MagickMax(const double x,const double y)
-{
- if (x > y)
- return(x);
- return(y);
-}
-
-static inline double MagickMin(const double x,const double y)
-{
- if (x < y)
- return(x);
- return(y);
-}
-
-
static MagickBooleanType resizeHorizontalFilter(cl_mem inputImage
, const unsigned int inputImageColumns, const unsigned int inputImageRows, const unsigned int matte
, cl_mem resizedImage, const unsigned int resizedColumns, const unsigned int resizedRows
/*
Apply filter to resize vertically from image to resize image.
*/
- scale=MagickMax(1.0/xFactor+MagickEpsilon,1.0);
+ scale=MAGICK_MAX(1.0/xFactor+MagickEpsilon,1.0);
support=scale*GetResizeFilterSupport(resizeFilter);
if (support < 0.5)
{
/*
Apply filter to resize vertically from image to resize image.
*/
- scale=MagickMax(1.0/yFactor+MagickEpsilon,1.0);
+ scale=MAGICK_MAX(1.0/yFactor+MagickEpsilon,1.0);
support=scale*GetResizeFilterSupport(resizeFilter);
if (support < 0.5)
{
return newImage;
}
+static Image* ComputeAddNoiseImage(const Image* inputImage,
+ const ChannelType channel, const NoiseType noise_type,
+ ExceptionInfo *exception)
+{
+ MagickBooleanType outputReady = MagickFalse;
+ MagickCLEnv clEnv = NULL;
+
+ cl_int clStatus;
+ size_t global_work_size[2];
+
+ const void *inputPixels = NULL;
+ Image* filteredImage = NULL;
+ void *filteredPixels = NULL;
+ void *hostPtr;
+ unsigned int inputColumns, inputRows;
+ float attenuate;
+ float *randomNumberBufferPtr = NULL;
+ MagickSizeType length;
+ unsigned int numRandomNumberPerPixel;
+ unsigned int numRowsPerKernelLaunch;
+ unsigned int numRandomNumberPerBuffer;
+ unsigned int r;
+ unsigned int k;
+ int i;
+
+ RandomInfo **restrict random_info;
+ const char *option;
+#if defined(MAGICKCORE_OPENMP_SUPPORT)
+ unsigned long key;
+#endif
+
+ cl_mem_flags mem_flags;
+ cl_context context = NULL;
+ cl_mem inputImageBuffer = NULL;
+ cl_mem randomNumberBuffer = NULL;
+ cl_mem filteredImageBuffer = NULL;
+ cl_command_queue queue = NULL;
+ cl_kernel addNoiseKernel = NULL;
+
+
+ clEnv = GetDefaultOpenCLEnv();
+ context = GetOpenCLContext(clEnv);
+ queue = AcquireOpenCLCommandQueue(clEnv);
+
+ inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
+ if (inputPixels == (void *) NULL)
+ {
+ (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
+ goto cleanup;
+ }
+
+ if (ALIGNED(inputPixels,CLPixelPacket))
+ {
+ mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
+ }
+ else
+ {
+ mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
+ }
+ /* create a CL buffer from image pixel buffer */
+ length = inputImage->columns * inputImage->rows;
+ inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
+ goto cleanup;
+ }
+
+
+ filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
+ assert(filteredImage != NULL);
+ if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
+ goto cleanup;
+ }
+ filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
+ if (filteredPixels == (void *) NULL)
+ {
+ (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
+ goto cleanup;
+ }
+
+ if (ALIGNED(filteredPixels,CLPixelPacket))
+ {
+ mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
+ hostPtr = filteredPixels;
+ }
+ else
+ {
+ mem_flags = CL_MEM_WRITE_ONLY;
+ hostPtr = NULL;
+ }
+ /* create a CL buffer from image pixel buffer */
+ length = inputImage->columns * inputImage->rows;
+ filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
+ goto cleanup;
+ }
+
+ /* find out how many random numbers needed by pixel */
+ numRandomNumberPerPixel = 0;
+ {
+ unsigned int numRandPerChannel = 0;
+ switch (noise_type)
+ {
+ case UniformNoise:
+ case ImpulseNoise:
+ case LaplacianNoise:
+ case RandomNoise:
+ default:
+ numRandPerChannel = 1;
+ break;
+ case GaussianNoise:
+ case MultiplicativeGaussianNoise:
+ case PoissonNoise:
+ numRandPerChannel = 2;
+ break;
+ };
+
+ if ((channel & RedChannel) != 0)
+ numRandomNumberPerPixel+=numRandPerChannel;
+ if ((channel & GreenChannel) != 0)
+ numRandomNumberPerPixel+=numRandPerChannel;
+ if ((channel & BlueChannel) != 0)
+ numRandomNumberPerPixel+=numRandPerChannel;
+ if ((channel & OpacityChannel) != 0)
+ numRandomNumberPerPixel+=numRandPerChannel;
+ }
+
+ numRowsPerKernelLaunch = 512;
+ /* create a buffer for random numbers */
+ numRandomNumberPerBuffer = (inputImage->columns*numRowsPerKernelLaunch)*numRandomNumberPerPixel;
+ randomNumberBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, numRandomNumberPerBuffer*sizeof(float)
+ , NULL, &clStatus);
+
+
+ /* set up the random number generators */
+ attenuate=1.0;
+ option=GetImageArtifact(inputImage,"attenuate");
+ if (option != (char *) NULL)
+ attenuate=StringToDouble(option,(char **) NULL);
+ random_info=AcquireRandomInfoThreadSet();
+#if defined(MAGICKCORE_OPENMP_SUPPORT)
+ key=GetRandomSecretKey(random_info[0]);
+#endif
+
+ addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"AddNoiseImage");
+
+ k = 0;
+ clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&inputImageBuffer);
+ clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer);
+ inputColumns = inputImage->columns;
+ clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputColumns);
+ inputRows = inputImage->rows;
+ clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputRows);
+ clSetKernelArg(addNoiseKernel,k++,sizeof(ChannelType),(void *)&channel);
+ clSetKernelArg(addNoiseKernel,k++,sizeof(NoiseType),(void *)&noise_type);
+ attenuate=1.0f;
+ option=GetImageArtifact(inputImage,"attenuate");
+ if (option != (char *) NULL)
+ attenuate=(float)StringToDouble(option,(char **) NULL);
+ clSetKernelArg(addNoiseKernel,k++,sizeof(float),(void *)&attenuate);
+ clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&randomNumberBuffer);
+ clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerPixel);
+
+ global_work_size[0] = inputColumns;
+ for (r = 0; r < inputRows; r+=numRowsPerKernelLaunch)
+ {
+ /* Generate random numbers in the buffer */
+ randomNumberBufferPtr = (float*)clEnqueueMapBuffer(queue, randomNumberBuffer, CL_TRUE, CL_MAP_WRITE, 0
+ , numRandomNumberPerBuffer*sizeof(float), 0, NULL, NULL, &clStatus);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.",".");
+ goto cleanup;
+ }
+
+#if defined(MAGICKCORE_OPENMP_SUPPORT)
+ #pragma omp parallel for schedule(static,4) \
+ num_threads((key == ~0UL) == 0 ? 1 : (size_t) GetMagickResourceLimit(ThreadResource))
+#endif
+ for (i = 0; i < numRandomNumberPerBuffer; i++)
+ {
+ const int id = GetOpenMPThreadId();
+ randomNumberBufferPtr[i] = (float)GetPseudoRandomValue(random_info[id]);
+ }
+
+ clStatus = clEnqueueUnmapMemObject(queue, randomNumberBuffer, randomNumberBufferPtr, 0, NULL, NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.",".");
+ goto cleanup;
+ }
+
+ /* set the row offset */
+ clSetKernelArg(addNoiseKernel,k,sizeof(unsigned int),(void *)&r);
+ global_work_size[1] = MAGICK_MIN(numRowsPerKernelLaunch, inputRows - r);
+ clEnqueueNDRangeKernel(queue,addNoiseKernel,2,NULL,global_work_size,NULL,0,NULL,NULL);
+ }
+
+ if (ALIGNED(filteredPixels,CLPixelPacket))
+ {
+ length = inputImage->columns * inputImage->rows;
+ clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
+ }
+ else
+ {
+ length = inputImage->columns * inputImage->rows;
+ clStatus = clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
+ }
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", ".");
+ goto cleanup;
+ }
+
+
+ outputReady = MagickTrue;
+cleanup:
+ if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
+ if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel);
+ if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer);
+ if (randomNumberBuffer!=NULL) clReleaseMemObject(randomNumberBuffer);
+ if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer);
+ if (outputReady == MagickFalse
+ && filteredImage != NULL)
+ {
+ DestroyImage(filteredImage);
+ filteredImage = NULL;
+ }
+ return filteredImage;
+}
+
+
+static Image* ComputeAddNoiseImageOptRandomNum(const Image* inputImage,
+ const ChannelType channel, const NoiseType noise_type,
+ ExceptionInfo *exception)
+{
+ MagickBooleanType outputReady = MagickFalse;
+ MagickCLEnv clEnv = NULL;
+
+ cl_int clStatus;
+ size_t global_work_size[2];
+ size_t random_work_size;
+
+ const void *inputPixels = NULL;
+ Image* filteredImage = NULL;
+ void *filteredPixels = NULL;
+ void *hostPtr;
+ unsigned int inputColumns, inputRows;
+ float attenuate;
+ MagickSizeType length;
+ unsigned int numRandomNumberPerPixel;
+ unsigned int numRowsPerKernelLaunch;
+ unsigned int numRandomNumberPerBuffer;
+ unsigned int numRandomNumberGenerators;
+ unsigned int initRandom;
+ float fNormalize;
+ unsigned int r;
+ unsigned int k;
+ int i;
+ const char *option;
+
+ cl_mem_flags mem_flags;
+ cl_context context = NULL;
+ cl_mem inputImageBuffer = NULL;
+ cl_mem randomNumberBuffer = NULL;
+ cl_mem filteredImageBuffer = NULL;
+ cl_mem randomNumberSeedsBuffer = NULL;
+ cl_command_queue queue = NULL;
+ cl_kernel addNoiseKernel = NULL;
+ cl_kernel randomNumberGeneratorKernel = NULL;
+
+
+ clEnv = GetDefaultOpenCLEnv();
+ context = GetOpenCLContext(clEnv);
+ queue = AcquireOpenCLCommandQueue(clEnv);
+
+ inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
+ if (inputPixels == (void *) NULL)
+ {
+ (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
+ goto cleanup;
+ }
+
+ if (ALIGNED(inputPixels,CLPixelPacket))
+ {
+ mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
+ }
+ else
+ {
+ mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
+ }
+ /* create a CL buffer from image pixel buffer */
+ length = inputImage->columns * inputImage->rows;
+ inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
+ goto cleanup;
+ }
+
+
+ filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
+ assert(filteredImage != NULL);
+ if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
+ goto cleanup;
+ }
+ filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
+ if (filteredPixels == (void *) NULL)
+ {
+ (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
+ goto cleanup;
+ }
+
+ if (ALIGNED(filteredPixels,CLPixelPacket))
+ {
+ mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
+ hostPtr = filteredPixels;
+ }
+ else
+ {
+ mem_flags = CL_MEM_WRITE_ONLY;
+ hostPtr = NULL;
+ }
+ /* create a CL buffer from image pixel buffer */
+ length = inputImage->columns * inputImage->rows;
+ filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
+ goto cleanup;
+ }
+
+ /* find out how many random numbers needed by pixel */
+ numRandomNumberPerPixel = 0;
+ {
+ unsigned int numRandPerChannel = 0;
+ switch (noise_type)
+ {
+ case UniformNoise:
+ case ImpulseNoise:
+ case LaplacianNoise:
+ case RandomNoise:
+ default:
+ numRandPerChannel = 1;
+ break;
+ case GaussianNoise:
+ case MultiplicativeGaussianNoise:
+ case PoissonNoise:
+ numRandPerChannel = 2;
+ break;
+ };
+
+ if ((channel & RedChannel) != 0)
+ numRandomNumberPerPixel+=numRandPerChannel;
+ if ((channel & GreenChannel) != 0)
+ numRandomNumberPerPixel+=numRandPerChannel;
+ if ((channel & BlueChannel) != 0)
+ numRandomNumberPerPixel+=numRandPerChannel;
+ if ((channel & OpacityChannel) != 0)
+ numRandomNumberPerPixel+=numRandPerChannel;
+ }
+
+ numRowsPerKernelLaunch = 512;
+
+ /* create a buffer for random numbers */
+ numRandomNumberPerBuffer = (inputImage->columns*numRowsPerKernelLaunch)*numRandomNumberPerPixel;
+ randomNumberBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, numRandomNumberPerBuffer*sizeof(float)
+ , NULL, &clStatus);
+
+ {
+ /* setup the random number generators */
+ unsigned long* seeds;
+ numRandomNumberGenerators = 512;
+ randomNumberSeedsBuffer = clCreateBuffer(context, CL_MEM_ALLOC_HOST_PTR|CL_MEM_READ_WRITE
+ , numRandomNumberGenerators * 4 * sizeof(unsigned long), NULL, &clStatus);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
+ goto cleanup;
+ }
+ seeds = (unsigned long*) clEnqueueMapBuffer(queue, randomNumberSeedsBuffer, CL_TRUE, CL_MAP_WRITE, 0
+ , numRandomNumberGenerators*4*sizeof(unsigned long), 0, NULL, NULL, &clStatus);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.",".");
+ goto cleanup;
+ }
+
+ for (i = 0; i < numRandomNumberGenerators; i++) {
+ RandomInfo* randomInfo = AcquireRandomInfo();
+ const unsigned long* s = GetRandomInfoSeed(randomInfo);
+
+ if (i == 0)
+ fNormalize = GetRandomInfoNormalize(randomInfo);
+
+ seeds[i*4] = s[0];
+ randomInfo = DestroyRandomInfo(randomInfo);
+ }
+
+ clStatus = clEnqueueUnmapMemObject(queue, randomNumberSeedsBuffer, seeds, 0, NULL, NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.",".");
+ goto cleanup;
+ }
+
+ randomNumberGeneratorKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE
+ ,"randomNumberGeneratorKernel");
+
+ k = 0;
+ clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(cl_mem),(void *)&randomNumberSeedsBuffer);
+ clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(float),(void *)&fNormalize);
+ clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(cl_mem),(void *)&randomNumberBuffer);
+ initRandom = 1;
+ clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(unsigned int),(void *)&initRandom);
+ clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerBuffer);
+
+ random_work_size = numRandomNumberGenerators;
+ }
+
+ addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"AddNoiseImage");
+ k = 0;
+ clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&inputImageBuffer);
+ clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer);
+ inputColumns = inputImage->columns;
+ clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputColumns);
+ inputRows = inputImage->rows;
+ clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputRows);
+ clSetKernelArg(addNoiseKernel,k++,sizeof(ChannelType),(void *)&channel);
+ clSetKernelArg(addNoiseKernel,k++,sizeof(NoiseType),(void *)&noise_type);
+ attenuate=1.0f;
+ option=GetImageArtifact(inputImage,"attenuate");
+ if (option != (char *) NULL)
+ attenuate=(float)StringToDouble(option,(char **) NULL);
+ clSetKernelArg(addNoiseKernel,k++,sizeof(float),(void *)&attenuate);
+ clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&randomNumberBuffer);
+ clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerPixel);
+
+ global_work_size[0] = inputColumns;
+ for (r = 0; r < inputRows; r+=numRowsPerKernelLaunch)
+ {
+ size_t generator_local_size = 64;
+ /* Generate random numbers in the buffer */
+ clEnqueueNDRangeKernel(queue,randomNumberGeneratorKernel,1,NULL
+ ,&random_work_size,&generator_local_size,0,NULL,NULL);
+ if (initRandom != 0)
+ {
+ /* make sure we only do init once */
+ initRandom = 0;
+ clSetKernelArg(randomNumberGeneratorKernel,3,sizeof(unsigned int),(void *)&initRandom);
+ }
+
+ /* set the row offset */
+ clSetKernelArg(addNoiseKernel,k,sizeof(unsigned int),(void *)&r);
+ global_work_size[1] = MAGICK_MIN(numRowsPerKernelLaunch, inputRows - r);
+ clEnqueueNDRangeKernel(queue,addNoiseKernel,2,NULL,global_work_size,NULL,0,NULL,NULL);
+ }
+
+ if (ALIGNED(filteredPixels,CLPixelPacket))
+ {
+ length = inputImage->columns * inputImage->rows;
+ clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
+ }
+ else
+ {
+ length = inputImage->columns * inputImage->rows;
+ clStatus = clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
+ }
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", ".");
+ goto cleanup;
+ }
+
+
+ outputReady = MagickTrue;
+cleanup:
+ if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
+ if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel);
+ if (randomNumberGeneratorKernel!=NULL) RelinquishOpenCLKernel(clEnv, randomNumberGeneratorKernel);
+ if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer);
+ if (randomNumberBuffer!=NULL) clReleaseMemObject(randomNumberBuffer);
+ if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer);
+ if (randomNumberSeedsBuffer!=NULL) clReleaseMemObject(randomNumberSeedsBuffer);
+ if (outputReady == MagickFalse
+ && filteredImage != NULL)
+ {
+ DestroyImage(filteredImage);
+ filteredImage = NULL;
+ }
+ return filteredImage;
+}
+
+
+
+MagickExport
+Image* AccelerateAddNoiseImage(const Image *image, const ChannelType channel,
+ const NoiseType noise_type,ExceptionInfo *exception)
+{
+ MagickBooleanType status;
+ Image* filteredImage = NULL;
+
+ assert(image != NULL);
+ assert(exception != NULL);
+
+ status = checkOpenCLEnvironment(exception);
+ if (status == MagickFalse)
+ return NULL;
+
+ status = checkAccelerateCondition(image, channel, exception);
+ if (status == MagickFalse)
+ return NULL;
+
+ if (sizeof(unsigned long) == 4)
+ filteredImage = ComputeAddNoiseImageOptRandomNum(image,channel,noise_type,exception);
+ else
+ filteredImage = ComputeAddNoiseImage(image,channel,noise_type,exception);
+
+ OpenCLLogException(__FUNCTION__,__LINE__,exception);
+ return filteredImage;
+}
+
+
#else /* MAGICKCORE_OPENCL_SUPPORT */
MagickExport Image *AccelerateConvolveImageChannel(
return(MagickFalse);
}
+MagickExport Image *AccelerateAddNoiseImage(const Image *image,
+ const ChannelType channel, const NoiseType noise_type,ExceptionInfo *exception)
+{
+ magick_unreferenced(image);
+ magick_unreferenced(channel);
+ magick_unreferenced(noise_type);
+ magick_unreferenced(exception);
+ return NULL;
+}
#endif /* MAGICKCORE_OPENCL_SUPPORT */