From: cristy Date: Wed, 27 Nov 2013 02:25:43 +0000 (+0000) Subject: (no commit message) X-Git-Tag: 7.0.1-0~3091 X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=e85d0f75864c9410b905ac836b58f48db47303cc;p=imagemagick --- diff --git a/MagickCore/accelerate-private.h b/MagickCore/accelerate-private.h index 772898d1a..56df71896 100644 --- a/MagickCore/accelerate-private.h +++ b/MagickCore/accelerate-private.h @@ -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) diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c index 6b5de7bc5..235de0f1c 100644 --- a/MagickCore/accelerate.c +++ b/MagickCore/accelerate.c @@ -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 */ diff --git a/MagickCore/accelerate.h b/MagickCore/accelerate.h index a6daae68f..e5199d218 100644 --- a/MagickCore/accelerate.h +++ b/MagickCore/accelerate.h @@ -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, diff --git a/MagickCore/opencl-private.h b/MagickCore/opencl-private.h index 4075d4e54..b58597616 100644 --- a/MagickCore/opencl-private.h +++ b/MagickCore/opencl-private.h @@ -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) diff --git a/MagickCore/opencl.c b/MagickCore/opencl.c index 4fad09da3..963172d2f 100644 --- a/MagickCore/opencl.c +++ b/MagickCore/opencl.c @@ -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 } diff --git a/MagickCore/random-private.h b/MagickCore/random-private.h index 8af2ea589..84f682c67 100644 --- a/MagickCore/random-private.h +++ b/MagickCore/random-private.h @@ -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); diff --git a/MagickCore/random.c b/MagickCore/random.c index 49a25faca..25b6fdc74 100644 --- a/MagickCore/random.c +++ b/MagickCore/random.c @@ -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; +}