← Back to branch summary

~ubuntu-branches/ubuntu/oneiric/imagemagick/oneiric-updates

« back to all changes in this revision

Viewing changes to magick/accelerate.c

  • Committer: Bazaar Package Importer
  • Author(s): Nelson A. de Oliveira
  • Date: 2010-06-20 19:59:55 UTC
  • mfrom: (6.2.8 sid)
  • Revision ID: james.westby@ubuntu.com-20100620195955-n3eq0yenhycw888i
Tags: 7:6.6.2.6-1
http://bugs.debian.org/579775
http://bugs.debian.org/582073
http://bugs.debian.org/583024
* New upstream release;
* Change Recommends on ufraw to ufraw-batch (Closes: #579775);
* Fix FTBFS when using dash to run the configure script, by setting
  CONFIG_SHELL=/bin/bash (Closes: #582073, #583024). Thank you, Niko Tyni!

Show diffs side-by-side

added added

removed removed

Lines of Context:
magick/accelerate.c
 
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
 
 
140
typedef 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
 
 
172
static 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
 
 
307
static 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
 
 
320
static 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
 
 
392
static 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
 
 
405
static 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
 
 
422
static 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
 
 
461
static 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
 
 
593
MagickExport 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
}