accelerate.c revision dc9c80d22d7aacc8604feb4fc1ffea13b230fb13
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 CLPixelInfo cl_float4 122#else 123#if (MAGICKCORE_QUANTUM_DEPTH == 8) 124#define CLOptions "-DCLQuantum=uchar -DCLPixelType=uchar4 " \ 125 "-DQuantumRange=%g -DMagickEpsilon=%g" 126#define CLPixelInfo cl_uchar4 127#elif (MAGICKCORE_QUANTUM_DEPTH == 16) 128#define CLOptions "-DCLQuantum=ushort -DCLPixelType=ushort4 " \ 129 "-DQuantumRange=%g -DMagickEpsilon=%g" 130#define CLPixelInfo cl_ushort4 131#elif (MAGICKCORE_QUANTUM_DEPTH == 32) 132#define CLOptions "-DCLQuantum=uint -DCLPixelType=uint4 " \ 133 "-DQuantumRange=%g -DMagickEpsilon=%g" 134#define CLPixelInfo cl_uint4 135#elif (MAGICKCORE_QUANTUM_DEPTH == 64) 136#define CLOptions "-DCLQuantum=ussize_t -DCLPixelType=ussize_t4 " \ 137 "-DQuantumRange=%g -DMagickEpsilon=%g" 138#define CLPixelInfo 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 const 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(CLPixelInfo), 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(CLPixelInfo),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 status=0; 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 (void) status; 407} 408 409static ConvolveInfo *DestroyConvolveInfo(ConvolveInfo *convolve_info) 410{ 411 cl_int 412 status; 413 414 status=0; 415 if (convolve_info->kernel != (cl_kernel) NULL) 416 status=clReleaseKernel(convolve_info->kernel); 417 if (convolve_info->program != (cl_program) NULL) 418 status=clReleaseProgram(convolve_info->program); 419 if (convolve_info->command_queue != (cl_command_queue) NULL) 420 status=clReleaseCommandQueue(convolve_info->command_queue); 421 if (convolve_info->context != (cl_context) NULL) 422 status=clReleaseContext(convolve_info->context); 423 (void) status; 424 convolve_info=(ConvolveInfo *) RelinquishMagickMemory(convolve_info); 425 return(convolve_info); 426} 427 428static MagickBooleanType EnqueueConvolveKernel(ConvolveInfo *convolve_info, 429 const Image *image,const void *pixels,float *filter,const size_t width, 430 const size_t height,void *convolve_pixels) 431{ 432 cl_int 433 status; 434 435 size_t 436 global_work_size[2], 437 length; 438 439 length=image->columns*image->rows; 440 status=clEnqueueWriteBuffer(convolve_info->command_queue, 441 convolve_info->pixels,CL_TRUE,0,length*sizeof(CLPixelInfo),pixels,0,NULL, 442 NULL); 443 length=width*height; 444 status=clEnqueueWriteBuffer(convolve_info->command_queue, 445 convolve_info->filter,CL_TRUE,0,length*sizeof(cl_float),filter,0,NULL, 446 NULL); 447 if (status != CL_SUCCESS) 448 return(MagickFalse); 449 global_work_size[0]=image->columns; 450 global_work_size[1]=image->rows; 451 status=clEnqueueNDRangeKernel(convolve_info->command_queue, 452 convolve_info->kernel,2,NULL,global_work_size,NULL,0,NULL,NULL); 453 if (status != CL_SUCCESS) 454 return(MagickFalse); 455 length=image->columns*image->rows; 456 status=clEnqueueReadBuffer(convolve_info->command_queue, 457 convolve_info->convolve_pixels,CL_TRUE,0,length*sizeof(CLPixelInfo), 458 convolve_pixels,0,NULL,NULL); 459 if (status != CL_SUCCESS) 460 return(MagickFalse); 461 status=clFinish(convolve_info->command_queue); 462 if (status != CL_SUCCESS) 463 return(MagickFalse); 464 return(MagickTrue); 465} 466 467static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name, 468 const char *source,ExceptionInfo *exception) 469{ 470 char 471 options[MaxTextExtent]; 472 473 cl_context_properties 474 context_properties[3]; 475 476 cl_int 477 status; 478 479 cl_platform_id 480 platforms[1]; 481 482 cl_uint 483 number_platforms; 484 485 ConvolveInfo 486 *convolve_info; 487 488 size_t 489 length, 490 lengths[] = { strlen(source) }; 491 492 /* 493 Create OpenCL info. 494 */ 495 convolve_info=(ConvolveInfo *) AcquireMagickMemory(sizeof(*convolve_info)); 496 if (convolve_info == (ConvolveInfo *) NULL) 497 { 498 (void) ThrowMagickException(exception,GetMagickModule(), 499 ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename); 500 return((ConvolveInfo *) NULL); 501 } 502 (void) ResetMagickMemory(convolve_info,0,sizeof(*convolve_info)); 503 /* 504 Create OpenCL context. 505 */ 506 status=clGetPlatformIDs(0,(cl_platform_id *) NULL,&number_platforms); 507 if ((status == CL_SUCCESS) && (number_platforms > 0)) 508 status=clGetPlatformIDs(1,platforms,NULL); 509 if (status != CL_SUCCESS) 510 { 511 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning, 512 "failed to create OpenCL context","`%s' (%d)",image->filename,status); 513 convolve_info=DestroyConvolveInfo(convolve_info); 514 return((ConvolveInfo *) NULL); 515 } 516 context_properties[0]=CL_CONTEXT_PLATFORM; 517 context_properties[1]=(cl_context_properties) platforms[0]; 518 context_properties[2]=0; 519 convolve_info->context=clCreateContextFromType(context_properties, 520 (cl_device_type) CL_DEVICE_TYPE_GPU,ConvolveNotify,exception,&status); 521 if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS)) 522 convolve_info->context=clCreateContextFromType(context_properties, 523 (cl_device_type) CL_DEVICE_TYPE_CPU,ConvolveNotify,exception,&status); 524 if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS)) 525 convolve_info->context=clCreateContextFromType(context_properties, 526 (cl_device_type) CL_DEVICE_TYPE_DEFAULT,ConvolveNotify,exception,&status); 527 if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS)) 528 { 529 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning, 530 "failed to create OpenCL context","`%s' (%d)",image->filename,status); 531 convolve_info=DestroyConvolveInfo(convolve_info); 532 return((ConvolveInfo *) NULL); 533 } 534 /* 535 Detect OpenCL devices. 536 */ 537 status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,0,NULL, 538 &length); 539 if ((status != CL_SUCCESS) || (length == 0)) 540 { 541 convolve_info=DestroyConvolveInfo(convolve_info); 542 return((ConvolveInfo *) NULL); 543 } 544 convolve_info->devices=(cl_device_id *) AcquireMagickMemory(length); 545 if (convolve_info->devices == (cl_device_id *) NULL) 546 { 547 (void) ThrowMagickException(exception,GetMagickModule(), 548 ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename); 549 convolve_info=DestroyConvolveInfo(convolve_info); 550 return((ConvolveInfo *) NULL); 551 } 552 status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,length, 553 convolve_info->devices,NULL); 554 if (status != CL_SUCCESS) 555 { 556 convolve_info=DestroyConvolveInfo(convolve_info); 557 return((ConvolveInfo *) NULL); 558 } 559 if (image->debug != MagickFalse) 560 { 561 char 562 attribute[MaxTextExtent]; 563 564 size_t 565 length; 566 567 clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_NAME, 568 sizeof(attribute),attribute,&length); 569 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Name: %s", 570 attribute); 571 clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_VENDOR, 572 sizeof(attribute),attribute,&length); 573 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Vendor: %s", 574 attribute); 575 clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_VERSION, 576 sizeof(attribute),attribute,&length); 577 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(), 578 "Driver Version: %s",attribute); 579 clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_PROFILE, 580 sizeof(attribute),attribute,&length); 581 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Profile: %s", 582 attribute); 583 clGetDeviceInfo(convolve_info->devices[0],CL_DRIVER_VERSION, 584 sizeof(attribute),attribute,&length); 585 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Driver: %s", 586 attribute); 587 clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_EXTENSIONS, 588 sizeof(attribute),attribute,&length); 589 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Extensions: %s", 590 attribute); 591 } 592 /* 593 Create OpenCL command queue. 594 */ 595 convolve_info->command_queue=clCreateCommandQueue(convolve_info->context, 596 convolve_info->devices[0],0,&status); 597 if ((convolve_info->command_queue == (cl_command_queue) NULL) || 598 (status != CL_SUCCESS)) 599 { 600 convolve_info=DestroyConvolveInfo(convolve_info); 601 return((ConvolveInfo *) NULL); 602 } 603 /* 604 Build OpenCL program. 605 */ 606 convolve_info->program=clCreateProgramWithSource(convolve_info->context,1, 607 &source,lengths,&status); 608 if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS)) 609 { 610 convolve_info=DestroyConvolveInfo(convolve_info); 611 return((ConvolveInfo *) NULL); 612 } 613 (void) FormatLocaleString(options,MaxTextExtent,CLOptions,(float) 614 QuantumRange,MagickEpsilon); 615 status=clBuildProgram(convolve_info->program,1,convolve_info->devices,options, 616 NULL,NULL); 617 if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS)) 618 { 619 char 620 *log; 621 622 status=clGetProgramBuildInfo(convolve_info->program, 623 convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,0,NULL,&length); 624 log=(char *) AcquireMagickMemory(length); 625 if (log == (char *) NULL) 626 { 627 convolve_info=DestroyConvolveInfo(convolve_info); 628 return((ConvolveInfo *) NULL); 629 } 630 status=clGetProgramBuildInfo(convolve_info->program, 631 convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,length,log,&length); 632 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning, 633 "failed to build OpenCL program","`%s' (%s)",image->filename,log); 634 log=DestroyString(log); 635 convolve_info=DestroyConvolveInfo(convolve_info); 636 return((ConvolveInfo *) NULL); 637 } 638 /* 639 Get a kernel object. 640 */ 641 convolve_info->kernel=clCreateKernel(convolve_info->program,name,&status); 642 if ((convolve_info->kernel == (cl_kernel) NULL) || (status != CL_SUCCESS)) 643 { 644 convolve_info=DestroyConvolveInfo(convolve_info); 645 return((ConvolveInfo *) NULL); 646 } 647 return(convolve_info); 648} 649 650#endif 651 652MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image, 653 const KernelInfo *kernel,Image *convolve_image,ExceptionInfo *exception) 654{ 655 assert(image != (Image *) NULL); 656 assert(image->signature == MagickSignature); 657 if (image->debug != MagickFalse) 658 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename); 659 assert(kernel != (KernelInfo *) NULL); 660 assert(kernel->signature == MagickSignature); 661 assert(convolve_image != (Image *) NULL); 662 assert(convolve_image->signature == MagickSignature); 663 assert(exception != (ExceptionInfo *) NULL); 664 assert(exception->signature == MagickSignature); 665 if ((image->storage_class != DirectClass) || 666 (image->colorspace == CMYKColorspace)) 667 return(MagickFalse); 668 if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) && 669 (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod)) 670 return(MagickFalse); 671 if (GetPixelChannels(image) != 4) 672 return(MagickFalse); 673#if !defined(MAGICKCORE_OPENCL_SUPPORT) 674 return(MagickFalse); 675#else 676 { 677 const void 678 *pixels; 679 680 float 681 *filter; 682 683 ConvolveInfo 684 *convolve_info; 685 686 MagickBooleanType 687 status; 688 689 MagickSizeType 690 length; 691 692 register ssize_t 693 i; 694 695 void 696 *convolve_pixels; 697 698 convolve_info=GetConvolveInfo(image,"Convolve",ConvolveKernel,exception); 699 if (convolve_info == (ConvolveInfo *) NULL) 700 return(MagickFalse); 701 pixels=AcquirePixelCachePixels(image,&length,exception); 702 if (pixels == (const void *) NULL) 703 { 704 convolve_info=DestroyConvolveInfo(convolve_info); 705 (void) ThrowMagickException(exception,GetMagickModule(),CacheError, 706 "UnableToReadPixelCache","`%s'",image->filename); 707 return(MagickFalse); 708 } 709 convolve_pixels=GetPixelCachePixels(convolve_image,&length,exception); 710 if (convolve_pixels == (void *) NULL) 711 { 712 convolve_info=DestroyConvolveInfo(convolve_info); 713 (void) ThrowMagickException(exception,GetMagickModule(),CacheError, 714 "UnableToReadPixelCache","`%s'",image->filename); 715 return(MagickFalse); 716 } 717 filter=(float *) AcquireQuantumMemory(kernel->width,kernel->height* 718 sizeof(*filter)); 719 if (filter == (float *) NULL) 720 { 721 DestroyConvolveBuffers(convolve_info); 722 convolve_info=DestroyConvolveInfo(convolve_info); 723 (void) ThrowMagickException(exception,GetMagickModule(), 724 ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename); 725 return(MagickFalse); 726 } 727 for (i=0; i < (ssize_t) (kernel->width*kernel->height); i++) 728 filter[i]=(float) kernel->values[i]; 729 status=BindConvolveParameters(convolve_info,image,pixels,filter, 730 kernel->width,kernel->height,convolve_pixels); 731 if (status == MagickFalse) 732 { 733 filter=(float *) RelinquishMagickMemory(filter); 734 DestroyConvolveBuffers(convolve_info); 735 convolve_info=DestroyConvolveInfo(convolve_info); 736 return(MagickFalse); 737 } 738 status=EnqueueConvolveKernel(convolve_info,image,pixels,filter, 739 kernel->width,kernel->height,convolve_pixels); 740 filter=(float *) RelinquishMagickMemory(filter); 741 if (status == MagickFalse) 742 { 743 DestroyConvolveBuffers(convolve_info); 744 convolve_info=DestroyConvolveInfo(convolve_info); 745 return(MagickFalse); 746 } 747 DestroyConvolveBuffers(convolve_info); 748 convolve_info=DestroyConvolveInfo(convolve_info); 749 return(MagickTrue); 750 } 751#endif 752} 753