accelerate.c revision d7a9cc48ee5c6c2003db9ed06ca88cbff3d97130
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% Cristy % 17% SiuChi Chan % 18% Guansong Zhang % 19% January 2010 % 20% % 21% % 22% Copyright 1999-2016 ImageMagick Studio LLC, a non-profit organization % 23% dedicated to making software imaging solutions freely available. % 24% % 25% You may not use this file except in compliance with the License. You may % 26% obtain a copy of the License at % 27% % 28% http://www.imagemagick.org/script/license.php % 29% % 30% Unless required by applicable law or agreed to in writing, software % 31% distributed under the License is distributed on an "AS IS" BASIS, % 32% WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. % 33% See the License for the specific language governing permissions and % 34% limitations under the License. % 35% % 36%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 37*/ 38 39/* 40Include declarations. 41*/ 42#include "MagickCore/studio.h" 43#include "MagickCore/accelerate.h" 44#include "MagickCore/accelerate-private.h" 45#include "MagickCore/artifact.h" 46#include "MagickCore/cache.h" 47#include "MagickCore/cache-private.h" 48#include "MagickCore/cache-view.h" 49#include "MagickCore/color-private.h" 50#include "MagickCore/delegate-private.h" 51#include "MagickCore/enhance.h" 52#include "MagickCore/exception.h" 53#include "MagickCore/exception-private.h" 54#include "MagickCore/gem.h" 55#include "MagickCore/image.h" 56#include "MagickCore/image-private.h" 57#include "MagickCore/linked-list.h" 58#include "MagickCore/list.h" 59#include "MagickCore/memory_.h" 60#include "MagickCore/monitor-private.h" 61#include "MagickCore/accelerate.h" 62#include "MagickCore/opencl.h" 63#include "MagickCore/opencl-private.h" 64#include "MagickCore/option.h" 65#include "MagickCore/pixel-accessor.h" 66#include "MagickCore/pixel-private.h" 67#include "MagickCore/prepress.h" 68#include "MagickCore/quantize.h" 69#include "MagickCore/quantum-private.h" 70#include "MagickCore/random_.h" 71#include "MagickCore/random-private.h" 72#include "MagickCore/registry.h" 73#include "MagickCore/resize.h" 74#include "MagickCore/resize-private.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#ifdef MAGICKCORE_CLPERFMARKER 83#include "CLPerfMarker.h" 84#endif 85 86#define MAGICK_MAX(x,y) (((x) >= (y))?(x):(y)) 87#define MAGICK_MIN(x,y) (((x) <= (y))?(x):(y)) 88 89#if defined(MAGICKCORE_OPENCL_SUPPORT) 90 91/* 92 Define declarations. 93*/ 94#define ALIGNED(pointer,type) ((((size_t)(pointer)) & (sizeof(type)-1)) == 0) 95 96/* 97 Static declarations. 98*/ 99static const ResizeWeightingFunctionType supportedResizeWeighting[] = 100{ 101 BoxWeightingFunction, 102 TriangleWeightingFunction, 103 HannWeightingFunction, 104 HammingWeightingFunction, 105 BlackmanWeightingFunction, 106 CubicBCWeightingFunction, 107 SincWeightingFunction, 108 SincFastWeightingFunction, 109 LastWeightingFunction 110}; 111 112/* 113 Forward declarations. 114*/ 115static Image *ComputeUnsharpMaskImageSingle(const Image *image, 116 const double radius,const double sigma,const double gain, 117 const double threshold,int blurOnly,ExceptionInfo *exception); 118 119/* 120 Helper functions. 121*/ 122static MagickBooleanType checkAccelerateCondition(const Image* image) 123{ 124 /* check if the image's colorspace is supported */ 125 if (image->colorspace != RGBColorspace && 126 image->colorspace != sRGBColorspace && 127 image->colorspace != GRAYColorspace) 128 return(MagickFalse); 129 130 /* check if the virtual pixel method is compatible with the OpenCL implementation */ 131 if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) && 132 (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod)) 133 return(MagickFalse); 134 135 /* check if the image has read / write mask */ 136 if (image->read_mask != MagickFalse || image->write_mask != MagickFalse) 137 return(MagickFalse); 138 139 if (image->number_channels > 4) 140 return(MagickFalse); 141 142 /* check if pixel order is R */ 143 if (GetPixelChannelOffset(image,RedPixelChannel) != 0) 144 return(MagickFalse); 145 146 if (image->number_channels == 1) 147 return(MagickTrue); 148 149 /* check if pixel order is RA */ 150 if ((image->number_channels == 2) && 151 (GetPixelChannelOffset(image,AlphaPixelChannel) == 1)) 152 return(MagickTrue); 153 154 if (image->number_channels == 2) 155 return(MagickFalse); 156 157 /* check if pixel order is RGB */ 158 if ((GetPixelChannelOffset(image,GreenPixelChannel) != 1) || 159 (GetPixelChannelOffset(image,BluePixelChannel) != 2)) 160 return(MagickFalse); 161 162 if (image->number_channels == 3) 163 return(MagickTrue); 164 165 /* check if pixel order is RGBA */ 166 if (GetPixelChannelOffset(image,AlphaPixelChannel) != 3) 167 return(MagickFalse); 168 169 return(MagickTrue); 170} 171 172static MagickBooleanType checkAccelerateConditionRGBA(const Image* image) 173{ 174 if (checkAccelerateCondition(image) == MagickFalse) 175 return(MagickFalse); 176 177 /* the order will be RGBA if the image has 4 channels */ 178 if (image->number_channels != 4) 179 return(MagickFalse); 180 181 if ((GetPixelRedTraits(image) == UndefinedPixelTrait) || 182 (GetPixelGreenTraits(image) == UndefinedPixelTrait) || 183 (GetPixelBlueTraits(image) == UndefinedPixelTrait) || 184 (GetPixelAlphaTraits(image) == UndefinedPixelTrait)) 185 return(MagickFalse); 186 187 return(MagickTrue); 188} 189 190static MagickBooleanType checkPixelIntensity(const Image *image, 191 const PixelIntensityMethod method) 192{ 193 /* EncodePixelGamma and DecodePixelGamma are not supported */ 194 if ((method == Rec601LumaPixelIntensityMethod) || 195 (method == Rec709LumaPixelIntensityMethod)) 196 { 197 if (image->colorspace == RGBColorspace) 198 return(MagickFalse); 199 } 200 201 if ((method == Rec601LuminancePixelIntensityMethod) || 202 (method == Rec709LuminancePixelIntensityMethod)) 203 { 204 if (image->colorspace == sRGBColorspace) 205 return(MagickFalse); 206 } 207 208 return(MagickTrue); 209} 210 211static MagickBooleanType checkHistogramCondition(const Image *image, 212 const PixelIntensityMethod method) 213{ 214 /* ensure this is the only pass get in for now. */ 215 if ((image->channel_mask & SyncChannels) == 0) 216 return MagickFalse; 217 218 return(checkPixelIntensity(image,method)); 219} 220 221static MagickBooleanType checkOpenCLEnvironment(ExceptionInfo* exception) 222{ 223 MagickBooleanType 224 flag; 225 226 MagickCLEnv 227 clEnv; 228 229 clEnv=GetDefaultOpenCLEnv(); 230 231 GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED, 232 sizeof(MagickBooleanType),&flag,exception); 233 if (flag != MagickFalse) 234 return(MagickFalse); 235 236 GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED, 237 sizeof(MagickBooleanType),&flag,exception); 238 if (flag == MagickFalse) 239 { 240 if (InitOpenCLEnv(clEnv,exception) == MagickFalse) 241 return(MagickFalse); 242 243 GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED, 244 sizeof(MagickBooleanType),&flag,exception); 245 if (flag != MagickFalse) 246 return(MagickFalse); 247 } 248 249 return(MagickTrue); 250} 251 252/* pad the global workgroup size to the next multiple of 253 the local workgroup size */ 254inline static unsigned int padGlobalWorkgroupSizeToLocalWorkgroupSize( 255 const unsigned int orgGlobalSize,const unsigned int localGroupSize) 256{ 257 return ((orgGlobalSize+(localGroupSize-1))/localGroupSize*localGroupSize); 258} 259 260static cl_mem createBuffer(const Image *image,CacheView *image_view, 261 MagickCLEnv clEnv,cl_context context,cl_mem_flags flags,void *pixels, 262 ExceptionInfo *exception) 263{ 264 cl_mem 265 buffer; 266 267 cl_mem_flags 268 mem_flags; 269 270 cl_int 271 status; 272 273 size_t 274 length; 275 276 void 277 *hostPtr; 278 279 pixels=(void *) GetCacheViewVirtualPixels(image_view,0,0,image->columns, 280 image->rows,exception); 281 if (pixels == (void *) NULL) 282 { 283 (void) OpenCLThrowMagickException(exception,GetMagickModule(), 284 CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); 285 return (cl_mem) NULL; 286 } 287 288 mem_flags=flags; 289 hostPtr=pixels; 290 if (ALIGNED(pixels,CLQuantum)) 291 mem_flags=mem_flags | CL_MEM_USE_HOST_PTR; 292 else if ((mem_flags == CL_MEM_READ_ONLY) || (mem_flags == CL_MEM_READ_WRITE)) 293 mem_flags=mem_flags | CL_MEM_COPY_HOST_PTR; 294 else if (mem_flags == CL_MEM_WRITE_ONLY) 295 hostPtr=NULL; 296 297 length=image->columns*image->rows*image->number_channels; 298 buffer=clEnv->library->clCreateBuffer(context,mem_flags,length* 299 sizeof(CLQuantum),hostPtr,&status); 300 if (status != CL_SUCCESS) 301 { 302 (void) OpenCLThrowMagickException(exception,GetMagickModule(), 303 ResourceLimitWarning,"clEnv->library->clCreateBuffer failed.","."); 304 } 305 306 return(buffer); 307} 308 309static inline cl_mem createReadBuffer(const Image *image,CacheView *image_view, 310 MagickCLEnv clEnv,cl_context context,ExceptionInfo *exception) 311{ 312 void 313 *pixels; 314 315 pixels=(void *) NULL; 316 return(createBuffer(image,image_view,clEnv,context,CL_MEM_READ_ONLY, 317 pixels,exception)); 318} 319 320static inline cl_mem createReadWriteBuffer(const Image *image, 321 CacheView *image_view,MagickCLEnv clEnv,cl_context context,void *pixels, 322 ExceptionInfo *exception) 323{ 324 return(createBuffer(image,image_view,clEnv,context,CL_MEM_READ_WRITE,pixels, 325 exception)); 326} 327 328static inline cl_mem createWriteBuffer(Image *image,CacheView *image_view, 329 MagickCLEnv clEnv,cl_context context,void *pixels,ExceptionInfo *exception) 330{ 331 return(createBuffer(image,image_view,clEnv,context,CL_MEM_WRITE_ONLY,pixels, 332 exception)); 333} 334 335static inline MagickBooleanType copyWriteBuffer(const Image *image, 336 MagickCLEnv clEnv,cl_command_queue queue,cl_mem buffer,void *pixels, 337 ExceptionInfo *exception) 338{ 339 cl_int 340 status; 341 342 size_t 343 length; 344 345 length=image->columns*image->rows*image->number_channels; 346 if (ALIGNED(pixels,CLQuantum)) 347 clEnv->library->clEnqueueMapBuffer(queue,buffer,CL_TRUE,CL_MAP_READ | 348 CL_MAP_WRITE,0,length*sizeof(CLQuantum),0,NULL,NULL,&status); 349 else 350 status=clEnv->library->clEnqueueReadBuffer(queue,buffer,CL_TRUE,0,length* 351 sizeof(CLQuantum),pixels,0,NULL,NULL); 352 if (status != CL_SUCCESS) 353 { 354 (void) OpenCLThrowMagickException(exception,GetMagickModule(), 355 ResourceLimitWarning,"Reading output image from CL buffer failed.", 356 "'%s'","."); 357 return(MagickFalse); 358 } 359 return(MagickTrue); 360} 361 362static cl_mem createKernelInfo(MagickCLEnv clEnv,cl_context context, 363 cl_command_queue queue,const double radius,const double sigma,cl_uint *width, 364 ExceptionInfo *exception) 365{ 366 char 367 geometry[MagickPathExtent]; 368 369 cl_int 370 status; 371 372 cl_mem 373 imageKernelBuffer; 374 375 float 376 *kernelBufferPtr; 377 378 KernelInfo 379 *kernel; 380 381 size_t 382 i; 383 384 (void) FormatLocaleString(geometry,MagickPathExtent, 385 "blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma); 386 kernel=AcquireKernelInfo(geometry,exception); 387 if (kernel == (KernelInfo *) NULL) 388 { 389 (void) OpenCLThrowMagickException(exception,GetMagickModule(), 390 ResourceLimitWarning,"AcquireKernelInfo failed.","."); 391 return((cl_mem) NULL); 392 } 393 394 imageKernelBuffer=clEnv->library->clCreateBuffer(context,CL_MEM_READ_ONLY, 395 kernel->width*sizeof(float),NULL,&status); 396 if (status != CL_SUCCESS) 397 { 398 kernel=DestroyKernelInfo(kernel); 399 (void) OpenCLThrowMagickException(exception,GetMagickModule(), 400 ResourceLimitWarning,"clEnv->library->clCreateBuffer failed.","."); 401 return((cl_mem) NULL); 402 } 403 404 kernelBufferPtr=(float*)clEnv->library->clEnqueueMapBuffer(queue, 405 imageKernelBuffer,CL_TRUE,CL_MAP_WRITE,0,kernel->width*sizeof(float),0, 406 NULL,NULL,&status); 407 if (status != CL_SUCCESS) 408 { 409 kernel=DestroyKernelInfo(kernel); 410 clEnv->library->clReleaseMemObject(imageKernelBuffer); 411 (void) OpenCLThrowMagickException(exception,GetMagickModule(), 412 ResourceLimitWarning,"clEnv->library->clEnqueueMapBuffer failed.","."); 413 return((cl_mem) NULL); 414 } 415 for (i = 0; i < kernel->width; i++) 416 { 417 kernelBufferPtr[i]=(float)kernel->values[i]; 418 } 419 420 *width=(cl_uint) kernel->width; 421 kernel=DestroyKernelInfo(kernel); 422 423 status=clEnv->library->clEnqueueUnmapMemObject(queue,imageKernelBuffer, 424 kernelBufferPtr,0,NULL,NULL); 425 if (status != CL_SUCCESS) 426 { 427 clEnv->library->clReleaseMemObject(imageKernelBuffer); 428 (void) OpenCLThrowMagickException(exception,GetMagickModule(), 429 ResourceLimitWarning,"clEnv->library->clEnqueueUnmapMemObject failed.", 430 "'%s'","."); 431 return((cl_mem) NULL); 432 } 433 return(imageKernelBuffer); 434} 435 436/* 437%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 438% % 439% % 440% % 441% A c c e l e r a t e A d d N o i s e I m a g e % 442% % 443% % 444% % 445%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 446*/ 447 448static Image *ComputeAddNoiseImage(const Image *image, 449 const NoiseType noise_type,ExceptionInfo *exception) 450{ 451 CacheView 452 *filteredImage_view, 453 *image_view; 454 455 cl_command_queue 456 queue; 457 458 cl_context 459 context; 460 461 cl_float 462 attenuate; 463 464 cl_int 465 clStatus; 466 467 cl_kernel 468 addNoiseKernel; 469 470 cl_event 471 event; 472 473 cl_mem 474 filteredImageBuffer, 475 imageBuffer; 476 477 cl_uint 478 bufferLength, 479 inputPixelCount, 480 number_channels, 481 numRandomNumberPerPixel, 482 pixelsPerWorkitem, 483 seed0, 484 seed1; 485 486 const char 487 *option; 488 489 MagickBooleanType 490 outputReady; 491 492 MagickCLEnv 493 clEnv; 494 495 Image 496 *filteredImage; 497 498 size_t 499 global_work_size[1], 500 local_work_size[1]; 501 502 unsigned int 503 k; 504 505 void 506 *filteredPixels; 507 508 outputReady = MagickFalse; 509 filteredImage = NULL; 510 filteredImage_view = NULL; 511 filteredPixels = NULL; 512 filteredImageBuffer = NULL; 513 addNoiseKernel = NULL; 514 515 clEnv = GetDefaultOpenCLEnv(); 516 context = GetOpenCLContext(clEnv); 517 queue = AcquireOpenCLCommandQueue(clEnv); 518 519 image_view=AcquireVirtualCacheView(image,exception); 520 imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); 521 if (imageBuffer == (cl_mem) NULL) 522 goto cleanup; 523 524 filteredImage=CloneImage(image,image->columns,image->rows,MagickTrue,exception); 525 if (filteredImage == (Image *) NULL) 526 goto cleanup; 527 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) 528 { 529 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); 530 goto cleanup; 531 } 532 533 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); 534 filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, 535 context,filteredPixels,exception); 536 if (filteredImageBuffer == (void *) NULL) 537 goto cleanup; 538 539 /* find out how many random numbers needed by pixel */ 540 numRandomNumberPerPixel = 0; 541 { 542 unsigned int numRandPerChannel = 0; 543 switch (noise_type) 544 { 545 case UniformNoise: 546 case ImpulseNoise: 547 case LaplacianNoise: 548 case RandomNoise: 549 default: 550 numRandPerChannel = 1; 551 break; 552 case GaussianNoise: 553 case MultiplicativeGaussianNoise: 554 case PoissonNoise: 555 numRandPerChannel = 2; 556 break; 557 }; 558 559 if (GetPixelRedTraits(image) != UndefinedPixelTrait) 560 numRandomNumberPerPixel+=numRandPerChannel; 561 if (GetPixelGreenTraits(image) != UndefinedPixelTrait) 562 numRandomNumberPerPixel+=numRandPerChannel; 563 if (GetPixelBlueTraits(image) != UndefinedPixelTrait) 564 numRandomNumberPerPixel+=numRandPerChannel; 565 if (GetPixelAlphaTraits(image) != UndefinedPixelTrait) 566 numRandomNumberPerPixel+=numRandPerChannel; 567 } 568 569 addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"AddNoise"); 570 if (addNoiseKernel == NULL) 571 { 572 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 573 goto cleanup; 574 } 575 576 { 577 cl_uint computeUnitCount; 578 cl_uint workItemCount; 579 clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &computeUnitCount, NULL); 580 workItemCount = computeUnitCount * 2 * 256; // 256 work items per group, 2 groups per CU 581 inputPixelCount = (cl_int) (image->columns * image->rows); 582 pixelsPerWorkitem = (inputPixelCount + workItemCount - 1) / workItemCount; 583 pixelsPerWorkitem = ((pixelsPerWorkitem + 3) / 4) * 4; 584 585 local_work_size[0] = 256; 586 global_work_size[0] = workItemCount; 587 } 588 { 589 RandomInfo* randomInfo = AcquireRandomInfo(); 590 const unsigned long* s = GetRandomInfoSeed(randomInfo); 591 seed0 = s[0]; 592 (void) GetPseudoRandomValue(randomInfo); 593 seed1 = s[0]; 594 randomInfo = DestroyRandomInfo(randomInfo); 595 } 596 597 number_channels = (cl_uint) image->number_channels; 598 bufferLength = (cl_uint)(image->columns * image->rows * image->number_channels); 599 attenuate=1.0f; 600 option=GetImageArtifact(image,"attenuate"); 601 if (option != (char *) NULL) 602 attenuate=(float)StringToDouble(option,(char **) NULL); 603 604 k = 0; 605 clStatus=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&imageBuffer); 606 clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&number_channels); 607 clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(ChannelType),(void *)&image->channel_mask); 608 clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&bufferLength); 609 clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&pixelsPerWorkitem); 610 clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(NoiseType),(void *)&noise_type); 611 clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_float),(void *)&attenuate); 612 clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&seed0); 613 clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&seed1); 614 clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&numRandomNumberPerPixel); 615 clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer); 616 if (clStatus != CL_SUCCESS) 617 { 618 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 619 goto cleanup; 620 } 621 622 clStatus=clEnv->library->clEnqueueNDRangeKernel(queue,addNoiseKernel,1,NULL,global_work_size,local_work_size,0,NULL,&event); 623 if (clStatus != CL_SUCCESS) 624 { 625 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 626 goto cleanup; 627 } 628 629 RecordProfileData(clEnv,AddNoiseKernel,event); 630 clEnv->library->clReleaseEvent(event); 631 if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) 632 goto cleanup; 633 634 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); 635 636cleanup: 637 OpenCLLogException(__FUNCTION__,__LINE__,exception); 638 639 image_view=DestroyCacheView(image_view); 640 if (filteredImage_view != NULL) 641 filteredImage_view=DestroyCacheView(filteredImage_view); 642 643 if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 644 if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel); 645 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); 646 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); 647 if (outputReady == MagickFalse && filteredImage != NULL) 648 filteredImage=DestroyImage(filteredImage); 649 650 return(filteredImage); 651} 652 653MagickExport Image *AccelerateAddNoiseImage(const Image *image, 654 const NoiseType noise_type,ExceptionInfo *exception) 655{ 656 Image 657 *filteredImage; 658 659 assert(image != NULL); 660 assert(exception != (ExceptionInfo *) NULL); 661 662 if ((checkAccelerateCondition(image) == MagickFalse) || 663 (checkOpenCLEnvironment(exception) == MagickFalse)) 664 return((Image *) NULL); 665 666 filteredImage=ComputeAddNoiseImage(image,noise_type,exception); 667 return(filteredImage); 668} 669 670/* 671%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 672% % 673% % 674% % 675% A c c e l e r a t e B l u r I m a g e % 676% % 677% % 678% % 679%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 680*/ 681 682static Image *ComputeBlurImage(const Image* image,const double radius, 683 const double sigma,ExceptionInfo *exception) 684{ 685 CacheView 686 *filteredImage_view, 687 *image_view; 688 689 cl_command_queue 690 queue; 691 692 cl_context 693 context; 694 695 cl_int 696 clStatus; 697 698 cl_kernel 699 blurColumnKernel, 700 blurRowKernel; 701 702 cl_event 703 event; 704 705 cl_mem 706 filteredImageBuffer, 707 imageBuffer, 708 imageKernelBuffer, 709 tempImageBuffer; 710 711 cl_uint 712 imageColumns, 713 imageRows, 714 kernelWidth, 715 number_channels; 716 717 Image 718 *filteredImage; 719 720 MagickBooleanType 721 outputReady; 722 723 MagickCLEnv 724 clEnv; 725 726 MagickSizeType 727 length; 728 729 unsigned int 730 i; 731 732 void 733 *filteredPixels; 734 735 context = NULL; 736 filteredImage = NULL; 737 filteredImage_view = NULL; 738 imageBuffer = NULL; 739 tempImageBuffer = NULL; 740 filteredImageBuffer = NULL; 741 filteredPixels = NULL; 742 imageKernelBuffer = NULL; 743 blurRowKernel = NULL; 744 blurColumnKernel = NULL; 745 queue = NULL; 746 747 outputReady = MagickFalse; 748 749 clEnv = GetDefaultOpenCLEnv(); 750 context = GetOpenCLContext(clEnv); 751 queue = AcquireOpenCLCommandQueue(clEnv); 752 753 image_view=AcquireVirtualCacheView(image,exception); 754 imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); 755 if (imageBuffer == (cl_mem) NULL) 756 goto cleanup; 757 758 filteredImage=CloneImage(image,image->columns,image->rows,MagickTrue,exception); 759 if (filteredImage == (Image *) NULL) 760 goto cleanup; 761 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) 762 { 763 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); 764 goto cleanup; 765 } 766 767 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); 768 filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, 769 context,filteredPixels,exception); 770 if (filteredImageBuffer == (void *) NULL) 771 goto cleanup; 772 773 imageKernelBuffer=createKernelInfo(clEnv,context,queue,radius,sigma, 774 &kernelWidth,exception); 775 776 { 777 /* create temp buffer */ 778 { 779 length = image->columns * image->rows; 780 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(cl_float4), NULL, &clStatus); 781 if (clStatus != CL_SUCCESS) 782 { 783 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 784 goto cleanup; 785 } 786 } 787 788 /* get the OpenCL kernels */ 789 { 790 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow"); 791 if (blurRowKernel == NULL) 792 { 793 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 794 goto cleanup; 795 }; 796 797 blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurColumn"); 798 if (blurColumnKernel == NULL) 799 { 800 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 801 goto cleanup; 802 }; 803 } 804 805 number_channels = (cl_uint) image->number_channels; 806 imageColumns = (cl_uint) image->columns; 807 imageRows = (cl_uint) image->rows; 808 809 { 810 /* need logic to decide this value */ 811 int chunkSize = 256; 812 813 { 814 /* set the kernel arguments */ 815 i = 0; 816 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 817 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),&number_channels); 818 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&image->channel_mask); 819 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer); 820 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&kernelWidth); 821 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageColumns); 822 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageRows); 823 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_float4)*(chunkSize+kernelWidth),(void *) NULL); 824 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer); 825 if (clStatus != CL_SUCCESS) 826 { 827 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 828 goto cleanup; 829 } 830 } 831 832 /* launch the kernel */ 833 { 834 size_t gsize[2]; 835 size_t wsize[2]; 836 837 gsize[0] = chunkSize*((image->columns+chunkSize-1)/chunkSize); 838 gsize[1] = image->rows; 839 wsize[0] = chunkSize; 840 wsize[1] = 1; 841 842 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, &event); 843 if (clStatus != CL_SUCCESS) 844 { 845 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 846 goto cleanup; 847 } 848 clEnv->library->clFlush(queue); 849 RecordProfileData(clEnv,BlurRowKernel,event); 850 clEnv->library->clReleaseEvent(event); 851 } 852 } 853 854 { 855 /* need logic to decide this value */ 856 int chunkSize = 256; 857 858 { 859 /* set the kernel arguments */ 860 i = 0; 861 clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer); 862 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_uint),&number_channels); 863 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(ChannelType),&image->channel_mask); 864 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer); 865 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_uint),(void *)&kernelWidth); 866 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_uint),(void *)&imageColumns); 867 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_uint),(void *)&imageRows); 868 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_float4)*(chunkSize+kernelWidth),(void *) NULL); 869 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); 870 if (clStatus != CL_SUCCESS) 871 { 872 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 873 goto cleanup; 874 } 875 } 876 877 /* launch the kernel */ 878 { 879 size_t gsize[2]; 880 size_t wsize[2]; 881 882 gsize[0] = image->columns; 883 gsize[1] = chunkSize*((image->rows+chunkSize-1)/chunkSize); 884 wsize[0] = 1; 885 wsize[1] = chunkSize; 886 887 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, &event); 888 if (clStatus != CL_SUCCESS) 889 { 890 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 891 goto cleanup; 892 } 893 clEnv->library->clFlush(queue); 894 RecordProfileData(clEnv,BlurColumnKernel,event); 895 clEnv->library->clReleaseEvent(event); 896 } 897 } 898 899 } 900 901 /* get result */ 902 if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) 903 goto cleanup; 904 905 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); 906 907cleanup: 908 OpenCLLogException(__FUNCTION__,__LINE__,exception); 909 910 image_view=DestroyCacheView(image_view); 911 if (filteredImage_view != NULL) 912 filteredImage_view=DestroyCacheView(filteredImage_view); 913 914 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); 915 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer); 916 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); 917 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer); 918 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel); 919 if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel); 920 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 921 if (outputReady == MagickFalse && filteredImage != NULL) 922 filteredImage=DestroyImage(filteredImage); 923 return(filteredImage); 924} 925 926static Image* ComputeBlurImageSingle(const Image* image, 927 const double radius,const double sigma,ExceptionInfo *exception) 928{ 929 return ComputeUnsharpMaskImageSingle(image,radius,sigma,0.0,0.0,1,exception); 930} 931 932MagickExport Image* AccelerateBlurImage(const Image *image, 933 const double radius,const double sigma,ExceptionInfo *exception) 934{ 935 Image 936 *filteredImage; 937 938 assert(image != NULL); 939 assert(exception != (ExceptionInfo *) NULL); 940 941 if ((checkAccelerateCondition(image) == MagickFalse) || 942 (checkOpenCLEnvironment(exception) == MagickFalse)) 943 return NULL; 944 945 if (radius < 12.1) 946 filteredImage=ComputeBlurImageSingle(image,radius,sigma,exception); 947 else 948 filteredImage=ComputeBlurImage(image,radius,sigma,exception); 949 return(filteredImage); 950} 951 952/* 953%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 954% % 955% % 956% % 957% A c c e l e r a t e C o m p o s i t e I m a g e % 958% % 959% % 960% % 961%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 962*/ 963 964static MagickBooleanType LaunchCompositeKernel(MagickCLEnv clEnv, 965 cl_command_queue queue,cl_mem imageBuffer,const unsigned int inputWidth, 966 const unsigned int inputHeight,const unsigned int matte, 967 const ChannelType channel,const CompositeOperator compose, 968 const cl_mem compositeImageBuffer,const unsigned int compositeWidth, 969 const unsigned int compositeHeight,const float destination_dissolve, 970 const float source_dissolve,ExceptionInfo *magick_unused(exception)) 971{ 972 cl_int 973 clStatus; 974 975 cl_kernel 976 compositeKernel; 977 978 cl_event 979 event; 980 981 int 982 k; 983 984 size_t 985 global_work_size[2], 986 local_work_size[2]; 987 988 unsigned int 989 composeOp; 990 991 magick_unreferenced(exception); 992 993 compositeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, 994 "Composite"); 995 996 k = 0; 997 clStatus=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(cl_mem),(void*)&imageBuffer); 998 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&inputWidth); 999 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&inputHeight); 1000 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(cl_mem),(void*)&compositeImageBuffer); 1001 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&compositeWidth); 1002 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&compositeHeight); 1003 composeOp = (unsigned int)compose; 1004 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&composeOp); 1005 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(ChannelType),(void*)&channel); 1006 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&matte); 1007 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(float),(void*)&destination_dissolve); 1008 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(float),(void*)&source_dissolve); 1009 1010 if (clStatus!=CL_SUCCESS) 1011 return MagickFalse; 1012 1013 local_work_size[0] = 64; 1014 local_work_size[1] = 1; 1015 1016 global_work_size[0] = padGlobalWorkgroupSizeToLocalWorkgroupSize(inputWidth, 1017 (unsigned int) local_work_size[0]); 1018 global_work_size[1] = inputHeight; 1019 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, compositeKernel, 2, NULL, 1020 global_work_size, local_work_size, 0, NULL, &event); 1021 1022 RecordProfileData(clEnv,CompositeKernel,event); 1023 clEnv->library->clReleaseEvent(event); 1024 1025 RelinquishOpenCLKernel(clEnv, compositeKernel); 1026 1027 return((clStatus==CL_SUCCESS) ? MagickTrue : MagickFalse); 1028} 1029 1030static MagickBooleanType ComputeCompositeImage(Image *image, 1031 const CompositeOperator compose,const Image *compositeImage, 1032 const float destination_dissolve,const float source_dissolve,ExceptionInfo *exception) 1033{ 1034 CacheView 1035 *image_view; 1036 1037 cl_command_queue 1038 queue; 1039 1040 cl_context 1041 context; 1042 1043 cl_int 1044 clStatus; 1045 1046 cl_mem_flags 1047 mem_flags; 1048 1049 cl_mem 1050 compositeImageBuffer, 1051 imageBuffer; 1052 1053 const void 1054 *composePixels; 1055 1056 MagickBooleanType 1057 outputReady, 1058 status; 1059 1060 MagickCLEnv 1061 clEnv; 1062 1063 MagickSizeType 1064 length; 1065 1066 void 1067 *inputPixels; 1068 1069 status = MagickFalse; 1070 outputReady = MagickFalse; 1071 composePixels = NULL; 1072 imageBuffer = NULL; 1073 compositeImageBuffer = NULL; 1074 1075 clEnv = GetDefaultOpenCLEnv(); 1076 context = GetOpenCLContext(clEnv); 1077 queue = AcquireOpenCLCommandQueue(clEnv); 1078 1079 /* Create and initialize OpenCL buffers. */ 1080 image_view=AcquireAuthenticCacheView(image,exception); 1081 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); 1082 if (inputPixels == (void *) NULL) 1083 { 1084 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, 1085 "UnableToReadPixelCache.","`%s'",image->filename); 1086 goto cleanup; 1087 } 1088 1089 /* If the host pointer is aligned to the size of CLPixelPacket, 1090 then use the host buffer directly from the GPU; otherwise, 1091 create a buffer on the GPU and copy the data over */ 1092 if (ALIGNED(inputPixels,CLPixelPacket)) 1093 { 1094 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR; 1095 } 1096 else 1097 { 1098 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; 1099 } 1100 /* create a CL buffer from image pixel buffer */ 1101 length = image->columns * image->rows; 1102 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, 1103 length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 1104 if (clStatus != CL_SUCCESS) 1105 { 1106 (void) OpenCLThrowMagickException(exception, GetMagickModule(), 1107 ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 1108 goto cleanup; 1109 } 1110 1111 1112 /* Create and initialize OpenCL buffers. */ 1113 composePixels = AcquirePixelCachePixels(compositeImage, &length, exception); 1114 if (composePixels == (void *) NULL) 1115 { 1116 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, 1117 "UnableToReadPixelCache.","`%s'",compositeImage->filename); 1118 goto cleanup; 1119 } 1120 1121 /* If the host pointer is aligned to the size of CLPixelPacket, 1122 then use the host buffer directly from the GPU; otherwise, 1123 create a buffer on the GPU and copy the data over */ 1124 if (ALIGNED(composePixels,CLPixelPacket)) 1125 { 1126 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; 1127 } 1128 else 1129 { 1130 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR; 1131 } 1132 /* create a CL buffer from image pixel buffer */ 1133 length = compositeImage->columns * compositeImage->rows; 1134 compositeImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, 1135 length * sizeof(CLPixelPacket), (void*)composePixels, &clStatus); 1136 if (clStatus != CL_SUCCESS) 1137 { 1138 (void) OpenCLThrowMagickException(exception, GetMagickModule(), 1139 ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 1140 goto cleanup; 1141 } 1142 1143 status = LaunchCompositeKernel(clEnv,queue,imageBuffer, 1144 (unsigned int) image->columns, 1145 (unsigned int) image->rows, 1146 (unsigned int) (image->alpha_trait > CopyPixelTrait) ? 1 : 0, 1147 image->channel_mask, compose, compositeImageBuffer, 1148 (unsigned int) compositeImage->columns, 1149 (unsigned int) compositeImage->rows, 1150 destination_dissolve,source_dissolve, 1151 exception); 1152 1153 if (status==MagickFalse) 1154 goto cleanup; 1155 1156 length = image->columns * image->rows; 1157 if (ALIGNED(inputPixels,CLPixelPacket)) 1158 { 1159 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, 1160 CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, 1161 NULL, &clStatus); 1162 } 1163 else 1164 { 1165 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, 1166 length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL); 1167 } 1168 if (clStatus==CL_SUCCESS) 1169 outputReady=SyncCacheViewAuthenticPixels(image_view,exception); 1170 1171cleanup: 1172 1173 image_view=DestroyCacheView(image_view); 1174 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); 1175 if (compositeImageBuffer!=NULL) clEnv->library->clReleaseMemObject(compositeImageBuffer); 1176 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 1177 1178 return(outputReady); 1179} 1180 1181MagickExport MagickBooleanType AccelerateCompositeImage(Image *image, 1182 const CompositeOperator compose,const Image *composite, 1183 const float destination_dissolve,const float source_dissolve, 1184 ExceptionInfo *exception) 1185{ 1186 MagickBooleanType 1187 status; 1188 1189 assert(image != NULL); 1190 assert(exception != (ExceptionInfo *) NULL); 1191 1192 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 1193 (checkOpenCLEnvironment(exception) == MagickFalse)) 1194 return(MagickFalse); 1195 1196 /* only support images with the size for now */ 1197 if ((image->columns != composite->columns) || 1198 (image->rows != composite->rows)) 1199 return MagickFalse; 1200 1201 switch(compose) 1202 { 1203 case ColorDodgeCompositeOp: 1204 case BlendCompositeOp: 1205 break; 1206 default: 1207 // unsupported compose operator, quit 1208 return MagickFalse; 1209 }; 1210 1211 status=ComputeCompositeImage(image,compose,composite,destination_dissolve, 1212 source_dissolve,exception); 1213 return(status); 1214} 1215 1216/* 1217%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 1218% % 1219% % 1220% % 1221% A c c e l e r a t e C o n t r a s t I m a g e % 1222% % 1223% % 1224% % 1225%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 1226*/ 1227 1228static MagickBooleanType ComputeContrastImage(Image *image, 1229 const MagickBooleanType sharpen,ExceptionInfo *exception) 1230{ 1231 CacheView 1232 *image_view; 1233 1234 cl_command_queue 1235 queue; 1236 1237 cl_context 1238 context; 1239 1240 cl_int 1241 clStatus; 1242 1243 cl_kernel 1244 filterKernel; 1245 1246 cl_event 1247 event; 1248 1249 cl_mem 1250 imageBuffer; 1251 1252 cl_mem_flags 1253 mem_flags; 1254 1255 MagickBooleanType 1256 outputReady; 1257 1258 MagickCLEnv 1259 clEnv; 1260 1261 MagickSizeType 1262 length; 1263 1264 size_t 1265 global_work_size[2]; 1266 1267 unsigned int 1268 i, 1269 uSharpen; 1270 1271 void 1272 *inputPixels; 1273 1274 outputReady = MagickFalse; 1275 clEnv = NULL; 1276 inputPixels = NULL; 1277 context = NULL; 1278 imageBuffer = NULL; 1279 filterKernel = NULL; 1280 queue = NULL; 1281 1282 clEnv = GetDefaultOpenCLEnv(); 1283 context = GetOpenCLContext(clEnv); 1284 1285 /* Create and initialize OpenCL buffers. */ 1286 image_view=AcquireAuthenticCacheView(image,exception); 1287 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); 1288 if (inputPixels == (void *) NULL) 1289 { 1290 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); 1291 goto cleanup; 1292 } 1293 1294 /* If the host pointer is aligned to the size of CLPixelPacket, 1295 then use the host buffer directly from the GPU; otherwise, 1296 create a buffer on the GPU and copy the data over */ 1297 if (ALIGNED(inputPixels,CLPixelPacket)) 1298 { 1299 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR; 1300 } 1301 else 1302 { 1303 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; 1304 } 1305 /* create a CL buffer from image pixel buffer */ 1306 length = image->columns * image->rows; 1307 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 1308 if (clStatus != CL_SUCCESS) 1309 { 1310 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 1311 goto cleanup; 1312 } 1313 1314 filterKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Contrast"); 1315 if (filterKernel == NULL) 1316 { 1317 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 1318 goto cleanup; 1319 } 1320 1321 i = 0; 1322 clStatus=clEnv->library->clSetKernelArg(filterKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 1323 1324 uSharpen = (sharpen == MagickFalse)?0:1; 1325 clStatus|=clEnv->library->clSetKernelArg(filterKernel,i++,sizeof(cl_uint),&uSharpen); 1326 if (clStatus != CL_SUCCESS) 1327 { 1328 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 1329 goto cleanup; 1330 } 1331 1332 global_work_size[0] = image->columns; 1333 global_work_size[1] = image->rows; 1334 /* launch the kernel */ 1335 queue = AcquireOpenCLCommandQueue(clEnv); 1336 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, filterKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); 1337 if (clStatus != CL_SUCCESS) 1338 { 1339 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 1340 goto cleanup; 1341 } 1342 clEnv->library->clFlush(queue); 1343 RecordProfileData(clEnv,ContrastKernel,event); 1344 clEnv->library->clReleaseEvent(event); 1345 1346 if (ALIGNED(inputPixels,CLPixelPacket)) 1347 { 1348 length = image->columns * image->rows; 1349 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); 1350 } 1351 else 1352 { 1353 length = image->columns * image->rows; 1354 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL); 1355 } 1356 if (clStatus != CL_SUCCESS) 1357 { 1358 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 1359 goto cleanup; 1360 } 1361 outputReady=SyncCacheViewAuthenticPixels(image_view,exception); 1362 1363cleanup: 1364 OpenCLLogException(__FUNCTION__,__LINE__,exception); 1365 1366 image_view=DestroyCacheView(image_view); 1367 1368 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); 1369 if (filterKernel!=NULL) RelinquishOpenCLKernel(clEnv, filterKernel); 1370 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 1371 return(outputReady); 1372} 1373 1374MagickExport MagickBooleanType AccelerateContrastImage(Image *image, 1375 const MagickBooleanType sharpen,ExceptionInfo *exception) 1376{ 1377 MagickBooleanType 1378 status; 1379 1380 assert(image != NULL); 1381 assert(exception != (ExceptionInfo *) NULL); 1382 1383 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 1384 (checkOpenCLEnvironment(exception) == MagickFalse)) 1385 return(MagickFalse); 1386 1387 status=ComputeContrastImage(image,sharpen,exception); 1388 return(status); 1389} 1390 1391/* 1392%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 1393% % 1394% % 1395% % 1396% A c c e l e r a t e C o n t r a s t S t r e t c h I m a g e % 1397% % 1398% % 1399% % 1400%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 1401*/ 1402 1403static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv, 1404 cl_command_queue queue,cl_mem imageBuffer,cl_mem histogramBuffer, 1405 Image *image,const ChannelType channel,ExceptionInfo *exception) 1406{ 1407 MagickBooleanType 1408 outputReady; 1409 1410 cl_int 1411 clStatus; 1412 1413 cl_kernel 1414 histogramKernel; 1415 1416 cl_event 1417 event; 1418 1419 cl_uint 1420 colorspace, 1421 method; 1422 1423 register ssize_t 1424 i; 1425 1426 size_t 1427 global_work_size[2]; 1428 1429 histogramKernel = NULL; 1430 1431 outputReady = MagickFalse; 1432 colorspace = image->colorspace; 1433 method = image->intensity; 1434 1435 /* get the OpenCL kernel */ 1436 histogramKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Histogram"); 1437 if (histogramKernel == NULL) 1438 { 1439 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 1440 goto cleanup; 1441 } 1442 1443 /* set the kernel arguments */ 1444 i = 0; 1445 clStatus=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 1446 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(ChannelType),&channel); 1447 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_uint),&colorspace); 1448 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_uint),&method); 1449 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&histogramBuffer); 1450 if (clStatus != CL_SUCCESS) 1451 { 1452 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 1453 goto cleanup; 1454 } 1455 1456 /* launch the kernel */ 1457 global_work_size[0] = image->columns; 1458 global_work_size[1] = image->rows; 1459 1460 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, histogramKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); 1461 1462 if (clStatus != CL_SUCCESS) 1463 { 1464 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 1465 goto cleanup; 1466 } 1467 clEnv->library->clFlush(queue); 1468 RecordProfileData(clEnv,HistogramKernel,event); 1469 clEnv->library->clReleaseEvent(event); 1470 1471 outputReady = MagickTrue; 1472 1473cleanup: 1474 OpenCLLogException(__FUNCTION__,__LINE__,exception); 1475 1476 if (histogramKernel!=NULL) 1477 RelinquishOpenCLKernel(clEnv, histogramKernel); 1478 1479 return(outputReady); 1480} 1481 1482static MagickBooleanType ComputeContrastStretchImage(Image *image, 1483 const double black_point,const double white_point,ExceptionInfo *exception) 1484{ 1485#define ContrastStretchImageTag "ContrastStretch/Image" 1486#define MaxRange(color) ((MagickRealType) ScaleQuantumToMap((Quantum) (color))) 1487 1488 CacheView 1489 *image_view; 1490 1491 cl_command_queue 1492 queue; 1493 1494 cl_context 1495 context; 1496 1497 cl_int 1498 clStatus; 1499 1500 cl_mem_flags 1501 mem_flags; 1502 1503 cl_mem 1504 histogramBuffer, 1505 imageBuffer, 1506 stretchMapBuffer; 1507 1508 cl_kernel 1509 histogramKernel, 1510 stretchKernel; 1511 1512 cl_event 1513 event; 1514 1515 cl_uint4 1516 *histogram; 1517 1518 double 1519 intensity; 1520 1521 FloatPixelPacket 1522 black, 1523 white; 1524 1525 MagickBooleanType 1526 outputReady, 1527 status; 1528 1529 MagickCLEnv 1530 clEnv; 1531 1532 MagickSizeType 1533 length; 1534 1535 PixelPacket 1536 *stretch_map; 1537 1538 register ssize_t 1539 i; 1540 1541 size_t 1542 global_work_size[2]; 1543 1544 void 1545 *hostPtr, 1546 *inputPixels; 1547 1548 histogram=NULL; 1549 stretch_map=NULL; 1550 inputPixels = NULL; 1551 imageBuffer = NULL; 1552 histogramBuffer = NULL; 1553 stretchMapBuffer = NULL; 1554 histogramKernel = NULL; 1555 stretchKernel = NULL; 1556 context = NULL; 1557 queue = NULL; 1558 outputReady = MagickFalse; 1559 1560 1561 assert(image != (Image *) NULL); 1562 assert(image->signature == MagickCoreSignature); 1563 if (image->debug != MagickFalse) 1564 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename); 1565 1566 //exception=(&image->exception); 1567 1568 /* 1569 * initialize opencl env 1570 */ 1571 clEnv = GetDefaultOpenCLEnv(); 1572 context = GetOpenCLContext(clEnv); 1573 queue = AcquireOpenCLCommandQueue(clEnv); 1574 1575 /* 1576 Allocate and initialize histogram arrays. 1577 */ 1578 histogram=(cl_uint4 *) AcquireQuantumMemory(MaxMap+1UL, sizeof(*histogram)); 1579 1580 if (histogram == (cl_uint4 *) NULL) 1581 ThrowBinaryException(ResourceLimitError,"MemoryAllocationFailed", image->filename); 1582 1583 /* reset histogram */ 1584 (void) ResetMagickMemory(histogram,0,(MaxMap+1)*sizeof(*histogram)); 1585 1586 /* 1587 if (IsGrayImage(image,exception) != MagickFalse) 1588 (void) SetImageColorspace(image,GRAYColorspace); 1589 */ 1590 1591 status=MagickTrue; 1592 1593 1594 /* 1595 Form histogram. 1596 */ 1597 /* Create and initialize OpenCL buffers. */ 1598 /* inputPixels = AcquirePixelCachePixels(image, &length, exception); */ 1599 /* assume this will get a writable image */ 1600 image_view=AcquireAuthenticCacheView(image,exception); 1601 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); 1602 1603 if (inputPixels == (void *) NULL) 1604 { 1605 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); 1606 goto cleanup; 1607 } 1608 /* If the host pointer is aligned to the size of CLPixelPacket, 1609 then use the host buffer directly from the GPU; otherwise, 1610 create a buffer on the GPU and copy the data over */ 1611 if (ALIGNED(inputPixels,CLPixelPacket)) 1612 { 1613 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; 1614 } 1615 else 1616 { 1617 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR; 1618 } 1619 /* create a CL buffer from image pixel buffer */ 1620 length = image->columns * image->rows; 1621 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 1622 if (clStatus != CL_SUCCESS) 1623 { 1624 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 1625 goto cleanup; 1626 } 1627 1628 /* If the host pointer is aligned to the size of cl_uint, 1629 then use the host buffer directly from the GPU; otherwise, 1630 create a buffer on the GPU and copy the data over */ 1631 if (ALIGNED(histogram,cl_uint4)) 1632 { 1633 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR; 1634 hostPtr = histogram; 1635 } 1636 else 1637 { 1638 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; 1639 hostPtr = histogram; 1640 } 1641 /* create a CL buffer for histogram */ 1642 length = (MaxMap+1); 1643 histogramBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus); 1644 if (clStatus != CL_SUCCESS) 1645 { 1646 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 1647 goto cleanup; 1648 } 1649 1650 status = LaunchHistogramKernel(clEnv, queue, imageBuffer, histogramBuffer, image, image->channel_mask, exception); 1651 if (status == MagickFalse) 1652 goto cleanup; 1653 1654 /* read from the kenel output */ 1655 if (ALIGNED(histogram,cl_uint4)) 1656 { 1657 length = (MaxMap+1); 1658 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(cl_uint4), 0, NULL, NULL, &clStatus); 1659 } 1660 else 1661 { 1662 length = (MaxMap+1); 1663 clStatus = clEnv->library->clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length * sizeof(cl_uint4), histogram, 0, NULL, NULL); 1664 } 1665 if (clStatus != CL_SUCCESS) 1666 { 1667 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 1668 goto cleanup; 1669 } 1670 1671 /* unmap, don't block gpu to use this buffer again. */ 1672 if (ALIGNED(histogram,cl_uint4)) 1673 { 1674 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL); 1675 if (clStatus != CL_SUCCESS) 1676 { 1677 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); 1678 goto cleanup; 1679 } 1680 } 1681 1682 /* recreate input buffer later, in case image updated */ 1683#ifdef RECREATEBUFFER 1684 if (imageBuffer!=NULL) 1685 clEnv->library->clReleaseMemObject(imageBuffer); 1686#endif 1687 1688 /* CPU stuff */ 1689 /* 1690 Find the histogram boundaries by locating the black/white levels. 1691 */ 1692 black.red=0.0; 1693 white.red=MaxRange(QuantumRange); 1694 if ((image->channel_mask & RedChannel) != 0) 1695 { 1696 intensity=0.0; 1697 for (i=0; i <= (ssize_t) MaxMap; i++) 1698 { 1699 intensity+=histogram[i].s[2]; 1700 if (intensity > black_point) 1701 break; 1702 } 1703 black.red=(MagickRealType) i; 1704 intensity=0.0; 1705 for (i=(ssize_t) MaxMap; i != 0; i--) 1706 { 1707 intensity+=histogram[i].s[2]; 1708 if (intensity > ((double) image->columns*image->rows-white_point)) 1709 break; 1710 } 1711 white.red=(MagickRealType) i; 1712 } 1713 black.green=0.0; 1714 white.green=MaxRange(QuantumRange); 1715 if ((image->channel_mask & GreenChannel) != 0) 1716 { 1717 intensity=0.0; 1718 for (i=0; i <= (ssize_t) MaxMap; i++) 1719 { 1720 intensity+=histogram[i].s[2]; 1721 if (intensity > black_point) 1722 break; 1723 } 1724 black.green=(MagickRealType) i; 1725 intensity=0.0; 1726 for (i=(ssize_t) MaxMap; i != 0; i--) 1727 { 1728 intensity+=histogram[i].s[2]; 1729 if (intensity > ((double) image->columns*image->rows-white_point)) 1730 break; 1731 } 1732 white.green=(MagickRealType) i; 1733 } 1734 black.blue=0.0; 1735 white.blue=MaxRange(QuantumRange); 1736 if ((image->channel_mask & BlueChannel) != 0) 1737 { 1738 intensity=0.0; 1739 for (i=0; i <= (ssize_t) MaxMap; i++) 1740 { 1741 intensity+=histogram[i].s[2]; 1742 if (intensity > black_point) 1743 break; 1744 } 1745 black.blue=(MagickRealType) i; 1746 intensity=0.0; 1747 for (i=(ssize_t) MaxMap; i != 0; i--) 1748 { 1749 intensity+=histogram[i].s[2]; 1750 if (intensity > ((double) image->columns*image->rows-white_point)) 1751 break; 1752 } 1753 white.blue=(MagickRealType) i; 1754 } 1755 black.alpha=0.0; 1756 white.alpha=MaxRange(QuantumRange); 1757 if ((image->channel_mask & AlphaChannel) != 0) 1758 { 1759 intensity=0.0; 1760 for (i=0; i <= (ssize_t) MaxMap; i++) 1761 { 1762 intensity+=histogram[i].s[2]; 1763 if (intensity > black_point) 1764 break; 1765 } 1766 black.alpha=(MagickRealType) i; 1767 intensity=0.0; 1768 for (i=(ssize_t) MaxMap; i != 0; i--) 1769 { 1770 intensity+=histogram[i].s[2]; 1771 if (intensity > ((double) image->columns*image->rows-white_point)) 1772 break; 1773 } 1774 white.alpha=(MagickRealType) i; 1775 } 1776 /* 1777 black.index=0.0; 1778 white.index=MaxRange(QuantumRange); 1779 if (((channel & IndexChannel) != 0) && (image->colorspace == CMYKColorspace)) 1780 { 1781 intensity=0.0; 1782 for (i=0; i <= (ssize_t) MaxMap; i++) 1783 { 1784 intensity+=histogram[i].index; 1785 if (intensity > black_point) 1786 break; 1787 } 1788 black.index=(MagickRealType) i; 1789 intensity=0.0; 1790 for (i=(ssize_t) MaxMap; i != 0; i--) 1791 { 1792 intensity+=histogram[i].index; 1793 if (intensity > ((double) image->columns*image->rows-white_point)) 1794 break; 1795 } 1796 white.index=(MagickRealType) i; 1797 } 1798 */ 1799 1800 1801 stretch_map=(PixelPacket *) AcquireQuantumMemory(MaxMap+1UL, 1802 sizeof(*stretch_map)); 1803 1804 if (stretch_map == (PixelPacket *) NULL) 1805 ThrowBinaryException(ResourceLimitError,"MemoryAllocationFailed", 1806 image->filename); 1807 1808 /* 1809 Stretch the histogram to create the stretched image mapping. 1810 */ 1811 (void) ResetMagickMemory(stretch_map,0,(MaxMap+1)*sizeof(*stretch_map)); 1812 for (i=0; i <= (ssize_t) MaxMap; i++) 1813 { 1814 if ((image->channel_mask & RedChannel) != 0) 1815 { 1816 if (i < (ssize_t) black.red) 1817 stretch_map[i].red=(Quantum) 0; 1818 else 1819 if (i > (ssize_t) white.red) 1820 stretch_map[i].red=QuantumRange; 1821 else 1822 if (black.red != white.red) 1823 stretch_map[i].red=ScaleMapToQuantum((MagickRealType) (MaxMap* 1824 (i-black.red)/(white.red-black.red))); 1825 } 1826 if ((image->channel_mask & GreenChannel) != 0) 1827 { 1828 if (i < (ssize_t) black.green) 1829 stretch_map[i].green=0; 1830 else 1831 if (i > (ssize_t) white.green) 1832 stretch_map[i].green=QuantumRange; 1833 else 1834 if (black.green != white.green) 1835 stretch_map[i].green=ScaleMapToQuantum((MagickRealType) (MaxMap* 1836 (i-black.green)/(white.green-black.green))); 1837 } 1838 if ((image->channel_mask & BlueChannel) != 0) 1839 { 1840 if (i < (ssize_t) black.blue) 1841 stretch_map[i].blue=0; 1842 else 1843 if (i > (ssize_t) white.blue) 1844 stretch_map[i].blue= QuantumRange; 1845 else 1846 if (black.blue != white.blue) 1847 stretch_map[i].blue=ScaleMapToQuantum((MagickRealType) (MaxMap* 1848 (i-black.blue)/(white.blue-black.blue))); 1849 } 1850 if ((image->channel_mask & AlphaChannel) != 0) 1851 { 1852 if (i < (ssize_t) black.alpha) 1853 stretch_map[i].alpha=0; 1854 else 1855 if (i > (ssize_t) white.alpha) 1856 stretch_map[i].alpha=QuantumRange; 1857 else 1858 if (black.alpha != white.alpha) 1859 stretch_map[i].alpha=ScaleMapToQuantum((MagickRealType) (MaxMap* 1860 (i-black.alpha)/(white.alpha-black.alpha))); 1861 } 1862 /* 1863 if (((channel & IndexChannel) != 0) && 1864 (image->colorspace == CMYKColorspace)) 1865 { 1866 if (i < (ssize_t) black.index) 1867 stretch_map[i].index=0; 1868 else 1869 if (i > (ssize_t) white.index) 1870 stretch_map[i].index=QuantumRange; 1871 else 1872 if (black.index != white.index) 1873 stretch_map[i].index=ScaleMapToQuantum((MagickRealType) (MaxMap* 1874 (i-black.index)/(white.index-black.index))); 1875 } 1876 */ 1877 } 1878 1879 /* 1880 Stretch the image. 1881 */ 1882 if (((image->channel_mask & AlphaChannel) != 0) || (((image->channel_mask & IndexChannel) != 0) && 1883 (image->colorspace == CMYKColorspace))) 1884 image->storage_class=DirectClass; 1885 if (image->storage_class == PseudoClass) 1886 { 1887 /* 1888 Stretch colormap. 1889 */ 1890 for (i=0; i < (ssize_t) image->colors; i++) 1891 { 1892 if ((image->channel_mask & RedChannel) != 0) 1893 { 1894 if (black.red != white.red) 1895 image->colormap[i].red=stretch_map[ 1896 ScaleQuantumToMap(image->colormap[i].red)].red; 1897 } 1898 if ((image->channel_mask & GreenChannel) != 0) 1899 { 1900 if (black.green != white.green) 1901 image->colormap[i].green=stretch_map[ 1902 ScaleQuantumToMap(image->colormap[i].green)].green; 1903 } 1904 if ((image->channel_mask & BlueChannel) != 0) 1905 { 1906 if (black.blue != white.blue) 1907 image->colormap[i].blue=stretch_map[ 1908 ScaleQuantumToMap(image->colormap[i].blue)].blue; 1909 } 1910 if ((image->channel_mask & AlphaChannel) != 0) 1911 { 1912 if (black.alpha != white.alpha) 1913 image->colormap[i].alpha=stretch_map[ 1914 ScaleQuantumToMap(image->colormap[i].alpha)].alpha; 1915 } 1916 } 1917 } 1918 1919 /* 1920 Stretch image. 1921 */ 1922 1923 1924 /* GPU can work on this again, image and equalize map as input 1925 image: uchar4 (CLPixelPacket) 1926 stretch_map: uchar4 (PixelPacket) 1927 black, white: float4 (FloatPixelPacket) */ 1928 1929#ifdef RECREATEBUFFER 1930 /* If the host pointer is aligned to the size of CLPixelPacket, 1931 then use the host buffer directly from the GPU; otherwise, 1932 create a buffer on the GPU and copy the data over */ 1933 if (ALIGNED(inputPixels,CLPixelPacket)) 1934 { 1935 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR; 1936 } 1937 else 1938 { 1939 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; 1940 } 1941 /* create a CL buffer from image pixel buffer */ 1942 length = image->columns * image->rows; 1943 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 1944 if (clStatus != CL_SUCCESS) 1945 { 1946 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 1947 goto cleanup; 1948 } 1949#endif 1950 1951 /* Create and initialize OpenCL buffers. */ 1952 if (ALIGNED(stretch_map, PixelPacket)) 1953 { 1954 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; 1955 hostPtr = stretch_map; 1956 } 1957 else 1958 { 1959 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; 1960 hostPtr = stretch_map; 1961 } 1962 /* create a CL buffer for stretch_map */ 1963 length = (MaxMap+1); 1964 stretchMapBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus); 1965 if (clStatus != CL_SUCCESS) 1966 { 1967 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 1968 goto cleanup; 1969 } 1970 1971 /* get the OpenCL kernel */ 1972 stretchKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ContrastStretch"); 1973 if (stretchKernel == NULL) 1974 { 1975 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 1976 goto cleanup; 1977 } 1978 1979 /* set the kernel arguments */ 1980 i = 0; 1981 clStatus=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 1982 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(ChannelType),&image->channel_mask); 1983 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_mem),(void *)&stretchMapBuffer); 1984 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(FloatPixelPacket),&white); 1985 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(FloatPixelPacket),&black); 1986 if (clStatus != CL_SUCCESS) 1987 { 1988 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 1989 goto cleanup; 1990 } 1991 1992 /* launch the kernel */ 1993 global_work_size[0] = image->columns; 1994 global_work_size[1] = image->rows; 1995 1996 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, stretchKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); 1997 1998 if (clStatus != CL_SUCCESS) 1999 { 2000 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 2001 goto cleanup; 2002 } 2003 clEnv->library->clFlush(queue); 2004 2005 RecordProfileData(clEnv,ContrastStretchKernel,event); 2006 clEnv->library->clReleaseEvent(event); 2007 2008 /* read the data back */ 2009 if (ALIGNED(inputPixels,CLPixelPacket)) 2010 { 2011 length = image->columns * image->rows; 2012 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); 2013 } 2014 else 2015 { 2016 length = image->columns * image->rows; 2017 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL); 2018 } 2019 if (clStatus != CL_SUCCESS) 2020 { 2021 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 2022 goto cleanup; 2023 } 2024 2025 outputReady=SyncCacheViewAuthenticPixels(image_view,exception); 2026 2027cleanup: 2028 OpenCLLogException(__FUNCTION__,__LINE__,exception); 2029 2030 image_view=DestroyCacheView(image_view); 2031 2032 if (imageBuffer!=NULL) 2033 clEnv->library->clReleaseMemObject(imageBuffer); 2034 2035 if (stretchMapBuffer!=NULL) 2036 clEnv->library->clReleaseMemObject(stretchMapBuffer); 2037 if (stretch_map!=NULL) 2038 stretch_map=(PixelPacket *) RelinquishMagickMemory(stretch_map); 2039 2040 2041 if (histogramBuffer!=NULL) 2042 clEnv->library->clReleaseMemObject(histogramBuffer); 2043 if (histogram!=NULL) 2044 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram); 2045 2046 2047 if (histogramKernel!=NULL) 2048 RelinquishOpenCLKernel(clEnv, histogramKernel); 2049 if (stretchKernel!=NULL) 2050 RelinquishOpenCLKernel(clEnv, stretchKernel); 2051 2052 if (queue != NULL) 2053 RelinquishOpenCLCommandQueue(clEnv, queue); 2054 2055 return(outputReady); 2056} 2057 2058MagickExport MagickBooleanType AccelerateContrastStretchImage( 2059 Image *image,const double black_point,const double white_point, 2060 ExceptionInfo *exception) 2061{ 2062 MagickBooleanType 2063 status; 2064 2065 assert(image != NULL); 2066 assert(exception != (ExceptionInfo *) NULL); 2067 2068 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 2069 (checkHistogramCondition(image,image->intensity) == MagickFalse) || 2070 (checkOpenCLEnvironment(exception) == MagickFalse)) 2071 return(MagickFalse); 2072 2073 status=ComputeContrastStretchImage(image,black_point,white_point,exception); 2074 return(status); 2075} 2076 2077/* 2078%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 2079% % 2080% % 2081% % 2082% A c c e l e r a t e C o n v o l v e I m a g e % 2083% % 2084% % 2085% % 2086%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 2087*/ 2088 2089static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, 2090 ExceptionInfo *exception) 2091{ 2092 CacheView 2093 *filteredImage_view, 2094 *image_view; 2095 2096 cl_command_queue 2097 queue; 2098 2099 cl_context 2100 context; 2101 2102 cl_kernel 2103 clkernel; 2104 2105 cl_event 2106 event; 2107 2108 cl_int 2109 clStatus; 2110 2111 cl_mem 2112 convolutionKernel, 2113 filteredImageBuffer, 2114 imageBuffer; 2115 2116 cl_mem_flags 2117 mem_flags; 2118 2119 cl_ulong 2120 deviceLocalMemorySize; 2121 2122 const void 2123 *inputPixels; 2124 2125 float 2126 *kernelBufferPtr; 2127 2128 Image 2129 *filteredImage; 2130 2131 MagickBooleanType 2132 outputReady; 2133 2134 MagickCLEnv 2135 clEnv; 2136 2137 MagickSizeType 2138 length; 2139 2140 size_t 2141 global_work_size[3], 2142 localGroupSize[3], 2143 localMemoryRequirement; 2144 2145 unsigned 2146 kernelSize; 2147 2148 unsigned int 2149 filterHeight, 2150 filterWidth, 2151 i, 2152 imageHeight, 2153 imageWidth, 2154 matte; 2155 2156 void 2157 *filteredPixels, 2158 *hostPtr; 2159 2160 /* intialize all CL objects to NULL */ 2161 context = NULL; 2162 imageBuffer = NULL; 2163 filteredImageBuffer = NULL; 2164 convolutionKernel = NULL; 2165 clkernel = NULL; 2166 queue = NULL; 2167 2168 filteredImage = NULL; 2169 filteredImage_view = NULL; 2170 outputReady = MagickFalse; 2171 2172 clEnv = GetDefaultOpenCLEnv(); 2173 context = GetOpenCLContext(clEnv); 2174 2175 image_view=AcquireVirtualCacheView(image,exception); 2176 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception); 2177 if (inputPixels == (const void *) NULL) 2178 { 2179 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); 2180 goto cleanup; 2181 } 2182 2183 /* Create and initialize OpenCL buffers. */ 2184 2185 /* If the host pointer is aligned to the size of CLPixelPacket, 2186 then use the host buffer directly from the GPU; otherwise, 2187 create a buffer on the GPU and copy the data over */ 2188 if (ALIGNED(inputPixels,CLPixelPacket)) 2189 { 2190 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; 2191 } 2192 else 2193 { 2194 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR; 2195 } 2196 /* create a CL buffer from image pixel buffer */ 2197 length = image->columns * image->rows; 2198 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 2199 if (clStatus != CL_SUCCESS) 2200 { 2201 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 2202 goto cleanup; 2203 } 2204 2205 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception); 2206 assert(filteredImage != NULL); 2207 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) 2208 { 2209 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); 2210 goto cleanup; 2211 } 2212 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); 2213 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception); 2214 if (filteredPixels == (void *) NULL) 2215 { 2216 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); 2217 goto cleanup; 2218 } 2219 2220 if (ALIGNED(filteredPixels,CLPixelPacket)) 2221 { 2222 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR; 2223 hostPtr = filteredPixels; 2224 } 2225 else 2226 { 2227 mem_flags = CL_MEM_WRITE_ONLY; 2228 hostPtr = NULL; 2229 } 2230 /* create a CL buffer from image pixel buffer */ 2231 length = image->columns * image->rows; 2232 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); 2233 if (clStatus != CL_SUCCESS) 2234 { 2235 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 2236 goto cleanup; 2237 } 2238 2239 kernelSize = (unsigned int) (kernel->width * kernel->height); 2240 convolutionKernel = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernelSize * sizeof(float), NULL, &clStatus); 2241 if (clStatus != CL_SUCCESS) 2242 { 2243 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 2244 goto cleanup; 2245 } 2246 2247 queue = AcquireOpenCLCommandQueue(clEnv); 2248 2249 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, convolutionKernel, CL_TRUE, CL_MAP_WRITE, 0, kernelSize * sizeof(float) 2250 , 0, NULL, NULL, &clStatus); 2251 if (clStatus != CL_SUCCESS) 2252 { 2253 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.","."); 2254 goto cleanup; 2255 } 2256 for (i = 0; i < kernelSize; i++) 2257 { 2258 kernelBufferPtr[i] = (float) kernel->values[i]; 2259 } 2260 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, convolutionKernel, kernelBufferPtr, 0, NULL, NULL); 2261 if (clStatus != CL_SUCCESS) 2262 { 2263 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); 2264 goto cleanup; 2265 } 2266 clEnv->library->clFlush(queue); 2267 2268 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv); 2269 2270 /* Compute the local memory requirement for a 16x16 workgroup. 2271 If it's larger than 16k, reduce the workgroup size to 8x8 */ 2272 localGroupSize[0] = 16; 2273 localGroupSize[1] = 16; 2274 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) * sizeof(CLPixelPacket) 2275 + kernel->width*kernel->height*sizeof(float); 2276 2277 if (localMemoryRequirement > deviceLocalMemorySize) 2278 { 2279 localGroupSize[0] = 8; 2280 localGroupSize[1] = 8; 2281 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) * sizeof(CLPixelPacket) 2282 + kernel->width*kernel->height*sizeof(float); 2283 } 2284 if (localMemoryRequirement <= deviceLocalMemorySize) 2285 { 2286 /* get the OpenCL kernel */ 2287 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ConvolveOptimized"); 2288 if (clkernel == NULL) 2289 { 2290 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 2291 goto cleanup; 2292 } 2293 2294 /* set the kernel arguments */ 2295 i = 0; 2296 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 2297 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); 2298 imageWidth = (unsigned int) image->columns; 2299 imageHeight = (unsigned int) image->rows; 2300 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageWidth); 2301 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageHeight); 2302 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&convolutionKernel); 2303 filterWidth = (unsigned int) kernel->width; 2304 filterHeight = (unsigned int) kernel->height; 2305 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterWidth); 2306 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterHeight); 2307 matte = (image->alpha_trait > CopyPixelTrait)?1:0; 2308 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&matte); 2309 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&image->channel_mask); 2310 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, (localGroupSize[0] + kernel->width-1)*(localGroupSize[1] + kernel->height-1)*sizeof(CLPixelPacket),NULL); 2311 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, kernel->width*kernel->height*sizeof(float),NULL); 2312 if (clStatus != CL_SUCCESS) 2313 { 2314 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 2315 goto cleanup; 2316 } 2317 2318 /* pad the global size to a multiple of the local work size dimension */ 2319 global_work_size[0] = ((image->columns + localGroupSize[0] - 1)/localGroupSize[0] ) * localGroupSize[0] ; 2320 global_work_size[1] = ((image->rows + localGroupSize[1] - 1)/localGroupSize[1]) * localGroupSize[1]; 2321 2322 /* launch the kernel */ 2323 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, &event); 2324 if (clStatus != CL_SUCCESS) 2325 { 2326 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 2327 goto cleanup; 2328 } 2329 RecordProfileData(clEnv,ConvolveKernel,event); 2330 clEnv->library->clReleaseEvent(event); 2331 } 2332 else 2333 { 2334 /* get the OpenCL kernel */ 2335 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Convolve"); 2336 if (clkernel == NULL) 2337 { 2338 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 2339 goto cleanup; 2340 } 2341 2342 /* set the kernel arguments */ 2343 i = 0; 2344 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 2345 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); 2346 imageWidth = (unsigned int) image->columns; 2347 imageHeight = (unsigned int) image->rows; 2348 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageWidth); 2349 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageHeight); 2350 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&convolutionKernel); 2351 filterWidth = (unsigned int) kernel->width; 2352 filterHeight = (unsigned int) kernel->height; 2353 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterWidth); 2354 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterHeight); 2355 matte = (image->alpha_trait > CopyPixelTrait)?1:0; 2356 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&matte); 2357 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&image->channel_mask); 2358 if (clStatus != CL_SUCCESS) 2359 { 2360 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 2361 goto cleanup; 2362 } 2363 2364 localGroupSize[0] = 8; 2365 localGroupSize[1] = 8; 2366 global_work_size[0] = (image->columns + (localGroupSize[0]-1))/localGroupSize[0] * localGroupSize[0]; 2367 global_work_size[1] = (image->rows + (localGroupSize[1]-1))/localGroupSize[1] * localGroupSize[1]; 2368 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, &event); 2369 2370 if (clStatus != CL_SUCCESS) 2371 { 2372 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 2373 goto cleanup; 2374 } 2375 } 2376 clEnv->library->clFlush(queue); 2377 RecordProfileData(clEnv,ConvolveKernel,event); 2378 clEnv->library->clReleaseEvent(event); 2379 2380 if (ALIGNED(filteredPixels,CLPixelPacket)) 2381 { 2382 length = image->columns * image->rows; 2383 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); 2384 } 2385 else 2386 { 2387 length = image->columns * image->rows; 2388 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL); 2389 } 2390 if (clStatus != CL_SUCCESS) 2391 { 2392 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 2393 goto cleanup; 2394 } 2395 2396 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); 2397 2398cleanup: 2399 OpenCLLogException(__FUNCTION__,__LINE__,exception); 2400 2401 image_view=DestroyCacheView(image_view); 2402 if (filteredImage_view != NULL) 2403 filteredImage_view=DestroyCacheView(filteredImage_view); 2404 2405 if (imageBuffer != NULL) 2406 clEnv->library->clReleaseMemObject(imageBuffer); 2407 2408 if (filteredImageBuffer != NULL) 2409 clEnv->library->clReleaseMemObject(filteredImageBuffer); 2410 2411 if (convolutionKernel != NULL) 2412 clEnv->library->clReleaseMemObject(convolutionKernel); 2413 2414 if (clkernel != NULL) 2415 RelinquishOpenCLKernel(clEnv, clkernel); 2416 2417 if (queue != NULL) 2418 RelinquishOpenCLCommandQueue(clEnv, queue); 2419 2420 if (outputReady == MagickFalse) 2421 { 2422 if (filteredImage != NULL) 2423 { 2424 DestroyImage(filteredImage); 2425 filteredImage = NULL; 2426 } 2427 } 2428 2429 return(filteredImage); 2430} 2431 2432MagickExport Image *AccelerateConvolveImage(const Image *image, 2433 const KernelInfo *kernel,ExceptionInfo *exception) 2434{ 2435 /* Temporary disabled due to access violation 2436 2437 Image 2438 *filteredImage; 2439 2440 assert(image != NULL); 2441 assert(kernel != (KernelInfo *) NULL); 2442 assert(exception != (ExceptionInfo *) NULL); 2443 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 2444 (checkOpenCLEnvironment(exception) == MagickFalse)) 2445 return((Image *) NULL); 2446 2447 filteredImage=ComputeConvolveImage(image,kernel,exception); 2448 return(filteredImage); 2449 */ 2450 magick_unreferenced(image); 2451 magick_unreferenced(kernel); 2452 magick_unreferenced(exception); 2453 return((Image *)NULL); 2454} 2455 2456/* 2457%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 2458% % 2459% % 2460% % 2461% A c c e l e r a t e D e s p e c k l e I m a g e % 2462% % 2463% % 2464% % 2465%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 2466*/ 2467 2468static Image *ComputeDespeckleImage(const Image *image, 2469 ExceptionInfo*exception) 2470{ 2471 static const int 2472 X[4] = {0, 1, 1,-1}, 2473 Y[4] = {1, 0, 1, 1}; 2474 2475 CacheView 2476 *filteredImage_view, 2477 *image_view; 2478 2479 cl_command_queue 2480 queue; 2481 2482 cl_context 2483 context; 2484 2485 cl_int 2486 clStatus; 2487 2488 cl_kernel 2489 hullPass1, 2490 hullPass2; 2491 2492 cl_event 2493 event; 2494 2495 cl_mem_flags 2496 mem_flags; 2497 2498 cl_mem 2499 filteredImageBuffer, 2500 imageBuffer, 2501 tempImageBuffer[2]; 2502 2503 const void 2504 *inputPixels; 2505 2506 Image 2507 *filteredImage; 2508 2509 int 2510 k, 2511 matte; 2512 2513 MagickBooleanType 2514 outputReady; 2515 2516 MagickCLEnv 2517 clEnv; 2518 2519 MagickSizeType 2520 length; 2521 2522 size_t 2523 global_work_size[2]; 2524 2525 unsigned int 2526 imageHeight, 2527 imageWidth; 2528 2529 void 2530 *filteredPixels, 2531 *hostPtr; 2532 2533 outputReady = MagickFalse; 2534 clEnv = NULL; 2535 inputPixels = NULL; 2536 filteredImage = NULL; 2537 filteredImage_view = NULL; 2538 filteredPixels = NULL; 2539 context = NULL; 2540 imageBuffer = NULL; 2541 filteredImageBuffer = NULL; 2542 hullPass1 = NULL; 2543 hullPass2 = NULL; 2544 queue = NULL; 2545 tempImageBuffer[0] = tempImageBuffer[1] = NULL; 2546 clEnv = GetDefaultOpenCLEnv(); 2547 context = GetOpenCLContext(clEnv); 2548 queue = AcquireOpenCLCommandQueue(clEnv); 2549 2550 image_view=AcquireVirtualCacheView(image,exception); 2551 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception); 2552 if (inputPixels == (void *) NULL) 2553 { 2554 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); 2555 goto cleanup; 2556 } 2557 2558 if (ALIGNED(inputPixels,CLPixelPacket)) 2559 { 2560 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; 2561 } 2562 else 2563 { 2564 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR; 2565 } 2566 /* create a CL buffer from image pixel buffer */ 2567 length = image->columns * image->rows; 2568 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 2569 if (clStatus != CL_SUCCESS) 2570 { 2571 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 2572 goto cleanup; 2573 } 2574 2575 mem_flags = CL_MEM_READ_WRITE; 2576 length = image->columns * image->rows; 2577 for (k = 0; k < 2; k++) 2578 { 2579 tempImageBuffer[k] = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), NULL, &clStatus); 2580 if (clStatus != CL_SUCCESS) 2581 { 2582 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 2583 goto cleanup; 2584 } 2585 } 2586 2587 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception); 2588 assert(filteredImage != NULL); 2589 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) 2590 { 2591 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); 2592 goto cleanup; 2593 } 2594 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); 2595 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception); 2596 if (filteredPixels == (void *) NULL) 2597 { 2598 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); 2599 goto cleanup; 2600 } 2601 2602 if (ALIGNED(filteredPixels,CLPixelPacket)) 2603 { 2604 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR; 2605 hostPtr = filteredPixels; 2606 } 2607 else 2608 { 2609 mem_flags = CL_MEM_WRITE_ONLY; 2610 hostPtr = NULL; 2611 } 2612 /* create a CL buffer from image pixel buffer */ 2613 length = image->columns * image->rows; 2614 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); 2615 if (clStatus != CL_SUCCESS) 2616 { 2617 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 2618 goto cleanup; 2619 } 2620 2621 hullPass1 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass1"); 2622 hullPass2 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass2"); 2623 2624 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)&imageBuffer); 2625 clStatus |=clEnv->library->clSetKernelArg(hullPass1,1,sizeof(cl_mem),(void *)(tempImageBuffer+1)); 2626 imageWidth = (unsigned int) image->columns; 2627 clStatus |=clEnv->library->clSetKernelArg(hullPass1,2,sizeof(unsigned int),(void *)&imageWidth); 2628 imageHeight = (unsigned int) image->rows; 2629 clStatus |=clEnv->library->clSetKernelArg(hullPass1,3,sizeof(unsigned int),(void *)&imageHeight); 2630 matte = (image->alpha_trait > CopyPixelTrait)?1:0; 2631 clStatus |=clEnv->library->clSetKernelArg(hullPass1,6,sizeof(int),(void *)&matte); 2632 if (clStatus != CL_SUCCESS) 2633 { 2634 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 2635 goto cleanup; 2636 } 2637 2638 clStatus = clEnv->library->clSetKernelArg(hullPass2,0,sizeof(cl_mem),(void *)(tempImageBuffer+1)); 2639 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)tempImageBuffer); 2640 imageWidth = (unsigned int) image->columns; 2641 clStatus |=clEnv->library->clSetKernelArg(hullPass2,2,sizeof(unsigned int),(void *)&imageWidth); 2642 imageHeight = (unsigned int) image->rows; 2643 clStatus |=clEnv->library->clSetKernelArg(hullPass2,3,sizeof(unsigned int),(void *)&imageHeight); 2644 matte = (image->alpha_trait > CopyPixelTrait)?1:0; 2645 clStatus |=clEnv->library->clSetKernelArg(hullPass2,6,sizeof(int),(void *)&matte); 2646 if (clStatus != CL_SUCCESS) 2647 { 2648 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 2649 goto cleanup; 2650 } 2651 2652 2653 global_work_size[0] = image->columns; 2654 global_work_size[1] = image->rows; 2655 2656 2657 for (k = 0; k < 4; k++) 2658 { 2659 cl_int2 offset; 2660 int polarity; 2661 2662 2663 offset.s[0] = X[k]; 2664 offset.s[1] = Y[k]; 2665 polarity = 1; 2666 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset); 2667 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity); 2668 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset); 2669 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity); 2670 if (clStatus != CL_SUCCESS) 2671 { 2672 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 2673 goto cleanup; 2674 } 2675 /* launch the kernel */ 2676 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event); 2677 if (clStatus != CL_SUCCESS) 2678 { 2679 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 2680 goto cleanup; 2681 } 2682 RecordProfileData(clEnv,HullPass1Kernel,event); 2683 clEnv->library->clReleaseEvent(event); 2684 2685 /* launch the kernel */ 2686 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event); 2687 if (clStatus != CL_SUCCESS) 2688 { 2689 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 2690 goto cleanup; 2691 } 2692 RecordProfileData(clEnv,HullPass2Kernel,event); 2693 clEnv->library->clReleaseEvent(event); 2694 2695 if (k == 0) 2696 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)(tempImageBuffer)); 2697 offset.s[0] = -X[k]; 2698 offset.s[1] = -Y[k]; 2699 polarity = 1; 2700 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset); 2701 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity); 2702 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset); 2703 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity); 2704 if (clStatus != CL_SUCCESS) 2705 { 2706 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 2707 goto cleanup; 2708 } 2709 /* launch the kernel */ 2710 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event); 2711 if (clStatus != CL_SUCCESS) 2712 { 2713 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 2714 goto cleanup; 2715 } 2716 RecordProfileData(clEnv,HullPass1Kernel,event); 2717 clEnv->library->clReleaseEvent(event); 2718 2719 /* launch the kernel */ 2720 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event); 2721 if (clStatus != CL_SUCCESS) 2722 { 2723 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 2724 goto cleanup; 2725 } 2726 RecordProfileData(clEnv,HullPass2Kernel,event); 2727 clEnv->library->clReleaseEvent(event); 2728 2729 offset.s[0] = -X[k]; 2730 offset.s[1] = -Y[k]; 2731 polarity = -1; 2732 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset); 2733 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity); 2734 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset); 2735 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity); 2736 if (clStatus != CL_SUCCESS) 2737 { 2738 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 2739 goto cleanup; 2740 } 2741 /* launch the kernel */ 2742 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event); 2743 if (clStatus != CL_SUCCESS) 2744 { 2745 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 2746 goto cleanup; 2747 } 2748 RecordProfileData(clEnv,HullPass1Kernel,event); 2749 clEnv->library->clReleaseEvent(event); 2750 2751 /* launch the kernel */ 2752 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event); 2753 if (clStatus != CL_SUCCESS) 2754 { 2755 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 2756 goto cleanup; 2757 } 2758 RecordProfileData(clEnv,HullPass2Kernel,event); 2759 clEnv->library->clReleaseEvent(event); 2760 2761 offset.s[0] = X[k]; 2762 offset.s[1] = Y[k]; 2763 polarity = -1; 2764 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset); 2765 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity); 2766 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset); 2767 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity); 2768 2769 if (k == 3) 2770 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)&filteredImageBuffer); 2771 2772 if (clStatus != CL_SUCCESS) 2773 { 2774 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 2775 goto cleanup; 2776 } 2777 /* launch the kernel */ 2778 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event); 2779 if (clStatus != CL_SUCCESS) 2780 { 2781 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 2782 goto cleanup; 2783 } 2784 RecordProfileData(clEnv,HullPass1Kernel,event); 2785 clEnv->library->clReleaseEvent(event); 2786 2787 /* launch the kernel */ 2788 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event); 2789 if (clStatus != CL_SUCCESS) 2790 { 2791 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 2792 goto cleanup; 2793 } 2794 RecordProfileData(clEnv,HullPass2Kernel,event); 2795 clEnv->library->clReleaseEvent(event); 2796 } 2797 2798 if (ALIGNED(filteredPixels,CLPixelPacket)) 2799 { 2800 length = image->columns * image->rows; 2801 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); 2802 } 2803 else 2804 { 2805 length = image->columns * image->rows; 2806 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL); 2807 } 2808 if (clStatus != CL_SUCCESS) 2809 { 2810 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 2811 goto cleanup; 2812 } 2813 2814 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); 2815 2816cleanup: 2817 OpenCLLogException(__FUNCTION__,__LINE__,exception); 2818 2819 image_view=DestroyCacheView(image_view); 2820 if (filteredImage_view != NULL) 2821 filteredImage_view=DestroyCacheView(filteredImage_view); 2822 2823 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 2824 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); 2825 for (k = 0; k < 2; k++) 2826 { 2827 if (tempImageBuffer[k]!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer[k]); 2828 } 2829 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); 2830 if (hullPass1!=NULL) RelinquishOpenCLKernel(clEnv, hullPass1); 2831 if (hullPass2!=NULL) RelinquishOpenCLKernel(clEnv, hullPass2); 2832 if (outputReady == MagickFalse && filteredImage != NULL) 2833 filteredImage=DestroyImage(filteredImage); 2834 return(filteredImage); 2835} 2836 2837MagickExport Image *AccelerateDespeckleImage(const Image* image, 2838 ExceptionInfo* exception) 2839{ 2840 Image 2841 *filteredImage; 2842 2843 assert(image != NULL); 2844 assert(exception != (ExceptionInfo *) NULL); 2845 2846 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 2847 (checkOpenCLEnvironment(exception) == MagickFalse)) 2848 return NULL; 2849 2850 filteredImage=ComputeDespeckleImage(image,exception); 2851 return(filteredImage); 2852} 2853 2854/* 2855%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 2856% % 2857% % 2858% % 2859% A c c e l e r a t e E q u a l i z e I m a g e % 2860% % 2861% % 2862% % 2863%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 2864*/ 2865 2866static MagickBooleanType ComputeEqualizeImage(Image *image, 2867 ExceptionInfo *exception) 2868{ 2869#define EqualizeImageTag "Equalize/Image" 2870 2871 CacheView 2872 *image_view; 2873 2874 cl_command_queue 2875 queue; 2876 2877 cl_context 2878 context; 2879 2880 cl_int 2881 clStatus; 2882 2883 cl_mem_flags 2884 mem_flags; 2885 2886 cl_mem 2887 equalizeMapBuffer, 2888 histogramBuffer, 2889 imageBuffer; 2890 2891 cl_kernel 2892 equalizeKernel, 2893 histogramKernel; 2894 2895 cl_event 2896 event; 2897 2898 cl_uint4 2899 *histogram; 2900 2901 FloatPixelPacket 2902 white, 2903 black, 2904 intensity, 2905 *map; 2906 2907 MagickBooleanType 2908 outputReady, 2909 status; 2910 2911 MagickCLEnv 2912 clEnv; 2913 2914 MagickSizeType 2915 length; 2916 2917 PixelPacket 2918 *equalize_map; 2919 2920 register ssize_t 2921 i; 2922 2923 size_t 2924 global_work_size[2]; 2925 2926 void 2927 *hostPtr, 2928 *inputPixels; 2929 2930 map=NULL; 2931 histogram=NULL; 2932 equalize_map=NULL; 2933 inputPixels = NULL; 2934 imageBuffer = NULL; 2935 histogramBuffer = NULL; 2936 equalizeMapBuffer = NULL; 2937 histogramKernel = NULL; 2938 equalizeKernel = NULL; 2939 context = NULL; 2940 queue = NULL; 2941 outputReady = MagickFalse; 2942 2943 assert(image != (Image *) NULL); 2944 assert(image->signature == MagickCoreSignature); 2945 if (image->debug != MagickFalse) 2946 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename); 2947 2948 /* 2949 * initialize opencl env 2950 */ 2951 clEnv = GetDefaultOpenCLEnv(); 2952 context = GetOpenCLContext(clEnv); 2953 queue = AcquireOpenCLCommandQueue(clEnv); 2954 2955 /* 2956 Allocate and initialize histogram arrays. 2957 */ 2958 histogram=(cl_uint4 *) AcquireQuantumMemory(MaxMap+1UL, sizeof(*histogram)); 2959 if (histogram == (cl_uint4 *) NULL) 2960 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename); 2961 2962 /* reset histogram */ 2963 (void) ResetMagickMemory(histogram,0,(MaxMap+1)*sizeof(*histogram)); 2964 2965 /* Create and initialize OpenCL buffers. */ 2966 /* inputPixels = AcquirePixelCachePixels(image, &length, exception); */ 2967 /* assume this will get a writable image */ 2968 image_view=AcquireAuthenticCacheView(image,exception); 2969 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); 2970 2971 if (inputPixels == (void *) NULL) 2972 { 2973 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); 2974 goto cleanup; 2975 } 2976 /* If the host pointer is aligned to the size of CLPixelPacket, 2977 then use the host buffer directly from the GPU; otherwise, 2978 create a buffer on the GPU and copy the data over */ 2979 if (ALIGNED(inputPixels,CLPixelPacket)) 2980 { 2981 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; 2982 } 2983 else 2984 { 2985 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR; 2986 } 2987 /* create a CL buffer from image pixel buffer */ 2988 length = image->columns * image->rows; 2989 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 2990 if (clStatus != CL_SUCCESS) 2991 { 2992 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 2993 goto cleanup; 2994 } 2995 2996 /* If the host pointer is aligned to the size of cl_uint, 2997 then use the host buffer directly from the GPU; otherwise, 2998 create a buffer on the GPU and copy the data over */ 2999 if (ALIGNED(histogram,cl_uint4)) 3000 { 3001 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR; 3002 hostPtr = histogram; 3003 } 3004 else 3005 { 3006 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; 3007 hostPtr = histogram; 3008 } 3009 /* create a CL buffer for histogram */ 3010 length = (MaxMap+1); 3011 histogramBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus); 3012 if (clStatus != CL_SUCCESS) 3013 { 3014 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 3015 goto cleanup; 3016 } 3017 3018 status = LaunchHistogramKernel(clEnv, queue, imageBuffer, histogramBuffer, image, image->channel_mask, exception); 3019 if (status == MagickFalse) 3020 goto cleanup; 3021 3022 /* read from the kenel output */ 3023 if (ALIGNED(histogram,cl_uint4)) 3024 { 3025 length = (MaxMap+1); 3026 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(cl_uint4), 0, NULL, NULL, &clStatus); 3027 } 3028 else 3029 { 3030 length = (MaxMap+1); 3031 clStatus = clEnv->library->clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length * sizeof(cl_uint4), histogram, 0, NULL, NULL); 3032 } 3033 if (clStatus != CL_SUCCESS) 3034 { 3035 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 3036 goto cleanup; 3037 } 3038 3039 /* unmap, don't block gpu to use this buffer again. */ 3040 if (ALIGNED(histogram,cl_uint4)) 3041 { 3042 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL); 3043 if (clStatus != CL_SUCCESS) 3044 { 3045 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); 3046 goto cleanup; 3047 } 3048 } 3049 3050 /* recreate input buffer later, in case image updated */ 3051#ifdef RECREATEBUFFER 3052 if (imageBuffer!=NULL) 3053 clEnv->library->clReleaseMemObject(imageBuffer); 3054#endif 3055 3056 /* CPU stuff */ 3057 equalize_map=(PixelPacket *) AcquireQuantumMemory(MaxMap+1UL, sizeof(*equalize_map)); 3058 if (equalize_map == (PixelPacket *) NULL) 3059 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename); 3060 3061 map=(FloatPixelPacket *) AcquireQuantumMemory(MaxMap+1UL,sizeof(*map)); 3062 if (map == (FloatPixelPacket *) NULL) 3063 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename); 3064 3065 /* 3066 Integrate the histogram to get the equalization map. 3067 */ 3068 (void) ResetMagickMemory(&intensity,0,sizeof(intensity)); 3069 for (i=0; i <= (ssize_t) MaxMap; i++) 3070 { 3071 if ((image->channel_mask & SyncChannels) != 0) 3072 { 3073 intensity.red+=histogram[i].s[2]; 3074 map[i]=intensity; 3075 continue; 3076 } 3077 if ((image->channel_mask & RedChannel) != 0) 3078 intensity.red+=histogram[i].s[2]; 3079 if ((image->channel_mask & GreenChannel) != 0) 3080 intensity.green+=histogram[i].s[1]; 3081 if ((image->channel_mask & BlueChannel) != 0) 3082 intensity.blue+=histogram[i].s[0]; 3083 if ((image->channel_mask & AlphaChannel) != 0) 3084 intensity.alpha+=histogram[i].s[3]; 3085 /* 3086 if (((channel & IndexChannel) != 0) && 3087 (image->colorspace == CMYKColorspace)) 3088 { 3089 intensity.index+=histogram[i].index; 3090 } 3091 */ 3092 map[i]=intensity; 3093 } 3094 black=map[0]; 3095 white=map[(int) MaxMap]; 3096 (void) ResetMagickMemory(equalize_map,0,(MaxMap+1)*sizeof(*equalize_map)); 3097 for (i=0; i <= (ssize_t) MaxMap; i++) 3098 { 3099 if ((image->channel_mask & SyncChannels) != 0) 3100 { 3101 if (white.red != black.red) 3102 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap* 3103 (map[i].red-black.red))/(white.red-black.red))); 3104 continue; 3105 } 3106 if (((image->channel_mask & RedChannel) != 0) && (white.red != black.red)) 3107 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap* 3108 (map[i].red-black.red))/(white.red-black.red))); 3109 if (((image->channel_mask & GreenChannel) != 0) && (white.green != black.green)) 3110 equalize_map[i].green=ScaleMapToQuantum((MagickRealType) ((MaxMap* 3111 (map[i].green-black.green))/(white.green-black.green))); 3112 if (((image->channel_mask & BlueChannel) != 0) && (white.blue != black.blue)) 3113 equalize_map[i].blue=ScaleMapToQuantum((MagickRealType) ((MaxMap* 3114 (map[i].blue-black.blue))/(white.blue-black.blue))); 3115 if (((image->channel_mask & AlphaChannel) != 0) && (white.alpha != black.alpha)) 3116 equalize_map[i].alpha=ScaleMapToQuantum((MagickRealType) ((MaxMap* 3117 (map[i].alpha-black.alpha))/(white.alpha-black.alpha))); 3118 /* 3119 if ((((channel & IndexChannel) != 0) && 3120 (image->colorspace == CMYKColorspace)) && 3121 (white.index != black.index)) 3122 equalize_map[i].index=ScaleMapToQuantum((MagickRealType) ((MaxMap* 3123 (map[i].index-black.index))/(white.index-black.index))); 3124 */ 3125 } 3126 3127 if (image->storage_class == PseudoClass) 3128 { 3129 /* 3130 Equalize colormap. 3131 */ 3132 for (i=0; i < (ssize_t) image->colors; i++) 3133 { 3134 if ((image->channel_mask & SyncChannels) != 0) 3135 { 3136 if (white.red != black.red) 3137 { 3138 image->colormap[i].red=equalize_map[ 3139 ScaleQuantumToMap(image->colormap[i].red)].red; 3140 image->colormap[i].green=equalize_map[ 3141 ScaleQuantumToMap(image->colormap[i].green)].red; 3142 image->colormap[i].blue=equalize_map[ 3143 ScaleQuantumToMap(image->colormap[i].blue)].red; 3144 image->colormap[i].alpha=equalize_map[ 3145 ScaleQuantumToMap(image->colormap[i].alpha)].red; 3146 } 3147 continue; 3148 } 3149 if (((image->channel_mask & RedChannel) != 0) && (white.red != black.red)) 3150 image->colormap[i].red=equalize_map[ 3151 ScaleQuantumToMap(image->colormap[i].red)].red; 3152 if (((image->channel_mask & GreenChannel) != 0) && (white.green != black.green)) 3153 image->colormap[i].green=equalize_map[ 3154 ScaleQuantumToMap(image->colormap[i].green)].green; 3155 if (((image->channel_mask & BlueChannel) != 0) && (white.blue != black.blue)) 3156 image->colormap[i].blue=equalize_map[ 3157 ScaleQuantumToMap(image->colormap[i].blue)].blue; 3158 if (((image->channel_mask & AlphaChannel) != 0) && 3159 (white.alpha != black.alpha)) 3160 image->colormap[i].alpha=equalize_map[ 3161 ScaleQuantumToMap(image->colormap[i].alpha)].alpha; 3162 } 3163 } 3164 3165 /* 3166 Equalize image. 3167 */ 3168 3169 /* GPU can work on this again, image and equalize map as input 3170 image: uchar4 (CLPixelPacket) 3171 equalize_map: uchar4 (PixelPacket) 3172 black, white: float4 (FloatPixelPacket) */ 3173 3174#ifdef RECREATEBUFFER 3175 /* If the host pointer is aligned to the size of CLPixelPacket, 3176 then use the host buffer directly from the GPU; otherwise, 3177 create a buffer on the GPU and copy the data over */ 3178 if (ALIGNED(inputPixels,CLPixelPacket)) 3179 { 3180 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR; 3181 } 3182 else 3183 { 3184 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; 3185 } 3186 /* create a CL buffer from image pixel buffer */ 3187 length = image->columns * image->rows; 3188 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 3189 if (clStatus != CL_SUCCESS) 3190 { 3191 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 3192 goto cleanup; 3193 } 3194#endif 3195 3196 /* Create and initialize OpenCL buffers. */ 3197 if (ALIGNED(equalize_map, PixelPacket)) 3198 { 3199 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; 3200 hostPtr = equalize_map; 3201 } 3202 else 3203 { 3204 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; 3205 hostPtr = equalize_map; 3206 } 3207 /* create a CL buffer for eqaulize_map */ 3208 length = (MaxMap+1); 3209 equalizeMapBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus); 3210 if (clStatus != CL_SUCCESS) 3211 { 3212 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 3213 goto cleanup; 3214 } 3215 3216 /* get the OpenCL kernel */ 3217 equalizeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Equalize"); 3218 if (equalizeKernel == NULL) 3219 { 3220 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 3221 goto cleanup; 3222 } 3223 3224 /* set the kernel arguments */ 3225 i = 0; 3226 clStatus=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 3227 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(ChannelType),&image->channel_mask); 3228 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(cl_mem),(void *)&equalizeMapBuffer); 3229 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(FloatPixelPacket),&white); 3230 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(FloatPixelPacket),&black); 3231 if (clStatus != CL_SUCCESS) 3232 { 3233 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 3234 goto cleanup; 3235 } 3236 3237 /* launch the kernel */ 3238 global_work_size[0] = image->columns; 3239 global_work_size[1] = image->rows; 3240 3241 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, equalizeKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); 3242 3243 if (clStatus != CL_SUCCESS) 3244 { 3245 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 3246 goto cleanup; 3247 } 3248 clEnv->library->clFlush(queue); 3249 RecordProfileData(clEnv,EqualizeKernel,event); 3250 clEnv->library->clReleaseEvent(event); 3251 3252 /* read the data back */ 3253 if (ALIGNED(inputPixels,CLPixelPacket)) 3254 { 3255 length = image->columns * image->rows; 3256 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); 3257 } 3258 else 3259 { 3260 length = image->columns * image->rows; 3261 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL); 3262 } 3263 if (clStatus != CL_SUCCESS) 3264 { 3265 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 3266 goto cleanup; 3267 } 3268 3269 outputReady=SyncCacheViewAuthenticPixels(image_view,exception); 3270 3271cleanup: 3272 OpenCLLogException(__FUNCTION__,__LINE__,exception); 3273 3274 image_view=DestroyCacheView(image_view); 3275 3276 if (imageBuffer!=NULL) 3277 clEnv->library->clReleaseMemObject(imageBuffer); 3278 3279 if (map!=NULL) 3280 map=(FloatPixelPacket *) RelinquishMagickMemory(map); 3281 3282 if (equalizeMapBuffer!=NULL) 3283 clEnv->library->clReleaseMemObject(equalizeMapBuffer); 3284 if (equalize_map!=NULL) 3285 equalize_map=(PixelPacket *) RelinquishMagickMemory(equalize_map); 3286 3287 if (histogramBuffer!=NULL) 3288 clEnv->library->clReleaseMemObject(histogramBuffer); 3289 if (histogram!=NULL) 3290 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram); 3291 3292 if (histogramKernel!=NULL) 3293 RelinquishOpenCLKernel(clEnv, histogramKernel); 3294 if (equalizeKernel!=NULL) 3295 RelinquishOpenCLKernel(clEnv, equalizeKernel); 3296 3297 if (queue != NULL) 3298 RelinquishOpenCLCommandQueue(clEnv, queue); 3299 3300 return(outputReady); 3301} 3302 3303MagickExport MagickBooleanType AccelerateEqualizeImage(Image *image, 3304 ExceptionInfo *exception) 3305{ 3306 MagickBooleanType 3307 status; 3308 3309 assert(image != NULL); 3310 assert(exception != (ExceptionInfo *) NULL); 3311 3312 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 3313 (checkHistogramCondition(image,image->intensity) == MagickFalse) || 3314 (checkOpenCLEnvironment(exception) == MagickFalse)) 3315 return(MagickFalse); 3316 3317 status=ComputeEqualizeImage(image,exception); 3318 return(status); 3319} 3320 3321/* 3322%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 3323% % 3324% % 3325% % 3326% A c c e l e r a t e F u n c t i o n I m a g e % 3327% % 3328% % 3329% % 3330%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 3331*/ 3332 3333static MagickBooleanType ComputeFunctionImage(Image *image, 3334 const MagickFunction function,const size_t number_parameters, 3335 const double *parameters,ExceptionInfo *exception) 3336{ 3337 CacheView 3338 *image_view; 3339 3340 cl_command_queue 3341 queue; 3342 3343 cl_context 3344 context; 3345 3346 cl_int 3347 clStatus; 3348 3349 cl_kernel 3350 clkernel; 3351 3352 cl_event 3353 event; 3354 3355 cl_mem 3356 imageBuffer, 3357 parametersBuffer; 3358 3359 cl_uint 3360 number_channels; 3361 3362 float 3363 *parametersBufferPtr; 3364 3365 MagickBooleanType 3366 status; 3367 3368 MagickCLEnv 3369 clEnv; 3370 3371 size_t 3372 globalWorkSize[2]; 3373 3374 unsigned int 3375 i; 3376 3377 void 3378 *pixels; 3379 3380 status = MagickFalse; 3381 3382 context = NULL; 3383 clkernel = NULL; 3384 queue = NULL; 3385 imageBuffer = NULL; 3386 parametersBuffer = NULL; 3387 pixels = NULL; 3388 3389 clEnv = GetDefaultOpenCLEnv(); 3390 context = GetOpenCLContext(clEnv); 3391 3392 image_view=AcquireAuthenticCacheView(image,exception); 3393 imageBuffer=createReadWriteBuffer(image,image_view,clEnv,context,pixels, 3394 exception); 3395 if (imageBuffer == (cl_mem) NULL) 3396 goto cleanup; 3397 3398 parametersBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, number_parameters * sizeof(float), NULL, &clStatus); 3399 if (clStatus != CL_SUCCESS) 3400 { 3401 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 3402 goto cleanup; 3403 } 3404 3405 queue = AcquireOpenCLCommandQueue(clEnv); 3406 3407 parametersBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, parametersBuffer, CL_TRUE, CL_MAP_WRITE, 0, number_parameters * sizeof(float) 3408 , 0, NULL, NULL, &clStatus); 3409 if (clStatus != CL_SUCCESS) 3410 { 3411 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.","."); 3412 goto cleanup; 3413 } 3414 for (i = 0; i < number_parameters; i++) 3415 { 3416 parametersBufferPtr[i] = (float)parameters[i]; 3417 } 3418 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, parametersBuffer, parametersBufferPtr, 0, NULL, NULL); 3419 if (clStatus != CL_SUCCESS) 3420 { 3421 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); 3422 goto cleanup; 3423 } 3424 clEnv->library->clFlush(queue); 3425 3426 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ComputeFunction"); 3427 if (clkernel == NULL) 3428 { 3429 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 3430 goto cleanup; 3431 } 3432 3433 number_channels = (cl_uint) image->number_channels; 3434 3435 /* set the kernel arguments */ 3436 i = 0; 3437 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 3438 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_uint),(void *)&number_channels); 3439 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&image->channel_mask); 3440 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(MagickFunction),(void *)&function); 3441 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&number_parameters); 3442 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)¶metersBuffer); 3443 if (clStatus != CL_SUCCESS) 3444 { 3445 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 3446 goto cleanup; 3447 } 3448 3449 globalWorkSize[0] = image->columns; 3450 globalWorkSize[1] = image->rows; 3451 /* launch the kernel */ 3452 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, globalWorkSize, NULL, 0, NULL, &event); 3453 if (clStatus != CL_SUCCESS) 3454 { 3455 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 3456 goto cleanup; 3457 } 3458 clEnv->library->clFlush(queue); 3459 RecordProfileData(clEnv,ComputeFunctionKernel,event); 3460 clEnv->library->clReleaseEvent(event); 3461 3462 if (copyWriteBuffer(image,clEnv,queue,imageBuffer,pixels,exception) == MagickFalse) 3463 goto cleanup; 3464 3465 status=SyncCacheViewAuthenticPixels(image_view,exception); 3466 3467cleanup: 3468 OpenCLLogException(__FUNCTION__,__LINE__,exception); 3469 3470 image_view=DestroyCacheView(image_view); 3471 3472 if (clkernel != NULL) RelinquishOpenCLKernel(clEnv, clkernel); 3473 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 3474 if (imageBuffer != NULL) clEnv->library->clReleaseMemObject(imageBuffer); 3475 if (parametersBuffer != NULL) clEnv->library->clReleaseMemObject(parametersBuffer); 3476 3477 return(status); 3478} 3479 3480MagickExport MagickBooleanType AccelerateFunctionImage(Image *image, 3481 const MagickFunction function,const size_t number_parameters, 3482 const double *parameters,ExceptionInfo *exception) 3483{ 3484 MagickBooleanType 3485 status; 3486 3487 assert(image != NULL); 3488 assert(exception != (ExceptionInfo *) NULL); 3489 3490 if ((checkAccelerateCondition(image) == MagickFalse) || 3491 (checkOpenCLEnvironment(exception) == MagickFalse)) 3492 return(MagickFalse); 3493 3494 status=ComputeFunctionImage(image,function,number_parameters,parameters, 3495 exception); 3496 return(status); 3497} 3498 3499/* 3500%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 3501% % 3502% % 3503% % 3504% A c c e l e r a t e G r a y s c a l e I m a g e % 3505% % 3506% % 3507% % 3508%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 3509*/ 3510 3511static MagickBooleanType ComputeGrayscaleImage(Image *image, 3512 const PixelIntensityMethod method,ExceptionInfo *exception) 3513{ 3514 CacheView 3515 *image_view; 3516 3517 cl_command_queue 3518 queue; 3519 3520 cl_context 3521 context; 3522 3523 cl_int 3524 clStatus; 3525 3526 cl_kernel 3527 grayscaleKernel; 3528 3529 cl_event 3530 event; 3531 3532 cl_mem 3533 imageBuffer; 3534 3535 cl_uint 3536 number_channels, 3537 colorspace, 3538 intensityMethod; 3539 3540 MagickBooleanType 3541 outputReady; 3542 3543 MagickCLEnv 3544 clEnv; 3545 3546 register ssize_t 3547 i; 3548 3549 void 3550 *inputPixels; 3551 3552 outputReady = MagickFalse; 3553 inputPixels = NULL; 3554 grayscaleKernel = NULL; 3555 3556 assert(image != (Image *) NULL); 3557 assert(image->signature == MagickCoreSignature); 3558 if (image->debug != MagickFalse) 3559 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename); 3560 3561 /* 3562 * initialize opencl env 3563 */ 3564 clEnv = GetDefaultOpenCLEnv(); 3565 context = GetOpenCLContext(clEnv); 3566 queue = AcquireOpenCLCommandQueue(clEnv); 3567 3568 /* Create and initialize OpenCL buffers. 3569 inputPixels = AcquirePixelCachePixels(image, &length, exception); 3570 assume this will get a writable image 3571 */ 3572 image_view=AcquireAuthenticCacheView(image,exception); 3573 imageBuffer=createReadWriteBuffer(image,image_view,clEnv,context,inputPixels, 3574 exception); 3575 if (imageBuffer == (cl_mem) NULL) 3576 goto cleanup; 3577 3578 grayscaleKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Grayscale"); 3579 if (grayscaleKernel == NULL) 3580 { 3581 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 3582 goto cleanup; 3583 } 3584 3585 number_channels = (cl_uint) image->number_channels; 3586 intensityMethod = (cl_uint) method; 3587 colorspace = (cl_uint) image->colorspace; 3588 3589 i = 0; 3590 clStatus=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 3591 clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_uint),&number_channels); 3592 clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_uint),&colorspace); 3593 clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_uint),&intensityMethod); 3594 if (clStatus != CL_SUCCESS) 3595 { 3596 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 3597 goto cleanup; 3598 } 3599 3600 { 3601 size_t global_work_size[2]; 3602 global_work_size[0] = image->columns; 3603 global_work_size[1] = image->rows; 3604 /* launch the kernel */ 3605 clStatus=clEnv->library->clEnqueueNDRangeKernel(queue, grayscaleKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); 3606 if (clStatus != CL_SUCCESS) 3607 { 3608 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 3609 goto cleanup; 3610 } 3611 clEnv->library->clFlush(queue); 3612 RecordProfileData(clEnv,GrayScaleKernel,event); 3613 clEnv->library->clReleaseEvent(event); 3614 } 3615 3616 if (copyWriteBuffer(image,clEnv,queue,imageBuffer,inputPixels,exception) == MagickFalse) 3617 goto cleanup; 3618 3619 outputReady=SyncCacheViewAuthenticPixels(image_view,exception); 3620 3621cleanup: 3622 OpenCLLogException(__FUNCTION__,__LINE__,exception); 3623 3624 image_view=DestroyCacheView(image_view); 3625 3626 if (imageBuffer!=NULL) 3627 clEnv->library->clReleaseMemObject(imageBuffer); 3628 if (grayscaleKernel!=NULL) 3629 RelinquishOpenCLKernel(clEnv, grayscaleKernel); 3630 if (queue != NULL) 3631 RelinquishOpenCLCommandQueue(clEnv, queue); 3632 3633 return( outputReady); 3634} 3635 3636MagickExport MagickBooleanType AccelerateGrayscaleImage(Image* image, 3637 const PixelIntensityMethod method,ExceptionInfo *exception) 3638{ 3639 MagickBooleanType 3640 status; 3641 3642 assert(image != NULL); 3643 assert(exception != (ExceptionInfo *) NULL); 3644 3645 if ((checkAccelerateCondition(image) == MagickFalse) || 3646 (checkPixelIntensity(image,method) == MagickFalse) || 3647 (checkOpenCLEnvironment(exception) == MagickFalse)) 3648 return(MagickFalse); 3649 3650 if (image->number_channels < 3) 3651 return(MagickFalse); 3652 3653 if ((GetPixelRedTraits(image) == UndefinedPixelTrait) || 3654 (GetPixelGreenTraits(image) == UndefinedPixelTrait) || 3655 (GetPixelBlueTraits(image) == UndefinedPixelTrait)) 3656 return(MagickFalse); 3657 3658 status=ComputeGrayscaleImage(image,method,exception); 3659 return(status); 3660} 3661 3662/* 3663%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 3664% % 3665% % 3666% % 3667% A c c e l e r a t e L o c a l C o n t r a s t I m a g e % 3668% % 3669% % 3670% % 3671%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 3672*/ 3673 3674static Image *ComputeLocalContrastImage(const Image *image, 3675 const double radius,const double strength,ExceptionInfo *exception) 3676{ 3677 CacheView 3678 *filteredImage_view, 3679 *image_view; 3680 3681 cl_command_queue 3682 queue; 3683 3684 cl_context 3685 context; 3686 3687 cl_int 3688 clStatus, 3689 iRadius; 3690 3691 cl_kernel 3692 blurRowKernel, 3693 blurColumnKernel; 3694 3695 cl_event 3696 event; 3697 3698 cl_mem 3699 filteredImageBuffer, 3700 imageBuffer, 3701 imageKernelBuffer, 3702 tempImageBuffer; 3703 3704 cl_mem_flags 3705 mem_flags; 3706 3707 const void 3708 *inputPixels; 3709 3710 Image 3711 *filteredImage; 3712 3713 MagickBooleanType 3714 outputReady; 3715 3716 MagickCLEnv 3717 clEnv; 3718 3719 MagickSizeType 3720 length; 3721 3722 void 3723 *filteredPixels, 3724 *hostPtr; 3725 3726 unsigned int 3727 i, 3728 imageColumns, 3729 imageRows, 3730 passes; 3731 3732 clEnv = NULL; 3733 filteredImage = NULL; 3734 filteredImage_view = NULL; 3735 context = NULL; 3736 imageBuffer = NULL; 3737 filteredImageBuffer = NULL; 3738 tempImageBuffer = NULL; 3739 imageKernelBuffer = NULL; 3740 blurRowKernel = NULL; 3741 blurColumnKernel = NULL; 3742 queue = NULL; 3743 outputReady = MagickFalse; 3744 3745 clEnv = GetDefaultOpenCLEnv(); 3746 context = GetOpenCLContext(clEnv); 3747 queue = AcquireOpenCLCommandQueue(clEnv); 3748 3749 /* Create and initialize OpenCL buffers. */ 3750 { 3751 image_view=AcquireVirtualCacheView(image,exception); 3752 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception); 3753 if (inputPixels == (const void *) NULL) 3754 { 3755 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); 3756 goto cleanup; 3757 } 3758 3759 /* If the host pointer is aligned to the size of CLPixelPacket, 3760 then use the host buffer directly from the GPU; otherwise, 3761 create a buffer on the GPU and copy the data over */ 3762 if (ALIGNED(inputPixels,CLPixelPacket)) 3763 { 3764 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; 3765 } 3766 else 3767 { 3768 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR; 3769 } 3770 /* create a CL buffer from image pixel buffer */ 3771 length = image->columns * image->rows; 3772 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 3773 if (clStatus != CL_SUCCESS) 3774 { 3775 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 3776 goto cleanup; 3777 } 3778 } 3779 3780 /* create output */ 3781 { 3782 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception); 3783 assert(filteredImage != NULL); 3784 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) 3785 { 3786 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); 3787 goto cleanup; 3788 } 3789 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); 3790 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception); 3791 if (filteredPixels == (void *) NULL) 3792 { 3793 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); 3794 goto cleanup; 3795 } 3796 3797 if (ALIGNED(filteredPixels,CLPixelPacket)) 3798 { 3799 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR; 3800 hostPtr = filteredPixels; 3801 } 3802 else 3803 { 3804 mem_flags = CL_MEM_WRITE_ONLY; 3805 hostPtr = NULL; 3806 } 3807 3808 /* create a CL buffer from image pixel buffer */ 3809 length = image->columns * image->rows; 3810 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); 3811 if (clStatus != CL_SUCCESS) 3812 { 3813 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 3814 goto cleanup; 3815 } 3816 } 3817 3818 { 3819 /* create temp buffer */ 3820 { 3821 length = image->columns * image->rows; 3822 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(float), NULL, &clStatus); 3823 if (clStatus != CL_SUCCESS) 3824 { 3825 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 3826 goto cleanup; 3827 } 3828 } 3829 3830 /* get the opencl kernel */ 3831 { 3832 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "LocalContrastBlurRow"); 3833 if (blurRowKernel == NULL) 3834 { 3835 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 3836 goto cleanup; 3837 }; 3838 3839 blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "LocalContrastBlurApplyColumn"); 3840 if (blurColumnKernel == NULL) 3841 { 3842 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 3843 goto cleanup; 3844 }; 3845 } 3846 3847 { 3848 imageColumns = (unsigned int) image->columns; 3849 imageRows = (unsigned int) image->rows; 3850 iRadius = (cl_int) (image->rows > image->columns ? image->rows : image->columns) * 0.002f * fabs(radius); // Normalized radius, 100% gives blur radius of 20% of the largest dimension 3851 3852 passes = ((1.0f * imageColumns) * imageColumns * iRadius) / 4000000000.0f; 3853 passes = (passes < 1) ? 1: passes; 3854 3855 /* set the kernel arguments */ 3856 i = 0; 3857 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 3858 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); 3859 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer); 3860 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_int),(void *)&iRadius); 3861 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns); 3862 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows); 3863 3864 if (clStatus != CL_SUCCESS) 3865 { 3866 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 3867 goto cleanup; 3868 } 3869 } 3870 3871 /* launch the kernel */ 3872 { 3873 int x; 3874 for (x = 0; x < passes; ++x) { 3875 size_t gsize[2]; 3876 size_t wsize[2]; 3877 size_t goffset[2]; 3878 3879 gsize[0] = 256; 3880 gsize[1] = image->rows / passes; 3881 wsize[0] = 256; 3882 wsize[1] = 1; 3883 goffset[0] = 0; 3884 goffset[1] = x * gsize[1]; 3885 3886 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, goffset, gsize, wsize, 0, NULL, &event); 3887 if (clStatus != CL_SUCCESS) 3888 { 3889 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 3890 goto cleanup; 3891 } 3892 RecordProfileData(clEnv,LocalContrastBlurRowKernel,event); 3893 clEnv->library->clReleaseEvent(event); 3894 } 3895 } 3896 3897 { 3898 cl_float FStrength = strength; 3899 i = 0; 3900 clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 3901 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); 3902 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer); 3903 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&iRadius); 3904 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_float),(void *)&FStrength); 3905 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns); 3906 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows); 3907 3908 if (clStatus != CL_SUCCESS) 3909 { 3910 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 3911 goto cleanup; 3912 } 3913 } 3914 3915 /* launch the kernel */ 3916 { 3917 int x; 3918 for (x = 0; x < passes; ++x) { 3919 size_t gsize[2]; 3920 size_t wsize[2]; 3921 size_t goffset[2]; 3922 3923 gsize[0] = ((image->columns + 3) / 4) * 4; 3924 gsize[1] = ((((image->rows + 63) / 64) + (passes + 1)) / passes) * 64; 3925 wsize[0] = 4; 3926 wsize[1] = 64; 3927 goffset[0] = 0; 3928 goffset[1] = x * gsize[1]; 3929 3930 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, goffset, gsize, wsize, 0, NULL, &event); 3931 if (clStatus != CL_SUCCESS) 3932 { 3933 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 3934 goto cleanup; 3935 } 3936 RecordProfileData(clEnv,LocalContrastBlurApplyColumnKernel,event); 3937 clEnv->library->clReleaseEvent(event); 3938 } 3939 } 3940 } 3941 3942 /* get result */ 3943 if (ALIGNED(filteredPixels,CLPixelPacket)) 3944 { 3945 length = image->columns * image->rows; 3946 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); 3947 } 3948 else 3949 { 3950 length = image->columns * image->rows; 3951 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL); 3952 } 3953 if (clStatus != CL_SUCCESS) 3954 { 3955 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 3956 goto cleanup; 3957 } 3958 3959 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); 3960 3961cleanup: 3962 OpenCLLogException(__FUNCTION__,__LINE__,exception); 3963 3964 image_view=DestroyCacheView(image_view); 3965 if (filteredImage_view != NULL) 3966 filteredImage_view=DestroyCacheView(filteredImage_view); 3967 3968 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); 3969 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); 3970 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer); 3971 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer); 3972 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel); 3973 if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel); 3974 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 3975 if (outputReady == MagickFalse) 3976 { 3977 if (filteredImage != NULL) 3978 { 3979 DestroyImage(filteredImage); 3980 filteredImage = NULL; 3981 } 3982 } 3983 return(filteredImage); 3984} 3985 3986MagickExport Image *AccelerateLocalContrastImage(const Image *image, 3987 const double radius,const double strength,ExceptionInfo *exception) 3988{ 3989 Image 3990 *filteredImage; 3991 3992 assert(image != NULL); 3993 assert(exception != (ExceptionInfo *) NULL); 3994 3995 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 3996 (checkOpenCLEnvironment(exception) == MagickFalse)) 3997 return NULL; 3998 3999 filteredImage=ComputeLocalContrastImage(image,radius,strength,exception); 4000 return(filteredImage); 4001} 4002 4003/* 4004%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 4005% % 4006% % 4007% % 4008% A c c e l e r a t e M o d u l a t e I m a g e % 4009% % 4010% % 4011% % 4012%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 4013*/ 4014 4015static MagickBooleanType ComputeModulateImage(Image *image, 4016 const double percent_brightness,const double percent_hue, 4017 const double percent_saturation,const ColorspaceType colorspace, 4018 ExceptionInfo *exception) 4019{ 4020 CacheView 4021 *image_view; 4022 4023 cl_float 4024 bright, 4025 hue, 4026 saturation; 4027 4028 cl_context 4029 context; 4030 4031 cl_command_queue 4032 queue; 4033 4034 cl_int 4035 color, 4036 clStatus; 4037 4038 cl_kernel 4039 modulateKernel; 4040 4041 cl_event 4042 event; 4043 4044 cl_mem 4045 imageBuffer; 4046 4047 cl_mem_flags 4048 mem_flags; 4049 4050 MagickBooleanType 4051 outputReady; 4052 4053 MagickCLEnv 4054 clEnv; 4055 4056 MagickSizeType 4057 length; 4058 4059 register ssize_t 4060 i; 4061 4062 void 4063 *inputPixels; 4064 4065 inputPixels = NULL; 4066 imageBuffer = NULL; 4067 modulateKernel = NULL; 4068 4069 assert(image != (Image *) NULL); 4070 assert(image->signature == MagickCoreSignature); 4071 if (image->debug != MagickFalse) 4072 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename); 4073 4074 /* 4075 * initialize opencl env 4076 */ 4077 clEnv = GetDefaultOpenCLEnv(); 4078 context = GetOpenCLContext(clEnv); 4079 queue = AcquireOpenCLCommandQueue(clEnv); 4080 4081 outputReady = MagickFalse; 4082 4083 /* Create and initialize OpenCL buffers. 4084 inputPixels = AcquirePixelCachePixels(image, &length, exception); 4085 assume this will get a writable image 4086 */ 4087 image_view=AcquireAuthenticCacheView(image,exception); 4088 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); 4089 if (inputPixels == (void *) NULL) 4090 { 4091 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); 4092 goto cleanup; 4093 } 4094 4095 /* If the host pointer is aligned to the size of CLPixelPacket, 4096 then use the host buffer directly from the GPU; otherwise, 4097 create a buffer on the GPU and copy the data over 4098 */ 4099 if (ALIGNED(inputPixels,CLPixelPacket)) 4100 { 4101 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR; 4102 } 4103 else 4104 { 4105 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; 4106 } 4107 /* create a CL buffer from image pixel buffer */ 4108 length = image->columns * image->rows; 4109 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 4110 if (clStatus != CL_SUCCESS) 4111 { 4112 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 4113 goto cleanup; 4114 } 4115 4116 modulateKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Modulate"); 4117 if (modulateKernel == NULL) 4118 { 4119 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 4120 goto cleanup; 4121 } 4122 4123 bright=percent_brightness; 4124 hue=percent_hue; 4125 saturation=percent_saturation; 4126 color=colorspace; 4127 4128 i = 0; 4129 clStatus=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 4130 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&bright); 4131 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&hue); 4132 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&saturation); 4133 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&color); 4134 if (clStatus != CL_SUCCESS) 4135 { 4136 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 4137 goto cleanup; 4138 } 4139 4140 { 4141 size_t global_work_size[2]; 4142 global_work_size[0] = image->columns; 4143 global_work_size[1] = image->rows; 4144 /* launch the kernel */ 4145 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, modulateKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); 4146 if (clStatus != CL_SUCCESS) 4147 { 4148 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 4149 goto cleanup; 4150 } 4151 clEnv->library->clFlush(queue); 4152 RecordProfileData(clEnv,ModulateKernel,event); 4153 clEnv->library->clReleaseEvent(event); 4154 } 4155 4156 if (ALIGNED(inputPixels,CLPixelPacket)) 4157 { 4158 length = image->columns * image->rows; 4159 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); 4160 } 4161 else 4162 { 4163 length = image->columns * image->rows; 4164 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL); 4165 } 4166 if (clStatus != CL_SUCCESS) 4167 { 4168 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 4169 goto cleanup; 4170 } 4171 4172 outputReady=SyncCacheViewAuthenticPixels(image_view,exception); 4173 4174cleanup: 4175 OpenCLLogException(__FUNCTION__,__LINE__,exception); 4176 4177 image_view=DestroyCacheView(image_view); 4178 4179 if (imageBuffer!=NULL) 4180 clEnv->library->clReleaseMemObject(imageBuffer); 4181 if (modulateKernel!=NULL) 4182 RelinquishOpenCLKernel(clEnv, modulateKernel); 4183 if (queue != NULL) 4184 RelinquishOpenCLCommandQueue(clEnv, queue); 4185 4186 return outputReady; 4187 4188} 4189 4190MagickExport MagickBooleanType AccelerateModulateImage(Image *image, 4191 const double percent_brightness,const double percent_hue, 4192 const double percent_saturation,const ColorspaceType colorspace, 4193 ExceptionInfo *exception) 4194{ 4195 MagickBooleanType 4196 status; 4197 4198 assert(image != NULL); 4199 assert(exception != (ExceptionInfo *) NULL); 4200 4201 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 4202 (checkOpenCLEnvironment(exception) == MagickFalse)) 4203 return(MagickFalse); 4204 4205 if ((colorspace != HSLColorspace) && (colorspace != UndefinedColorspace)) 4206 return(MagickFalse); 4207 4208 status=ComputeModulateImage(image,percent_brightness,percent_hue, 4209 percent_saturation,colorspace,exception); 4210 return(status); 4211} 4212 4213/* 4214%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 4215% % 4216% % 4217% % 4218% A c c e l e r a t e M o t i o n B l u r I m a g e % 4219% % 4220% % 4221% % 4222%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 4223*/ 4224 4225static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, 4226 const size_t width,const OffsetInfo *offset,ExceptionInfo *exception) 4227{ 4228 CacheView 4229 *filteredImage_view, 4230 *image_view; 4231 4232 cl_command_queue 4233 queue; 4234 4235 cl_context 4236 context; 4237 4238 cl_float4 4239 biasPixel; 4240 4241 cl_int 4242 clStatus; 4243 4244 cl_kernel 4245 motionBlurKernel; 4246 4247 cl_event 4248 event; 4249 4250 cl_mem 4251 filteredImageBuffer, 4252 imageBuffer, 4253 imageKernelBuffer, 4254 offsetBuffer; 4255 4256 cl_mem_flags 4257 mem_flags; 4258 4259 const void 4260 *inputPixels; 4261 4262 float 4263 *kernelBufferPtr; 4264 4265 Image 4266 *filteredImage; 4267 4268 int 4269 *offsetBufferPtr; 4270 4271 MagickBooleanType 4272 outputReady; 4273 4274 MagickCLEnv 4275 clEnv; 4276 4277 PixelInfo 4278 bias; 4279 4280 MagickSizeType 4281 length; 4282 4283 size_t 4284 global_work_size[2], 4285 local_work_size[2]; 4286 4287 unsigned int 4288 i, 4289 imageHeight, 4290 imageWidth, 4291 matte; 4292 4293 void 4294 *filteredPixels, 4295 *hostPtr; 4296 4297 outputReady = MagickFalse; 4298 context = NULL; 4299 filteredImage = NULL; 4300 filteredImage_view = NULL; 4301 imageBuffer = NULL; 4302 filteredImageBuffer = NULL; 4303 imageKernelBuffer = NULL; 4304 motionBlurKernel = NULL; 4305 queue = NULL; 4306 4307 clEnv = GetDefaultOpenCLEnv(); 4308 context = GetOpenCLContext(clEnv); 4309 4310 /* Create and initialize OpenCL buffers. */ 4311 4312 image_view=AcquireVirtualCacheView(image,exception); 4313 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception); 4314 if (inputPixels == (const void *) NULL) 4315 { 4316 (void) ThrowMagickException(exception,GetMagickModule(),CacheError, 4317 "UnableToReadPixelCache.","`%s'",image->filename); 4318 goto cleanup; 4319 } 4320 4321 // If the host pointer is aligned to the size of CLPixelPacket, 4322 // then use the host buffer directly from the GPU; otherwise, 4323 // create a buffer on the GPU and copy the data over 4324 if (ALIGNED(inputPixels,CLPixelPacket)) 4325 { 4326 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; 4327 } 4328 else 4329 { 4330 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR; 4331 } 4332 // create a CL buffer from image pixel buffer 4333 length = image->columns * image->rows; 4334 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, 4335 length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 4336 if (clStatus != CL_SUCCESS) 4337 { 4338 (void) ThrowMagickException(exception, GetMagickModule(), 4339 ResourceLimitError, "clEnv->library->clCreateBuffer failed.","."); 4340 goto cleanup; 4341 } 4342 4343 4344 filteredImage = CloneImage(image,image->columns,image->rows, 4345 MagickTrue,exception); 4346 assert(filteredImage != NULL); 4347 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) 4348 { 4349 (void) ThrowMagickException(exception, GetMagickModule(), 4350 ResourceLimitError, "CloneImage failed.", "'%s'", "."); 4351 goto cleanup; 4352 } 4353 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); 4354 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception); 4355 if (filteredPixels == (void *) NULL) 4356 { 4357 (void) ThrowMagickException(exception,GetMagickModule(),CacheError, 4358 "UnableToReadPixelCache.","`%s'",filteredImage->filename); 4359 goto cleanup; 4360 } 4361 4362 if (ALIGNED(filteredPixels,CLPixelPacket)) 4363 { 4364 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR; 4365 hostPtr = filteredPixels; 4366 } 4367 else 4368 { 4369 mem_flags = CL_MEM_WRITE_ONLY; 4370 hostPtr = NULL; 4371 } 4372 // create a CL buffer from image pixel buffer 4373 length = image->columns * image->rows; 4374 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, 4375 length * sizeof(CLPixelPacket), hostPtr, &clStatus); 4376 if (clStatus != CL_SUCCESS) 4377 { 4378 (void) ThrowMagickException(exception, GetMagickModule(), 4379 ResourceLimitError, "clEnv->library->clCreateBuffer failed.","."); 4380 goto cleanup; 4381 } 4382 4383 4384 imageKernelBuffer = clEnv->library->clCreateBuffer(context, 4385 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width * sizeof(float), NULL, 4386 &clStatus); 4387 if (clStatus != CL_SUCCESS) 4388 { 4389 (void) ThrowMagickException(exception, GetMagickModule(), 4390 ResourceLimitError, "clEnv->library->clCreateBuffer failed.","."); 4391 goto cleanup; 4392 } 4393 4394 queue = AcquireOpenCLCommandQueue(clEnv); 4395 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, 4396 CL_TRUE, CL_MAP_WRITE, 0, width * sizeof(float), 0, NULL, NULL, &clStatus); 4397 if (clStatus != CL_SUCCESS) 4398 { 4399 (void) ThrowMagickException(exception, GetMagickModule(), 4400 ResourceLimitError, "clEnv->library->clEnqueueMapBuffer failed.","."); 4401 goto cleanup; 4402 } 4403 for (i = 0; i < width; i++) 4404 { 4405 kernelBufferPtr[i] = (float) kernel[i]; 4406 } 4407 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 4408 0, NULL, NULL); 4409 if (clStatus != CL_SUCCESS) 4410 { 4411 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, 4412 "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); 4413 goto cleanup; 4414 } 4415 4416 offsetBuffer = clEnv->library->clCreateBuffer(context, 4417 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width * sizeof(cl_int2), NULL, 4418 &clStatus); 4419 if (clStatus != CL_SUCCESS) 4420 { 4421 (void) ThrowMagickException(exception, GetMagickModule(), 4422 ResourceLimitError, "clEnv->library->clCreateBuffer failed.","."); 4423 goto cleanup; 4424 } 4425 4426 offsetBufferPtr = (int*)clEnv->library->clEnqueueMapBuffer(queue, offsetBuffer, CL_TRUE, 4427 CL_MAP_WRITE, 0, width * sizeof(cl_int2), 0, NULL, NULL, &clStatus); 4428 if (clStatus != CL_SUCCESS) 4429 { 4430 (void) ThrowMagickException(exception, GetMagickModule(), 4431 ResourceLimitError, "clEnv->library->clEnqueueMapBuffer failed.","."); 4432 goto cleanup; 4433 } 4434 for (i = 0; i < width; i++) 4435 { 4436 offsetBufferPtr[2*i] = (int)offset[i].x; 4437 offsetBufferPtr[2*i+1] = (int)offset[i].y; 4438 } 4439 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, offsetBuffer, offsetBufferPtr, 0, 4440 NULL, NULL); 4441 if (clStatus != CL_SUCCESS) 4442 { 4443 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, 4444 "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); 4445 goto cleanup; 4446 } 4447 4448 4449 // get the OpenCL kernel 4450 motionBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, 4451 "MotionBlur"); 4452 if (motionBlurKernel == NULL) 4453 { 4454 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, 4455 "AcquireOpenCLKernel failed.", "'%s'", "."); 4456 goto cleanup; 4457 } 4458 4459 // set the kernel arguments 4460 i = 0; 4461 clStatus=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem), 4462 (void *)&imageBuffer); 4463 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem), 4464 (void *)&filteredImageBuffer); 4465 imageWidth = (unsigned int) image->columns; 4466 imageHeight = (unsigned int) image->rows; 4467 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int), 4468 &imageWidth); 4469 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int), 4470 &imageHeight); 4471 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem), 4472 (void *)&imageKernelBuffer); 4473 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int), 4474 &width); 4475 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem), 4476 (void *)&offsetBuffer); 4477 4478 GetPixelInfo(image,&bias); 4479 biasPixel.s[0] = bias.red; 4480 biasPixel.s[1] = bias.green; 4481 biasPixel.s[2] = bias.blue; 4482 biasPixel.s[3] = bias.alpha; 4483 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_float4), &biasPixel); 4484 4485 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(ChannelType), &image->channel_mask); 4486 matte = (image->alpha_trait > CopyPixelTrait)?1:0; 4487 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int), &matte); 4488 if (clStatus != CL_SUCCESS) 4489 { 4490 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, 4491 "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 4492 goto cleanup; 4493 } 4494 4495 // launch the kernel 4496 local_work_size[0] = 16; 4497 local_work_size[1] = 16; 4498 global_work_size[0] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize( 4499 (unsigned int) image->columns,(unsigned int) local_work_size[0]); 4500 global_work_size[1] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize( 4501 (unsigned int) image->rows,(unsigned int) local_work_size[1]); 4502 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, motionBlurKernel, 2, NULL, 4503 global_work_size, local_work_size, 0, NULL, &event); 4504 4505 if (clStatus != CL_SUCCESS) 4506 { 4507 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, 4508 "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 4509 goto cleanup; 4510 } 4511 clEnv->library->clFlush(queue); 4512 RecordProfileData(clEnv,MotionBlurKernel,event); 4513 clEnv->library->clReleaseEvent(event); 4514 4515 if (ALIGNED(filteredPixels,CLPixelPacket)) 4516 { 4517 length = image->columns * image->rows; 4518 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, 4519 CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, 4520 NULL, &clStatus); 4521 } 4522 else 4523 { 4524 length = image->columns * image->rows; 4525 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, 4526 length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL); 4527 } 4528 if (clStatus != CL_SUCCESS) 4529 { 4530 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, 4531 "Reading output image from CL buffer failed.", "'%s'", "."); 4532 goto cleanup; 4533 } 4534 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); 4535 4536cleanup: 4537 4538 image_view=DestroyCacheView(image_view); 4539 if (filteredImage_view != NULL) 4540 filteredImage_view=DestroyCacheView(filteredImage_view); 4541 4542 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); 4543 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); 4544 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer); 4545 if (motionBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, motionBlurKernel); 4546 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 4547 if (outputReady == MagickFalse && filteredImage != NULL) 4548 filteredImage=DestroyImage(filteredImage); 4549 4550 return(filteredImage); 4551} 4552 4553MagickExport Image *AccelerateMotionBlurImage(const Image *image, 4554 const double* kernel,const size_t width,const OffsetInfo *offset, 4555 ExceptionInfo *exception) 4556{ 4557 Image 4558 *filteredImage; 4559 4560 assert(image != NULL); 4561 assert(kernel != (double *) NULL); 4562 assert(offset != (OffsetInfo *) NULL); 4563 assert(exception != (ExceptionInfo *) NULL); 4564 4565 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 4566 (checkOpenCLEnvironment(exception) == MagickFalse)) 4567 return NULL; 4568 4569 filteredImage=ComputeMotionBlurImage(image,kernel,width,offset,exception); 4570 return(filteredImage); 4571} 4572 4573/* 4574%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 4575% % 4576% % 4577% % 4578% A c c e l e r a t e R a n d o m I m a g e % 4579% % 4580% % 4581% % 4582%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 4583*/ 4584 4585static MagickBooleanType LaunchRandomImageKernel(MagickCLEnv clEnv, 4586 cl_command_queue queue,cl_mem imageBuffer,const unsigned int imageColumns, 4587 const unsigned int imageRows,cl_mem seedBuffer, 4588 const unsigned int numGenerators,ExceptionInfo *exception) 4589{ 4590 int 4591 k; 4592 4593 cl_int 4594 clStatus; 4595 4596 cl_kernel 4597 randomImageKernel; 4598 4599 cl_event 4600 event; 4601 4602 MagickBooleanType 4603 status; 4604 4605 size_t 4606 global_work_size, 4607 local_work_size; 4608 4609 status = MagickFalse; 4610 randomImageKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "RandomNumberGenerator"); 4611 4612 k = 0; 4613 clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_mem),(void*)&imageBuffer); 4614 clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_uint),(void*)&imageColumns); 4615 clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_uint),(void*)&imageRows); 4616 clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_mem),(void*)&seedBuffer); 4617 { 4618 const float randNormNumerator = 1.0f; 4619 const unsigned int randNormDenominator = (unsigned int)(~0UL); 4620 clEnv->library->clSetKernelArg(randomImageKernel,k++, 4621 sizeof(float),(void*)&randNormNumerator); 4622 clEnv->library->clSetKernelArg(randomImageKernel,k++, 4623 sizeof(cl_uint),(void*)&randNormDenominator); 4624 } 4625 4626 4627 global_work_size = numGenerators; 4628 local_work_size = 64; 4629 4630 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue,randomImageKernel,1,NULL,&global_work_size, 4631 &local_work_size, 0, NULL, &event); 4632 4633 if (clStatus != CL_SUCCESS) 4634 { 4635 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, 4636 "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 4637 goto cleanup; 4638 } 4639 RecordProfileData(clEnv,RandomNumberGeneratorKernel,event); 4640 clEnv->library->clReleaseEvent(event); 4641 4642 status = MagickTrue; 4643 4644cleanup: 4645 if (randomImageKernel!=NULL) RelinquishOpenCLKernel(clEnv, randomImageKernel); 4646 return(status); 4647} 4648 4649static MagickBooleanType ComputeRandomImage(Image* image, 4650 ExceptionInfo* exception) 4651{ 4652 CacheView 4653 *image_view; 4654 4655 cl_command_queue 4656 queue; 4657 4658 cl_context 4659 context; 4660 4661 cl_int 4662 clStatus; 4663 4664 /* Don't release this buffer in this function !!! */ 4665 cl_mem 4666 randomNumberSeedsBuffer; 4667 4668 cl_mem_flags 4669 mem_flags; 4670 4671 cl_mem 4672 imageBuffer; 4673 4674 MagickBooleanType 4675 outputReady, 4676 status; 4677 4678 MagickCLEnv 4679 clEnv; 4680 4681 MagickSizeType 4682 length; 4683 4684 void 4685 *inputPixels; 4686 4687 status = MagickFalse; 4688 outputReady = MagickFalse; 4689 inputPixels = NULL; 4690 context = NULL; 4691 imageBuffer = NULL; 4692 queue = NULL; 4693 4694 clEnv = GetDefaultOpenCLEnv(); 4695 context = GetOpenCLContext(clEnv); 4696 4697 /* Create and initialize OpenCL buffers. */ 4698 image_view=AcquireAuthenticCacheView(image,exception); 4699 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); 4700 if (inputPixels == (void *) NULL) 4701 { 4702 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); 4703 goto cleanup; 4704 } 4705 4706 /* If the host pointer is aligned to the size of CLPixelPacket, 4707 then use the host buffer directly from the GPU; otherwise, 4708 create a buffer on the GPU and copy the data over */ 4709 if (ALIGNED(inputPixels,CLPixelPacket)) 4710 { 4711 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR; 4712 } 4713 else 4714 { 4715 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; 4716 } 4717 /* create a CL buffer from image pixel buffer */ 4718 length = image->columns * image->rows; 4719 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 4720 if (clStatus != CL_SUCCESS) 4721 { 4722 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 4723 goto cleanup; 4724 } 4725 4726 queue = AcquireOpenCLCommandQueue(clEnv); 4727 4728 randomNumberSeedsBuffer = GetAndLockRandSeedBuffer(clEnv); 4729 if (randomNumberSeedsBuffer==NULL) 4730 { 4731 (void) OpenCLThrowMagickException(exception, GetMagickModule(), 4732 ResourceLimitWarning, "Failed to get GPU random number generators.", 4733 "'%s'", "."); 4734 goto cleanup; 4735 } 4736 4737 status = LaunchRandomImageKernel(clEnv,queue, 4738 imageBuffer, 4739 (unsigned int) image->columns, 4740 (unsigned int) image->rows, 4741 randomNumberSeedsBuffer, 4742 GetNumRandGenerators(clEnv), 4743 exception); 4744 if (status==MagickFalse) 4745 { 4746 goto cleanup; 4747 } 4748 4749 if (ALIGNED(inputPixels,CLPixelPacket)) 4750 { 4751 length = image->columns * image->rows; 4752 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); 4753 } 4754 else 4755 { 4756 length = image->columns * image->rows; 4757 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL); 4758 } 4759 if (clStatus != CL_SUCCESS) 4760 { 4761 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 4762 goto cleanup; 4763 } 4764 outputReady=SyncCacheViewAuthenticPixels(image_view,exception); 4765 4766cleanup: 4767 OpenCLLogException(__FUNCTION__,__LINE__,exception); 4768 4769 image_view=DestroyCacheView(image_view); 4770 4771 UnlockRandSeedBuffer(clEnv); 4772 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); 4773 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 4774 return outputReady; 4775} 4776 4777MagickExport MagickBooleanType AccelerateRandomImage(Image *image, 4778 ExceptionInfo* exception) 4779{ 4780 MagickBooleanType 4781 status; 4782 4783 assert(image != NULL); 4784 assert(exception != (ExceptionInfo *) NULL); 4785 4786 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 4787 (checkOpenCLEnvironment(exception) == MagickFalse)) 4788 return(MagickFalse); 4789 4790 status=ComputeRandomImage(image,exception); 4791 return(status); 4792} 4793 4794/* 4795%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 4796% % 4797% % 4798% % 4799% A c c e l e r a t e R e s i z e I m a g e % 4800% % 4801% % 4802% % 4803%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 4804*/ 4805 4806static MagickBooleanType resizeHorizontalFilter(cl_mem image, 4807 const unsigned int imageColumns,const unsigned int imageRows, 4808 const unsigned int matte,cl_mem resizedImage, 4809 const unsigned int resizedColumns,const unsigned int resizedRows, 4810 const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients, 4811 const float xFactor,MagickCLEnv clEnv,cl_command_queue queue, 4812 ExceptionInfo *exception) 4813{ 4814 cl_kernel 4815 horizontalKernel; 4816 4817 cl_event 4818 event; 4819 4820 cl_int clStatus; 4821 4822 const unsigned int 4823 workgroupSize = 256; 4824 4825 float 4826 resizeFilterScale, 4827 resizeFilterSupport, 4828 resizeFilterWindowSupport, 4829 resizeFilterBlur, 4830 scale, 4831 support; 4832 4833 int 4834 cacheRangeStart, 4835 cacheRangeEnd, 4836 numCachedPixels, 4837 resizeFilterType, 4838 resizeWindowType; 4839 4840 MagickBooleanType 4841 status = MagickFalse; 4842 4843 size_t 4844 deviceLocalMemorySize, 4845 gammaAccumulatorLocalMemorySize, 4846 global_work_size[2], 4847 imageCacheLocalMemorySize, 4848 pixelAccumulatorLocalMemorySize, 4849 local_work_size[2], 4850 totalLocalMemorySize, 4851 weightAccumulatorLocalMemorySize; 4852 4853 unsigned int 4854 chunkSize, 4855 i, 4856 pixelPerWorkgroup; 4857 4858 horizontalKernel = NULL; 4859 status = MagickFalse; 4860 4861 /* 4862 Apply filter to resize vertically from image to resize image. 4863 */ 4864 scale=MAGICK_MAX(1.0/xFactor+MagickEpsilon,1.0); 4865 support=scale*GetResizeFilterSupport(resizeFilter); 4866 if (support < 0.5) 4867 { 4868 /* 4869 Support too small even for nearest neighbour: Reduce to point 4870 sampling. 4871 */ 4872 support=(MagickRealType) 0.5; 4873 scale=1.0; 4874 } 4875 scale=PerceptibleReciprocal(scale); 4876 4877 if (resizedColumns < workgroupSize) 4878 { 4879 chunkSize = 32; 4880 pixelPerWorkgroup = 32; 4881 } 4882 else 4883 { 4884 chunkSize = workgroupSize; 4885 pixelPerWorkgroup = workgroupSize; 4886 } 4887 4888 /* get the local memory size supported by the device */ 4889 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv); 4890 4891DisableMSCWarning(4127) 4892 while(1) 4893RestoreMSCWarning 4894 { 4895 /* calculate the local memory size needed per workgroup */ 4896 cacheRangeStart = (int) (((0 + 0.5)/xFactor+MagickEpsilon)-support+0.5); 4897 cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/xFactor+MagickEpsilon)+support+0.5); 4898 numCachedPixels = cacheRangeEnd - cacheRangeStart + 1; 4899 imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket); 4900 totalLocalMemorySize = imageCacheLocalMemorySize; 4901 4902 /* local size for the pixel accumulator */ 4903 pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4); 4904 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize; 4905 4906 /* local memory size for the weight accumulator */ 4907 weightAccumulatorLocalMemorySize = chunkSize * sizeof(float); 4908 totalLocalMemorySize+=weightAccumulatorLocalMemorySize; 4909 4910 /* local memory size for the gamma accumulator */ 4911 if (matte == 0) 4912 gammaAccumulatorLocalMemorySize = sizeof(float); 4913 else 4914 gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float); 4915 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize; 4916 4917 if (totalLocalMemorySize <= deviceLocalMemorySize) 4918 break; 4919 else 4920 { 4921 pixelPerWorkgroup = pixelPerWorkgroup/2; 4922 chunkSize = chunkSize/2; 4923 if (pixelPerWorkgroup == 0 4924 || chunkSize == 0) 4925 { 4926 /* quit, fallback to CPU */ 4927 goto cleanup; 4928 } 4929 } 4930 } 4931 4932 resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter); 4933 resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter); 4934 4935 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilter"); 4936 if (horizontalKernel == NULL) 4937 { 4938 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 4939 goto cleanup; 4940 } 4941 4942 i = 0; 4943 clStatus = clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&image); 4944 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageColumns); 4945 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageRows); 4946 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&matte); 4947 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&xFactor); 4948 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizedImage); 4949 4950 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedColumns); 4951 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedRows); 4952 4953 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeFilterType); 4954 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeWindowType); 4955 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients); 4956 4957 resizeFilterScale = (float) GetResizeFilterScale(resizeFilter); 4958 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterScale); 4959 4960 resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter); 4961 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterSupport); 4962 4963 resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter); 4964 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport); 4965 4966 resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter); 4967 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterBlur); 4968 4969 4970 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL); 4971 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), &numCachedPixels); 4972 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup); 4973 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &chunkSize); 4974 4975 4976 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL); 4977 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL); 4978 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL); 4979 4980 if (clStatus != CL_SUCCESS) 4981 { 4982 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 4983 goto cleanup; 4984 } 4985 4986 global_work_size[0] = (resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize; 4987 global_work_size[1] = resizedRows; 4988 4989 local_work_size[0] = workgroupSize; 4990 local_work_size[1] = 1; 4991 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, &event); 4992 (void) local_work_size; 4993 if (clStatus != CL_SUCCESS) 4994 { 4995 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 4996 goto cleanup; 4997 } 4998 clEnv->library->clFlush(queue); 4999 RecordProfileData(clEnv,ResizeHorizontalKernel,event); 5000 clEnv->library->clReleaseEvent(event); 5001 status = MagickTrue; 5002 5003 5004cleanup: 5005 OpenCLLogException(__FUNCTION__,__LINE__,exception); 5006 5007 if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel); 5008 5009 return(status); 5010} 5011 5012static MagickBooleanType resizeVerticalFilter(cl_mem image, 5013 const unsigned int imageColumns,const unsigned int imageRows, 5014 const unsigned int matte,cl_mem resizedImage, 5015 const unsigned int resizedColumns,const unsigned int resizedRows, 5016 const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients, 5017 const float yFactor,MagickCLEnv clEnv,cl_command_queue queue, 5018 ExceptionInfo *exception) 5019{ 5020 cl_kernel 5021 verticalKernel; 5022 5023 cl_event 5024 event; 5025 5026 cl_int clStatus; 5027 5028 const unsigned int 5029 workgroupSize = 256; 5030 5031 float 5032 resizeFilterScale, 5033 resizeFilterSupport, 5034 resizeFilterWindowSupport, 5035 resizeFilterBlur, 5036 scale, 5037 support; 5038 5039 int 5040 cacheRangeStart, 5041 cacheRangeEnd, 5042 numCachedPixels, 5043 resizeFilterType, 5044 resizeWindowType; 5045 5046 MagickBooleanType 5047 status = MagickFalse; 5048 5049 size_t 5050 deviceLocalMemorySize, 5051 gammaAccumulatorLocalMemorySize, 5052 global_work_size[2], 5053 imageCacheLocalMemorySize, 5054 pixelAccumulatorLocalMemorySize, 5055 local_work_size[2], 5056 totalLocalMemorySize, 5057 weightAccumulatorLocalMemorySize; 5058 5059 unsigned int 5060 chunkSize, 5061 i, 5062 pixelPerWorkgroup; 5063 5064 verticalKernel = NULL; 5065 status = MagickFalse; 5066 5067 /* 5068 Apply filter to resize vertically from image to resize image. 5069 */ 5070 scale=MAGICK_MAX(1.0/yFactor+MagickEpsilon,1.0); 5071 support=scale*GetResizeFilterSupport(resizeFilter); 5072 if (support < 0.5) 5073 { 5074 /* 5075 Support too small even for nearest neighbour: Reduce to point 5076 sampling. 5077 */ 5078 support=(MagickRealType) 0.5; 5079 scale=1.0; 5080 } 5081 scale=PerceptibleReciprocal(scale); 5082 5083 if (resizedRows < workgroupSize) 5084 { 5085 chunkSize = 32; 5086 pixelPerWorkgroup = 32; 5087 } 5088 else 5089 { 5090 chunkSize = workgroupSize; 5091 pixelPerWorkgroup = workgroupSize; 5092 } 5093 5094 /* get the local memory size supported by the device */ 5095 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv); 5096 5097DisableMSCWarning(4127) 5098 while(1) 5099RestoreMSCWarning 5100 { 5101 /* calculate the local memory size needed per workgroup */ 5102 cacheRangeStart = (int) (((0 + 0.5)/yFactor+MagickEpsilon)-support+0.5); 5103 cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/yFactor+MagickEpsilon)+support+0.5); 5104 numCachedPixels = cacheRangeEnd - cacheRangeStart + 1; 5105 imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket); 5106 totalLocalMemorySize = imageCacheLocalMemorySize; 5107 5108 /* local size for the pixel accumulator */ 5109 pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4); 5110 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize; 5111 5112 /* local memory size for the weight accumulator */ 5113 weightAccumulatorLocalMemorySize = chunkSize * sizeof(float); 5114 totalLocalMemorySize+=weightAccumulatorLocalMemorySize; 5115 5116 /* local memory size for the gamma accumulator */ 5117 if (matte == 0) 5118 gammaAccumulatorLocalMemorySize = sizeof(float); 5119 else 5120 gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float); 5121 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize; 5122 5123 if (totalLocalMemorySize <= deviceLocalMemorySize) 5124 break; 5125 else 5126 { 5127 pixelPerWorkgroup = pixelPerWorkgroup/2; 5128 chunkSize = chunkSize/2; 5129 if (pixelPerWorkgroup == 0 5130 || chunkSize == 0) 5131 { 5132 /* quit, fallback to CPU */ 5133 goto cleanup; 5134 } 5135 } 5136 } 5137 5138 resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter); 5139 resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter); 5140 5141 verticalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilter"); 5142 if (verticalKernel == NULL) 5143 { 5144 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 5145 goto cleanup; 5146 } 5147 5148 i = 0; 5149 clStatus = clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(cl_mem), (void*)&image); 5150 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(unsigned int), (void*)&imageColumns); 5151 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(unsigned int), (void*)&imageRows); 5152 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(unsigned int), (void*)&matte); 5153 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(float), (void*)&yFactor); 5154 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(cl_mem), (void*)&resizedImage); 5155 5156 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(unsigned int), (void*)&resizedColumns); 5157 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(unsigned int), (void*)&resizedRows); 5158 5159 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(int), (void*)&resizeFilterType); 5160 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(int), (void*)&resizeWindowType); 5161 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients); 5162 5163 resizeFilterScale = (float) GetResizeFilterScale(resizeFilter); 5164 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(float), (void*)&resizeFilterScale); 5165 5166 resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter); 5167 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(float), (void*)&resizeFilterSupport); 5168 5169 resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter); 5170 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport); 5171 5172 resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter); 5173 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(float), (void*)&resizeFilterBlur); 5174 5175 5176 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, imageCacheLocalMemorySize, NULL); 5177 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(int), &numCachedPixels); 5178 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup); 5179 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(unsigned int), &chunkSize); 5180 5181 5182 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, pixelAccumulatorLocalMemorySize, NULL); 5183 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, weightAccumulatorLocalMemorySize, NULL); 5184 clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, gammaAccumulatorLocalMemorySize, NULL); 5185 5186 if (clStatus != CL_SUCCESS) 5187 { 5188 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 5189 goto cleanup; 5190 } 5191 5192 global_work_size[0] = resizedColumns; 5193 global_work_size[1] = (resizedRows+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize; 5194 5195 local_work_size[0] = 1; 5196 local_work_size[1] = workgroupSize; 5197 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, verticalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, &event); 5198 if (clStatus != CL_SUCCESS) 5199 { 5200 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 5201 goto cleanup; 5202 } 5203 clEnv->library->clFlush(queue); 5204 RecordProfileData(clEnv,ResizeVerticalKernel,event); 5205 clEnv->library->clReleaseEvent(event); 5206 status = MagickTrue; 5207 5208 5209cleanup: 5210 OpenCLLogException(__FUNCTION__,__LINE__,exception); 5211 5212 if (verticalKernel != NULL) RelinquishOpenCLKernel(clEnv, verticalKernel); 5213 5214 return(status); 5215} 5216 5217static Image *ComputeResizeImage(const Image* image, 5218 const size_t resizedColumns,const size_t resizedRows, 5219 const ResizeFilter *resizeFilter,ExceptionInfo *exception) 5220{ 5221 CacheView 5222 *filteredImage_view, 5223 *image_view; 5224 5225 cl_command_queue 5226 queue; 5227 5228 cl_int 5229 clStatus; 5230 5231 cl_context 5232 context; 5233 5234 cl_mem 5235 cubicCoefficientsBuffer, 5236 filteredImageBuffer, 5237 imageBuffer, 5238 tempImageBuffer; 5239 5240 cl_mem_flags 5241 mem_flags; 5242 5243 const double 5244 *resizeFilterCoefficient; 5245 5246 const void 5247 *inputPixels; 5248 5249 float 5250 *mappedCoefficientBuffer, 5251 xFactor, 5252 yFactor; 5253 5254 MagickBooleanType 5255 outputReady, 5256 status; 5257 5258 MagickCLEnv 5259 clEnv; 5260 5261 MagickSizeType 5262 length; 5263 5264 Image 5265 *filteredImage; 5266 5267 unsigned int 5268 i, 5269 matte; 5270 5271 void 5272 *filteredPixels, 5273 *hostPtr; 5274 5275 outputReady = MagickFalse; 5276 filteredImage = NULL; 5277 filteredImage_view = NULL; 5278 clEnv = NULL; 5279 context = NULL; 5280 imageBuffer = NULL; 5281 tempImageBuffer = NULL; 5282 filteredImageBuffer = NULL; 5283 cubicCoefficientsBuffer = NULL; 5284 queue = NULL; 5285 5286 clEnv = GetDefaultOpenCLEnv(); 5287 context = GetOpenCLContext(clEnv); 5288 5289 /* Create and initialize OpenCL buffers. */ 5290 image_view=AcquireVirtualCacheView(image,exception); 5291 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception); 5292 if (inputPixels == (const void *) NULL) 5293 { 5294 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); 5295 goto cleanup; 5296 } 5297 5298 /* If the host pointer is aligned to the size of CLPixelPacket, 5299 then use the host buffer directly from the GPU; otherwise, 5300 create a buffer on the GPU and copy the data over */ 5301 if (ALIGNED(inputPixels,CLPixelPacket)) 5302 { 5303 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; 5304 } 5305 else 5306 { 5307 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR; 5308 } 5309 /* create a CL buffer from image pixel buffer */ 5310 length = image->columns * image->rows; 5311 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 5312 if (clStatus != CL_SUCCESS) 5313 { 5314 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 5315 goto cleanup; 5316 } 5317 5318 cubicCoefficientsBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY, 7 * sizeof(float), NULL, &clStatus); 5319 if (clStatus != CL_SUCCESS) 5320 { 5321 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 5322 goto cleanup; 5323 } 5324 queue = AcquireOpenCLCommandQueue(clEnv); 5325 mappedCoefficientBuffer = (float*)clEnv->library->clEnqueueMapBuffer(queue, cubicCoefficientsBuffer, CL_TRUE, CL_MAP_WRITE, 0, 7 * sizeof(float) 5326 , 0, NULL, NULL, &clStatus); 5327 if (clStatus != CL_SUCCESS) 5328 { 5329 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.","."); 5330 goto cleanup; 5331 } 5332 resizeFilterCoefficient = GetResizeFilterCoefficient(resizeFilter); 5333 for (i = 0; i < 7; i++) 5334 { 5335 mappedCoefficientBuffer[i] = (float) resizeFilterCoefficient[i]; 5336 } 5337 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, cubicCoefficientsBuffer, mappedCoefficientBuffer, 0, NULL, NULL); 5338 if (clStatus != CL_SUCCESS) 5339 { 5340 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); 5341 goto cleanup; 5342 } 5343 5344 filteredImage = CloneImage(image,resizedColumns,resizedRows,MagickTrue,exception); 5345 if (filteredImage == NULL) 5346 goto cleanup; 5347 5348 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) 5349 { 5350 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); 5351 goto cleanup; 5352 } 5353 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); 5354 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception); 5355 if (filteredPixels == (void *) NULL) 5356 { 5357 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); 5358 goto cleanup; 5359 } 5360 5361 if (ALIGNED(filteredPixels,CLPixelPacket)) 5362 { 5363 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR; 5364 hostPtr = filteredPixels; 5365 } 5366 else 5367 { 5368 mem_flags = CL_MEM_WRITE_ONLY; 5369 hostPtr = NULL; 5370 } 5371 5372 /* create a CL buffer from image pixel buffer */ 5373 length = filteredImage->columns * filteredImage->rows; 5374 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); 5375 if (clStatus != CL_SUCCESS) 5376 { 5377 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 5378 goto cleanup; 5379 } 5380 5381 xFactor=(float) resizedColumns/(float) image->columns; 5382 yFactor=(float) resizedRows/(float) image->rows; 5383 matte=(image->alpha_trait > CopyPixelTrait)?1:0; 5384 if (xFactor > yFactor) 5385 { 5386 5387 length = resizedColumns*image->rows; 5388 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLPixelPacket), NULL, &clStatus); 5389 if (clStatus != CL_SUCCESS) 5390 { 5391 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 5392 goto cleanup; 5393 } 5394 5395 status = resizeHorizontalFilter(imageBuffer, (unsigned int) image->columns, (unsigned int) image->rows, matte 5396 , tempImageBuffer, (unsigned int) resizedColumns, (unsigned int) image->rows 5397 , resizeFilter, cubicCoefficientsBuffer 5398 , xFactor, clEnv, queue, exception); 5399 if (status != MagickTrue) 5400 goto cleanup; 5401 5402 status = resizeVerticalFilter(tempImageBuffer, (unsigned int) resizedColumns, (unsigned int) image->rows, matte 5403 , filteredImageBuffer, (unsigned int) resizedColumns, (unsigned int) resizedRows 5404 , resizeFilter, cubicCoefficientsBuffer 5405 , yFactor, clEnv, queue, exception); 5406 if (status != MagickTrue) 5407 goto cleanup; 5408 } 5409 else 5410 { 5411 length = image->columns*resizedRows; 5412 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLPixelPacket), NULL, &clStatus); 5413 if (clStatus != CL_SUCCESS) 5414 { 5415 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 5416 goto cleanup; 5417 } 5418 5419 status = resizeVerticalFilter(imageBuffer, (unsigned int) image->columns, (unsigned int) image->rows, matte 5420 , tempImageBuffer, (unsigned int) image->columns, (unsigned int) resizedRows 5421 , resizeFilter, cubicCoefficientsBuffer 5422 , yFactor, clEnv, queue, exception); 5423 if (status != MagickTrue) 5424 goto cleanup; 5425 5426 status = resizeHorizontalFilter(tempImageBuffer, (unsigned int) image->columns, (unsigned int) resizedRows, matte 5427 , filteredImageBuffer, (unsigned int) resizedColumns, (unsigned int) resizedRows 5428 , resizeFilter, cubicCoefficientsBuffer 5429 , xFactor, clEnv, queue, exception); 5430 if (status != MagickTrue) 5431 goto cleanup; 5432 } 5433 length = resizedColumns*resizedRows; 5434 if (ALIGNED(filteredPixels,CLPixelPacket)) 5435 { 5436 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); 5437 } 5438 else 5439 { 5440 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL); 5441 } 5442 if (clStatus != CL_SUCCESS) 5443 { 5444 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 5445 goto cleanup; 5446 } 5447 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); 5448 5449cleanup: 5450 OpenCLLogException(__FUNCTION__,__LINE__,exception); 5451 5452 image_view=DestroyCacheView(image_view); 5453 if (filteredImage_view != NULL) 5454 filteredImage_view=DestroyCacheView(filteredImage_view); 5455 5456 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); 5457 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer); 5458 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); 5459 if (cubicCoefficientsBuffer!=NULL) clEnv->library->clReleaseMemObject(cubicCoefficientsBuffer); 5460 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 5461 if (outputReady == MagickFalse && filteredImage != NULL) 5462 filteredImage=DestroyImage(filteredImage); 5463 return(filteredImage); 5464} 5465 5466static MagickBooleanType gpuSupportedResizeWeighting( 5467 ResizeWeightingFunctionType f) 5468{ 5469 unsigned int 5470 i; 5471 5472 for (i = 0; ;i++) 5473 { 5474 if (supportedResizeWeighting[i] == LastWeightingFunction) 5475 break; 5476 if (supportedResizeWeighting[i] == f) 5477 return(MagickTrue); 5478 } 5479 return(MagickFalse); 5480} 5481 5482MagickExport Image *AccelerateResizeImage(const Image *image, 5483 const size_t resizedColumns,const size_t resizedRows, 5484 const ResizeFilter *resizeFilter,ExceptionInfo *exception) 5485{ 5486 Image 5487 *filteredImage; 5488 5489 assert(image != NULL); 5490 assert(exception != (ExceptionInfo *) NULL); 5491 5492 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 5493 (checkOpenCLEnvironment(exception) == MagickFalse)) 5494 return NULL; 5495 5496 if ((gpuSupportedResizeWeighting(GetResizeFilterWeightingType( 5497 resizeFilter)) == MagickFalse) || 5498 (gpuSupportedResizeWeighting(GetResizeFilterWindowWeightingType( 5499 resizeFilter)) == MagickFalse)) 5500 return NULL; 5501 5502 filteredImage=ComputeResizeImage(image,resizedColumns,resizedRows, 5503 resizeFilter,exception); 5504 return(filteredImage); 5505} 5506 5507/* 5508%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 5509% % 5510% % 5511% % 5512% A c c e l e r a t e R o t a t i o n a l B l u r I m a g e % 5513% % 5514% % 5515% % 5516%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 5517*/ 5518 5519static Image* ComputeRotationalBlurImage(const Image *image,const double angle, 5520 ExceptionInfo *exception) 5521{ 5522 CacheView 5523 *image_view, 5524 *filteredImage_view; 5525 5526 cl_command_queue 5527 queue; 5528 5529 cl_context 5530 context; 5531 5532 cl_float2 5533 blurCenter; 5534 5535 cl_float4 5536 biasPixel; 5537 5538 cl_int 5539 clStatus; 5540 5541 cl_mem 5542 cosThetaBuffer, 5543 filteredImageBuffer, 5544 imageBuffer, 5545 sinThetaBuffer; 5546 5547 cl_mem_flags 5548 mem_flags; 5549 5550 cl_kernel 5551 rotationalBlurKernel; 5552 5553 cl_event 5554 event; 5555 5556 const void 5557 *inputPixels; 5558 5559 float 5560 blurRadius, 5561 *cosThetaPtr, 5562 offset, 5563 *sinThetaPtr, 5564 theta; 5565 5566 Image 5567 *filteredImage; 5568 5569 MagickBooleanType 5570 outputReady; 5571 5572 MagickCLEnv 5573 clEnv; 5574 5575 PixelInfo 5576 bias; 5577 5578 MagickSizeType 5579 length; 5580 5581 size_t 5582 global_work_size[2]; 5583 5584 unsigned int 5585 cossin_theta_size, 5586 i, 5587 matte; 5588 5589 void 5590 *filteredPixels, 5591 *hostPtr; 5592 5593 outputReady = MagickFalse; 5594 context = NULL; 5595 filteredImage = NULL; 5596 filteredImage_view = NULL; 5597 imageBuffer = NULL; 5598 filteredImageBuffer = NULL; 5599 sinThetaBuffer = NULL; 5600 cosThetaBuffer = NULL; 5601 queue = NULL; 5602 rotationalBlurKernel = NULL; 5603 5604 5605 clEnv = GetDefaultOpenCLEnv(); 5606 context = GetOpenCLContext(clEnv); 5607 5608 5609 /* Create and initialize OpenCL buffers. */ 5610 5611 image_view=AcquireVirtualCacheView(image,exception); 5612 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception); 5613 if (inputPixels == (const void *) NULL) 5614 { 5615 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); 5616 goto cleanup; 5617 } 5618 5619 /* If the host pointer is aligned to the size of CLPixelPacket, 5620 then use the host buffer directly from the GPU; otherwise, 5621 create a buffer on the GPU and copy the data over */ 5622 if (ALIGNED(inputPixels,CLPixelPacket)) 5623 { 5624 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; 5625 } 5626 else 5627 { 5628 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR; 5629 } 5630 /* create a CL buffer from image pixel buffer */ 5631 length = image->columns * image->rows; 5632 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 5633 if (clStatus != CL_SUCCESS) 5634 { 5635 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 5636 goto cleanup; 5637 } 5638 5639 5640 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception); 5641 assert(filteredImage != NULL); 5642 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) 5643 { 5644 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); 5645 goto cleanup; 5646 } 5647 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); 5648 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception); 5649 if (filteredPixels == (void *) NULL) 5650 { 5651 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); 5652 goto cleanup; 5653 } 5654 5655 if (ALIGNED(filteredPixels,CLPixelPacket)) 5656 { 5657 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR; 5658 hostPtr = filteredPixels; 5659 } 5660 else 5661 { 5662 mem_flags = CL_MEM_WRITE_ONLY; 5663 hostPtr = NULL; 5664 } 5665 /* create a CL buffer from image pixel buffer */ 5666 length = image->columns * image->rows; 5667 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); 5668 if (clStatus != CL_SUCCESS) 5669 { 5670 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 5671 goto cleanup; 5672 } 5673 5674 blurCenter.s[0] = (float) (image->columns-1)/2.0; 5675 blurCenter.s[1] = (float) (image->rows-1)/2.0; 5676 blurRadius=hypot(blurCenter.s[0],blurCenter.s[1]); 5677 cossin_theta_size=(unsigned int) fabs(4.0*DegreesToRadians(angle)*sqrt((double)blurRadius)+2UL); 5678 5679 /* create a buffer for sin_theta and cos_theta */ 5680 sinThetaBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus); 5681 if (clStatus != CL_SUCCESS) 5682 { 5683 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 5684 goto cleanup; 5685 } 5686 cosThetaBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus); 5687 if (clStatus != CL_SUCCESS) 5688 { 5689 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 5690 goto cleanup; 5691 } 5692 5693 5694 queue = AcquireOpenCLCommandQueue(clEnv); 5695 sinThetaPtr = (float*) clEnv->library->clEnqueueMapBuffer(queue, sinThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus); 5696 if (clStatus != CL_SUCCESS) 5697 { 5698 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.","."); 5699 goto cleanup; 5700 } 5701 5702 cosThetaPtr = (float*) clEnv->library->clEnqueueMapBuffer(queue, cosThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus); 5703 if (clStatus != CL_SUCCESS) 5704 { 5705 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.","."); 5706 goto cleanup; 5707 } 5708 5709 theta=DegreesToRadians(angle)/(MagickRealType) (cossin_theta_size-1); 5710 offset=theta*(MagickRealType) (cossin_theta_size-1)/2.0; 5711 for (i=0; i < (ssize_t) cossin_theta_size; i++) 5712 { 5713 cosThetaPtr[i]=(float)cos((double) (theta*i-offset)); 5714 sinThetaPtr[i]=(float)sin((double) (theta*i-offset)); 5715 } 5716 5717 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, sinThetaBuffer, sinThetaPtr, 0, NULL, NULL); 5718 clStatus |= clEnv->library->clEnqueueUnmapMemObject(queue, cosThetaBuffer, cosThetaPtr, 0, NULL, NULL); 5719 if (clStatus != CL_SUCCESS) 5720 { 5721 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); 5722 goto cleanup; 5723 } 5724 5725 /* get the OpenCL kernel */ 5726 rotationalBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "RotationalBlur"); 5727 if (rotationalBlurKernel == NULL) 5728 { 5729 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 5730 goto cleanup; 5731 } 5732 5733 5734 /* set the kernel arguments */ 5735 i = 0; 5736 clStatus=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 5737 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); 5738 5739 GetPixelInfo(image,&bias); 5740 biasPixel.s[0] = bias.red; 5741 biasPixel.s[1] = bias.green; 5742 biasPixel.s[2] = bias.blue; 5743 biasPixel.s[3] = bias.alpha; 5744 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_float4), &biasPixel); 5745 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(ChannelType), &image->channel_mask); 5746 5747 matte = (image->alpha_trait > CopyPixelTrait)?1:0; 5748 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(unsigned int), &matte); 5749 5750 clStatus=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_float2), &blurCenter); 5751 5752 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&cosThetaBuffer); 5753 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&sinThetaBuffer); 5754 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(unsigned int), &cossin_theta_size); 5755 if (clStatus != CL_SUCCESS) 5756 { 5757 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 5758 goto cleanup; 5759 } 5760 5761 5762 global_work_size[0] = image->columns; 5763 global_work_size[1] = image->rows; 5764 /* launch the kernel */ 5765 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, rotationalBlurKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); 5766 if (clStatus != CL_SUCCESS) 5767 { 5768 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 5769 goto cleanup; 5770 } 5771 clEnv->library->clFlush(queue); 5772 RecordProfileData(clEnv,RotationalBlurKernel,event); 5773 clEnv->library->clReleaseEvent(event); 5774 5775 if (ALIGNED(filteredPixels,CLPixelPacket)) 5776 { 5777 length = image->columns * image->rows; 5778 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); 5779 } 5780 else 5781 { 5782 length = image->columns * image->rows; 5783 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL); 5784 } 5785 if (clStatus != CL_SUCCESS) 5786 { 5787 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 5788 goto cleanup; 5789 } 5790 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); 5791 5792cleanup: 5793 OpenCLLogException(__FUNCTION__,__LINE__,exception); 5794 5795 image_view=DestroyCacheView(image_view); 5796 if (filteredImage_view != NULL) 5797 filteredImage_view=DestroyCacheView(filteredImage_view); 5798 5799 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); 5800 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); 5801 if (sinThetaBuffer!=NULL) clEnv->library->clReleaseMemObject(sinThetaBuffer); 5802 if (cosThetaBuffer!=NULL) clEnv->library->clReleaseMemObject(cosThetaBuffer); 5803 if (rotationalBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, rotationalBlurKernel); 5804 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 5805 if (outputReady == MagickFalse) 5806 { 5807 if (filteredImage != NULL) 5808 { 5809 DestroyImage(filteredImage); 5810 filteredImage = NULL; 5811 } 5812 } 5813 return filteredImage; 5814} 5815 5816MagickExport Image* AccelerateRotationalBlurImage(const Image *image, 5817 const double angle,ExceptionInfo *exception) 5818{ 5819 Image 5820 *filteredImage; 5821 5822 assert(image != NULL); 5823 assert(exception != (ExceptionInfo *) NULL); 5824 5825 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 5826 (checkOpenCLEnvironment(exception) == MagickFalse)) 5827 return NULL; 5828 5829 filteredImage=ComputeRotationalBlurImage(image,angle,exception); 5830 return filteredImage; 5831} 5832 5833/* 5834%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 5835% % 5836% % 5837% % 5838% A c c e l e r a t e U n s h a r p M a s k I m a g e % 5839% % 5840% % 5841% % 5842%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% 5843*/ 5844 5845static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, 5846 const double sigma,const double gain,const double threshold, 5847 ExceptionInfo *exception) 5848{ 5849 CacheView 5850 *filteredImage_view, 5851 *image_view; 5852 5853 char 5854 geometry[MagickPathExtent]; 5855 5856 cl_command_queue 5857 queue; 5858 5859 cl_context 5860 context; 5861 5862 cl_int 5863 clStatus; 5864 5865 cl_kernel 5866 blurRowKernel, 5867 unsharpMaskBlurColumnKernel; 5868 5869 cl_event 5870 event; 5871 5872 cl_mem 5873 filteredImageBuffer, 5874 imageBuffer, 5875 imageKernelBuffer, 5876 tempImageBuffer; 5877 5878 cl_mem_flags 5879 mem_flags; 5880 5881 const void 5882 *inputPixels; 5883 5884 float 5885 fGain, 5886 fThreshold, 5887 *kernelBufferPtr; 5888 5889 Image 5890 *filteredImage; 5891 5892 int 5893 chunkSize; 5894 5895 KernelInfo 5896 *kernel; 5897 5898 MagickBooleanType 5899 outputReady; 5900 5901 MagickCLEnv 5902 clEnv; 5903 5904 MagickSizeType 5905 length; 5906 5907 void 5908 *filteredPixels, 5909 *hostPtr; 5910 5911 unsigned int 5912 i, 5913 imageColumns, 5914 imageRows, 5915 kernelWidth; 5916 5917 clEnv = NULL; 5918 filteredImage = NULL; 5919 filteredImage_view = NULL; 5920 kernel = NULL; 5921 context = NULL; 5922 imageBuffer = NULL; 5923 filteredImageBuffer = NULL; 5924 tempImageBuffer = NULL; 5925 imageKernelBuffer = NULL; 5926 blurRowKernel = NULL; 5927 unsharpMaskBlurColumnKernel = NULL; 5928 queue = NULL; 5929 outputReady = MagickFalse; 5930 5931 clEnv = GetDefaultOpenCLEnv(); 5932 context = GetOpenCLContext(clEnv); 5933 queue = AcquireOpenCLCommandQueue(clEnv); 5934 5935 /* Create and initialize OpenCL buffers. */ 5936 { 5937 image_view=AcquireVirtualCacheView(image,exception); 5938 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception); 5939 if (inputPixels == (const void *) NULL) 5940 { 5941 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); 5942 goto cleanup; 5943 } 5944 5945 /* If the host pointer is aligned to the size of CLPixelPacket, 5946 then use the host buffer directly from the GPU; otherwise, 5947 create a buffer on the GPU and copy the data over */ 5948 if (ALIGNED(inputPixels,CLPixelPacket)) 5949 { 5950 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; 5951 } 5952 else 5953 { 5954 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR; 5955 } 5956 /* create a CL buffer from image pixel buffer */ 5957 length = image->columns * image->rows; 5958 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); 5959 if (clStatus != CL_SUCCESS) 5960 { 5961 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 5962 goto cleanup; 5963 } 5964 } 5965 5966 /* create output */ 5967 { 5968 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception); 5969 assert(filteredImage != NULL); 5970 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) 5971 { 5972 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); 5973 goto cleanup; 5974 } 5975 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); 5976 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception); 5977 if (filteredPixels == (void *) NULL) 5978 { 5979 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); 5980 goto cleanup; 5981 } 5982 5983 if (ALIGNED(filteredPixels,CLPixelPacket)) 5984 { 5985 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR; 5986 hostPtr = filteredPixels; 5987 } 5988 else 5989 { 5990 mem_flags = CL_MEM_WRITE_ONLY; 5991 hostPtr = NULL; 5992 } 5993 5994 /* create a CL buffer from image pixel buffer */ 5995 length = image->columns * image->rows; 5996 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); 5997 if (clStatus != CL_SUCCESS) 5998 { 5999 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 6000 goto cleanup; 6001 } 6002 } 6003 6004 /* create the blur kernel */ 6005 { 6006 (void) FormatLocaleString(geometry,MagickPathExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma); 6007 kernel=AcquireKernelInfo(geometry,exception); 6008 if (kernel == (KernelInfo *) NULL) 6009 { 6010 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.","."); 6011 goto cleanup; 6012 } 6013 6014 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY, kernel->width * sizeof(float), NULL, &clStatus); 6015 if (clStatus != CL_SUCCESS) 6016 { 6017 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 6018 goto cleanup; 6019 } 6020 6021 6022 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus); 6023 if (clStatus != CL_SUCCESS) 6024 { 6025 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.","."); 6026 goto cleanup; 6027 } 6028 for (i = 0; i < kernel->width; i++) 6029 { 6030 kernelBufferPtr[i] = (float) kernel->values[i]; 6031 } 6032 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL); 6033 if (clStatus != CL_SUCCESS) 6034 { 6035 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); 6036 goto cleanup; 6037 } 6038 } 6039 6040 { 6041 /* create temp buffer */ 6042 { 6043 length = image->columns * image->rows; 6044 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus); 6045 if (clStatus != CL_SUCCESS) 6046 { 6047 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); 6048 goto cleanup; 6049 } 6050 } 6051 6052 /* get the opencl kernel */ 6053 { 6054 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow"); 6055 if (blurRowKernel == NULL) 6056 { 6057 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 6058 goto cleanup; 6059 }; 6060 6061 unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMaskBlurColumn"); 6062 if (unsharpMaskBlurColumnKernel == NULL) 6063 { 6064 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 6065 goto cleanup; 6066 }; 6067 } 6068 6069 { 6070 chunkSize = 256; 6071 6072 imageColumns = (unsigned int) image->columns; 6073 imageRows = (unsigned int) image->rows; 6074 6075 kernelWidth = (unsigned int) kernel->width; 6076 6077 /* set the kernel arguments */ 6078 i = 0; 6079 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 6080 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer); 6081 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&image->channel_mask); 6082 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer); 6083 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth); 6084 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns); 6085 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows); 6086 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *) NULL); 6087 if (clStatus != CL_SUCCESS) 6088 { 6089 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 6090 goto cleanup; 6091 } 6092 } 6093 6094 /* launch the kernel */ 6095 { 6096 size_t gsize[2]; 6097 size_t wsize[2]; 6098 6099 gsize[0] = chunkSize*((image->columns+chunkSize-1)/chunkSize); 6100 gsize[1] = image->rows; 6101 wsize[0] = chunkSize; 6102 wsize[1] = 1; 6103 6104 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, &event); 6105 if (clStatus != CL_SUCCESS) 6106 { 6107 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 6108 goto cleanup; 6109 } 6110 clEnv->library->clFlush(queue); 6111 RecordProfileData(clEnv,BlurRowKernel,event); 6112 clEnv->library->clReleaseEvent(event); 6113 } 6114 6115 6116 { 6117 chunkSize = 256; 6118 imageColumns = (unsigned int) image->columns; 6119 imageRows = (unsigned int) image->rows; 6120 kernelWidth = (unsigned int) kernel->width; 6121 fGain = (float) gain; 6122 fThreshold = (float) threshold; 6123 6124 i = 0; 6125 clStatus=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 6126 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer); 6127 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); 6128 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns); 6129 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows); 6130 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, (chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL); 6131 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, kernelWidth*sizeof(float),NULL); 6132 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(ChannelType),&image->channel_mask); 6133 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer); 6134 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth); 6135 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fGain); 6136 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fThreshold); 6137 6138 if (clStatus != CL_SUCCESS) 6139 { 6140 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 6141 goto cleanup; 6142 } 6143 } 6144 6145 /* launch the kernel */ 6146 { 6147 size_t gsize[2]; 6148 size_t wsize[2]; 6149 6150 gsize[0] = image->columns; 6151 gsize[1] = chunkSize*((image->rows+chunkSize-1)/chunkSize); 6152 wsize[0] = 1; 6153 wsize[1] = chunkSize; 6154 6155 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, &event); 6156 if (clStatus != CL_SUCCESS) 6157 { 6158 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 6159 goto cleanup; 6160 } 6161 clEnv->library->clFlush(queue); 6162 RecordProfileData(clEnv,UnsharpMaskBlurColumnKernel,event); 6163 clEnv->library->clReleaseEvent(event); 6164 } 6165 6166 } 6167 6168 /* get result */ 6169 if (ALIGNED(filteredPixels,CLPixelPacket)) 6170 { 6171 length = image->columns * image->rows; 6172 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); 6173 } 6174 else 6175 { 6176 length = image->columns * image->rows; 6177 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL); 6178 } 6179 if (clStatus != CL_SUCCESS) 6180 { 6181 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 6182 goto cleanup; 6183 } 6184 6185 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); 6186 6187cleanup: 6188 OpenCLLogException(__FUNCTION__,__LINE__,exception); 6189 6190 image_view=DestroyCacheView(image_view); 6191 if (filteredImage_view != NULL) 6192 filteredImage_view=DestroyCacheView(filteredImage_view); 6193 6194 if (kernel != NULL) kernel=DestroyKernelInfo(kernel); 6195 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); 6196 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); 6197 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer); 6198 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer); 6199 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel); 6200 if (unsharpMaskBlurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskBlurColumnKernel); 6201 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 6202 if (outputReady == MagickFalse) 6203 { 6204 if (filteredImage != NULL) 6205 { 6206 DestroyImage(filteredImage); 6207 filteredImage = NULL; 6208 } 6209 } 6210 return(filteredImage); 6211} 6212 6213static Image *ComputeUnsharpMaskImageSingle(const Image *image, 6214 const double radius,const double sigma,const double gain, 6215 const double threshold,int blurOnly,ExceptionInfo *exception) 6216{ 6217 CacheView 6218 *filteredImage_view, 6219 *image_view; 6220 6221 cl_command_queue 6222 queue; 6223 6224 cl_context 6225 context; 6226 6227 cl_int 6228 justBlur, 6229 clStatus; 6230 6231 cl_kernel 6232 unsharpMaskKernel; 6233 6234 cl_event 6235 event; 6236 6237 cl_mem 6238 filteredImageBuffer, 6239 imageBuffer, 6240 imageKernelBuffer; 6241 6242 cl_uint 6243 i, 6244 imageColumns, 6245 imageRows, 6246 kernelWidth, 6247 number_channels; 6248 6249 float 6250 fGain, 6251 fThreshold; 6252 6253 Image 6254 *filteredImage; 6255 6256 MagickBooleanType 6257 outputReady; 6258 6259 MagickCLEnv 6260 clEnv; 6261 6262 void 6263 *filteredPixels; 6264 6265 clEnv = NULL; 6266 filteredImage = NULL; 6267 filteredImage_view = NULL; 6268 filteredPixels = NULL; 6269 context = NULL; 6270 imageBuffer = NULL; 6271 filteredImageBuffer = NULL; 6272 imageKernelBuffer = NULL; 6273 unsharpMaskKernel = NULL; 6274 queue = NULL; 6275 outputReady = MagickFalse; 6276 6277 clEnv = GetDefaultOpenCLEnv(); 6278 context = GetOpenCLContext(clEnv); 6279 queue = AcquireOpenCLCommandQueue(clEnv); 6280 6281 image_view=AcquireVirtualCacheView(image,exception); 6282 imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); 6283 if (imageBuffer == (cl_mem) NULL) 6284 goto cleanup; 6285 6286 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception); 6287 if (filteredImage == (Image *) NULL) 6288 goto cleanup; 6289 6290 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); 6291 filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, 6292 context,filteredPixels,exception); 6293 if (filteredImageBuffer == (void *) NULL) 6294 goto cleanup; 6295 6296 imageKernelBuffer=createKernelInfo(clEnv,context,queue,radius,sigma, 6297 &kernelWidth,exception); 6298 6299 { 6300 /* get the opencl kernel */ 6301 { 6302 unsharpMaskKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMask"); 6303 if (unsharpMaskKernel == NULL) 6304 { 6305 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 6306 goto cleanup; 6307 }; 6308 } 6309 6310 { 6311 imageColumns = (cl_uint) image->columns; 6312 imageRows = (cl_uint) image->rows; 6313 number_channels = (cl_uint) image->number_channels; 6314 fGain = (float) gain; 6315 fThreshold = (float) threshold; 6316 justBlur = blurOnly; 6317 6318 /* set the kernel arguments */ 6319 i = 0; 6320 clStatus=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); 6321 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&number_channels); 6322 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(ChannelType),(void *)&image->channel_mask); 6323 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer); 6324 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&kernelWidth); 6325 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&imageColumns); 6326 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&imageRows); 6327 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_float4)*(8 * (32 + kernelWidth)),(void *) NULL); 6328 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fGain); 6329 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fThreshold); 6330 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&justBlur); 6331 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); 6332 if (clStatus != CL_SUCCESS) 6333 { 6334 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 6335 goto cleanup; 6336 } 6337 } 6338 6339 /* launch the kernel */ 6340 { 6341 size_t gsize[2]; 6342 size_t wsize[2]; 6343 6344 gsize[0] = ((image->columns + 7) / 8) * 8; 6345 gsize[1] = ((image->rows + 31) / 32) * 32; 6346 wsize[0] = 8; 6347 wsize[1] = 32; 6348 6349 clStatus = clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskKernel, 2, NULL, gsize, wsize, 0, NULL, &event); 6350 if (clStatus != CL_SUCCESS) 6351 { 6352 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 6353 goto cleanup; 6354 } 6355 clEnv->library->clFlush(queue); 6356 RecordProfileData(clEnv,UnsharpMaskKernel,event); 6357 clEnv->library->clReleaseEvent(event); 6358 } 6359 } 6360 6361 if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) 6362 goto cleanup; 6363 6364 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); 6365 6366cleanup: 6367 OpenCLLogException(__FUNCTION__,__LINE__,exception); 6368 6369 image_view=DestroyCacheView(image_view); 6370 if (filteredImage_view != NULL) 6371 filteredImage_view=DestroyCacheView(filteredImage_view); 6372 6373 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); 6374 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); 6375 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer); 6376 if (unsharpMaskKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskKernel); 6377 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 6378 if (outputReady == MagickFalse) 6379 { 6380 if (filteredImage != NULL) 6381 { 6382 DestroyImage(filteredImage); 6383 filteredImage = NULL; 6384 } 6385 } 6386 return(filteredImage); 6387} 6388 6389MagickExport Image *AccelerateUnsharpMaskImage(const Image *image, 6390 const double radius,const double sigma,const double gain, 6391 const double threshold,ExceptionInfo *exception) 6392{ 6393 Image 6394 *filteredImage; 6395 6396 assert(image != NULL); 6397 assert(exception != (ExceptionInfo *) NULL); 6398 6399 if ((checkAccelerateConditionRGBA(image) == MagickFalse) || 6400 (checkOpenCLEnvironment(exception) == MagickFalse)) 6401 return NULL; 6402 6403 if (radius < 12.1) 6404 filteredImage=ComputeUnsharpMaskImageSingle(image,radius,sigma,gain, 6405 threshold,0,exception); 6406 else 6407 filteredImage=ComputeUnsharpMaskImage(image,radius,sigma,gain,threshold, 6408 exception); 6409 return(filteredImage); 6410} 6411 6412static Image *ComputeWaveletDenoiseImage(const Image *image, 6413 const double threshold,ExceptionInfo *exception) 6414{ 6415 CacheView 6416 *filteredImage_view, 6417 *image_view; 6418 6419 cl_command_queue 6420 queue; 6421 6422 cl_context 6423 context; 6424 6425 cl_int 6426 clStatus; 6427 6428 cl_kernel 6429 denoiseKernel; 6430 6431 cl_event 6432 event; 6433 6434 cl_mem 6435 filteredImageBuffer, 6436 imageBuffer; 6437 6438 Image 6439 *filteredImage; 6440 6441 MagickBooleanType 6442 outputReady; 6443 6444 MagickCLEnv 6445 clEnv; 6446 6447 void 6448 *filteredPixels; 6449 6450 unsigned int 6451 i; 6452 6453 filteredImage = NULL; 6454 filteredImage_view = NULL; 6455 filteredImageBuffer = NULL; 6456 filteredPixels = NULL; 6457 denoiseKernel = NULL; 6458 outputReady = MagickFalse; 6459 6460 clEnv = GetDefaultOpenCLEnv(); 6461 context = GetOpenCLContext(clEnv); 6462 queue = AcquireOpenCLCommandQueue(clEnv); 6463 6464 /* Create and initialize OpenCL buffers. */ 6465 image_view = AcquireVirtualCacheView(image, exception); 6466 imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); 6467 if (imageBuffer == (cl_mem) NULL) 6468 goto cleanup; 6469 6470 /* create output */ 6471 filteredImage=CloneImage(image,0,0,MagickTrue,exception); 6472 if (filteredImage == (Image *) NULL) 6473 goto cleanup; 6474 if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue) 6475 { 6476 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); 6477 goto cleanup; 6478 } 6479 filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception); 6480 filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, 6481 context,filteredPixels,exception); 6482 if (filteredImageBuffer == (cl_mem) NULL) 6483 goto cleanup; 6484 6485 /* get the opencl kernel */ 6486 denoiseKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "WaveletDenoise"); 6487 if (denoiseKernel == NULL) 6488 { 6489 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); 6490 goto cleanup; 6491 } 6492 6493 // Process image 6494 { 6495 const int PASSES = 5; 6496 cl_uint number_channels = (cl_uint)image->number_channels; 6497 cl_uint width = (cl_uint)image->columns; 6498 cl_uint height = (cl_uint)image->rows; 6499 cl_uint max_channels = number_channels; 6500 if ((max_channels == 4) || (max_channels == 2)) 6501 max_channels=max_channels-1; 6502 cl_float thresh = threshold; 6503 6504 /* set the kernel arguments */ 6505 i = 0; 6506 clStatus = clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_mem), (void *)&imageBuffer); 6507 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_mem), (void *)&filteredImageBuffer); 6508 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_uint), (void *)&number_channels); 6509 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_uint), (void *)&max_channels); 6510 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_float), (void *)&thresh); 6511 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_int), (void *)&PASSES); 6512 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_uint), (void *)&width); 6513 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_uint), (void *)&height); 6514 if (clStatus != CL_SUCCESS) 6515 { 6516 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); 6517 goto cleanup; 6518 } 6519 6520 { 6521 const int TILESIZE = 64; 6522 const int PAD = 1 << (PASSES - 1); 6523 const int SIZE = TILESIZE - 2 * PAD; 6524 6525 size_t gsize[2]; 6526 size_t wsize[2]; 6527 6528 gsize[0] = ((width + (SIZE - 1)) / SIZE) * TILESIZE; 6529 gsize[1] = ((height + (SIZE - 1)) / SIZE) * 4; 6530 wsize[0] = TILESIZE; 6531 wsize[1] = 4; 6532 6533 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, denoiseKernel, 2, NULL, gsize, wsize, 0, NULL, &event); 6534 if (clStatus != CL_SUCCESS) 6535 { 6536 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); 6537 goto cleanup; 6538 } 6539 } 6540 RecordProfileData(clEnv, WaveletDenoiseKernel, event); 6541 clEnv->library->clReleaseEvent(event); 6542 } 6543 6544 if (copyWriteBuffer(image,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) 6545 { 6546 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); 6547 goto cleanup; 6548 } 6549 6550 outputReady = SyncCacheViewAuthenticPixels(filteredImage_view, exception); 6551 6552cleanup: 6553 OpenCLLogException(__FUNCTION__, __LINE__, exception); 6554 6555 image_view = DestroyCacheView(image_view); 6556 if (filteredImage_view != NULL) 6557 filteredImage_view = DestroyCacheView(filteredImage_view); 6558 6559 if (imageBuffer != NULL) clEnv->library->clReleaseMemObject(imageBuffer); 6560 if (filteredImageBuffer != NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); 6561 if (denoiseKernel != NULL) RelinquishOpenCLKernel(clEnv, denoiseKernel); 6562 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); 6563 if (outputReady == MagickFalse) 6564 { 6565 if (filteredImage != NULL) 6566 { 6567 DestroyImage(filteredImage); 6568 filteredImage = NULL; 6569 } 6570 } 6571 return(filteredImage); 6572} 6573 6574MagickExport Image *AccelerateWaveletDenoiseImage(const Image *image, 6575 const double threshold,ExceptionInfo *exception) 6576{ 6577 Image 6578 *filteredImage; 6579 6580 assert(image != NULL); 6581 assert(exception != (ExceptionInfo *)NULL); 6582 6583 if ((checkAccelerateCondition(image) == MagickFalse) || 6584 (checkOpenCLEnvironment(exception) == MagickFalse)) 6585 return (Image *) NULL; 6586 6587 filteredImage=ComputeWaveletDenoiseImage(image,threshold,exception); 6588 6589 return(filteredImage); 6590} 6591 6592#else /* MAGICKCORE_OPENCL_SUPPORT */ 6593 6594MagickExport Image *AccelerateAddNoiseImage(const Image *magick_unused(image), 6595 const NoiseType magick_unused(noise_type), 6596 ExceptionInfo *magick_unused(exception)) 6597{ 6598 magick_unreferenced(image); 6599 magick_unreferenced(noise_type); 6600 magick_unreferenced(exception); 6601 return((Image *) NULL); 6602} 6603 6604MagickExport Image *AccelerateBlurImage(const Image *magick_unused(image), 6605 const double magick_unused(radius),const double magick_unused(sigma), 6606 ExceptionInfo *magick_unused(exception)) 6607{ 6608 magick_unreferenced(image); 6609 magick_unreferenced(radius); 6610 magick_unreferenced(sigma); 6611 magick_unreferenced(exception); 6612 6613 return((Image *) NULL); 6614} 6615 6616MagickExport MagickBooleanType AccelerateCompositeImage( 6617 Image *magick_unused(image),const CompositeOperator magick_unused(compose), 6618 const Image *magick_unused(composite), 6619 const float magick_unused(destination_dissolve), 6620 const float magick_unused(source_dissolve), 6621 ExceptionInfo *magick_unused(exception)) 6622{ 6623 magick_unreferenced(image); 6624 magick_unreferenced(compose); 6625 magick_unreferenced(composite); 6626 magick_unreferenced(destination_dissolve); 6627 magick_unreferenced(source_dissolve); 6628 magick_unreferenced(exception); 6629 6630 return(MagickFalse); 6631} 6632 6633MagickExport MagickBooleanType AccelerateContrastImage( 6634 Image* magick_unused(image),const MagickBooleanType magick_unused(sharpen), 6635 ExceptionInfo* magick_unused(exception)) 6636{ 6637 magick_unreferenced(image); 6638 magick_unreferenced(sharpen); 6639 magick_unreferenced(exception); 6640 6641 return(MagickFalse); 6642} 6643 6644MagickExport MagickBooleanType AccelerateContrastStretchImage( 6645 Image *magick_unused(image),const double magick_unused(black_point), 6646 const double magick_unused(white_point), 6647 ExceptionInfo* magick_unused(exception)) 6648{ 6649 magick_unreferenced(image); 6650 magick_unreferenced(black_point); 6651 magick_unreferenced(white_point); 6652 magick_unreferenced(exception); 6653 6654 return(MagickFalse); 6655} 6656 6657MagickExport Image *AccelerateConvolveImage(const Image *magick_unused(image), 6658 const KernelInfo *magick_unused(kernel), 6659 ExceptionInfo *magick_unused(exception)) 6660{ 6661 magick_unreferenced(image); 6662 magick_unreferenced(kernel); 6663 magick_unreferenced(exception); 6664 6665 return((Image *) NULL); 6666} 6667 6668MagickExport MagickBooleanType AccelerateEqualizeImage( 6669 Image* magick_unused(image),ExceptionInfo* magick_unused(exception)) 6670{ 6671 magick_unreferenced(image); 6672 magick_unreferenced(exception); 6673 6674 return(MagickFalse); 6675} 6676 6677MagickExport Image *AccelerateDespeckleImage(const Image* magick_unused(image), 6678 ExceptionInfo* magick_unused(exception)) 6679{ 6680 magick_unreferenced(image); 6681 magick_unreferenced(exception); 6682 6683 return((Image *) NULL); 6684} 6685 6686MagickExport MagickBooleanType AccelerateFunctionImage( 6687 Image *magick_unused(image), 6688 const MagickFunction magick_unused(function), 6689 const size_t magick_unused(number_parameters), 6690 const double *magick_unused(parameters), 6691 ExceptionInfo *magick_unused(exception)) 6692{ 6693 magick_unreferenced(image); 6694 magick_unreferenced(function); 6695 magick_unreferenced(number_parameters); 6696 magick_unreferenced(parameters); 6697 magick_unreferenced(exception); 6698 6699 return(MagickFalse); 6700} 6701 6702MagickExport MagickBooleanType AccelerateGrayscaleImage( 6703 Image *magick_unused(image),const PixelIntensityMethod magick_unused(method), 6704 ExceptionInfo *magick_unused(exception)) 6705{ 6706 magick_unreferenced(image); 6707 magick_unreferenced(method); 6708 magick_unreferenced(exception); 6709 6710 return(MagickFalse); 6711} 6712 6713MagickExport Image *AccelerateLocalContrastImage( 6714 const Image *magick_unused(image),const double magick_unused(radius), 6715 const double magick_unused(strength),ExceptionInfo *magick_unused(exception)) 6716{ 6717 magick_unreferenced(image); 6718 magick_unreferenced(radius); 6719 magick_unreferenced(strength); 6720 magick_unreferenced(exception); 6721 6722 return((Image *) NULL); 6723} 6724 6725MagickExport MagickBooleanType AccelerateModulateImage( 6726 Image *magick_unused(image),const double magick_unused(percent_brightness), 6727 const double magick_unused(percent_hue), 6728 const double magick_unused(percent_saturation), 6729 ColorspaceType magick_unused(colorspace), 6730 ExceptionInfo *magick_unused(exception)) 6731{ 6732 magick_unreferenced(image); 6733 magick_unreferenced(percent_brightness); 6734 magick_unreferenced(percent_hue); 6735 magick_unreferenced(percent_saturation); 6736 magick_unreferenced(colorspace); 6737 magick_unreferenced(exception); 6738 6739 return(MagickFalse); 6740} 6741 6742MagickExport Image *AccelerateMotionBlurImage( 6743 const Image *magick_unused(image),const double *magick_unused(kernel), 6744 const size_t magick_unused(width),const OffsetInfo *magick_unused(offset), 6745 ExceptionInfo *magick_unused(exception)) 6746{ 6747 magick_unreferenced(image); 6748 magick_unreferenced(kernel); 6749 magick_unreferenced(width); 6750 magick_unreferenced(offset); 6751 magick_unreferenced(exception); 6752 6753 return((Image *) NULL); 6754} 6755 6756MagickExport MagickBooleanType AccelerateRandomImage( 6757 Image *magick_unused(image),ExceptionInfo *magick_unused(exception)) 6758{ 6759 magick_unreferenced(image); 6760 magick_unreferenced(exception); 6761 6762 return MagickFalse; 6763} 6764 6765MagickExport Image *AccelerateResizeImage(const Image *magick_unused(image), 6766 const size_t magick_unused(resizedColumns), 6767 const size_t magick_unused(resizedRows), 6768 const ResizeFilter *magick_unused(resizeFilter), 6769 ExceptionInfo *magick_unused(exception)) 6770{ 6771 magick_unreferenced(image); 6772 magick_unreferenced(resizedColumns); 6773 magick_unreferenced(resizedRows); 6774 magick_unreferenced(resizeFilter); 6775 magick_unreferenced(exception); 6776 6777 return((Image *) NULL); 6778} 6779 6780MagickExport Image *AccelerateRotationalBlurImage( 6781 const Image *magick_unused(image),const double magick_unused(angle), 6782 ExceptionInfo *magick_unused(exception)) 6783{ 6784 magick_unreferenced(image); 6785 magick_unreferenced(angle); 6786 magick_unreferenced(exception); 6787 6788 return((Image *) NULL); 6789} 6790 6791MagickExport Image *AccelerateUnsharpMaskImage( 6792 const Image *magick_unused(image),const double magick_unused(radius), 6793 const double magick_unused(sigma),const double magick_unused(gain), 6794 const double magick_unused(threshold), 6795 ExceptionInfo *magick_unused(exception)) 6796{ 6797 magick_unreferenced(image); 6798 magick_unreferenced(radius); 6799 magick_unreferenced(sigma); 6800 magick_unreferenced(gain); 6801 magick_unreferenced(threshold); 6802 magick_unreferenced(exception); 6803 6804 return((Image *) NULL); 6805} 6806 6807MagickExport Image *AccelerateWaveletDenoiseImage( 6808 const Image *magick_unused(image),const double magick_unused(threshold), 6809 ExceptionInfo *magick_unused(exception)) 6810{ 6811 magick_unreferenced(image); 6812 magick_unreferenced(threshold); 6813 magick_unreferenced(exception); 6814 6815 return((Image *)NULL); 6816} 6817#endif /* MAGICKCORE_OPENCL_SUPPORT */ 6818