accelerate.c revision 1a2e276c9e1be7859ac1c933ffdc00f9d91bbc02
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-2011 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 "MagickCore/studio.h" 53#include "MagickCore/accelerate.h" 54#include "MagickCore/artifact.h" 55#include "MagickCore/cache.h" 56#include "MagickCore/cache-private.h" 57#include "MagickCore/cache-view.h" 58#include "MagickCore/color-private.h" 59#include "MagickCore/enhance.h" 60#include "MagickCore/exception.h" 61#include "MagickCore/exception-private.h" 62#include "MagickCore/gem.h" 63#include "MagickCore/hashmap.h" 64#include "MagickCore/image.h" 65#include "MagickCore/image-private.h" 66#include "MagickCore/list.h" 67#include "MagickCore/memory_.h" 68#include "MagickCore/monitor-private.h" 69#include "MagickCore/accelerate.h" 70#include "MagickCore/option.h" 71#include "MagickCore/pixel-accessor.h" 72#include "MagickCore/prepress.h" 73#include "MagickCore/quantize.h" 74#include "MagickCore/registry.h" 75#include "MagickCore/semaphore.h" 76#include "MagickCore/splay-tree.h" 77#include "MagickCore/statistic.h" 78#include "MagickCore/string_.h" 79#include "MagickCore/string-private.h" 80#include "MagickCore/token.h" 81 82/* 83%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 84% % 85% % 86% % 87% A c c e l e r a t e C o n v o l v e I m a g e % 88% % 89% % 90% % 91%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 92% 93% AccelerateConvolveImage() applies a custom convolution kernel to the image. 94% It is accelerated by taking advantage of speed-ups offered by executing in 95% concert across heterogeneous platforms consisting of CPUs, GPUs, and other 96% processors. 97% 98% The format of the AccelerateConvolveImage method is: 99% 100% Image *AccelerateConvolveImage(const Image *image, 101% const KernelInfo *kernel,Image *convolve_image, 102% ExceptionInfo *exception) 103% 104% A description of each parameter follows: 105% 106% o image: the image. 107% 108% o kernel: the convolution kernel. 109% 110% o convole_image: the convoleed image. 111% 112% o exception: return any errors or warnings in this structure. 113% 114*/ 115 116#if defined(MAGICKCORE_OPENCL_SUPPORT) 117 118#if defined(MAGICKCORE_HDRI_SUPPORT) 119#define CLOptions "-DMAGICKCORE_HDRI_SUPPORT=1 -DCLQuantum=float " \ 120 "-DCLPixelType=float4 -DQuantumRange=%g -DMagickEpsilon=%g" 121#define CLPixelPacket cl_float4 122#else 123#if (MAGICKCORE_QUANTUM_DEPTH == 8) 124#define CLOptions "-DCLQuantum=uchar -DCLPixelType=uchar4 " \ 125 "-DQuantumRange=%g -DMagickEpsilon=%g" 126#define CLPixelPacket cl_uchar4 127#elif (MAGICKCORE_QUANTUM_DEPTH == 16) 128#define CLOptions "-DCLQuantum=ushort -DCLPixelType=ushort4 " \ 129 "-DQuantumRange=%g -DMagickEpsilon=%g" 130#define CLPixelPacket cl_ushort4 131#elif (MAGICKCORE_QUANTUM_DEPTH == 32) 132#define CLOptions "-DCLQuantum=uint -DCLPixelType=uint4 " \ 133 "-DQuantumRange=%g -DMagickEpsilon=%g" 134#define CLPixelPacket cl_uint4 135#elif (MAGICKCORE_QUANTUM_DEPTH == 64) 136#define CLOptions "-DCLQuantum=ussize_t -DCLPixelType=ussize_t4 " \ 137 "-DQuantumRange=%g -DMagickEpsilon=%g" 138#define CLPixelPacket cl_ulong4 139#endif 140#endif 141 142typedef struct _ConvolveInfo 143{ 144 cl_context 145 context; 146 147 cl_device_id 148 *devices; 149 150 cl_command_queue 151 command_queue; 152 153 cl_kernel 154 kernel; 155 156 cl_program 157 program; 158 159 cl_mem 160 pixels, 161 convolve_pixels; 162 163 cl_ulong 164 width, 165 height; 166 167 cl_uint 168 matte; 169 170 cl_mem 171 filter; 172} ConvolveInfo; 173 174static char 175 *ConvolveKernel = 176 "static inline long ClampToCanvas(const long offset,const unsigned long range)\n" 177 "{\n" 178 " if (offset < 0L)\n" 179 " return(0L);\n" 180 " if (offset >= range)\n" 181 " return((long) (range-1L));\n" 182 " return(offset);\n" 183 "}\n" 184 "\n" 185 "static inline CLQuantum ClampToQuantum(const float value)\n" 186 "{\n" 187 "#if defined(MAGICKCORE_HDRI_SUPPORT)\n" 188 " return((CLQuantum) value)\n" 189 "#else\n" 190 " if (value < 0.0)\n" 191 " return((CLQuantum) 0);\n" 192 " if (value >= (float) QuantumRange)\n" 193 " return((CLQuantum) QuantumRange);\n" 194 " return((CLQuantum) (value+0.5));\n" 195 "#endif\n" 196 "}\n" 197 "\n" 198 "__kernel void Convolve(const __global CLPixelType *input,\n" 199 " __constant float *filter,const unsigned long width,const unsigned long height,\n" 200 " const unsigned int matte,__global CLPixelType *output)\n" 201 "{\n" 202 " const unsigned long columns = get_global_size(0);\n" 203 " const unsigned long rows = get_global_size(1);\n" 204 "\n" 205 " const long x = get_global_id(0);\n" 206 " const long y = get_global_id(1);\n" 207 "\n" 208 " const float scale = (1.0/QuantumRange);\n" 209 " const long mid_width = (width-1)/2;\n" 210 " const long mid_height = (height-1)/2;\n" 211 " float4 sum = { 0.0, 0.0, 0.0, 0.0 };\n" 212 " float gamma = 0.0;\n" 213 " register unsigned long i = 0;\n" 214 "\n" 215 " int method = 0;\n" 216 " if (matte != false)\n" 217 " method=1;\n" 218 " if ((x >= width) && (x < (columns-width-1)) &&\n" 219 " (y >= height) && (y < (rows-height-1)))\n" 220 " {\n" 221 " method=2;\n" 222 " if (matte != false)\n" 223 " method=3;\n" 224 " }\n" 225 " switch (method)\n" 226 " {\n" 227 " case 0:\n" 228 " {\n" 229 " for (long v=(-mid_height); v <= mid_height; v++)\n" 230 " {\n" 231 " for (long u=(-mid_width); u <= mid_width; u++)\n" 232 " {\n" 233 " const long index=ClampToCanvas(y+v,rows)*columns+\n" 234 " ClampToCanvas(x+u,columns);\n" 235 " sum.x+=filter[i]*input[index].x;\n" 236 " sum.y+=filter[i]*input[index].y;\n" 237 " sum.z+=filter[i]*input[index].z;\n" 238 " gamma+=filter[i];\n" 239 " i++;\n" 240 " }\n" 241 " }\n" 242 " break;\n" 243 " }\n" 244 " case 1:\n" 245 " {\n" 246 " for (long v=(-mid_height); v <= mid_height; v++)\n" 247 " {\n" 248 " for (long u=(-mid_width); u <= mid_width; u++)\n" 249 " {\n" 250 " const unsigned long index=ClampToCanvas(y+v,rows)*columns+\n" 251 " ClampToCanvas(x+u,columns);\n" 252 " const float alpha=scale*input[index].w;\n" 253 " sum.x+=alpha*filter[i]*input[index].x;\n" 254 " sum.y+=alpha*filter[i]*input[index].y;\n" 255 " sum.z+=alpha*filter[i]*input[index].z;\n" 256 " sum.w+=filter[i]*input[index].w;\n" 257 " gamma+=alpha*filter[i];\n" 258 " i++;\n" 259 " }\n" 260 " }\n" 261 " break;\n" 262 " }\n" 263 " case 2:\n" 264 " {\n" 265 " for (long v=(-mid_height); v <= mid_height; v++)\n" 266 " {\n" 267 " for (long u=(-mid_width); u <= mid_width; u++)\n" 268 " {\n" 269 " const unsigned long index=(y+v)*columns+(x+u);\n" 270 " sum.x+=filter[i]*input[index].x;\n" 271 " sum.y+=filter[i]*input[index].y;\n" 272 " sum.z+=filter[i]*input[index].z;\n" 273 " gamma+=filter[i];\n" 274 " i++;\n" 275 " }\n" 276 " }\n" 277 " break;\n" 278 " }\n" 279 " case 3:\n" 280 " {\n" 281 " for (long v=(-mid_height); v <= mid_height; v++)\n" 282 " {\n" 283 " for (long u=(-mid_width); u <= mid_width; u++)\n" 284 " {\n" 285 " const unsigned long index=(y+v)*columns+(x+u);\n" 286 " const float alpha=scale*input[index].w;\n" 287 " sum.x+=alpha*filter[i]*input[index].x;\n" 288 " sum.y+=alpha*filter[i]*input[index].y;\n" 289 " sum.z+=alpha*filter[i]*input[index].z;\n" 290 " sum.w+=filter[i]*input[index].w;\n" 291 " gamma+=alpha*filter[i];\n" 292 " i++;\n" 293 " }\n" 294 " }\n" 295 " break;\n" 296 " }\n" 297 " }\n" 298 " gamma=1.0/(fabs(gamma) <= MagickEpsilon ? 1.0 : gamma);\n" 299 " const unsigned long index = y*columns+x;\n" 300 " output[index].x=ClampToQuantum(gamma*sum.x);\n" 301 " output[index].y=ClampToQuantum(gamma*sum.y);\n" 302 " output[index].z=ClampToQuantum(gamma*sum.z);\n" 303 " if (matte == false)\n" 304 " output[index].w=input[index].w;\n" 305 " else\n" 306 " output[index].w=ClampToQuantum(sum.w);\n" 307 "}\n"; 308 309static void ConvolveNotify(const char *message,const void *data,size_t length, 310 void *user_context) 311{ 312 ExceptionInfo 313 *exception; 314 315 (void) data; 316 (void) length; 317 exception=(ExceptionInfo *) user_context; 318 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning, 319 "DelegateFailed","`%s'",message); 320} 321 322static MagickBooleanType BindConvolveParameters(ConvolveInfo *convolve_info, 323 const Image *image,const void *pixels,float *filter,const size_t width, 324 const size_t height,void *convolve_pixels) 325{ 326 cl_int 327 status; 328 329 register cl_uint 330 i; 331 332 size_t 333 length; 334 335 /* 336 Allocate OpenCL buffers. 337 */ 338 length=image->columns*image->rows; 339 convolve_info->pixels=clCreateBuffer(convolve_info->context,(cl_mem_flags) 340 (CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR),length*sizeof(CLPixelPacket), 341 (void *) pixels,&status); 342 if ((convolve_info->pixels == (cl_mem) NULL) || (status != CL_SUCCESS)) 343 return(MagickFalse); 344 length=width*height; 345 convolve_info->filter=clCreateBuffer(convolve_info->context,(cl_mem_flags) 346 (CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR),length*sizeof(cl_float),filter, 347 &status); 348 if ((convolve_info->filter == (cl_mem) NULL) || (status != CL_SUCCESS)) 349 return(MagickFalse); 350 length=image->columns*image->rows; 351 convolve_info->convolve_pixels=clCreateBuffer(convolve_info->context, 352 (cl_mem_flags) (CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR),length* 353 sizeof(CLPixelPacket),convolve_pixels,&status); 354 if ((convolve_info->convolve_pixels == (cl_mem) NULL) || 355 (status != CL_SUCCESS)) 356 return(MagickFalse); 357 /* 358 Bind OpenCL buffers. 359 */ 360 i=0; 361 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *) 362 &convolve_info->pixels); 363 if (status != CL_SUCCESS) 364 return(MagickFalse); 365 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *) 366 &convolve_info->filter); 367 if (status != CL_SUCCESS) 368 return(MagickFalse); 369 convolve_info->width=(cl_ulong) width; 370 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_ulong),(void *) 371 &convolve_info->width); 372 if (status != CL_SUCCESS) 373 return(MagickFalse); 374 convolve_info->height=(cl_ulong) height; 375 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_ulong),(void *) 376 &convolve_info->height); 377 if (status != CL_SUCCESS) 378 return(MagickFalse); 379 convolve_info->matte=(cl_uint) image->matte; 380 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_uint),(void *) 381 &convolve_info->matte); 382 if (status != CL_SUCCESS) 383 return(MagickFalse); 384 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *) 385 &convolve_info->convolve_pixels); 386 if (status != CL_SUCCESS) 387 return(MagickFalse); 388 status=clFinish(convolve_info->command_queue); 389 if (status != CL_SUCCESS) 390 return(MagickFalse); 391 return(MagickTrue); 392} 393 394static void DestroyConvolveBuffers(ConvolveInfo *convolve_info) 395{ 396 cl_int 397 status; 398 399 (void) status; 400 if (convolve_info->convolve_pixels != (cl_mem) NULL) 401 status=clReleaseMemObject(convolve_info->convolve_pixels); 402 if (convolve_info->pixels != (cl_mem) NULL) 403 status=clReleaseMemObject(convolve_info->pixels); 404 if (convolve_info->filter != (cl_mem) NULL) 405 status=clReleaseMemObject(convolve_info->filter); 406} 407 408static ConvolveInfo *DestroyConvolveInfo(ConvolveInfo *convolve_info) 409{ 410 cl_int 411 status; 412 413 (void) status; 414 if (convolve_info->kernel != (cl_kernel) NULL) 415 status=clReleaseKernel(convolve_info->kernel); 416 if (convolve_info->program != (cl_program) NULL) 417 status=clReleaseProgram(convolve_info->program); 418 if (convolve_info->command_queue != (cl_command_queue) NULL) 419 status=clReleaseCommandQueue(convolve_info->command_queue); 420 if (convolve_info->context != (cl_context) NULL) 421 status=clReleaseContext(convolve_info->context); 422 convolve_info=(ConvolveInfo *) RelinquishMagickMemory(convolve_info); 423 return(convolve_info); 424} 425 426static MagickBooleanType EnqueueConvolveKernel(ConvolveInfo *convolve_info, 427 const Image *image,const void *pixels,float *filter,const size_t width, 428 const size_t height,void *convolve_pixels) 429{ 430 cl_int 431 status; 432 433 size_t 434 global_work_size[2], 435 length; 436 437 length=image->columns*image->rows; 438 status=clEnqueueWriteBuffer(convolve_info->command_queue, 439 convolve_info->pixels,CL_TRUE,0,length*sizeof(CLPixelPacket),pixels,0,NULL, 440 NULL); 441 length=width*height; 442 status=clEnqueueWriteBuffer(convolve_info->command_queue, 443 convolve_info->filter,CL_TRUE,0,length*sizeof(cl_float),filter,0,NULL, 444 NULL); 445 if (status != CL_SUCCESS) 446 return(MagickFalse); 447 global_work_size[0]=image->columns; 448 global_work_size[1]=image->rows; 449 status=clEnqueueNDRangeKernel(convolve_info->command_queue, 450 convolve_info->kernel,2,NULL,global_work_size,NULL,0,NULL,NULL); 451 if (status != CL_SUCCESS) 452 return(MagickFalse); 453 length=image->columns*image->rows; 454 status=clEnqueueReadBuffer(convolve_info->command_queue, 455 convolve_info->convolve_pixels,CL_TRUE,0,length*sizeof(CLPixelPacket), 456 convolve_pixels,0,NULL,NULL); 457 if (status != CL_SUCCESS) 458 return(MagickFalse); 459 status=clFinish(convolve_info->command_queue); 460 if (status != CL_SUCCESS) 461 return(MagickFalse); 462 return(MagickTrue); 463} 464 465static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name, 466 const char *source,ExceptionInfo *exception) 467{ 468 char 469 options[MaxTextExtent]; 470 471 cl_context_properties 472 context_properties[3]; 473 474 cl_int 475 status; 476 477 cl_platform_id 478 platforms[1]; 479 480 cl_uint 481 number_platforms; 482 483 ConvolveInfo 484 *convolve_info; 485 486 size_t 487 length, 488 lengths[] = { strlen(source) }; 489 490 /* 491 Create OpenCL info. 492 */ 493 convolve_info=(ConvolveInfo *) AcquireMagickMemory(sizeof(*convolve_info)); 494 if (convolve_info == (ConvolveInfo *) NULL) 495 { 496 (void) ThrowMagickException(exception,GetMagickModule(), 497 ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename); 498 return((ConvolveInfo *) NULL); 499 } 500 (void) ResetMagickMemory(convolve_info,0,sizeof(*convolve_info)); 501 /* 502 Create OpenCL context. 503 */ 504 status=clGetPlatformIDs(0,NULL,&number_platforms); 505 if (status == CL_SUCCESS) 506 status=clGetPlatformIDs(1,platforms,NULL); 507 if (status != CL_SUCCESS) 508 { 509 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning, 510 "failed to create OpenCL context","`%s' (%d)",image->filename,status); 511 convolve_info=DestroyConvolveInfo(convolve_info); 512 return((ConvolveInfo *) NULL); 513 } 514 context_properties[0]=CL_CONTEXT_PLATFORM; 515 context_properties[1]=(cl_context_properties) platforms[0]; 516 context_properties[2]=0; 517 convolve_info->context=clCreateContextFromType(context_properties, 518 (cl_device_type) CL_DEVICE_TYPE_GPU,ConvolveNotify,exception,&status); 519 if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS)) 520 convolve_info->context=clCreateContextFromType(context_properties, 521 (cl_device_type) CL_DEVICE_TYPE_CPU,ConvolveNotify,exception,&status); 522 if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS)) 523 convolve_info->context=clCreateContextFromType(context_properties, 524 (cl_device_type) CL_DEVICE_TYPE_DEFAULT,ConvolveNotify,exception,&status); 525 if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS)) 526 { 527 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning, 528 "failed to create OpenCL context","`%s' (%d)",image->filename,status); 529 convolve_info=DestroyConvolveInfo(convolve_info); 530 return((ConvolveInfo *) NULL); 531 } 532 /* 533 Detect OpenCL devices. 534 */ 535 status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,0,NULL, 536 &length); 537 if ((status != CL_SUCCESS) || (length == 0)) 538 { 539 convolve_info=DestroyConvolveInfo(convolve_info); 540 return((ConvolveInfo *) NULL); 541 } 542 convolve_info->devices=(cl_device_id *) AcquireMagickMemory(length); 543 if (convolve_info->devices == (cl_device_id *) NULL) 544 { 545 (void) ThrowMagickException(exception,GetMagickModule(), 546 ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename); 547 convolve_info=DestroyConvolveInfo(convolve_info); 548 return((ConvolveInfo *) NULL); 549 } 550 status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,length, 551 convolve_info->devices,NULL); 552 if (status != CL_SUCCESS) 553 { 554 convolve_info=DestroyConvolveInfo(convolve_info); 555 return((ConvolveInfo *) NULL); 556 } 557 /* 558 Create OpenCL command queue. 559 */ 560 convolve_info->command_queue=clCreateCommandQueue(convolve_info->context, 561 convolve_info->devices[0],0,&status); 562 if ((convolve_info->command_queue == (cl_command_queue) NULL) || 563 (status != CL_SUCCESS)) 564 { 565 convolve_info=DestroyConvolveInfo(convolve_info); 566 return((ConvolveInfo *) NULL); 567 } 568 /* 569 Build OpenCL program. 570 */ 571 convolve_info->program=clCreateProgramWithSource(convolve_info->context,1, 572 &source,lengths,&status); 573 if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS)) 574 { 575 convolve_info=DestroyConvolveInfo(convolve_info); 576 return((ConvolveInfo *) NULL); 577 } 578 (void) FormatLocaleString(options,MaxTextExtent,CLOptions,(float) 579 QuantumRange,MagickEpsilon); 580 status=clBuildProgram(convolve_info->program,1,convolve_info->devices,options, 581 NULL,NULL); 582 if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS)) 583 { 584 char 585 *log; 586 587 status=clGetProgramBuildInfo(convolve_info->program, 588 convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,0,NULL,&length); 589 log=(char *) AcquireMagickMemory(length); 590 if (log == (char *) NULL) 591 { 592 convolve_info=DestroyConvolveInfo(convolve_info); 593 return((ConvolveInfo *) NULL); 594 } 595 status=clGetProgramBuildInfo(convolve_info->program, 596 convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,length,log,&length); 597 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning, 598 "failed to build OpenCL program","`%s' (%s)",image->filename,log); 599 log=DestroyString(log); 600 convolve_info=DestroyConvolveInfo(convolve_info); 601 return((ConvolveInfo *) NULL); 602 } 603 /* 604 Get a kernel object. 605 */ 606 convolve_info->kernel=clCreateKernel(convolve_info->program,name,&status); 607 if ((convolve_info->kernel == (cl_kernel) NULL) || (status != CL_SUCCESS)) 608 { 609 convolve_info=DestroyConvolveInfo(convolve_info); 610 return((ConvolveInfo *) NULL); 611 } 612 return(convolve_info); 613} 614 615#endif 616 617MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image, 618 const KernelInfo *kernel,Image *convolve_image,ExceptionInfo *exception) 619{ 620 assert(image != (Image *) NULL); 621 assert(image->signature == MagickSignature); 622 if (image->debug != MagickFalse) 623 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename); 624 assert(kernel != (KernelInfo *) NULL); 625 assert(kernel->signature == MagickSignature); 626 assert(convolve_image != (Image *) NULL); 627 assert(convolve_image->signature == MagickSignature); 628 assert(exception != (ExceptionInfo *) NULL); 629 assert(exception->signature == MagickSignature); 630 if ((image->storage_class != DirectClass) || 631 (image->colorspace == CMYKColorspace)) 632 if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) && 633 (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod)) 634 return(MagickFalse); 635#if !defined(MAGICKCORE_OPENCL_SUPPORT) 636 return(MagickFalse); 637#else 638 { 639 const void 640 *pixels; 641 642 float 643 *filter; 644 645 ConvolveInfo 646 *convolve_info; 647 648 MagickBooleanType 649 status; 650 651 MagickSizeType 652 length; 653 654 register ssize_t 655 i; 656 657 void 658 *convolve_pixels; 659 660 convolve_info=GetConvolveInfo(image,"Convolve",ConvolveKernel,exception); 661 if (convolve_info == (ConvolveInfo *) NULL) 662 return(MagickFalse); 663 pixels=AcquirePixelCachePixels(image,&length,exception); 664 if (pixels == (const void *) NULL) 665 { 666 convolve_info=DestroyConvolveInfo(convolve_info); 667 (void) ThrowMagickException(exception,GetMagickModule(),CacheError, 668 "UnableToReadPixelCache","`%s'",image->filename); 669 return(MagickFalse); 670 } 671 convolve_pixels=GetPixelCachePixels(convolve_image,&length,exception); 672 if (convolve_pixels == (void *) NULL) 673 { 674 convolve_info=DestroyConvolveInfo(convolve_info); 675 (void) ThrowMagickException(exception,GetMagickModule(),CacheError, 676 "UnableToReadPixelCache","`%s'",image->filename); 677 return(MagickFalse); 678 } 679 filter=(float *) AcquireQuantumMemory(kernel->width,kernel->height* 680 sizeof(*filter)); 681 if (filter == (float *) NULL) 682 { 683 DestroyConvolveBuffers(convolve_info); 684 convolve_info=DestroyConvolveInfo(convolve_info); 685 (void) ThrowMagickException(exception,GetMagickModule(), 686 ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename); 687 return(MagickFalse); 688 } 689 for (i=0; i < (ssize_t) (kernel->width*kernel->height); i++) 690 filter[i]=(float) kernel->values[i]; 691 status=BindConvolveParameters(convolve_info,image,pixels,filter, 692 kernel->width,kernel->height,convolve_pixels); 693 if (status == MagickFalse) 694 { 695 filter=(float *) RelinquishMagickMemory(filter); 696 DestroyConvolveBuffers(convolve_info); 697 convolve_info=DestroyConvolveInfo(convolve_info); 698 return(MagickFalse); 699 } 700 status=EnqueueConvolveKernel(convolve_info,image,pixels,filter, 701 kernel->width,kernel->height,convolve_pixels); 702 filter=(float *) RelinquishMagickMemory(filter); 703 if (status == MagickFalse) 704 { 705 DestroyConvolveBuffers(convolve_info); 706 convolve_info=DestroyConvolveInfo(convolve_info); 707 return(MagickFalse); 708 } 709 DestroyConvolveBuffers(convolve_info); 710 convolve_info=DestroyConvolveInfo(convolve_info); 711 return(MagickTrue); 712 } 713#endif 714} 715