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