From a14003ee00fae6a256442e5846f721d6e2a5d262 Mon Sep 17 00:00:00 2001 From: dirk Date: Fri, 17 Jun 2016 21:20:30 +0200 Subject: [PATCH] Removed AccelerateCompositeImage. --- MagickCore/accelerate-kernels-private.h | 183 ---------------- MagickCore/accelerate-private.h | 2 - MagickCore/accelerate.c | 266 ------------------------ 3 files changed, 451 deletions(-) diff --git a/MagickCore/accelerate-kernels-private.h b/MagickCore/accelerate-kernels-private.h index c6dbfc094..effb440e5 100644 --- a/MagickCore/accelerate-kernels-private.h +++ b/MagickCore/accelerate-kernels-private.h @@ -974,189 +974,6 @@ OPENCL_ENDIF() } ) -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -% C o m p o s i t e % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -*/ - - STRINGIFY( - inline float ColorDodge(const float Sca, - const float Sa,const float Dca,const float Da) - { - /* - Oct 2004 SVG specification. - */ - if ((Sca*Da+Dca*Sa) >= Sa*Da) - return(Sa*Da+Sca*(1.0-Da)+Dca*(1.0-Sa)); - return(Dca*Sa*Sa/(Sa-Sca)+Sca*(1.0-Da)+Dca*(1.0-Sa)); - - - /* - New specification, March 2009 SVG specification. This specification was - also wrong of non-overlap cases. - */ - /* - if ((fabs(Sca-Sa) < MagickEpsilon) && (fabs(Dca) < MagickEpsilon)) - return(Sca*(1.0-Da)); - if (fabs(Sca-Sa) < MagickEpsilon) - return(Sa*Da+Sca*(1.0-Da)+Dca*(1.0-Sa)); - return(Sa*MagickMin(Da,Dca*Sa/(Sa-Sca))); - */ - - /* - Working from first principles using the original formula: - - f(Sc,Dc) = Dc/(1-Sc) - - This works correctly! Looks like the 2004 model was right but just - required a extra condition for correct handling. - */ - - /* - if ((fabs(Sca-Sa) < MagickEpsilon) && (fabs(Dca) < MagickEpsilon)) - return(Sca*(1.0-Da)+Dca*(1.0-Sa)); - if (fabs(Sca-Sa) < MagickEpsilon) - return(Sa*Da+Sca*(1.0-Da)+Dca*(1.0-Sa)); - return(Dca*Sa*Sa/(Sa-Sca)+Sca*(1.0-Da)+Dca*(1.0-Sa)); - */ - } - - inline void CompositeColorDodge(const float4 *p, - const float4 *q,float4 *composite) { - - float - Da, - gamma, - Sa; - - Sa=QuantumScale*getAlphaF4(*p); /* simplify and speed up equations */ - Da=QuantumScale*getAlphaF4(*q); - gamma=RoundToUnity(Sa+Da-Sa*Da); /* over blend, as per SVG doc */ - setAlphaF4(composite,QuantumRange*gamma); - gamma=QuantumRange/(fabs(gamma) < MagickEpsilon ? MagickEpsilon : gamma); - setRedF4(composite,gamma*ColorDodge(QuantumScale*getRedF4(*p)*Sa,Sa,QuantumScale* - getRedF4(*q)*Da,Da)); - setGreenF4(composite,gamma*ColorDodge(QuantumScale*getGreenF4(*p)*Sa,Sa,QuantumScale* - getGreenF4(*q)*Da,Da)); - setBlueF4(composite,gamma*ColorDodge(QuantumScale*getBlueF4(*p)*Sa,Sa,QuantumScale* - getBlueF4(*q)*Da,Da)); - } - ) - - STRINGIFY( - inline void MagickPixelCompositePlus(const float4 *p, - const float alpha,const float4 *q, - const float beta,float4 *composite) - { - float - gamma; - - float - Da, - Sa; - /* - Add two pixels with the given opacities. - */ - Sa=QuantumScale*alpha; - Da=QuantumScale*beta; - gamma=RoundToUnity(Sa+Da); /* 'Plus' blending -- not 'Over' blending */ - setAlphaF4(composite,QuantumRange*gamma); - gamma=PerceptibleReciprocal(gamma); - setRedF4(composite,gamma*(Sa*getRedF4(*p)+Da*getRedF4(*q))); - setGreenF4(composite,gamma*(Sa*getGreenF4(*p)+Da*getGreenF4(*q))); - setBlueF4(composite,gamma*(Sa*getBlueF4(*p)+Da*getBlueF4(*q))); - } - ) - - STRINGIFY( - inline void MagickPixelCompositeBlend(const float4 *p, - const float alpha,const float4 *q, - const float beta,float4 *composite) - { - MagickPixelCompositePlus(p,(float) (alpha* - (getAlphaF4(*p))),q,(float) (beta* - (getAlphaF4(*q))),composite); - } - ) - - STRINGIFY( - __kernel - void Composite(__global CLPixelType *image, - const unsigned int imageWidth, - const unsigned int imageHeight, - const __global CLPixelType *compositeImage, - const unsigned int compositeWidth, - const unsigned int compositeHeight, - const unsigned int compose, - const ChannelType channel, - const unsigned int matte, - const float destination_dissolve, - const float source_dissolve) { - - uint2 index; - index.x = get_global_id(0); - index.y = get_global_id(1); - - - if (index.x >= imageWidth - || index.y >= imageHeight) { - return; - } - const CLPixelType inputPixel = image[index.y*imageWidth+index.x]; - float4 destination; - setRedF4(&destination,getRed(inputPixel)); - setGreenF4(&destination,getGreen(inputPixel)); - setBlueF4(&destination,getBlue(inputPixel)); - - - const CLPixelType compositePixel - = compositeImage[index.y*imageWidth+index.x]; - float4 source; - setRedF4(&source,getRed(compositePixel)); - setGreenF4(&source,getGreen(compositePixel)); - setBlueF4(&source,getBlue(compositePixel)); - - if (matte != 0) { - setAlphaF4(&destination,getAlpha(inputPixel)); - setAlphaF4(&source,getAlpha(compositePixel)); - } - else { - setAlphaF4(&destination,1.0f); - setAlphaF4(&source,1.0f); - } - - float4 composite=destination; - - CompositeOperator op = (CompositeOperator)compose; - switch (op) { - case ColorDodgeCompositeOp: - CompositeColorDodge(&source,&destination,&composite); - break; - case BlendCompositeOp: - MagickPixelCompositeBlend(&source,source_dissolve,&destination, - destination_dissolve,&composite); - break; - default: - // unsupported operators - break; - }; - - CLPixelType outputPixel; - setRed(&outputPixel, ClampToQuantum(getRedF4(composite))); - setGreen(&outputPixel, ClampToQuantum(getGreenF4(composite))); - setBlue(&outputPixel, ClampToQuantum(getBlueF4(composite))); - setAlpha(&outputPixel, ClampToQuantum(getAlphaF4(composite))); - image[index.y*imageWidth+index.x] = outputPixel; - } - ) - /* %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % % diff --git a/MagickCore/accelerate-private.h b/MagickCore/accelerate-private.h index dd4303b55..d4b67a7dc 100644 --- a/MagickCore/accelerate-private.h +++ b/MagickCore/accelerate-private.h @@ -48,8 +48,6 @@ extern MagickPrivate Image *AccelerateWaveletDenoiseImage(const Image *,const double,ExceptionInfo *); extern MagickPrivate MagickBooleanType - AccelerateCompositeImage(Image *,const CompositeOperator,const Image *, - const float,const float,ExceptionInfo *), AccelerateContrastImage(Image *,const MagickBooleanType,ExceptionInfo *), AccelerateContrastStretchImage(Image *,const double,const double, ExceptionInfo*), diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c index 631e1dafd..4c70efe97 100644 --- a/MagickCore/accelerate.c +++ b/MagickCore/accelerate.c @@ -739,272 +739,6 @@ MagickPrivate Image* AccelerateBlurImage(const Image *image, return(filteredImage); } -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -% A c c e l e r a t e C o m p o s i t e I m a g e % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -*/ - -static MagickBooleanType LaunchCompositeKernel(MagickCLEnv clEnv, - MagickCLDevice device,cl_command_queue queue,cl_mem imageBuffer, - const unsigned int inputWidth,const unsigned int inputHeight, - const unsigned int matte,const ChannelType channel, - const CompositeOperator compose,const cl_mem compositeImageBuffer, - const unsigned int compositeWidth,const unsigned int compositeHeight, - const float destination_dissolve,const float source_dissolve) -{ - cl_int - clStatus; - - cl_kernel - compositeKernel; - - cl_event - event; - - int - k; - - size_t - global_work_size[2], - local_work_size[2]; - - unsigned int - composeOp; - - compositeKernel = AcquireOpenCLKernel(device,"Composite"); - - k = 0; - clStatus=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(cl_mem),(void*)&imageBuffer); - clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&inputWidth); - clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&inputHeight); - clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(cl_mem),(void*)&compositeImageBuffer); - clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&compositeWidth); - clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&compositeHeight); - composeOp = (unsigned int)compose; - clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&composeOp); - clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(ChannelType),(void*)&channel); - clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&matte); - clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(float),(void*)&destination_dissolve); - clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(float),(void*)&source_dissolve); - - if (clStatus!=CL_SUCCESS) - return MagickFalse; - - local_work_size[0] = 64; - local_work_size[1] = 1; - - global_work_size[0] = padGlobalWorkgroupSizeToLocalWorkgroupSize(inputWidth, - (unsigned int) local_work_size[0]); - global_work_size[1] = inputHeight; - clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, compositeKernel, 2, NULL, - global_work_size, local_work_size, 0, NULL, &event); - - RecordProfileData(device,compositeKernel,event); - - ReleaseOpenCLKernel(compositeKernel); - - return((clStatus==CL_SUCCESS) ? MagickTrue : MagickFalse); -} - -static MagickBooleanType ComputeCompositeImage(Image *image,MagickCLEnv clEnv, - const CompositeOperator compose,const Image *compositeImage, - const float destination_dissolve,const float source_dissolve,ExceptionInfo *exception) -{ - CacheView - *image_view; - - cl_command_queue - queue; - - cl_int - clStatus; - - cl_mem_flags - mem_flags; - - cl_mem - compositeImageBuffer, - imageBuffer; - - const void - *composePixels; - - MagickBooleanType - outputReady, - status; - - MagickCLDevice - device; - - MagickSizeType - length; - - void - *inputPixels; - - status = MagickFalse; - outputReady = MagickFalse; - composePixels = NULL; - imageBuffer = NULL; - compositeImageBuffer = NULL; - - device = RequestOpenCLDevice(clEnv); - queue = AcquireOpenCLCommandQueue(device); - - /* Create and initialize OpenCL buffers. */ - image_view=AcquireAuthenticCacheView(image,exception); - inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); - if (inputPixels == (void *) NULL) - { - (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning, - "UnableToReadPixelCache.","`%s'",image->filename); - goto cleanup; - } - - /* If the host pointer is aligned to the size of CLPixelPacket, - then use the host buffer directly from the GPU; otherwise, - create a buffer on the GPU and copy the data over */ - if (ALIGNED(inputPixels,CLPixelPacket)) - { - mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR; - } - else - { - mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; - } - /* create a CL buffer from image pixel buffer */ - length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, - length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), - ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); - goto cleanup; - } - - - /* Create and initialize OpenCL buffers. */ - composePixels = AcquirePixelCachePixels(compositeImage, &length, exception); - if (composePixels == (void *) NULL) - { - (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning, - "UnableToReadPixelCache.","`%s'",compositeImage->filename); - goto cleanup; - } - - /* If the host pointer is aligned to the size of CLPixelPacket, - then use the host buffer directly from the GPU; otherwise, - create a buffer on the GPU and copy the data over */ - if (ALIGNED(composePixels,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 = compositeImage->columns * compositeImage->rows; - compositeImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, - length * sizeof(CLPixelPacket), (void*)composePixels, &clStatus); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), - ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); - goto cleanup; - } - - status = LaunchCompositeKernel(clEnv,device,queue,imageBuffer, - (unsigned int) image->columns, - (unsigned int) image->rows, - (unsigned int) (image->alpha_trait > CopyPixelTrait) ? 1 : 0, - image->channel_mask, compose, compositeImageBuffer, - (unsigned int) compositeImage->columns, - (unsigned int) compositeImage->rows, - destination_dissolve,source_dissolve); - - if (status==MagickFalse) - goto cleanup; - - length = image->columns * image->rows; - if (ALIGNED(inputPixels,CLPixelPacket)) - { - clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, - CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, - NULL, &clStatus); - } - else - { - clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, - length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL); - } - if (clStatus==CL_SUCCESS) - outputReady=SyncCacheViewAuthenticPixels(image_view,exception); - -cleanup: - - image_view=DestroyCacheView(image_view); - if (imageBuffer!=NULL) - clEnv->library->clReleaseMemObject(imageBuffer); - if (compositeImageBuffer!=NULL) - clEnv->library->clReleaseMemObject(compositeImageBuffer); - if (queue != NULL) - RelinquishOpenCLCommandQueue(device,queue); - if (device != NULL) - ReleaseOpenCLDevice(device); - - return(outputReady); -} - -MagickPrivate MagickBooleanType AccelerateCompositeImage(Image *image, - const CompositeOperator compose,const Image *composite, - const float destination_dissolve,const float source_dissolve, - ExceptionInfo *exception) -{ - MagickBooleanType - status; - - MagickCLEnv - clEnv; - - assert(image != NULL); - assert(exception != (ExceptionInfo *) NULL); - - /* only support images with the size for now */ - if ((image->columns != composite->columns) || - (image->rows != composite->rows)) - return(MagickFalse); - - switch(compose) - { - case ColorDodgeCompositeOp: - case BlendCompositeOp: - break; - default: - // unsupported compose operator, quit - return(MagickFalse); - }; - - if (checkAccelerateConditionRGBA(image) == MagickFalse) - return(MagickFalse); - - clEnv=getOpenCLEnvironment(exception); - if (clEnv == (MagickCLEnv) NULL) - return(MagickFalse); - - status=ComputeCompositeImage(image,clEnv,compose,composite, - destination_dissolve,source_dissolve,exception); - return(status); -} - /* %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % % -- 2.40.0