return(status);
}
-/*
-%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
-% %
-% %
-% %
-% N e g a t e I m a g e w i t h O p e n C L %
-% %
-% %
-% %
-%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
-%
-%
-% A description of each parameter follows:
-%
-% o image: the image.
-%
-% o channel: the channel.
-%
-% o grayscale: If MagickTrue, only negate grayscale pixels within the image.
-%
-*/
-
-MagickBooleanType ComputeNegateImageChannel(Image *image,
- const ChannelType channel,const MagickBooleanType magick_unused(grayscale),
- ExceptionInfo* exception)
-{
- CacheView
- *image_view;
-
- cl_context
- context;
-
- cl_command_queue
- queue;
-
- cl_int
- clStatus;
-
- cl_kernel
- negateKernel;
-
- cl_mem
- imageBuffer;
-
- cl_mem_flags
- mem_flags;
-
- MagickBooleanType
- outputReady;
-
- MagickCLEnv
- clEnv;
-
- MagickSizeType
- length;
-
- register ssize_t
- i;
-
- void
- *inputPixels;
-
- magick_unreferenced(grayscale);
-
- inputPixels = NULL;
- imageBuffer = NULL;
- negateKernel = NULL;
-
- assert(image != (Image *) NULL);
- assert(image->signature == MagickSignature);
- if (image->debug != MagickFalse)
- (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
-
- /*
- * initialize opencl env
- */
- clEnv = GetDefaultOpenCLEnv();
- context = GetOpenCLContext(clEnv);
- queue = AcquireOpenCLCommandQueue(clEnv);
-
- outputReady = MagickFalse;
-
- /* Create and initialize OpenCL buffers.
- inputPixels = AcquirePixelCachePixels(image, &length, exception);
- assume this will get a writable image
- */
- image_view=AcquireAuthenticCacheView(image,exception);
- inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
- if (inputPixels == (void *) NULL)
- {
- (void) OpenCLThrowMagickException(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(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
- goto cleanup;
- }
-
- negateKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Negate");
- if (negateKernel == NULL)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
- goto cleanup;
- }
-
- i = 0;
- clStatus=clEnv->library->clSetKernelArg(negateKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
- clStatus=clEnv->library->clSetKernelArg(negateKernel,i++,sizeof(ChannelType),(void *)&channel);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
- printf("no kernel\n");
- goto cleanup;
- }
-
- {
- size_t global_work_size[2];
- global_work_size[0] = image->columns;
- global_work_size[1] = image->rows;
- /* launch the kernel */
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, negateKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
- goto cleanup;
- }
- clEnv->library->clFlush(queue);
- }
-
- if (ALIGNED(inputPixels,CLPixelPacket))
- {
- length = image->columns * image->rows;
- clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
- }
- else
- {
- length = image->columns * image->rows;
- clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
- }
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
- goto cleanup;
- }
-
- outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
-
-cleanup:
- OpenCLLogException(__FUNCTION__,__LINE__,exception);
-
- image_view=DestroyCacheView(image_view);
-
- if (imageBuffer!=NULL)
- clEnv->library->clReleaseMemObject(imageBuffer);
- if (negateKernel!=NULL)
- RelinquishOpenCLKernel(clEnv, negateKernel);
- if (queue != NULL)
- RelinquishOpenCLCommandQueue(clEnv, queue);
-
- return(outputReady);
-}
-
-MagickExport MagickBooleanType AccelerateNegateImageChannel(Image *image,
- const ChannelType channel,const MagickBooleanType grayscale,
- ExceptionInfo* exception)
-{
- MagickBooleanType
- status;
-
- assert(image != NULL);
- assert(exception != (ExceptionInfo *) NULL);
-
- if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
- (checkAccelerateCondition(image, channel) == MagickFalse))
- return(MagickFalse);
-
- status=ComputeNegateImageChannel(image,channel,grayscale,exception);
- return(status);
-}
-
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
return(MagickFalse);
}
-MagickExport
-MagickBooleanType AccelerateNegateImageChannel(
- Image* image, const ChannelType channel, const MagickBooleanType grayscale, ExceptionInfo* exception)
-{
- magick_unreferenced(image);
- magick_unreferenced(channel);
- magick_unreferenced(grayscale);
- magick_unreferenced(exception);
- return(MagickFalse);
-}
-
MagickExport
MagickBooleanType AccelerateGrayscaleImage(
Image* image, const PixelIntensityMethod method, ExceptionInfo* exception)
return NULL;
}
-#endif /* MAGICKCORE_OPENCL_SUPPORT */
-
-MagickExport MagickBooleanType AccelerateConvolveImage(
- const Image *magick_unused(image),const KernelInfo *magick_unused(kernel),
- Image *magick_unused(convolve_image),ExceptionInfo *magick_unused(exception))
-{
- magick_unreferenced(image);
- magick_unreferenced(kernel);
- magick_unreferenced(convolve_image);
- magick_unreferenced(exception);
-
- /* legacy, do not use */
- return(MagickFalse);
-}
-
+#endif /* MAGICKCORE_OPENCL_SUPPORT */
\ No newline at end of file
projects / imagemagick / commitdiff