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