}
)
-/*
-%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
-% %
-% %
-% %
-% 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;
- }
- )
-
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
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);
-}
-
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
projects / imagemagick / commitdiff