root/ImageMagick/trunk/magick/accelerate.c

Revision 2003, 23.7 KB (checked in by cristy, 3 months ago)
Line 
1/*
2%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3%                                                                             %
4%                                                                             %
5%                                                                             %
6%     AAA     CCCC    CCCC  EEEEE  L      EEEEE  RRRR    AAA   TTTTT  EEEEE   %
7%    A   A   C       C      E      L      E      R   R  A   A    T    E       %
8%    AAAAA   C       C      EEE    L      EEE    RRRR   AAAAA    T    EEE     %
9%    A   A   C       C      E      L      E      R R    A   A    T    E       %
10%    A   A    CCCC    CCCC  EEEEE  LLLLL  EEEEE  R  R   A   A    T    EEEEE   %
11%                                                                             %
12%                                                                             %
13%                       MagickCore Acceleration Methods                       %
14%                                                                             %
15%                              Software Design                                %
16%                               John Cristy                                   %
17%                               January 2010                                  %
18%                                                                             %
19%                                                                             %
20%  Copyright 1999-2010 ImageMagick Studio LLC, a non-profit organization      %
21%  dedicated to making software imaging solutions freely available.           %
22%                                                                             %
23%  You may not use this file except in compliance with the License.  You may  %
24%  obtain a copy of the License at                                            %
25%                                                                             %
26%    http://www.imagemagick.org/script/license.php                            %
27%                                                                             %
28%  Unless required by applicable law or agreed to in writing, software        %
29%  distributed under the License is distributed on an "AS IS" BASIS,          %
30%  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.   %
31%  See the License for the specific language governing permissions and        %
32%  limitations under the License.                                             %
33%                                                                             %
34%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
35%
36% Morphology is the the application of various kernals, of any size and even
37% shape, to a image in various ways (typically binary, but not always).
38%
39% Convolution (weighted sum or average) is just one specific type of
40% accelerate. Just one that is very common for image bluring and sharpening
41% effects.  Not only 2D Gaussian blurring, but also 2-pass 1D Blurring.
42%
43% This module provides not only a general accelerate function, and the ability
44% to apply more advanced or iterative morphologies, but also functions for the
45% generation of many different types of kernel arrays from user supplied
46% arguments. Prehaps even the generation of a kernel from a small image.
47*/
48
49/*
50  Include declarations.
51*/
52#include "magick/studio.h"
53#include "magick/accelerate.h"
54#include "magick/artifact.h"
55#include "magick/cache-view.h"
56#include "magick/color-private.h"
57#include "magick/enhance.h"
58#include "magick/exception.h"
59#include "magick/exception-private.h"
60#include "magick/gem.h"
61#include "magick/hashmap.h"
62#include "magick/image.h"
63#include "magick/image-private.h"
64#include "magick/list.h"
65#include "magick/memory_.h"
66#include "magick/monitor-private.h"
67#include "magick/accelerate.h"
68#include "magick/option.h"
69#include "magick/pixel-private.h"
70#include "magick/prepress.h"
71#include "magick/quantize.h"
72#include "magick/registry.h"
73#include "magick/semaphore.h"
74#include "magick/splay-tree.h"
75#include "magick/statistic.h"
76#include "magick/string_.h"
77#include "magick/string-private.h"
78#include "magick/token.h"
79
80/*
81%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
82%                                                                             %
83%                                                                             %
84%                                                                             %
85%     A c c e l e r a t e C o n v o l v e I m a g e                           %
86%                                                                             %
87%                                                                             %
88%                                                                             %
89%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
90%
91%  AccelerateConvolveImage() applies a custom convolution kernel to the image.
92%  It is accelerated by taking advantage of speed-ups offered by executing in
93%  concert across heterogeneous platforms consisting of CPUs, GPUs, and other
94%  processors.
95%
96%  The format of the AccelerateConvolveImage method is:
97%
98%      Image *AccelerateConvolveImage(const Image *image,
99%        const KernelInfo *kernel,Image *convolve_image,
100%        ExceptionInfo *exception)
101%
102%  A description of each parameter follows:
103%
104%    o image: the image.
105%
106%    o kernel: the convolution kernel.
107%
108%    o convole_image: the convoleed image.
109%
110%    o exception: return any errors or warnings in this structure.
111%
112*/
113
114#if defined(MAGICKCORE_OPENCL_SUPPORT)
115
116#if defined(MAGICKCORE_HDRI_SUPPORT)
117#define CLOptions "-DMAGICKCORE_HDRI_SUPPORT=1 -DCLQuantum=float " \
118  "-DCLPixelType=float4 -DQuantumRange=%g -DMagickEpsilon=%g"
119#define CLPixelPacket  cl_float4
120#else
121#if (MAGICKCORE_QUANTUM_DEPTH == 8)
122#define CLOptions "-DCLQuantum=uchar -DCLPixelType=uchar4 " \
123  "-DQuantumRange=%g -DMagickEpsilon=%g"
124#define CLPixelPacket  cl_uchar4
125#elif (MAGICKCORE_QUANTUM_DEPTH == 16)
126#define CLOptions "-DCLQuantum=ushort -DCLPixelType=ushort4 " \
127  "-DQuantumRange=%g -DMagickEpsilon=%g"
128#define CLPixelPacket  cl_ushort4
129#elif (MAGICKCORE_QUANTUM_DEPTH == 32)
130#define CLOptions "-DCLQuantum=uint -DCLPixelType=uint4 " \
131  "-DQuantumRange=%g -DMagickEpsilon=%g"
132#define CLPixelPacket  cl_uint4
133#elif (MAGICKCORE_QUANTUM_DEPTH == 32)
134#define CLOptions "-DCLQuantum=ussize_t -DCLPixelType=ussize_t4 " \
135  "-DQuantumRange=%g -DMagickEpsilon=%g"
136#define CLPixelPacket  cl_ulong4
137#endif
138#endif
139
140typedef struct _ConvolveInfo
141{
142  cl_context
143    context;
144
145  cl_device_id
146    *devices;
147
148  cl_command_queue
149    command_queue;
150
151  cl_kernel
152    kernel;
153
154  cl_program
155    program;
156
157  cl_mem
158    pixels,
159    convolve_pixels;
160
161  cl_ulong
162    width,
163    height;
164
165  cl_bool
166    matte;
167
168  cl_mem
169    filter;
170} ConvolveInfo;
171
172static char
173  *ConvolveKernel =
174    "static inline long ClampToCanvas(const long offset,const unsigned long range)\n"
175    "{\n"
176    "  if (offset < 0L)\n"
177    "    return(0L);\n"
178    "  if (offset >= range)\n"
179    "    return((long) (range-1L));\n"
180    "  return(offset);\n"
181    "}\n"
182    "\n"
183    "static inline CLQuantum ClampToQuantum(const double value)\n"
184    "{\n"
185    "#if defined(MAGICKCORE_HDRI_SUPPORT)\n"
186    "  return((CLQuantum) value)\n"
187    "#else\n"
188    "  if (value < 0.0)\n"
189    "    return((CLQuantum) 0);\n"
190    "  if (value >= (double) QuantumRange)\n"
191    "    return((CLQuantum) QuantumRange);\n"
192    "  return((CLQuantum) (value+0.5));\n"
193    "#endif\n"
194    "}\n"
195    "\n"
196    "__kernel void Convolve(const __global CLPixelType *input,\n"
197    "  __constant double *filter,const unsigned long width,const unsigned long height,\n"
198    "  const bool matte,__global CLPixelType *output)\n"
199    "{\n"
200    "  const unsigned long columns = get_global_size(0);\n"
201    "  const unsigned long rows = get_global_size(1);\n"
202    "\n"
203    "  const long x = get_global_id(0);\n"
204    "  const long y = get_global_id(1);\n"
205    "\n"
206    "  const double scale = (1.0/QuantumRange);\n"
207    "  const long mid_width = (width-1)/2;\n"
208    "  const long mid_height = (height-1)/2;\n"
209    "  double4 sum = { 0.0, 0.0, 0.0, 0.0 };\n"
210    "  double gamma = 0.0;\n"
211    "  register unsigned long i = 0;\n"
212    "\n"
213    "  int method = 0;\n"
214    "  if (matte != false)\n"
215    "    method=1;\n"
216    "  if ((x >= width) && (x < (columns-width-1)) &&\n"
217    "      (y >= height) && (y < (rows-height-1)))\n"
218    "    {\n"
219    "      method=2;\n"
220    "      if (matte != false)\n"
221    "        method=3;\n"
222    "    }\n"
223    "  switch (method)\n"
224    "  {\n"
225    "    case 0:\n"
226    "    {\n"
227    "      for (long v=(-mid_height); v <= mid_height; v++)\n"
228    "      {\n"
229    "        for (long u=(-mid_width); u <= mid_width; u++)\n"
230    "        {\n"
231    "          const long index=ClampToCanvas(y+v,rows)*columns+\n"
232    "            ClampToCanvas(x+u,columns);\n"
233    "          sum.x+=filter[i]*input[index].x;\n"
234    "          sum.y+=filter[i]*input[index].y;\n"
235    "          sum.z+=filter[i]*input[index].z;\n"
236    "          gamma+=filter[i];\n"
237    "          i++;\n"
238    "        }\n"
239    "      }\n"
240    "      break;\n"
241    "    }\n"
242    "    case 1:\n"
243    "    {\n"
244    "      for (long v=(-mid_height); v <= mid_height; v++)\n"
245    "      {\n"
246    "        for (long u=(-mid_width); u <= mid_width; u++)\n"
247    "        {\n"
248    "          const unsigned long index=ClampToCanvas(y+v,rows)*columns+\n"
249    "            ClampToCanvas(x+u,columns);\n"
250    "          const double alpha=scale*(QuantumRange-input[index].w);\n"
251    "          sum.x+=alpha*filter[i]*input[index].x;\n"
252    "          sum.y+=alpha*filter[i]*input[index].y;\n"
253    "          sum.z+=alpha*filter[i]*input[index].z;\n"
254    "          sum.w+=filter[i]*input[index].w;\n"
255    "          gamma+=alpha*filter[i];\n"
256    "          i++;\n"
257    "        }\n"
258    "      }\n"
259    "      break;\n"
260    "    }\n"
261    "    case 2:\n"
262    "    {\n"
263    "      for (long v=(-mid_height); v <= mid_height; v++)\n"
264    "      {\n"
265    "        for (long u=(-mid_width); u <= mid_width; u++)\n"
266    "        {\n"
267    "          const unsigned long index=(y+v)*columns+(x+u);\n"
268    "          sum.x+=filter[i]*input[index].x;\n"
269    "          sum.y+=filter[i]*input[index].y;\n"
270    "          sum.z+=filter[i]*input[index].z;\n"
271    "          gamma+=filter[i];\n"
272    "          i++;\n"
273    "        }\n"
274    "      }\n"
275    "      break;\n"
276    "    }\n"
277    "    case 3:\n"
278    "    {\n"
279    "      for (long v=(-mid_height); v <= mid_height; v++)\n"
280    "      {\n"
281    "        for (long u=(-mid_width); u <= mid_width; u++)\n"
282    "        {\n"
283    "          const unsigned long index=(y+v)*columns+(x+u);\n"
284    "          const double alpha=scale*(QuantumRange-input[index].w);\n"
285    "          sum.x+=alpha*filter[i]*input[index].x;\n"
286    "          sum.y+=alpha*filter[i]*input[index].y;\n"
287    "          sum.z+=alpha*filter[i]*input[index].z;\n"
288    "          sum.w+=filter[i]*input[index].w;\n"
289    "          gamma+=alpha*filter[i];\n"
290    "          i++;\n"
291    "        }\n"
292    "      }\n"
293    "      break;\n"
294    "    }\n"
295    "  }\n"
296    "  gamma=1.0/(fabs(gamma) <= MagickEpsilon ? 1.0 : gamma);\n"
297    "  const unsigned long index = y*columns+x;\n"
298    "  output[index].x=ClampToQuantum(gamma*sum.x);\n"
299    "  output[index].y=ClampToQuantum(gamma*sum.y);\n"
300    "  output[index].z=ClampToQuantum(gamma*sum.z);\n"
301    "  if (matte == false)\n"
302    "    output[index].w=input[index].w;\n"
303    "  else\n"
304    "    output[index].w=ClampToQuantum(sum.w);\n"
305    "}\n";
306
307static void ConvolveNotify(const char *message,const void *data,size_t length,
308  void *user_context)
309{
310  ExceptionInfo
311    *exception;
312
313  (void) data;
314  (void) length;
315  exception=(ExceptionInfo *) user_context;
316  (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
317    "DelegateFailed","`%s'",message);
318}
319
320static MagickBooleanType BindConvolveParameters(ConvolveInfo *convolve_info,
321  const Image *image,const void *pixels,double *filter,
322  const size_t width,const size_t height,void *convolve_pixels)
323{
324  cl_int
325    status;
326
327  register cl_uint
328    i;
329
330  size_t
331    length;
332
333  /*
334    Allocate OpenCL buffers.
335  */
336  length=image->columns*image->rows;
337  convolve_info->pixels=clCreateBuffer(convolve_info->context,(cl_mem_flags)
338    (CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR),length*sizeof(CLPixelPacket),
339    (void *) pixels,&status);
340  if ((convolve_info->pixels == (cl_mem) NULL) || (status != CL_SUCCESS))
341    return(MagickFalse);
342  length=width*height;
343  convolve_info->filter=clCreateBuffer(convolve_info->context,(cl_mem_flags)
344    (CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR),length*sizeof(cl_double),filter,
345    &status);
346  if ((convolve_info->filter == (cl_mem) NULL) || (status != CL_SUCCESS))
347    return(MagickFalse);
348  length=image->columns*image->rows;
349  convolve_info->convolve_pixels=clCreateBuffer(convolve_info->context,
350    (cl_mem_flags) (CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR),length*
351    sizeof(CLPixelPacket),convolve_pixels,&status);
352  if ((convolve_info->convolve_pixels == (cl_mem) NULL) ||
353      (status != CL_SUCCESS))
354    return(MagickFalse);
355  /*
356    Bind OpenCL buffers.
357  */
358  i=0;
359  status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
360    &convolve_info->pixels);
361  if (status != CL_SUCCESS)
362    return(MagickFalse);
363  status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
364    &convolve_info->filter);
365  if (status != CL_SUCCESS)
366    return(MagickFalse);
367  convolve_info->width=(cl_ulong) width;
368  status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_ulong),(void *)
369    &convolve_info->width);
370  if (status != CL_SUCCESS)
371    return(MagickFalse);
372  convolve_info->height=(cl_ulong) height;
373  status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_ulong),(void *)
374    &convolve_info->height);
375  if (status != CL_SUCCESS)
376    return(MagickFalse);
377  convolve_info->matte=(cl_bool) image->matte;
378  status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_bool),(void *)
379    &convolve_info->matte);
380  if (status != CL_SUCCESS)
381    return(MagickFalse);
382  status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
383    &convolve_info->convolve_pixels);
384  if (status != CL_SUCCESS)
385    return(MagickFalse);
386  status=clFinish(convolve_info->command_queue);
387  if (status != CL_SUCCESS)
388    return(MagickFalse);
389  return(MagickTrue);
390}
391
392static void DestroyConvolveBuffers(ConvolveInfo *convolve_info)
393{
394  cl_int
395    status;
396
397  if (convolve_info->convolve_pixels != (cl_mem) NULL)
398    status=clReleaseMemObject(convolve_info->convolve_pixels);
399  if (convolve_info->pixels != (cl_mem) NULL)
400    status=clReleaseMemObject(convolve_info->pixels);
401  if (convolve_info->filter != (cl_mem) NULL)
402    status=clReleaseMemObject(convolve_info->filter);
403}
404
405static ConvolveInfo *DestroyConvolveInfo(ConvolveInfo *convolve_info)
406{
407  cl_int
408    status;
409
410  if (convolve_info->kernel != (cl_kernel) NULL)
411    status=clReleaseKernel(convolve_info->kernel);
412  if (convolve_info->program != (cl_program) NULL)
413    status=clReleaseProgram(convolve_info->program);
414  if (convolve_info->command_queue != (cl_command_queue) NULL)
415    status=clReleaseCommandQueue(convolve_info->command_queue);
416  if (convolve_info->context != (cl_context) NULL)
417    status=clReleaseContext(convolve_info->context);
418  convolve_info=(ConvolveInfo *) RelinquishMagickMemory(convolve_info);
419  return(convolve_info);
420}
421
422static MagickBooleanType EnqueueConvolveKernel(ConvolveInfo *convolve_info,
423  const Image *image,const void *pixels,double *filter,
424  const size_t width,const size_t height,void *convolve_pixels)
425{
426  cl_int
427    status;
428
429  size_t
430    global_work_size[2],
431    length;
432
433  length=image->columns*image->rows;
434  status=clEnqueueWriteBuffer(convolve_info->command_queue,
435    convolve_info->pixels,CL_TRUE,0,length*sizeof(CLPixelPacket),pixels,0,NULL,
436    NULL);
437  length=width*height;
438  status=clEnqueueWriteBuffer(convolve_info->command_queue,
439    convolve_info->filter,CL_TRUE,0,length*sizeof(cl_double),filter,0,NULL,
440    NULL);
441  if (status != CL_SUCCESS)
442    return(MagickFalse);
443  global_work_size[0]=image->columns;
444  global_work_size[1]=image->rows;
445  status=clEnqueueNDRangeKernel(convolve_info->command_queue,
446    convolve_info->kernel,2,NULL,global_work_size,NULL,0,NULL,NULL);
447  if (status != CL_SUCCESS)
448    return(MagickFalse);
449  length=image->columns*image->rows;
450  status=clEnqueueReadBuffer(convolve_info->command_queue,
451    convolve_info->convolve_pixels,CL_TRUE,0,length*sizeof(CLPixelPacket),
452    convolve_pixels,0,NULL,NULL);
453  if (status != CL_SUCCESS)
454    return(MagickFalse);
455  status=clFinish(convolve_info->command_queue);
456  if (status != CL_SUCCESS)
457    return(MagickFalse);
458  return(MagickTrue);
459}
460
461static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name,
462  const char *source,ExceptionInfo *exception)
463{
464  char
465    options[MaxTextExtent];
466
467  cl_int
468    status;
469
470  ConvolveInfo
471    *convolve_info;
472
473  size_t
474    length,
475    lengths[] = { strlen(source) };
476
477  /*
478    Create OpenCL info.
479  */
480  convolve_info=(ConvolveInfo *) AcquireAlignedMemory(1,sizeof(*convolve_info));
481  if (convolve_info == (ConvolveInfo *) NULL)
482    {
483      (void) ThrowMagickException(exception,GetMagickModule(),
484        ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename);
485      return((ConvolveInfo *) NULL);
486    }
487  (void) ResetMagickMemory(convolve_info,0,sizeof(*convolve_info));
488  /*
489    Create OpenCL context.
490  */
491  convolve_info->context=clCreateContextFromType((cl_context_properties *)
492    NULL,(cl_device_type) CL_DEVICE_TYPE_GPU,ConvolveNotify,exception,&status);
493  if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
494    convolve_info->context=clCreateContextFromType((cl_context_properties *)
495      NULL,(cl_device_type) CL_DEVICE_TYPE_CPU,ConvolveNotify,exception,
496      &status);
497  if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
498    convolve_info->context=clCreateContextFromType((cl_context_properties *)
499      NULL,(cl_device_type) CL_DEVICE_TYPE_DEFAULT,ConvolveNotify,exception,
500      &status);
501  if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
502    {
503      (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
504        "failed to create OpenCL context","`%s' (%d)",image->filename,status);
505      convolve_info=DestroyConvolveInfo(convolve_info);
506      return((ConvolveInfo *) NULL);
507    }
508  /*
509    Detect OpenCL devices.
510  */
511  status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,0,NULL,
512    &length);
513  if ((status != CL_SUCCESS) || (length == 0))
514    {
515      convolve_info=DestroyConvolveInfo(convolve_info);
516      return((ConvolveInfo *) NULL);
517    }
518  convolve_info->devices=(cl_device_id *) AcquireMagickMemory(length);
519  if (convolve_info->devices == (cl_device_id *) NULL)
520    {
521      (void) ThrowMagickException(exception,GetMagickModule(),
522        ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename);
523      convolve_info=DestroyConvolveInfo(convolve_info);
524      return((ConvolveInfo *) NULL);
525    }
526  status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,length,
527    convolve_info->devices,NULL);
528  if (status != CL_SUCCESS)
529    {
530      convolve_info=DestroyConvolveInfo(convolve_info);
531      return((ConvolveInfo *) NULL);
532    }
533  /*
534    Create OpenCL command queue.
535  */
536  convolve_info->command_queue=clCreateCommandQueue(convolve_info->context,
537    convolve_info->devices[0],0,&status);
538  if ((convolve_info->command_queue == (cl_command_queue) NULL) ||
539      (status != CL_SUCCESS))
540    {
541      convolve_info=DestroyConvolveInfo(convolve_info);
542      return((ConvolveInfo *) NULL);
543    }
544  /*
545    Build OpenCL program.
546  */
547  convolve_info->program=clCreateProgramWithSource(convolve_info->context,1,
548    &source,lengths,&status);
549  if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS))
550    {
551      convolve_info=DestroyConvolveInfo(convolve_info);
552      return((ConvolveInfo *) NULL);
553    }
554  (void) FormatMagickString(options,MaxTextExtent,CLOptions,(double)
555    QuantumRange,MagickEpsilon);
556  status=clBuildProgram(convolve_info->program,1,convolve_info->devices,options,
557    NULL,NULL);
558  if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS))
559    {
560      char
561        *log;
562
563      status=clGetProgramBuildInfo(convolve_info->program,
564        convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,0,NULL,&length);
565      log=(char *) AcquireMagickMemory(length);
566      if (log == (char *) NULL)
567        {
568          convolve_info=DestroyConvolveInfo(convolve_info);
569          return((ConvolveInfo *) NULL);
570        }
571      status=clGetProgramBuildInfo(convolve_info->program,
572        convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,length,log,&length);
573      (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
574        "failed to build OpenCL program","`%s' (%s)",image->filename,log);
575      log=DestroyString(log);
576      convolve_info=DestroyConvolveInfo(convolve_info);
577      return((ConvolveInfo *) NULL);
578    }
579  /*
580    Get a kernel object.
581  */
582  convolve_info->kernel=clCreateKernel(convolve_info->program,name,&status);
583  if ((convolve_info->kernel == (cl_kernel) NULL) || (status != CL_SUCCESS))
584    {
585      convolve_info=DestroyConvolveInfo(convolve_info);
586      return((ConvolveInfo *) NULL);
587    }
588  return(convolve_info);
589}
590
591#endif
592
593MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image,
594  const KernelInfo *kernel,Image *convolve_image,ExceptionInfo *exception)
595{
596  assert(image != (Image *) NULL);
597  assert(image->signature == MagickSignature);
598  if (image->debug != MagickFalse)
599    (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
600  assert(kernel != (KernelInfo *) NULL);
601  assert(kernel->signature == MagickSignature);
602  assert(convolve_image != (Image *) NULL);
603  assert(convolve_image->signature == MagickSignature);
604  assert(exception != (ExceptionInfo *) NULL);
605  assert(exception->signature == MagickSignature);
606  if ((image->storage_class != DirectClass) ||
607      (image->colorspace == CMYKColorspace))
608  if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) &&
609      (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod))
610    return(MagickFalse);
611#if !defined(MAGICKCORE_OPENCL_SUPPORT)
612  return(MagickFalse);
613#else
614  {
615    const void
616      *pixels;
617
618    ConvolveInfo
619      *convolve_info;
620
621    MagickBooleanType
622      status;
623
624    MagickSizeType
625      length;
626
627    void
628      *convolve_pixels;
629
630    convolve_info=GetConvolveInfo(image,"Convolve",ConvolveKernel,exception);
631    if (convolve_info == (ConvolveInfo *) NULL)
632      return(MagickFalse);
633    pixels=AcquirePixelCachePixels(image,&length,exception);
634    if (pixels == (const void *) NULL)
635      {
636        (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
637          "UnableToReadPixelCache","`%s'",image->filename);
638        convolve_info=DestroyConvolveInfo(convolve_info);
639        return(MagickFalse);
640      }
641    convolve_pixels=GetPixelCachePixels(convolve_image,&length,exception);
642    if (convolve_pixels == (void *) NULL)
643      {
644        (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
645          "UnableToReadPixelCache","`%s'",image->filename);
646        convolve_info=DestroyConvolveInfo(convolve_info);
647        return(MagickFalse);
648      }
649    status=BindConvolveParameters(convolve_info,image,pixels,kernel->values,
650      kernel->width,kernel->height,convolve_pixels);
651    if (status == MagickFalse)
652      {
653        DestroyConvolveBuffers(convolve_info);
654        convolve_info=DestroyConvolveInfo(convolve_info);
655        return(MagickFalse);
656      }
657    status=EnqueueConvolveKernel(convolve_info,image,pixels,kernel->values,
658      kernel->width,kernel->height,convolve_pixels);
659    if (status == MagickFalse)
660      {
661        DestroyConvolveBuffers(convolve_info);
662        convolve_info=DestroyConvolveInfo(convolve_info);
663        return(MagickFalse);
664      }
665    DestroyConvolveBuffers(convolve_info);
666    convolve_info=DestroyConvolveInfo(convolve_info);
667    return(MagickTrue);
668  }
669#endif
670}
Note: See TracBrowser for help on using the browser.