]> granicus.if.org Git - imagemagick/commitdiff
(no commit message)
authorcristy <urban-warrior@git.imagemagick.org>
Wed, 27 Nov 2013 02:25:43 +0000 (02:25 +0000)
committercristy <urban-warrior@git.imagemagick.org>
Wed, 27 Nov 2013 02:25:43 +0000 (02:25 +0000)
MagickCore/accelerate-private.h
MagickCore/accelerate.c
MagickCore/accelerate.h
MagickCore/opencl-private.h
MagickCore/opencl.c
MagickCore/random-private.h
MagickCore/random.c

index 772898d1ae338660528b276e66ac409b4e626ce3..56df718965e77ca9b8c0f093523a0c151b8c7cac 100644 (file)
@@ -506,10 +506,7 @@ const char* accelerateKernels =
         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);  
@@ -614,10 +611,7 @@ const char* accelerateKernels =
         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);  
@@ -1784,7 +1778,7 @@ const char* accelerateKernels =
   )
 
   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
@@ -1977,7 +1971,7 @@ const char* accelerateKernels =
 
 
   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
@@ -2148,7 +2142,7 @@ const char* accelerateKernels =
 
 
   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
@@ -2164,10 +2158,234 @@ const char* accelerateKernels =
       ,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)
index 6b5de7bc5c46273357fb54069b381e65d4cb984e..235de0f1cc1beb91b5271080d08ff2c12f441eec 100644 (file)
@@ -65,6 +65,8 @@ Include declarations.
 #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"
@@ -79,6 +81,9 @@ Include declarations.
 #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)
@@ -2391,22 +2396,6 @@ Image* AccelerateUnsharpMaskImage(const Image *image, const ChannelType channel,
 
 }
 
-
-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
@@ -2436,7 +2425,7 @@ static MagickBooleanType resizeHorizontalFilter(cl_mem inputImage
   /*
   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)
   {
@@ -2616,7 +2605,7 @@ static MagickBooleanType resizeVerticalFilter(cl_mem inputImage
   /*
   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)
   {
@@ -4304,6 +4293,536 @@ Image* AccelerateDespeckleImage(const Image* image, ExceptionInfo* exception)
   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(
@@ -4441,6 +4960,15 @@ MagickBooleanType AccelerateModulateImage(
   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 */
 
index a6daae68fc43d4b80f75a853c8a14439b5bffe67..e5199d218bead472a850481d7b15b75256584d49 100644 (file)
@@ -22,6 +22,7 @@
 extern "C" {
 #endif
 
+#include "MagickCore/fx.h"
 #include "MagickCore/morphology.h"
 #include "MagickCore/resample.h"
 #include "MagickCore/resize.h"
@@ -34,12 +35,13 @@ extern MagickExport MagickBooleanType
   AccelerateEqualizeImage(Image *,const ChannelType,ExceptionInfo *),
   AccelerateFunctionImage(Image *,const ChannelType,const MagickFunction,
     const size_t,const double *,ExceptionInfo *),
-  AccelerateModulateImage(Image* image, double percent_brightness, 
-    double percent_hue, double percent_saturation, 
-    ColorspaceType colorspace, ExceptionInfo* exception);
+  AccelerateModulateImage(Image*, double, double, double, 
+    ColorspaceType, ExceptionInfo*);
 
 
 extern MagickExport Image
+  *AccelerateAddNoiseImage(const Image*,const ChannelType,const NoiseType,
+    ExceptionInfo *),
   *AccelerateBlurImage(const Image *,const ChannelType,const double,
     const double,ExceptionInfo *),
   *AccelerateConvolveImageChannel(const Image *,const ChannelType,
index 4075d4e542ccad148a610929b29320e6b85d10fb..b5859761648d2446206db77e144a6dc739ac2360 100644 (file)
@@ -22,6 +22,7 @@ MagickCore OpenCL private methods.
 Include declarations.
 */
 #include "MagickCore/studio.h"
+#include "MagickCore/opencl.h"
 
 #if defined(__cplusplus) || defined(c_plusplus)
 extern "C" {
@@ -98,19 +99,23 @@ extern MagickExport const char*
 extern MagickExport void 
   OpenCLLog(const char*);
 
-/* #define ACCELERATE_LOG_EXCEPTION 1 */
+/* #define OPENCLLOG_ENABLED 1 */
 static inline void OpenCLLogException(const char* function, 
                         const unsigned int line, 
                         ExceptionInfo* exception) {
+#ifdef OPENCLLOG_ENABLED
   if (exception->severity!=0) {
     char message[MaxTextExtent];
     /*  dump the source into a file */
     (void) FormatLocaleString(message,MaxTextExtent,"%s:%d Exception(%d)"
       ,function,line,exception->severity);
-#ifdef ACCELERATE_LOG_EXCEPTION
     OpenCLLog(message);
-#endif
   }
+#else
+  magick_unreferenced(function);
+  magick_unreferenced(line);
+  magick_unreferenced(exception);
+#endif
 }
 
 #if defined(__cplusplus) || defined(c_plusplus)
index 4fad09da30862dbc541f94a978de3c75821d5876..963172d2f624890186625416630604fe81e401a1 100644 (file)
@@ -457,10 +457,22 @@ cl_context GetOpenCLContext(MagickCLEnv clEnv) {
 static char* getBinaryCLProgramName(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature)
 {
   char* name;
+  char* ptr;
   char path[MaxTextExtent];
   char deviceName[MaxTextExtent];
   const char* prefix = "magick_opencl";
   clGetDeviceInfo(clEnv->device, CL_DEVICE_NAME, MaxTextExtent, deviceName, NULL);
+  ptr=deviceName;
+  /* strip out illegal characters for file names */
+  while (*ptr != '\0')
+  {
+    if ( *ptr == ' ' || *ptr == '\\' || *ptr == '/' || *ptr == ':' || *ptr == '*' 
+        || *ptr == '?' || *ptr == '"' || *ptr == '<' || *ptr == '>' || *ptr == '|')
+    {
+      *ptr = '_';
+    }
+    ptr++;
+  }
   (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s_%s_%02d_%08x.bin"
          ,GetOpenCLCachedFilesDirectory()
          ,DirectorySeparator,prefix,deviceName, (unsigned int)prog, signature);
@@ -2426,27 +2438,33 @@ const char* GetOpenCLCachedFilesDirectory() {
   return openclCachedFilesDirectory;
 }
 
-
-/* create a loggin function */
+/* create a function for OpenCL log */
 MagickExport
 void OpenCLLog(const char* message) {
 
+#ifdef OPENCLLOG_ENABLED
 #define OPENCL_LOG_FILE "ImageMagickOpenCL.log"
 
   FILE* log;
-  if (message) {
-    char path[MaxTextExtent];
+  if (getenv("MAGICK_OCL_LOG"))
+  {
+    if (message) {
+      char path[MaxTextExtent];
 
-    /*  dump the source into a file */
-    (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
-      ,GetOpenCLCachedFilesDirectory()
-      ,DirectorySeparator,OPENCL_LOG_FILE);
+      /*  dump the source into a file */
+      (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
+        ,GetOpenCLCachedFilesDirectory()
+        ,DirectorySeparator,OPENCL_LOG_FILE);
 
 
-    log = fopen(path, "ab");
-    fwrite(message, sizeof(char), strlen(message), log);
-    fwrite("\n", sizeof(char), 1, log);
-    fclose(log);
+      log = fopen(path, "ab");
+      fwrite(message, sizeof(char), strlen(message), log);
+      fwrite("\n", sizeof(char), 1, log);
+      fclose(log);
+    }
   }
+#else
+  magick_unreferenced(message);
+#endif
 }
 
index 8af2ea58963bc4def5633723aeb6b6e9884825b0..84f682c671f07491df254d1e7f5f928979c721d2 100644 (file)
@@ -24,9 +24,15 @@ extern "C" {
 
 #include "MagickCore/thread-private.h"
 
+extern MagickPrivate double
+  GetRandomInfoNormalize(RandomInfo *);
+
 extern MagickPrivate MagickBooleanType
   RandomComponentGenesis(void);
 
+extern MagickPrivate unsigned long
+  *GetRandomInfoSeed(RandomInfo *);
+
 extern MagickPrivate void
   RandomComponentTerminus(void);
 
index 49a25faca3b72cf541d3f2c4f4a1998e520a4a79..25b6fdc742a7830393fb4c6915195a84715aabf2 100644 (file)
@@ -922,3 +922,15 @@ MagickExport void SetRandomTrueRandom(const MagickBooleanType true_random)
 {
   gather_true_random=true_random;
 }
+
+MagickPrivate unsigned long *GetRandomInfoSeed(RandomInfo *random_info)
+{
+  assert(random_info != (RandomInfo *) NULL);
+  return random_info->seed;
+}
+
+MagickPrivate double GetRandomInfoNormalize(RandomInfo *random_info)
+{
+  assert(random_info != (RandomInfo *) NULL);
+  return random_info->normalize;
+}